diff options
author | Jeroen Bakker <j.bakker@atmind.nl> | 2012-06-08 13:17:07 +0400 |
---|---|---|
committer | Jeroen Bakker <j.bakker@atmind.nl> | 2012-06-08 13:17:07 +0400 |
commit | de7fe937ff24121ce8c66af902639cd96244a55f (patch) | |
tree | 625821742e2a136e10fe1cd76a81303766b8633f /source/blender/compositor/intern/COM_NodeOperation.cpp | |
parent | 95641388471d178552fea26bb477c13536bd58ef (diff) |
* Added OpenCL kernel for bokeh blur
* Uncomment COM_OPENCL_ENABLED from COM_defines.h to test
Diffstat (limited to 'source/blender/compositor/intern/COM_NodeOperation.cpp')
-rw-r--r-- | source/blender/compositor/intern/COM_NodeOperation.cpp | 108 |
1 files changed, 108 insertions, 0 deletions
diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp index fae652e39d7..650e4af5ae0 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.cpp +++ b/source/blender/compositor/intern/COM_NodeOperation.cpp @@ -124,3 +124,111 @@ bool NodeOperation::determineDependingAreaOfInterest(rcti * input, ReadBufferOpe return false; } } + +cl_mem NodeOperation::COM_clAttachMemoryBufferToKernelParameter(cl_context context, cl_kernel kernel, int parameterIndex, int offsetIndex, list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader* reader) +{ + cl_int error; + MemoryBuffer* result = (MemoryBuffer*)reader->initializeTileData(NULL, inputMemoryBuffers); + + const cl_image_format imageFormat = { + CL_RGBA, + CL_FLOAT + }; + + cl_mem clBuffer = clCreateImage2D(context, CL_MEM_READ_ONLY|CL_MEM_USE_HOST_PTR, &imageFormat, result->getWidth(), + result->getHeight(), 0, result->getBuffer(), &error); + + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + if (error == CL_SUCCESS) cleanup->push_back(clBuffer); + + error = clSetKernelArg(kernel, parameterIndex, sizeof(cl_mem), &clBuffer); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + + COM_clAttachMemoryBufferOffsetToKernelParameter(kernel, offsetIndex, result); + return clBuffer; +} + +void NodeOperation::COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffer) +{ + if (offsetIndex != -1) { + cl_int error; + rcti* rect = memoryBuffer->getRect(); + cl_int2 offset = {rect->xmin, rect->ymin}; + + error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + } +} + +void NodeOperation::COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex) +{ + if (offsetIndex != -1) { + cl_int error; + cl_int2 offset = {this->getWidth(), this->getHeight()}; + + error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + } +} + +void NodeOperation::COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer) +{ + cl_int error; + error = clSetKernelArg(kernel, parameterIndex, sizeof(cl_mem), &clOutputMemoryBuffer); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } +} + +void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer) { + cl_int error; + const size_t size[] = {outputMemoryBuffer->getWidth(),outputMemoryBuffer->getHeight()}; + + error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } +} + +void NodeOperation::COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex) { + cl_int error; + const int width = outputMemoryBuffer->getWidth(); + const int height = outputMemoryBuffer->getHeight(); + int offsetx; + int offsety; + const int localSize = 32; + size_t size[2]; + cl_int2 offset; + + for (offsety = 0 ; offsety < height; offsety+=localSize) { + offset[1] = offsety; + if (offsety+localSize < height) { + size[1] = localSize; + } else { + size[1] = height - offsety; + } + for (offsetx = 0 ; offsetx < width ; offsetx+=localSize) { + if (offsetx+localSize < width) { + size[0] = localSize; + } else { + size[0] = width - offsetx; + } + offset[0] = offsetx; + + error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + clFlush(queue); + } + } +} + +cl_kernel NodeOperation::COM_clCreateKernel(cl_program program, const char *kernelname, list<cl_kernel> *clKernelsToCleanUp) +{ + cl_int error; + cl_kernel kernel = clCreateKernel(program, kernelname, &error) ; + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); + } + else { + if (clKernelsToCleanUp) clKernelsToCleanUp->push_back(kernel); + } + return kernel; + +} |