apache-tuscany/sca-cpp/trunk/modules/opencl/eval.hpp

/*
 * 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);
}

/**
 * OpenCL profiling counters.
 */
#ifdef WANT_MAINTAINER_OPENCL_PROF

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.
 */
const bool resetOpenCLCounters() {
    memtime = kernelqtime = kerneltime = preptime = evaltime = 0;
    return true;
}

/**
 * Print the OpenCL profiling counters.
 */
const 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.
 */
const 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.
 */
const 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.
 */
const failable<cl_ulong> profileMemEvents(const cl_uint n, const cl_event* const 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_OPENCL_PROF
            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(const OpenCLContext& c) : dev(c.dev), ndevs(c.ndevs), ctx(c.ctx) {
        for (cl_uint i = 0; i < ndevs; i++) {
            devid[i] = c.devid[i];
            cq[i] = c.cq[i];
            if (cq[i] != 0) {
                const cl_int rcqerr = clRetainCommandQueue(cq[i]);
                if (rcqerr != CL_SUCCESS)
                    mkfailure<bool>(string("Couldn't retain OpenCL command queue: ") + clError(rcqerr));
            }
        }
        const cl_int rcerr = clRetainContext(ctx);
        if (rcerr != CL_SUCCESS)
            mkfailure<bool>(string("Couldn't retain OpenCL context: ") + clError(rcerr));
    }

    ~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* const 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* const evt) {
    if (isNull(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* const 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("'") + (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 (isNull(v))
        return list<OpenCLBuffer>();
    const failable<OpenCLBuffer> a = valueToKernelArg(car(v), i, kernel, cl, cq);
    if (!hasContent(a))
        return mkfailure<list<OpenCLBuffer>>(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* const p, const value::ValueType type) {
    switch(type) {
    case value::Symbol: {
        const char* const s = (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* const s = (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)(*(const cl_float*)p);
    case value::Bool:
        return (bool)(*(const cl_int*)p);
    default:
        return *(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(lvvlambda());

        return mkfailure<OpenCLKernel>(string("Couldn't find function: ") + (string)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_OPENCL_PROF
    const cl_uint estart = (cl_uint)timens();
    const cl_uint pstart = estart;
#endif

    if (!hasContent(fkernel))
        return mkfailure<value>(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>(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>(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>(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_OPENCL_PROF
    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_OPENCL_PROF
    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_OPENCL_PROF
    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>(clprog);
    return evalKernel(createKernel(expr, content(clprog)), expr, 1, value::Nil, 0, cl);
}

}
}
#endif /* tuscany_opencl_eval_hpp */