diff options
Diffstat (limited to 'sca-cpp/branches/lightweight-sca/modules/opencl/eval.hpp')
| -rw-r--r-- | sca-cpp/branches/lightweight-sca/modules/opencl/eval.hpp | 724 |
1 files changed, 724 insertions, 0 deletions
diff --git a/sca-cpp/branches/lightweight-sca/modules/opencl/eval.hpp b/sca-cpp/branches/lightweight-sca/modules/opencl/eval.hpp new file mode 100644 index 0000000000..5606b2f57a --- /dev/null +++ b/sca-cpp/branches/lightweight-sca/modules/opencl/eval.hpp @@ -0,0 +1,724 @@ +/* + * 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. + */ +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_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() { + 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>>(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_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 */ |