Implement an OpenCL component implementation evaluator. Add OpenCL tests, working on Ubuntu and MacOS with both CPU and NVIDIA and ATI GPUs.

git-svn-id: http://svn.us.apache.org/repos/asf/tuscany@1162473 13f79535-47bb-0310-9956-ffa450edef68

11 files changed, 1380 insertions, 1 deletions

diff --git a/sca-cpp/trunk/modules/Makefile.am b/sca-cpp/trunk/modules/Makefile.am
index 278d8b40cc..f4195dc7e8 100644
--- a/sca-cpp/trunk/modules/Makefile.am
+++ b/sca-cpp/trunk/modules/Makefile.am
@@ -15,5 +15,5 @@
 #  specific language governing permissions and limitations
 #  under the License.
 
-SUBDIRS = scheme atom rss js json scdl http server python java openid oauth wsgi edit
+SUBDIRS = scheme atom rss js json scdl http server python opencl java openid oauth wsgi edit
 
diff --git a/sca-cpp/trunk/modules/opencl/Makefile.am b/sca-cpp/trunk/modules/opencl/Makefile.am
new file mode 100644
index 0000000000..95e30d20b1
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/Makefile.am
@@ -0,0 +1,57 @@
+#  Licensed to the Apache Software Foundation (ASF) under one
+#  or more contributor license agreements.  See the NOTICE file
+#  distributed with this work for additional information
+#  regarding copyright ownership.  The ASF licenses this file
+#  to you under the Apache License, Version 2.0 (the
+#  "License"); you may not use this file except in compliance
+#  with the License.  You may obtain a copy of the License at
+#  
+#    http://www.apache.org/licenses/LICENSE-2.0
+#    
+#  Unless required by applicable law or agreed to in writing,
+#  software distributed under the License is distributed on an
+#  "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+#  KIND, either express or implied.  See the License for the
+#  specific language governing permissions and limitations
+#  under the License.
+
+
+if WANT_OPENCL
+
+INCLUDES = -I${OPENCL_INCLUDE}
+
+incl_HEADERS = *.hpp
+incldir = $(prefix)/include/modules/opencl
+
+dist_mod_SCRIPTS = opencl-conf
+moddir = $(prefix)/modules/opencl
+
+EXTRA_DIST = domain-test.composite server-test.cl
+
+if DARWIN
+OPENCL_FLAGS = -framework OpenCL
+else
+OPENCL_FLAGS = -L${OPENCL_LIB} -R${OPENCL_LIB} -lOpenCL
+endif
+
+#mod_LTLIBRARIES = libmod_tuscany_opencl.la
+#libmod_tuscany_opencl_la_SOURCES = mod-opencl.cpp
+#libmod_tuscany_opencl_la_LDFLAGS = -lxml2 -lcurl -lmozjs -framework OpenCL
+#noinst_DATA = libmod_tuscany_opencl${libsuffix}
+#libmod_tuscany_opencl${libsuffix}:
+#	ln -s .libs/libmod_tuscany_opencl${libsuffix}
+
+opencl_test_SOURCES = opencl-test.cpp
+opencl_test_LDFLAGS = ${OPENCL_FLAGS}
+
+opencl_shell_SOURCES = opencl-shell.cpp
+opencl_shell_LDFLAGS = ${OPENCL_FLAGS}
+
+client_test_SOURCES = client-test.cpp
+client_test_LDFLAGS = -lxml2 -lcurl -lmozjs
+
+dist_noinst_SCRIPTS = server-test
+noinst_PROGRAMS = opencl-test opencl-shell client-test
+TESTS = opencl-test
+
+endif
diff --git a/sca-cpp/trunk/modules/opencl/client-test.cpp b/sca-cpp/trunk/modules/opencl/client-test.cpp
new file mode 100644
index 0000000000..7af3cc73d2
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/client-test.cpp
@@ -0,0 +1,39 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/* $Rev$ $Date$ */
+
+/**
+ * Test HTTP client functions.
+ */
+
+#include "stream.hpp"
+#include "string.hpp"
+#include "../server/client-test.hpp"
+
+int main() {
+    tuscany::cout << "Testing..." << tuscany::endl;
+    tuscany::server::testURI = "http://localhost:8090/opencl";
+
+    tuscany::server::testServer();
+
+    tuscany::cout << "OK" << tuscany::endl;
+
+    return 0;
+}
diff --git a/sca-cpp/trunk/modules/opencl/domain-test.composite b/sca-cpp/trunk/modules/opencl/domain-test.composite
new file mode 100644
index 0000000000..e69399d581
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/domain-test.composite
@@ -0,0 +1,31 @@
+<?xml version="1.0" encoding="UTF-8"?>
+<!--
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ * 
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ * 
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.    
+-->
+<composite xmlns="http://docs.oasis-open.org/ns/opencsa/sca/200912"
+  targetNamespace="http://domain/test"
+  name="domain-test">
+        
+    <component name="opencl-test">
+        <implementation.opencl kernel="server-test.cl"/>
+        <service name="test">
+            <binding.http uri="opencl"/>
+        </service>
+    </component>     
+
+</composite>
diff --git a/sca-cpp/trunk/modules/opencl/driver.hpp b/sca-cpp/trunk/modules/opencl/driver.hpp
new file mode 100644
index 0000000000..b4b6c2845b
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/driver.hpp
@@ -0,0 +1,62 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/* $Rev$ $Date$ */
+
+#ifndef tuscany_opencl_driver_hpp
+#define tuscany_opencl_driver_hpp
+
+/**
+ * OpenCL evaluator main driver loop.
+ */
+
+#include "string.hpp"
+#include "stream.hpp"
+#include "monad.hpp"
+#include "../scheme/driver.hpp"
+#include "eval.hpp"
+
+namespace tuscany {
+namespace opencl {
+
+const value evalDriverLoop(const OpenCLProgram& clprog, istream& in, ostream& out, const OpenCLContext& cl) {
+    scheme::promptForInput(scheme::evalInputPrompt, out);
+    value input = scheme::readValue(in);
+    if (isNil(input))
+        return input;
+    const failable<value> output = evalKernel(createKernel(input, clprog), input, 1, value::String, 512, cl);
+    scheme::announceOutput(scheme::evalOutputPrompt, out);
+    scheme::userPrint(content(output), out);
+    return evalDriverLoop(clprog, in, out, cl);
+}
+
+const bool evalDriverRun(const char* path, istream& in, ostream& out) {
+    OpenCLContext cl(OpenCLContext::DEFAULT);
+    scheme::setupDisplay(out);
+    ifstream is(path);
+    failable<OpenCLProgram> clprog = readProgram(path, is, cl);
+    if (!hasContent(clprog))
+        return true;
+    evalDriverLoop(content(clprog), in, out, cl);
+    return true;
+}
+
+}
+}
+#endif /* tuscany_opencl_driver_hpp */
diff --git a/sca-cpp/trunk/modules/opencl/eval.hpp b/sca-cpp/trunk/modules/opencl/eval.hpp
new file mode 100644
index 0000000000..ef0e028e71
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/eval.hpp
@@ -0,0 +1,725 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/* $Rev$ $Date$ */
+
+#ifndef tuscany_opencl_eval_hpp
+#define tuscany_opencl_eval_hpp
+
+/**
+ * OpenCL kernel evaluation logic.
+ */
+#ifdef IS_DARWIN
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#endif
+
+#include "list.hpp"
+#include "value.hpp"
+#include "perf.hpp"
+
+namespace tuscany {
+namespace opencl {
+
+/**
+ * Convert an OpenCL error code to a string.
+ */
+const string clError(const cl_int e) {
+    ostringstream s;
+    s << "error " << e;
+    return str(s);
+}
+
+#ifdef WANT_MAINTAINER_MODE
+
+/**
+ * OpenCL profiling counters.
+ */
+cl_ulong memtime = 0;
+cl_ulong kernelqtime = 0;
+cl_ulong kerneltime = 0;
+cl_ulong preptime = 0;
+cl_ulong evaltime = 0;
+
+/**
+ * Reset the OpenCL profiling counters.
+ */
+bool resetOpenCLCounters() {
+    memtime = kernelqtime = kerneltime = preptime = evaltime = 0;
+    return true;
+}
+
+/**
+ * Print the OpenCL profiling counters.
+ */
+bool printOpenCLCounters(const long n) {
+    cout << " kernelq " << ((double)kernelqtime / 1000000.0) / (double)n << " ms kernel " << ((double)kerneltime / 1000000.0) / (double)n << " ms memory " << ((double)memtime / 1000000.0) / (double)n << " ms prep " << ((double)preptime / 1000000.0) / (double)n << " ms eval " << ((double)evaltime / 1000000.0) / (double)n << " ms";
+    return true;
+}
+
+/**
+ * Profile a memory event.
+ */
+failable<cl_ulong> profileMemEvent(const cl_event evt) {
+    cl_ulong start = 0;
+    const cl_int serr = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
+    if (serr != CL_SUCCESS)
+        return mkfailure<cl_ulong>("Couldn't profile memory event start: " + clError(serr));
+    cl_ulong end = 0;
+    const cl_int eerr = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
+    if (eerr != CL_SUCCESS)
+        return mkfailure<cl_ulong>("Couldn't profile memory event end: " + clError(eerr));
+    const cl_ulong t = end - start;
+    memtime += t;
+    return t;
+}
+
+/**
+ * Profile a kernel event.
+ */
+failable<cl_ulong> profileKernelEvent(const cl_event evt) {
+    cl_ulong queued = 0;
+    const cl_int qerr = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_QUEUED, sizeof(cl_ulong), &queued, NULL);
+    if (qerr != CL_SUCCESS)
+        return mkfailure<cl_ulong>("Couldn't profile kernel event queue: " + clError(qerr));
+    cl_ulong start = 0;
+    const cl_int serr = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
+    if (serr != CL_SUCCESS)
+        return mkfailure<cl_ulong>("Couldn't profile kernel event start: " + clError(serr));
+    cl_ulong end = 0;
+    const cl_int eerr = clGetEventProfilingInfo(evt, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
+    if (eerr != CL_SUCCESS)
+        return mkfailure<cl_ulong>("Couldn't profile kernel event end: " + clError(eerr));
+    const cl_ulong q = start - queued;
+    kernelqtime += q;
+    const cl_ulong t = end - start;
+    kerneltime += t;
+    return t;
+}
+
+/**
+ * Profile an array of memory events.
+ */
+failable<cl_ulong> profileMemEvents(const cl_uint n, const cl_event* evt) {
+    if (n == 0)
+        return 0;
+    const failable<cl_ulong> t = profileMemEvent(*evt);
+    if (!hasContent(t))
+        return t;
+    const failable<cl_ulong> r = profileMemEvents(n - 1, evt + 1);
+    if (!hasContent(r))
+        return r;
+    return content(t) + content(r);
+}
+
+#else
+
+#define resetOpenCLCounters()
+#define printOpenCLCounters(n)
+
+#endif
+
+class OpenCLContext;
+class OpenCLProgram;
+class OpenCLKernel;
+class OpenCLBuffer;
+
+/**
+ * Represent an OpenCL context.
+ */
+class OpenCLContext {
+public:
+#define OPENCL_MAX_DEVICES 64
+
+    enum DeviceType {
+        DEFAULT = 0, CPU = 1, GPU = 2
+    };
+
+    OpenCLContext(const OpenCLContext::DeviceType devtype) : dev(OpenCLContext::DEFAULT), ndevs(0), ctx(0) {
+        debug("opencl::OpenCLContext");
+        for (int i = 0; i < OPENCL_MAX_DEVICES; i++)
+            cq[i] = 0;
+
+        // Get the available platforms
+        cl_uint nplatforms;
+        cl_platform_id platforms[16];
+        const cl_int gperr = clGetPlatformIDs(16, platforms, &nplatforms);
+        if(nplatforms == 0 || gperr != CL_SUCCESS) {
+            mkfailure<bool>("Couldn't get OpenCL platforms: " + clError(gperr));
+            return;
+        }
+        for(cl_uint i = 0; i < nplatforms; ++i) {
+            char vendor[256];
+            const cl_int gverr = clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, NULL);
+            if(gverr != CL_SUCCESS) {
+                mkfailure<bool>("Couldn't get OpenCL platform: " + clError(gverr));
+                return;
+            }
+            debug(vendor, "opencl::OpenCLContext::vendor");
+        }
+
+        // Get the available devices of the requested type
+        if (devtype == OpenCLContext::DEFAULT || devtype == OpenCLContext::GPU) {
+            for(cl_uint i = 0; i < nplatforms; ++i) {
+                const cl_int err = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_GPU, OPENCL_MAX_DEVICES, devid, &ndevs);
+                if (err == CL_SUCCESS) {
+                    debug(ndevs, "opencl::OpenCLContext::gcpus");
+                    dev = OpenCLContext::GPU;
+                    break;
+                }
+            }
+        }
+        if ((devtype == OpenCLContext::DEFAULT && ndevs == 0) || devtype == OpenCLContext::CPU) {
+            for(cl_uint i = 0; i < nplatforms; ++i) {
+                const cl_int err = clGetDeviceIDs(platforms[i], CL_DEVICE_TYPE_CPU, OPENCL_MAX_DEVICES, devid, &ndevs);
+                if (err == CL_SUCCESS) {
+                    debug(ndevs, "opencl::OpenCLContext::ncpus");
+                    dev = OpenCLContext::CPU;
+                    break;
+                }
+            }
+        }
+        if (ndevs == 0)
+            return;
+
+        // Initialize OpenCL context and command queues
+        cl_int ccerr;
+        ctx = clCreateContext(0, ndevs, devid, NULL, NULL, &ccerr);
+        if(!ctx || ccerr != CL_SUCCESS) {
+            mkfailure<bool>("Couldn't create OpenCL context: " + clError(ccerr));
+            return;
+        }
+
+        for (cl_uint i = 0; i < ndevs; i++) {
+            cl_int cqerr;
+#ifdef WANT_MAINTAINER_MODE
+            cq[i] = clCreateCommandQueue(ctx, devid[i], CL_QUEUE_PROFILING_ENABLE, &cqerr);
+#else
+            cq[i] = clCreateCommandQueue(ctx, devid[i], 0, &cqerr);
+#endif
+            if (!cq[i] || cqerr != CL_SUCCESS) {
+                mkfailure<bool>("Couldn't create OpenCL command queue: " + clError(cqerr));
+                return;
+            }
+        }
+    }
+
+    ~OpenCLContext() {
+        debug("opencl::~OpenCLContext");
+        for (cl_uint i = 0; i < ndevs; i++) {
+            if (cq[i] != 0)
+                clReleaseCommandQueue(cq[i]);
+        }
+        if (ctx != 0)
+            clReleaseContext(ctx);
+    }
+
+private:
+    OpenCLContext::DeviceType dev;
+    cl_uint ndevs;
+    cl_device_id devid[OPENCL_MAX_DEVICES];
+    cl_context ctx;
+    cl_command_queue cq[OPENCL_MAX_DEVICES];
+
+    friend const cl_uint devices(const OpenCLContext& cl);
+    friend const cl_command_queue commandq(const OpenCLContext& cl);
+    friend const failable<OpenCLBuffer> readOnlyBuffer(const size_t size, const void* p, const OpenCLContext& cl, cl_command_queue cq);
+    friend const failable<OpenCLBuffer> writeOnlyBuffer(const size_t size, const OpenCLContext& cl);
+    friend const failable<value> evalKernel(const failable<OpenCLKernel>& kernel, const value& expr, const size_t gwsize, const value::ValueType type, const size_t n, const OpenCLContext& cl);
+    friend const failable<OpenCLProgram> readProgram(const string& path, istream& is, const OpenCLContext& cl);
+};
+
+/**
+ * Return the number of computing devices available in a context.
+ */
+const cl_uint devices(const OpenCLContext& cl) {
+    return cl.ndevs;
+}
+
+/**
+ * Return a command queue from a context.
+ */
+const cl_command_queue commandq(const OpenCLContext& cl) {
+    return cl.cq[0];
+}
+
+/**
+ * Represents an OpenCL program.
+ */
+class OpenCLProgram {
+public:
+    OpenCLProgram() : prog(0) {
+    }
+
+    OpenCLProgram(const cl_program prog) : prog(prog) {
+    }
+
+    OpenCLProgram(const OpenCLProgram& c) : prog(c.prog) {
+        const cl_int rperr = clRetainProgram(prog);
+        if (rperr != CL_SUCCESS)
+            mkfailure<bool>(string("Couldn't retain OpenCL program: ") + clError(rperr));
+    }
+
+    ~OpenCLProgram() {
+        if (!prog)
+            return;
+        const cl_int rperr = clReleaseProgram(prog);
+        if (rperr != CL_SUCCESS)
+            mkfailure<bool>(string("Couldn't release OpenCL program: ") + clError(rperr));
+    }
+
+private:
+    const cl_program prog;
+
+    friend const failable<OpenCLKernel> createKernel(const value& expr, const OpenCLProgram& clprog);
+    friend const failable<value> evalKernel(const failable<OpenCLKernel>& kernel, const value& expr, const size_t gwsize, const value::ValueType type, const size_t n, const OpenCLContext& cl);
+};
+
+/**
+ * Represents an OpenCL kernel.
+ */
+class OpenCLKernel {
+public:
+    OpenCLKernel() : k(0) {
+    }
+
+    OpenCLKernel(const cl_kernel k) : k(k) {
+    }
+
+    OpenCLKernel(const OpenCLKernel& c) : k(c.k) {
+        const cl_int rkerr = clRetainKernel(k);
+        if (rkerr != CL_SUCCESS)
+            mkfailure<bool>(string("Couldn't retain OpenCL kernel: ") + clError(rkerr));
+    }
+
+    ~OpenCLKernel() {
+        if (!k)
+            return;
+        const cl_int rkerr = clReleaseKernel(k);
+        if (rkerr != CL_SUCCESS)
+            mkfailure<bool>(string("Couldn't release OpenCL kernel: ") + clError(rkerr));
+    }
+
+private:
+    const cl_kernel k;
+
+    friend const failable<OpenCLBuffer> valueToKernelArg(const cl_uint i, const size_t size, const void* p, const failable<OpenCLBuffer>& buf, const OpenCLKernel& kernel);
+    friend const failable<value> evalKernel(const failable<OpenCLKernel>& kernel, const value& expr, const size_t gwsize, const value::ValueType type, const size_t n, const OpenCLContext& cl);
+};
+
+/**
+ * Represents an OpenCL buffer.
+ */
+class OpenCLBuffer {
+public:
+    OpenCLBuffer() : mem(0), evt(0) {
+    }
+
+    OpenCLBuffer(const cl_mem mem, const cl_event evt) : mem(mem), evt(evt) {
+    }
+
+    OpenCLBuffer(const OpenCLBuffer& c) : mem(c.mem), evt(c.evt) {
+        if (mem != 0) {
+            const cl_int rmerr = clRetainMemObject(mem);
+            if (rmerr != CL_SUCCESS)
+                mkfailure<bool>(string("Couldn't retain OpenCL buffer: ") + clError(rmerr));
+        }
+        if (evt != 0) {
+            const cl_int reerr = clRetainEvent(evt);
+            if (reerr != CL_SUCCESS)
+                mkfailure<bool>(string("Couldn't retain OpenCL event: ") + clError(reerr));
+        }
+    }
+
+    ~OpenCLBuffer() {
+        if (mem != 0) {
+            const cl_int rmerr = clReleaseMemObject(mem);
+            if (rmerr != CL_SUCCESS)
+                mkfailure<bool>(string("Couldn't release OpenCL buffer: ") + clError(rmerr));
+        }
+        if (evt != 0) {
+            const cl_int reerr = clReleaseEvent(evt);
+            if (reerr != CL_SUCCESS)
+                mkfailure<bool>(string("Couldn't release OpenCL event: ") + clError(reerr));
+        }
+    }
+
+private:
+    const cl_mem mem;
+    const cl_event evt;
+
+    friend const cl_uint writeBufferEvents(const list<OpenCLBuffer>& buf, cl_event* evt);
+    friend const failable<OpenCLBuffer> valueToKernelArg(const cl_uint i, const size_t size, const void* p, const failable<OpenCLBuffer>& buf, const OpenCLKernel& kernel);
+    friend const failable<value> evalKernel(const failable<OpenCLKernel>& kernel, const value& expr, const size_t gwsize, const value::ValueType type, const size_t n, const OpenCLContext& cl);
+};
+
+/**
+ * Return a read-only memory buffer.
+ */
+const failable<OpenCLBuffer> readOnlyBuffer(const size_t size, const void* p, const OpenCLContext& cl, const cl_command_queue cq) {
+    if (cl.dev == OpenCLContext::CPU) {
+        cl_int err;
+        const cl_mem buf = clCreateBuffer(cl.ctx, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, size, const_cast<void*>(p), &err);
+        if (!buf || err != CL_SUCCESS)
+            return mkfailure<OpenCLBuffer>(string("Couldn't map OpenCL host memory: ") + clError(err));
+        return OpenCLBuffer(buf, 0);
+    }
+    cl_int berr;
+    const cl_mem buf = clCreateBuffer(cl.ctx, CL_MEM_READ_ONLY, size, NULL, &berr);
+    if (!buf || berr != CL_SUCCESS)
+        return mkfailure<OpenCLBuffer>(string("Couldn't allocate OpenCL device memory: ") + clError(berr));
+    cl_event wevt;
+    const cl_int werr = clEnqueueWriteBuffer(cq, buf, CL_FALSE, 0, size, p, 0, NULL, &wevt);
+    if (werr != CL_SUCCESS) {
+        clReleaseMemObject(buf);
+        return mkfailure<OpenCLBuffer>(string("Couldn't enqueue write to device memory: ") + clError(werr));
+    }
+    return OpenCLBuffer(buf, wevt);
+}
+
+/**
+ * Fill an array of write events for a given list of buffers.
+ */
+const cl_uint writeBufferEvents(const list<OpenCLBuffer>& buf, cl_event* evt) {
+    if (isNil(buf))
+        return 0;
+    const cl_event e = car(buf).evt;
+    if (e == 0)
+        return writeBufferEvents(cdr(buf), evt);
+    *evt = e;
+    return 1 + writeBufferEvents(cdr(buf), evt + 1);
+}
+
+/**
+ * Return a write-only memory buffer.
+ */
+const failable<OpenCLBuffer> writeOnlyBuffer(const size_t size, const OpenCLContext& cl) {
+    if (cl.dev == OpenCLContext::CPU) {
+        cl_int err;
+        const cl_mem buf = clCreateBuffer(cl.ctx, CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR, size, NULL, &err);
+        if (!buf || err != CL_SUCCESS)
+            return mkfailure<OpenCLBuffer>(string("Couldn't map OpenCL host memory: ") + clError(err));
+        return OpenCLBuffer(buf, 0);
+    }
+    cl_int err;
+    const cl_mem buf = clCreateBuffer(cl.ctx, CL_MEM_WRITE_ONLY, size, NULL, &err);
+    if (!buf || err != CL_SUCCESS)
+        return mkfailure<OpenCLBuffer>(string("Couldn't allocate OpenCL device memory: ") + clError(err));
+    return OpenCLBuffer(buf, 0);
+}
+
+/**
+ * Convert a value to a kernel arg.
+ */
+const failable<OpenCLBuffer> valueToKernelArg(const cl_uint i, const size_t size, const void* p, const failable<OpenCLBuffer>& buf, const OpenCLKernel& kernel) {
+    if (!hasContent(buf))
+        return buf;
+    if (p != NULL) {
+        const cl_int err = clSetKernelArg(kernel.k, (cl_uint)i, size, p);
+        if (err != CL_SUCCESS)
+            return mkfailure<OpenCLBuffer>(string("Couldn't set OpenCL simple kernel arg: ") + clError(err));
+        return buf;
+    }
+    const OpenCLBuffer b = content(buf);
+    const cl_int err = clSetKernelArg(kernel.k, i, sizeof(cl_mem), &b.mem);
+    if (err != CL_SUCCESS)
+        return mkfailure<OpenCLBuffer>(string("Couldn't set OpenCL buffer kernel arg: ") + clError(err));
+    return buf;
+}
+
+const failable<OpenCLBuffer> valueToKernelArg(const value& v, const cl_uint i, const OpenCLKernel& kernel, const OpenCLContext& cl, const cl_command_queue cq) {
+    switch (type(v)) {
+    case value::Symbol: {
+        const string s = string("'") + v;
+        return valueToKernelArg(i, 0, NULL, readOnlyBuffer(length(s) + 1, c_str(s), cl, cq), kernel);
+    }
+    case value::String: {
+        const string s = (string)v;
+        return valueToKernelArg(i, 0, NULL, readOnlyBuffer(length(s) + 1, c_str(s), cl, cq), kernel);
+    }
+    case value::Number: {
+        const cl_float d = (cl_float)((double)v);
+        return valueToKernelArg(i, sizeof(cl_float), &d, OpenCLBuffer(), kernel);
+    }
+    case value::Bool: {
+        const cl_int b = (cl_int)((bool)v);
+        return valueToKernelArg(i, sizeof(cl_int), &b, OpenCLBuffer(), kernel);
+    }
+    default: {
+        return valueToKernelArg(i, sizeof(cl_mem), NULL, readOnlyBuffer(sizeof(value), &v, cl, cq), kernel);
+    }
+    }
+}
+
+/**
+ * Convert a list of values to kernel args.
+ */
+const failable<list<OpenCLBuffer>> valuesToKernelArgsListHelper(const list<value>& v, const cl_uint i, const OpenCLKernel& kernel, const OpenCLContext& cl, const cl_command_queue cq) {
+    if (isNil(v))
+        return list<OpenCLBuffer>();
+    const failable<OpenCLBuffer> a = valueToKernelArg(car(v), i, kernel, cl, cq);
+    if (!hasContent(a))
+        return mkfailure<list<OpenCLBuffer>>(reason(a));
+    const failable<list<OpenCLBuffer>> al = valuesToKernelArgsListHelper(cdr(v), i + 1, kernel, cl, cq);
+    if (!hasContent(al))
+        return al;
+    return cons<OpenCLBuffer>(content(a), content(al));
+}
+
+const failable<list<OpenCLBuffer>> valuesToKernelArgs(const list<value>& v, const OpenCLKernel& kernel, const OpenCLContext& cl, const cl_command_queue cq) {
+    return valuesToKernelArgsListHelper(v, 0, kernel, cl, cq);
+}
+
+/**
+ * Convert a kernel result to a value.
+ */
+const value kernelResultToValue(const void* p, const value::ValueType type) {
+    switch(type) {
+    case value::Symbol: {
+        const char* s = static_cast<const char*>(p);
+        const size_t l = strlen(s);
+        if (l != 0 && *s == '\'')
+            return value(s + 1);
+        return value(s);
+    }
+    case value::String: {
+        const char* s = static_cast<const char*>(p);
+        const size_t l = strlen(s);
+        if (l != 0 && *s == '\'')
+            return value(s + 1);
+        return value(string(s, l));
+    }
+    case value::Number:
+        return (double)(*(static_cast<const cl_float*>(p)));
+    case value::Bool:
+        return (bool)(*(static_cast<const cl_int*>(p)));
+    default:
+        return *(static_cast<const value*>(p));
+    }
+}
+
+/**
+ * Return the value type corresponding to a C99 type name.
+ */
+const value::ValueType valueType(const string& t) {
+    if (t == "float")
+        return value::Number;
+    if (t == "int")
+        return value::Bool;
+    if (t == "char")
+        return value::String;
+    return value::Nil;
+}
+
+/**
+ * Return the size of a C99 type corresponding to a value type.
+ */
+const size_t valueSize(const value::ValueType type) {
+    switch(type) {
+    case value::Number:
+        return sizeof(cl_float);
+    case value::Bool:
+        return sizeof(cl_int);
+    case value::Symbol:
+        return sizeof(cl_char);
+    case value::String:
+        return sizeof(cl_char);
+    default:
+        return sizeof(value);
+    }
+}
+
+/**
+ * Return the result type of a kernel.
+ */
+class OpenCLResultType {
+public:
+    OpenCLResultType(const value::ValueType type, const size_t n, const size_t size) : type(type), n(n), size(size) {}
+    const value::ValueType type;
+    const size_t n;
+    const size_t size;
+};
+
+const OpenCLResultType kernelResultType(const string& fn, value::ValueType type, const size_t n) {
+    if (type != value::Nil)
+        return OpenCLResultType(type, n, valueSize(type));
+    const string s = car(tokenize("_", fn));
+    const size_t d = find_first_of(s, "0123456789");
+    if (d == length(s)) {
+        const value::ValueType vt = valueType(s);
+        return OpenCLResultType(vt, 1, valueSize(vt));
+    }
+    const value::ValueType vt = valueType(substr(s, 0, d));
+    return OpenCLResultType(vt, atoi(c_str(substr(s, d))), valueSize(vt));
+}
+
+/**
+ * Create the kernel implementing an expression.
+ */
+const failable<OpenCLKernel> createKernel(const value& expr, const OpenCLProgram& clprog) {
+
+    // Create an OpenCL kernel for the requested function
+    const value fn = car<value>(expr);
+    cl_int ckerr;
+    const cl_kernel k = clCreateKernel(clprog.prog, c_str(fn), &ckerr);
+    if (k == NULL || ckerr != CL_SUCCESS) {
+
+        // The start, stop, and restart functions are optional
+        //if (fn == "start" || fn == "stop")
+            //return value(lambda<value(const list<value>&)>());
+
+        return mkfailure<OpenCLKernel>(string("Couldn't find function: ") + car<value>(expr) + " : " + clError(ckerr));
+    }
+    return OpenCLKernel(k);
+}
+
+/**
+ * Evaluate an expression implemented by a kernel.
+ */
+const failable<value> evalKernel(const failable<OpenCLKernel>& fkernel, const value& expr, const size_t gwsize, const value::ValueType type, const size_t n, const OpenCLContext& cl) {
+
+#ifdef WANT_MAINTAINER_MODE
+    const cl_uint estart = (cl_uint)timens();
+    const cl_uint pstart = estart;
+#endif
+
+    if (!hasContent(fkernel))
+        return mkfailure<value>(reason(fkernel));
+    const OpenCLKernel kernel = content(fkernel);
+
+    // Get a command queue for the specified device type
+    const cl_command_queue cq = commandq(cl);
+
+    // Set the kernel input args
+    const failable<list<OpenCLBuffer>> args = valuesToKernelArgs(cdr<value>(expr), kernel, cl, cq);
+    if (!hasContent(args)) {
+        return mkfailure<value>(reason(args));
+    }
+
+    // Allocate result buffer in device memory
+    const value fn = car<value>(expr);
+    const OpenCLResultType rtype = kernelResultType(fn, type, n);
+    const size_t rsize = rtype.n * rtype.size;
+    const failable<OpenCLBuffer> rbuf = writeOnlyBuffer(rsize, cl);
+    if (!hasContent(rbuf))
+        return mkfailure<value>(reason(rbuf));
+
+    // Set it as a kernel output arg
+    const cl_mem rmem = content(rbuf).mem;
+    const failable<OpenCLBuffer> rarg = valueToKernelArg((cl_uint)length(cdr<value>(expr)), sizeof(cl_mem), &rmem, rbuf, kernel);
+    if (!hasContent(rarg))
+        return mkfailure<value>(reason(rarg));
+
+    // Enqueue the kernel, to be executed after all the writes complete
+    cl_event wevt[32];
+    const cl_uint nwevt = writeBufferEvents(content(args), wevt);
+    cl_event kevt;
+    const cl_int qerr = clEnqueueNDRangeKernel(cq, kernel.k, 1, NULL, &gwsize, NULL, nwevt, nwevt != 0? wevt : NULL, &kevt);
+    if (qerr != CL_SUCCESS)
+        return mkfailure<value>(string("Couldn't enqueue kernel task: ") + clError(qerr));
+
+    // Enqueue result buffer read, to be executed after the kernel completes
+    char res[rsize];
+    cl_event revt;
+    const cl_int rerr = clEnqueueReadBuffer(cq, rmem, CL_FALSE, 0, rsize, res, 1, &kevt, &revt);  
+    if (rerr != CL_SUCCESS) {
+        clReleaseEvent(kevt);
+        return mkfailure<value>(string("Couldn't read from OpenCL device memory: ") + clError(rerr));
+    }
+
+#ifdef WANT_MAINTAINER_MODE
+    const cl_uint pend = (cl_uint)timens();
+    preptime += (pend - pstart);
+#endif
+
+    // Wait for completion
+    const cl_int werr = clWaitForEvents(1, &revt);
+    if (werr != CL_SUCCESS) {
+        clReleaseEvent(revt);
+        clReleaseEvent(kevt);
+        return mkfailure<value>(string("Couldn't wait for kernel completion: ") + clError(werr));
+    }
+
+#ifdef WANT_MAINTAINER_MODE
+    profileMemEvents(nwevt, wevt);
+    profileKernelEvent(kevt);
+    profileMemEvent(revt);
+#endif
+
+    // Convert the result to a value
+    const value v = kernelResultToValue(res, rtype.type);
+
+    // Release OpenCL resources
+    clReleaseEvent(revt);
+    clReleaseEvent(kevt);
+
+#ifdef WANT_MAINTAINER_MODE
+    const cl_uint eend = (cl_uint)timens();
+    evaltime += (eend - estart);
+#endif
+
+    return v;
+}
+
+const failable<value> evalKernel(const failable<OpenCLKernel>& kernel, const value& expr, const OpenCLContext& cl) {
+    return evalKernel(kernel, expr, 1, value::Nil, 0, cl);
+}
+
+/**
+ * Read an opencl program from an input stream.
+ */
+const failable<OpenCLProgram> readProgram(const string& path, istream& is, const OpenCLContext& cl) {
+
+    // Read the program source
+    const list<string> ls = streamList(is);
+    ostringstream os;
+    write(ls, os);
+    const char* cs = c_str(str(os));
+
+    // Create the OpenCL program
+    cl_int cperr;
+    const cl_program prog = clCreateProgramWithSource(cl.ctx, 1, (const char **)&cs, NULL, &cperr);
+    if (!prog || cperr != CL_SUCCESS)
+        return mkfailure<OpenCLProgram>(string("Couldn't create OpenCL program from source: ") + path + " : " + clError(cperr));
+
+    // Built it
+    const cl_int bperr = clBuildProgram(prog, 0, NULL, NULL, NULL, NULL);
+    if(bperr != CL_SUCCESS) {
+        size_t l;
+        char b[2048];
+        clGetProgramBuildInfo(prog, cl.devid[0], CL_PROGRAM_BUILD_LOG, sizeof(b), b, &l);
+        return mkfailure<OpenCLProgram>(string("Couldn't build OpenCL program: ") + path + " : " + clError(bperr) + "\n" + string(b));
+    }
+    return OpenCLProgram(prog);
+}
+
+/**
+ * Evaluate an expression against an OpenCL program provided as an input stream.
+ */
+const failable<value> evalKernel(const value& expr, istream& is, const OpenCLContext& cl) {
+    failable<OpenCLProgram> clprog = readProgram("program.cl", is, cl);
+    if (!hasContent(clprog))
+        return mkfailure<value>(reason(clprog));
+    return evalKernel(createKernel(expr, content(clprog)), expr, 1, value::Nil, 0, cl);
+}
+
+}
+}
+#endif /* tuscany_opencl_eval_hpp */
diff --git a/sca-cpp/trunk/modules/opencl/opencl-conf b/sca-cpp/trunk/modules/opencl/opencl-conf
new file mode 100755
index 0000000000..1ba2c336a3
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/opencl-conf
@@ -0,0 +1,37 @@
+#!/bin/sh
+
+#  Licensed to the Apache Software Foundation (ASF) under one
+#  or more contributor license agreements.  See the NOTICE file
+#  distributed with this work for additional information
+#  regarding copyright ownership.  The ASF licenses this file
+#  to you under the Apache License, Version 2.0 (the
+#  "License"); you may not use this file except in compliance
+#  with the License.  You may obtain a copy of the License at
+#  
+#    http://www.apache.org/licenses/LICENSE-2.0
+#    
+#  Unless required by applicable law or agreed to in writing,
+#  software distributed under the License is distributed on an
+#  "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+#  KIND, either express or implied.  See the License for the
+#  specific language governing permissions and limitations
+#  under the License.
+
+# Generate an OpenCL server conf
+here=`echo "import os; print os.path.realpath('$0')" | python`; here=`dirname $here`
+mkdir -p $1
+root=`echo "import os; print os.path.realpath('$1')" | python`
+
+uname=`uname -s`
+if [ $uname = "Darwin" ]; then
+    libsuffix=".dylib"
+else
+    libsuffix=".so"
+fi
+
+cat >>$root/conf/modules.conf <<EOF
+# Generated by: opencl-conf $*
+# Support for OpenCL SCA components
+LoadModule mod_tuscany_eval $here/libmod_tuscany_opencl$libsuffix
+
+EOF
diff --git a/sca-cpp/trunk/modules/opencl/opencl-shell.cpp b/sca-cpp/trunk/modules/opencl/opencl-shell.cpp
new file mode 100644
index 0000000000..1dfeaeea1d
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/opencl-shell.cpp
@@ -0,0 +1,40 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/* $Rev$ $Date$ */
+
+/**
+ * OpenCL script evaluator shell, used for interactive testing of scripts.
+ */
+
+#include <assert.h>
+#include "gc.hpp"
+#include "stream.hpp"
+#include "string.hpp"
+#include "driver.hpp"
+
+int main(const int argc, char** argv) {
+    tuscany::gc_scoped_pool pool;
+    if (argc != 2) {
+        tuscany::cerr << "Usage: opencl-shell <kernel.cl>" << tuscany::endl;
+        return 1;
+    }
+    tuscany::opencl::evalDriverRun(argv[1], tuscany::cin, tuscany::cout);
+    return 0;
+}
diff --git a/sca-cpp/trunk/modules/opencl/opencl-test.cpp b/sca-cpp/trunk/modules/opencl/opencl-test.cpp
new file mode 100644
index 0000000000..17bc5ccfa6
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/opencl-test.cpp
@@ -0,0 +1,332 @@
+/*
+ * Licensed to the Apache Software Foundation (ASF) under one
+ * or more contributor license agreements.  See the NOTICE file
+ * distributed with this work for additional information
+ * regarding copyright ownership.  The ASF licenses this file
+ * to you under the Apache License, Version 2.0 (the
+ * "License"); you may not use this file except in compliance
+ * with the License.  You may obtain a copy of the License at
+ *
+ *   http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing,
+ * software distributed under the License is distributed on an
+ * "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+ * KIND, either express or implied.  See the License for the
+ * specific language governing permissions and limitations
+ * under the License.
+ */
+
+/* $Rev$ $Date$ */
+
+/**
+ * Test OpenCL kernel evaluator.
+ */
+
+#include <assert.h>
+#include "stream.hpp"
+#include "string.hpp"
+#include "driver.hpp"
+#include "parallel.hpp"
+#include "perf.hpp"
+
+namespace tuscany {
+namespace opencl {
+
+const string testFloatsCPU =
+        "kernel void add(const float x, const float y, global float* r) {\n"
+        "    float l_x;\n"
+        "    float l_y;\n"
+        "    l_x = x;\n"
+        "    l_y = y;\n"
+        "    *r = l_x + l_y;\n"
+        "}\n"
+        "kernel void float_add(const float x, const float y, global float* r) {\n"
+        "    add(x, y, r);\n"
+        "}\n";
+
+const string testFloatsGPU =
+        "kernel void add(const float x, const float y, global float* r) {\n"
+        "    local float l_x;\n"
+        "    local float l_y;\n"
+        "    l_x = x;\n"
+        "    l_y = y;\n"
+        "    barrier(CLK_LOCAL_MEM_FENCE);\n"
+        "    *r = l_x + l_y;\n"
+        "}\n"
+        "kernel void float_add(const float x, const float y, global float* r) {\n"
+        "    add(x, y, r);\n"
+        "}\n";
+
+const string testBoolsCPU =
+        "kernel void int_or(const int x, const int y, global int* r) {\n"
+        "    int l_x;\n"
+        "    int l_y;\n"
+        "    l_x = x;\n"
+        "    l_y = y;\n"
+        "    *r = l_x || l_y;\n"
+        "}\n";
+
+const string testBoolsGPU =
+        "kernel void int_or(const int x, const int y, global int* r) {\n"
+        "    local int l_x;\n"
+        "    local int l_y;\n"
+        "    l_x = x;\n"
+        "    l_y = y;\n"
+        "    barrier(CLK_LOCAL_MEM_FENCE);\n"
+        "    *r = l_x || l_y;\n"
+        "}\n";
+
+const string testCharsCPU =
+        "kernel void add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    const int ixl = (int)xl;\n"
+        "    const int iyl = (int)yl;\n"
+        "    for (int i = 0; i < ixl; i++)\n"
+        "        r[i] = x[i];\n"
+        "    for (int i = 0; i < iyl; i++)\n"
+        "        r[ixl + i] = y[i];\n"
+        "    r[ixl + iyl] = '\\0';\n"
+        "}\n"
+        "kernel void char128_add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    add(xl, x, yl, y, r);\n"
+        "}\n";
+
+const string testCharsCopyGPU =
+        "kernel void add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    const int ixl = (int)xl;\n"
+        "    const int iyl = (int)yl;\n"
+        "    local char l_x[128];\n"
+        "    local char l_y[128];\n"
+        "    event_t re = async_work_group_copy(l_x, x, ixl, 0);\n"
+        "    async_work_group_copy(l_y, y, iyl, re);\n"
+        "    wait_group_events(1, &re);\n"
+        "    local char l_r[128];\n"
+        "    for (int i = 0; i < ixl; i++)\n"
+        "        l_r[i] = l_x[i];\n"
+        "    for (int i = 0; i < iyl; i++)\n"
+        "        l_r[ixl + i] = l_y[i];\n"
+        "    l_r[ixl + iyl] = '\\0';\n"
+        "    event_t we = async_work_group_copy(r, l_r, ixl + iyl + 1, 0);\n"
+        "    wait_group_events(1, &we);\n"
+        "}\n"
+        "kernel void char128(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    add(xl, x, yl, y, r);\n"
+        "}\n";
+
+const string testCharsGPU =
+        "kernel void add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    const int ixl = (int)xl;\n"
+        "    const int iyl = (int)yl;\n"
+        "    for (int i = 0; i < ixl; i++)\n"
+        "        r[i] = x[i];\n"
+        "    for (int i = 0; i < iyl; i++)\n"
+        "        r[ixl + i] = y[i];\n"
+        "    r[ixl + iyl] = '\\0';\n"
+        "}\n"
+        "kernel void char128(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    add(xl, x, yl, y, r);\n"
+        "}\n";
+
+const string testCharsParallelCPU =
+        "kernel void add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    const int i = get_global_id(0);\n"
+        "    const int ixl = (int)xl;\n"
+        "    const int iyl = (int)yl;\n"
+        "    r[i] = i < ixl? x[i] : i < ixl + iyl? y[i - ixl] : '\\0';\n"
+        "}\n";
+
+const string testCharsParallelCopyGPU =
+        "kernel void add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    const int i = get_global_id(0);\n"
+        "    const int ixl = (int)xl;\n"
+        "    const int iyl = (int)yl;\n"
+        "    local char l_x[128];\n"
+        "    local char l_y[128];\n"
+        "    event_t re = async_work_group_copy(l_x, x, ixl, 0);\n"
+        "    async_work_group_copy(l_y, y, iyl, re);\n"
+        "    wait_group_events(1, &re);\n"
+        "    local char l_r[128];\n"
+        "    l_r[i] = i < ixl? l_x[i] : i < ixl + iyl? l_y[i - ixl] : '\\0';\n"
+        "    event_t we = async_work_group_copy(r, l_r, ixl + iyl + 1, 0);\n"
+        "    wait_group_events(1, &we);\n"
+        "}\n";
+
+const string testCharsParallelGPU =
+        "kernel void add(const float xl, global const char* x, const float yl, global const char* y, global char* r) {\n"
+        "    const int i = get_global_id(0);\n"
+        "    const int ixl = (int)xl;\n"
+        "    const int iyl = (int)yl;\n"
+        "    r[i] = i < ixl? x[i] : i < ixl + iyl? y[i - ixl] : '\\0';\n"
+        "}\n";
+
+bool testTaskParallel(const OpenCLContext::DeviceType dev) {
+    gc_scoped_pool pool;
+    OpenCLContext cl(dev);
+    if (!devices(cl) != 0)
+        return true;
+
+    {
+        istringstream is(dev == OpenCLContext::CPU? testFloatsCPU : testFloatsGPU);
+        failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+        assert(hasContent(clprog));
+        const value exp = mklist<value>("float_add", 2, 3);
+        const failable<value> r = evalKernel(createKernel(exp, content(clprog)), exp, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(5));
+    }
+    if (true) return true;
+    {
+        istringstream is(dev == OpenCLContext::CPU? testFloatsCPU : testFloatsGPU);
+        failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+        assert(hasContent(clprog));
+        const value exp = mklist<value>("add", 2, 3);
+        const failable<value> r = evalKernel(createKernel(exp, content(clprog)), exp, 1, value::Number, 1, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(5));
+    }
+    {
+        istringstream is(dev == OpenCLContext::CPU? testBoolsCPU : testBoolsGPU);
+        failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+        assert(hasContent(clprog));
+
+        const value exp = mklist<value>("int_or", true, false);
+        const failable<value> r = evalKernel(createKernel(exp, content(clprog)), exp, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(true));
+
+        const value exp2 = mklist<value>("int_or", false, false);
+        const failable<value> r2 = evalKernel(createKernel(exp2, content(clprog)), exp2, cl);
+        assert(hasContent(r2));
+        assert(content(r2) == value(false));
+    }
+    {
+        istringstream is(dev == OpenCLContext::CPU? testCharsCPU : testCharsGPU);
+        failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+        assert(hasContent(clprog));
+
+        const value exp = mklist<value>("char128", 60, string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello "), 60, string("World World World World World World World World World World "));
+        const failable<value> r = evalKernel(createKernel(exp, content(clprog)), exp, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello World World World World World World World World World World ")));
+    }
+    {
+        istringstream is(dev == OpenCLContext::CPU? testCharsCPU : testCharsGPU);
+        failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+        assert(hasContent(clprog));
+
+        const value exp = mklist<value>("add", 60, string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello "), 60, string("World World World World World World World World World World "));
+        const failable<value> r = evalKernel(createKernel(exp, content(clprog)), exp, 1, value::String, 128, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello World World World World World World World World World World ")));
+    }
+
+    return true;
+}
+
+struct evalTaskParallelLoop {
+    evalTaskParallelLoop(const OpenCLContext& cl, const OpenCLProgram& clprog) : cl(cl), clprog(clprog), exp(mklist<value>("add", 60, string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello "), 60, string("World World World World World World World World World World "))) {
+    }
+    const bool operator()() const {
+        const failable<value> r = evalKernel(createKernel(exp, clprog), exp, 1, value::String, 128, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello World World World World World World World World World World ")));
+        return true;
+    }
+
+    const OpenCLContext& cl;
+    const OpenCLProgram& clprog;
+    const value exp;
+};
+
+const bool testTaskParallelPerf(const OpenCLContext::DeviceType dev, const bool copy) {
+    gc_scoped_pool pool;
+    OpenCLContext cl(dev);
+    if (!devices(cl) != 0)
+        return true;
+
+    istringstream is(dev == OpenCLContext::CPU? testCharsCPU : copy? testCharsCopyGPU : testCharsGPU);
+    failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+    assert(hasContent(clprog));
+
+    resetOpenCLCounters();
+    const lambda<bool()> el = evalTaskParallelLoop(cl, content(clprog));
+    cout << "OpenCL task-parallel eval " << (dev == OpenCLContext::CPU? "CPU" : "GPU") << (copy? " copy" : "") << " test " << time(el, 5, 500) << " ms";
+    printOpenCLCounters(500);
+    cout << endl;
+    return true;
+}
+
+bool testDataParallel(const OpenCLContext::DeviceType dev) {
+    gc_scoped_pool pool;
+    OpenCLContext cl(dev);
+    if (!devices(cl) != 0)
+        return true;
+
+    {
+        istringstream is(dev == OpenCLContext::CPU? testCharsParallelCPU : testCharsParallelGPU);
+        failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+        assert(hasContent(clprog));
+
+        const value exp = mklist<value>("add", 6, string("Hello "), 5, string("World"));
+        const failable<value> r = evalKernel(createKernel(exp, content(clprog)), exp, 121, value::String, 128, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(string("Hello World")));
+    }
+
+    return true;
+}
+
+struct evalDataParallelLoop {
+    evalDataParallelLoop(const OpenCLContext& cl, const OpenCLProgram& clprog) : cl(cl), clprog(clprog), exp(mklist<value>("add", 60, string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello "), 60, string("World World World World World World World World World World "))) {
+    }
+    const bool operator()() const {
+        const failable<value> r = evalKernel(createKernel(exp, clprog), exp, 121, value::String, 128, cl);
+        assert(hasContent(r));
+        assert(content(r) == value(string("Hello Hello Hello Hello Hello Hello Hello Hello Hello Hello World World World World World World World World World World ")));
+        return true;
+    }
+
+    const OpenCLContext& cl;
+    const OpenCLProgram& clprog;
+    const value exp;
+};
+
+const bool testDataParallelPerf(const OpenCLContext::DeviceType dev, const bool copy) {
+    gc_scoped_pool pool;
+    OpenCLContext cl(dev);
+    if (!devices(cl) != 0)
+        return true;
+
+    istringstream is(dev == OpenCLContext::CPU? testCharsParallelCPU : copy? testCharsParallelCopyGPU : testCharsParallelGPU);
+    failable<OpenCLProgram> clprog = readProgram("kernel.cl", is, cl);
+    assert(hasContent(clprog));
+
+    resetOpenCLCounters();
+    const lambda<bool()> el = evalDataParallelLoop(cl, content(clprog));
+    cout << "OpenCL data-parallel eval " << (dev == OpenCLContext::CPU? "CPU" : "GPU") << (copy? " copy" : "") << " test " << time(el, 5, 500) << " ms";
+    printOpenCLCounters(500);
+    cout << endl;
+    return true;
+}
+
+}
+}
+
+int main() {
+    tuscany::cout << "Testing..." << tuscany::endl;
+
+    tuscany::opencl::testTaskParallel(tuscany::opencl::OpenCLContext::CPU);
+    tuscany::opencl::testTaskParallelPerf(tuscany::opencl::OpenCLContext::CPU, false);
+    tuscany::opencl::testDataParallel(tuscany::opencl::OpenCLContext::CPU);
+    tuscany::opencl::testDataParallelPerf(tuscany::opencl::OpenCLContext::CPU, false);
+
+    tuscany::opencl::testTaskParallel(tuscany::opencl::OpenCLContext::GPU);
+    tuscany::opencl::testTaskParallelPerf(tuscany::opencl::OpenCLContext::GPU, false);
+    tuscany::opencl::testTaskParallelPerf(tuscany::opencl::OpenCLContext::GPU, true);
+    tuscany::opencl::testDataParallel(tuscany::opencl::OpenCLContext::GPU);
+    tuscany::opencl::testDataParallelPerf(tuscany::opencl::OpenCLContext::GPU, false);
+    tuscany::opencl::testDataParallelPerf(tuscany::opencl::OpenCLContext::GPU, true);
+
+    tuscany::cout << "OK" << tuscany::endl;
+    return 0;
+}
diff --git a/sca-cpp/trunk/modules/opencl/server-test b/sca-cpp/trunk/modules/opencl/server-test
new file mode 100755
index 0000000000..e623599b39
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/server-test
@@ -0,0 +1,39 @@
+#!/bin/sh
+
+#  Licensed to the Apache Software Foundation (ASF) under one
+#  or more contributor license agreements.  See the NOTICE file
+#  distributed with this work for additional information
+#  regarding copyright ownership.  The ASF licenses this file
+#  to you under the Apache License, Version 2.0 (the
+#  "License"); you may not use this file except in compliance
+#  with the License.  You may obtain a copy of the License at
+#  
+#    http://www.apache.org/licenses/LICENSE-2.0
+#    
+#  Unless required by applicable law or agreed to in writing,
+#  software distributed under the License is distributed on an
+#  "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+#  KIND, either express or implied.  See the License for the
+#  specific language governing permissions and limitations
+#  under the License.
+
+# Setup
+../http/httpd-conf tmp localhost 8090 ../server/htdocs
+../server/server-conf tmp
+./opencl-conf tmp
+cat >>tmp/conf/httpd.conf <<EOF
+SCAContribution `pwd`/
+SCAComposite domain-test.composite
+EOF
+
+../http/httpd-start tmp
+sleep 2
+
+# Test
+./client-test 2>/dev/null
+rc=$?
+
+# Cleanup
+../http/httpd-stop tmp
+sleep 2
+exit $rc
diff --git a/sca-cpp/trunk/modules/opencl/server-test.cl b/sca-cpp/trunk/modules/opencl/server-test.cl
new file mode 100644
index 0000000000..de5c2d1b1e
--- /dev/null
+++ b/sca-cpp/trunk/modules/opencl/server-test.cl
@@ -0,0 +1,17 @@
+#  Licensed to the Apache Software Foundation (ASF) under one
+#  or more contributor license agreements.  See the NOTICE file
+#  distributed with this work for additional information
+#  regarding copyright ownership.  The ASF licenses this file
+#  to you under the Apache License, Version 2.0 (the
+#  "License"); you may not use this file except in compliance
+#  with the License.  You may obtain a copy of the License at
+#  
+#    http://www.apache.org/licenses/LICENSE-2.0
+#    
+#  Unless required by applicable law or agreed to in writing,
+#  software distributed under the License is distributed on an
+#  "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
+#  KIND, either express or implied.  See the License for the
+#  specific language governing permissions and limitations
+#  under the License.
+