From 82bad4bd6cb9f285ad541ccc0b91f16784f73c6b Mon Sep 17 00:00:00 2001 From: Monique Dewanchand Date: Wed, 20 Jun 2012 20:05:21 +0000 Subject: Refactoring of tiles opencl implementation: - Moved methods from NodeOperation to OpenCLDevice - Added check on Nvidia for local size --- source/blender/compositor/intern/COM_Device.h | 4 - source/blender/compositor/intern/COM_Node.h | 1 + .../compositor/intern/COM_NodeOperation.cpp | 115 ------------------- .../blender/compositor/intern/COM_NodeOperation.h | 17 +-- .../blender/compositor/intern/COM_OpenCLDevice.cpp | 122 ++++++++++++++++++++- .../blender/compositor/intern/COM_OpenCLDevice.h | 23 +++- source/blender/compositor/intern/COM_WorkPackage.h | 2 +- .../compositor/intern/COM_WorkScheduler.cpp | 5 +- .../operations/COM_BokehBlurOperation.cpp | 19 ++-- .../compositor/operations/COM_BokehBlurOperation.h | 2 +- .../operations/COM_DilateErodeOperation.cpp | 29 ++--- .../operations/COM_DilateErodeOperation.h | 4 +- .../operations/COM_WriteBufferOperation.cpp | 11 +- .../operations/COM_WriteBufferOperation.h | 2 +- 14 files changed, 185 insertions(+), 171 deletions(-) diff --git a/source/blender/compositor/intern/COM_Device.h b/source/blender/compositor/intern/COM_Device.h index 08fdb5bb578..2a86382a191 100644 --- a/source/blender/compositor/intern/COM_Device.h +++ b/source/blender/compositor/intern/COM_Device.h @@ -23,11 +23,7 @@ #ifndef _COM_Device_h #define _COM_Device_h -#include "COM_ExecutionSystem.h" #include "COM_WorkPackage.h" -#include "COM_NodeOperation.h" -#include "BLI_rect.h" -#include "COM_MemoryBuffer.h" /** * @brief Abstract class for device implementations to be used by the Compositor. diff --git a/source/blender/compositor/intern/COM_Node.h b/source/blender/compositor/intern/COM_Node.h index 12baa26cd6e..090b1455440 100644 --- a/source/blender/compositor/intern/COM_Node.h +++ b/source/blender/compositor/intern/COM_Node.h @@ -29,6 +29,7 @@ #include "COM_CompositorContext.h" #include "DNA_node_types.h" #include "BKE_text.h" +#include "COM_ExecutionSystem.h" #include #include diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp index b39b1758051..33989fa5787 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.cpp +++ b/source/blender/compositor/intern/COM_NodeOperation.cpp @@ -140,118 +140,3 @@ bool NodeOperation::determineDependingAreaOfInterest(rcti *input, ReadBufferOper return false; } } - -cl_mem NodeOperation::COM_clAttachMemoryBufferToKernelParameter(cl_context context, cl_kernel kernel, int parameterIndex, int offsetIndex, list *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; - - bool breaked = false; - for (offsety = 0; offsety < height && (!breaked); offsety += localSize) { - offset[1] = offsety; - if (offsety + localSize < height) { - size[1] = localSize; - } - else { - size[1] = height - offsety; - } - for (offsetx = 0; offsetx < width && (!breaked); 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); - if (isBreaked()) { - breaked = false; - } - } - } -} - -cl_kernel NodeOperation::COM_clCreateKernel(cl_program program, const char *kernelname, list *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; - -} diff --git a/source/blender/compositor/intern/COM_NodeOperation.h b/source/blender/compositor/intern/COM_NodeOperation.h index 30731572712..f96b994685a 100644 --- a/source/blender/compositor/intern/COM_NodeOperation.h +++ b/source/blender/compositor/intern/COM_NodeOperation.h @@ -22,9 +22,7 @@ #ifndef _COM_NodeOperation_h #define _COM_NodeOperation_h - -class NodeOperation; - +class OpenCLDevice; #include "COM_Node.h" #include #include @@ -150,7 +148,7 @@ public: * @param memoryBuffers all input MemoryBuffer's needed * @param outputBuffer the outputbuffer to write to */ - virtual void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, + virtual void executeOpenCLRegion(OpenCLDevice* device, rcti *rect, unsigned int chunkNumber, MemoryBuffer **memoryBuffers, MemoryBuffer *outputBuffer) {} /** @@ -165,7 +163,7 @@ public: * @param clMemToCleanUp all created cl_mem references must be added to this list. Framework will clean this after execution * @param clKernelsToCleanUp all created cl_kernel references must be added to this list. Framework will clean this after execution */ - virtual void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp) {} + virtual void executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp) {} virtual void deinitExecution(); bool isResolutionSet() { @@ -272,15 +270,6 @@ protected: * @brief set if this NodeOperation can be scheduled on a OpenCLDevice */ void setOpenCL(bool openCL) { this->openCL = openCL; } - - static cl_mem COM_clAttachMemoryBufferToKernelParameter(cl_context context, cl_kernel kernel, int parameterIndex, int offsetIndex, list *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader *reader); - static void COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffers); - static void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer); - void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex); - static void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer); - void COM_clEnqueueRange(cl_command_queue queue, cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex); - cl_kernel COM_clCreateKernel(cl_program program, const char *kernelname, list *clKernelsToCleanUp); - }; #endif diff --git a/source/blender/compositor/intern/COM_OpenCLDevice.cpp b/source/blender/compositor/intern/COM_OpenCLDevice.cpp index 9d005804098..c9d27b8543c 100644 --- a/source/blender/compositor/intern/COM_OpenCLDevice.cpp +++ b/source/blender/compositor/intern/COM_OpenCLDevice.cpp @@ -23,13 +23,15 @@ #include "COM_OpenCLDevice.h" #include "COM_WorkScheduler.h" +typedef enum COM_VendorID {NVIDIA=0x10DE, AMD=0x1002} COM_VendorID; -OpenCLDevice::OpenCLDevice(cl_context context, cl_device_id device, cl_program program) +OpenCLDevice::OpenCLDevice(cl_context context, cl_device_id device, cl_program program, cl_int vendorId) { this->device = device; this->context = context; this->program = program; this->queue = NULL; + this->vendorID = vendorId; } bool OpenCLDevice::initialize() @@ -56,10 +58,126 @@ void OpenCLDevice::execute(WorkPackage *work) MemoryBuffer **inputBuffers = executionGroup->getInputBuffersOpenCL(chunkNumber); MemoryBuffer *outputBuffer = executionGroup->allocateOutputBuffer(chunkNumber, &rect); - executionGroup->getOutputNodeOperation()->executeOpenCLRegion(this->context, this->program, this->queue, &rect, + executionGroup->getOutputNodeOperation()->executeOpenCLRegion(this, &rect, chunkNumber, inputBuffers, outputBuffer); delete outputBuffer; executionGroup->finalizeChunkExecution(chunkNumber, inputBuffers); } + +cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex, list *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(this->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 OpenCLDevice::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 OpenCLDevice::COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex, NodeOperation* operation) +{ + if (offsetIndex != -1) { + cl_int error; + cl_int2 offset = {operation->getWidth(), operation->getHeight()}; + + error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + } +} + +void OpenCLDevice::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 OpenCLDevice::COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer) +{ + cl_int error; + const size_t size[] = {outputMemoryBuffer->getWidth(), outputMemoryBuffer->getHeight()}; + + error = clEnqueueNDRangeKernel(this->queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } +} + +void OpenCLDevice::COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex, NodeOperation* operation) +{ + cl_int error; + const int width = outputMemoryBuffer->getWidth(); + const int height = outputMemoryBuffer->getHeight(); + int offsetx; + int offsety; + int localSize = 1024; + size_t size[2]; + cl_int2 offset; + + if (this->vendorID == NVIDIA){localSize = 32;} + bool breaked = false; + for (offsety = 0; offsety < height && (!breaked); offsety += localSize) { + offset[1] = offsety; + if (offsety + localSize < height) { + size[1] = localSize; + } + else { + size[1] = height - offsety; + } + for (offsetx = 0; offsetx < width && (!breaked); 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(this->queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + clFlush(this->queue); + if (operation->isBreaked()) { + breaked = false; + } + } + } +} + +cl_kernel OpenCLDevice::COM_clCreateKernel(const char *kernelname, list *clKernelsToCleanUp) +{ + cl_int error; + cl_kernel kernel = clCreateKernel(this->program, kernelname, &error); + if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + else { + if (clKernelsToCleanUp) clKernelsToCleanUp->push_back(kernel); + } + return kernel; + +} diff --git a/source/blender/compositor/intern/COM_OpenCLDevice.h b/source/blender/compositor/intern/COM_OpenCLDevice.h index 83ce8cec811..d132f330651 100644 --- a/source/blender/compositor/intern/COM_OpenCLDevice.h +++ b/source/blender/compositor/intern/COM_OpenCLDevice.h @@ -29,7 +29,6 @@ class OpenCLDevice; #include "OCL_opencl.h" #include "COM_WorkScheduler.h" - /** * @brief device representing an GPU OpenCL device. * an instance of this class represents a single cl_device @@ -55,13 +54,21 @@ private: * @brief opencl command queue */ cl_command_queue queue; + + /** + * @brief opencl vendor ID + */ + cl_int vendorID; + public: /** * @brief constructor with opencl device * @param context * @param device + * @param program + * @param vendorID */ - OpenCLDevice(cl_context context, cl_device_id device, cl_program program); + OpenCLDevice(cl_context context, cl_device_id device, cl_program program, cl_int vendorId); /** @@ -83,6 +90,18 @@ public: * @param work the WorkPackage to execute */ void execute(WorkPackage *work); + + cl_context getContext(){return this->context;} + + cl_command_queue getQueue(){return this->queue;} + + cl_mem COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex, list *cleanup, MemoryBuffer **inputMemoryBuffers, SocketReader *reader); + void COM_clAttachMemoryBufferOffsetToKernelParameter(cl_kernel kernel, int offsetIndex, MemoryBuffer *memoryBuffers); + void COM_clAttachOutputMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, cl_mem clOutputMemoryBuffer); + void COM_clAttachSizeToKernelParameter(cl_kernel kernel, int offsetIndex, NodeOperation* operation); + void COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer); + void COM_clEnqueueRange(cl_kernel kernel, MemoryBuffer *outputMemoryBuffer, int offsetIndex, NodeOperation* operation); + cl_kernel COM_clCreateKernel(const char *kernelname, list *clKernelsToCleanUp); }; #endif diff --git a/source/blender/compositor/intern/COM_WorkPackage.h b/source/blender/compositor/intern/COM_WorkPackage.h index 18d83cc151c..fed87186d20 100644 --- a/source/blender/compositor/intern/COM_WorkPackage.h +++ b/source/blender/compositor/intern/COM_WorkPackage.h @@ -24,7 +24,7 @@ class WorkPackage; #ifndef _COM_WorkPackage_h_ #define _COM_WorkPackage_h_ - +class ExecutionGroup; #include "COM_ExecutionGroup.h" /** diff --git a/source/blender/compositor/intern/COM_WorkScheduler.cpp b/source/blender/compositor/intern/COM_WorkScheduler.cpp index a410c28f47d..12c0f28ec9b 100644 --- a/source/blender/compositor/intern/COM_WorkScheduler.cpp +++ b/source/blender/compositor/intern/COM_WorkScheduler.cpp @@ -257,7 +257,10 @@ void WorkScheduler::initialize() unsigned int indexDevices; for (indexDevices = 0; indexDevices < totalNumberOfDevices; indexDevices++) { cl_device_id device = cldevices[indexDevices]; - OpenCLDevice *clDevice = new OpenCLDevice(context, device, program); + cl_int vendorID = 0; + cl_int error = clGetDeviceInfo(device, CL_DEVICE_VENDOR_ID, sizeof(cl_int), &vendorID, NULL); + if (error!= CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + OpenCLDevice *clDevice = new OpenCLDevice(context, device, program, vendorID); clDevice->initialize(), gpudevices.push_back(clDevice); if (G.f & G_DEBUG) { diff --git a/source/blender/compositor/operations/COM_BokehBlurOperation.cpp b/source/blender/compositor/operations/COM_BokehBlurOperation.cpp index e2fce504791..9fe5abcb075 100644 --- a/source/blender/compositor/operations/COM_BokehBlurOperation.cpp +++ b/source/blender/compositor/operations/COM_BokehBlurOperation.cpp @@ -22,6 +22,7 @@ #include "COM_BokehBlurOperation.h" #include "BLI_math.h" +#include "COM_OpenCLDevice.h" extern "C" { #include "RE_pipeline.h" @@ -160,25 +161,25 @@ bool BokehBlurOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe } static cl_kernel kernel = 0; -void BokehBlurOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, +void BokehBlurOperation::executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp) { if (!kernel) { - kernel = COM_clCreateKernel(program, "bokehBlurKernel", NULL); + kernel = device->COM_clCreateKernel("bokehBlurKernel", NULL); } cl_int radius = this->getWidth() * this->size / 100.0f; cl_int step = this->getStep(); - COM_clAttachMemoryBufferToKernelParameter(context, kernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->inputBoundingBoxReader); - COM_clAttachMemoryBufferToKernelParameter(context, kernel, 1, 4, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); - COM_clAttachMemoryBufferToKernelParameter(context, kernel, 2, -1, clMemToCleanUp, inputMemoryBuffers, this->inputBokehProgram); - COM_clAttachOutputMemoryBufferToKernelParameter(kernel, 3, clOutputBuffer); - COM_clAttachMemoryBufferOffsetToKernelParameter(kernel, 5, outputMemoryBuffer); + device->COM_clAttachMemoryBufferToKernelParameter(kernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->inputBoundingBoxReader); + device->COM_clAttachMemoryBufferToKernelParameter(kernel, 1, 4, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); + device->COM_clAttachMemoryBufferToKernelParameter(kernel, 2, -1, clMemToCleanUp, inputMemoryBuffers, this->inputBokehProgram); + device->COM_clAttachOutputMemoryBufferToKernelParameter(kernel, 3, clOutputBuffer); + device->COM_clAttachMemoryBufferOffsetToKernelParameter(kernel, 5, outputMemoryBuffer); clSetKernelArg(kernel, 6, sizeof(cl_int), &radius); clSetKernelArg(kernel, 7, sizeof(cl_int), &step); - COM_clAttachSizeToKernelParameter(kernel, 8); + device->COM_clAttachSizeToKernelParameter(kernel, 8, this); - COM_clEnqueueRange(queue, kernel, outputMemoryBuffer, 9); + device->COM_clEnqueueRange(kernel, outputMemoryBuffer, 9, this); } diff --git a/source/blender/compositor/operations/COM_BokehBlurOperation.h b/source/blender/compositor/operations/COM_BokehBlurOperation.h index 3ec61c5ce01..853855d5c34 100644 --- a/source/blender/compositor/operations/COM_BokehBlurOperation.h +++ b/source/blender/compositor/operations/COM_BokehBlurOperation.h @@ -57,6 +57,6 @@ public: void setSize(float size) { this->size = size; } - void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp); + void executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp); }; #endif diff --git a/source/blender/compositor/operations/COM_DilateErodeOperation.cpp b/source/blender/compositor/operations/COM_DilateErodeOperation.cpp index 306a2d96985..80d1c6444eb 100644 --- a/source/blender/compositor/operations/COM_DilateErodeOperation.cpp +++ b/source/blender/compositor/operations/COM_DilateErodeOperation.cpp @@ -22,6 +22,7 @@ #include "COM_DilateErodeOperation.h" #include "BLI_math.h" +#include "COM_OpenCLDevice.h" // DilateErode Distance Threshold DilateErodeThresholdOperation::DilateErodeThresholdOperation() : NodeOperation() @@ -234,24 +235,24 @@ bool DilateDistanceOperation::determineDependingAreaOfInterest(rcti *input, Read } static cl_kernel dilateKernel = 0; -void DilateDistanceOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, +void DilateDistanceOperation::executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp) { if (!dilateKernel) { - dilateKernel = COM_clCreateKernel(program, "dilateKernel", NULL); + dilateKernel = device->COM_clCreateKernel("dilateKernel", NULL); } cl_int distanceSquared = this->distance * this->distance; cl_int scope = this->scope; - COM_clAttachMemoryBufferToKernelParameter(context, dilateKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); - COM_clAttachOutputMemoryBufferToKernelParameter(dilateKernel, 1, clOutputBuffer); - COM_clAttachMemoryBufferOffsetToKernelParameter(dilateKernel, 3, outputMemoryBuffer); + device->COM_clAttachMemoryBufferToKernelParameter(dilateKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); + device->COM_clAttachOutputMemoryBufferToKernelParameter(dilateKernel, 1, clOutputBuffer); + device->COM_clAttachMemoryBufferOffsetToKernelParameter(dilateKernel, 3, outputMemoryBuffer); clSetKernelArg(dilateKernel, 4, sizeof(cl_int), &scope); clSetKernelArg(dilateKernel, 5, sizeof(cl_int), &distanceSquared); - COM_clAttachSizeToKernelParameter(dilateKernel, 6); - COM_clEnqueueRange(queue, dilateKernel, outputMemoryBuffer, 7); + device->COM_clAttachSizeToKernelParameter(dilateKernel, 6, this); + device->COM_clEnqueueRange(dilateKernel, outputMemoryBuffer, 7, this); } // Erode Distance @@ -293,24 +294,24 @@ void ErodeDistanceOperation::executePixel(float *color, int x, int y, MemoryBuff } static cl_kernel erodeKernel = 0; -void ErodeDistanceOperation::executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, +void ErodeDistanceOperation::executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp) { if (!erodeKernel) { - erodeKernel = COM_clCreateKernel(program, "erodeKernel", NULL); + erodeKernel = device->COM_clCreateKernel("erodeKernel", NULL); } cl_int distanceSquared = this->distance * this->distance; cl_int scope = this->scope; - COM_clAttachMemoryBufferToKernelParameter(context, erodeKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); - COM_clAttachOutputMemoryBufferToKernelParameter(erodeKernel, 1, clOutputBuffer); - COM_clAttachMemoryBufferOffsetToKernelParameter(erodeKernel, 3, outputMemoryBuffer); + device->COM_clAttachMemoryBufferToKernelParameter(erodeKernel, 0, 2, clMemToCleanUp, inputMemoryBuffers, this->inputProgram); + device->COM_clAttachOutputMemoryBufferToKernelParameter(erodeKernel, 1, clOutputBuffer); + device->COM_clAttachMemoryBufferOffsetToKernelParameter(erodeKernel, 3, outputMemoryBuffer); clSetKernelArg(erodeKernel, 4, sizeof(cl_int), &scope); clSetKernelArg(erodeKernel, 5, sizeof(cl_int), &distanceSquared); - COM_clAttachSizeToKernelParameter(erodeKernel, 6); - COM_clEnqueueRange(queue, erodeKernel, outputMemoryBuffer, 7); + device->COM_clAttachSizeToKernelParameter(erodeKernel, 6, this); + device->COM_clEnqueueRange(erodeKernel, outputMemoryBuffer, 7, this); } // Dilate step diff --git a/source/blender/compositor/operations/COM_DilateErodeOperation.h b/source/blender/compositor/operations/COM_DilateErodeOperation.h index b11356129b4..4d0bf9de0ec 100644 --- a/source/blender/compositor/operations/COM_DilateErodeOperation.h +++ b/source/blender/compositor/operations/COM_DilateErodeOperation.h @@ -99,7 +99,7 @@ public: void setDistance(float distance) { this->distance = distance; } bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); - void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, + void executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp); @@ -113,7 +113,7 @@ public: */ void executePixel(float *color, int x, int y, MemoryBuffer * inputBuffers[], void *data); - void executeOpenCL(cl_context context, cl_program program, cl_command_queue queue, + void executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list *clMemToCleanUp, list *clKernelsToCleanUp); diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp index 4fff3fdcc31..356ba452185 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp @@ -23,6 +23,7 @@ #include "COM_WriteBufferOperation.h" #include "COM_defines.h" #include +#include "COM_OpenCLDevice.h" WriteBufferOperation::WriteBufferOperation() : NodeOperation() { @@ -110,7 +111,7 @@ void WriteBufferOperation::executeRegion(rcti *rect, unsigned int tileNumber, Me memoryBuffer->setCreatedState(); } -void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer **inputMemoryBuffers, MemoryBuffer *outputBuffer) +void WriteBufferOperation::executeOpenCLRegion(OpenCLDevice* device, rcti *rect, unsigned int chunkNumber, MemoryBuffer **inputMemoryBuffers, MemoryBuffer *outputBuffer) { float *outputFloatBuffer = outputBuffer->getBuffer(); cl_int error; @@ -131,7 +132,7 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr CL_FLOAT }; - cl_mem clOutputBuffer = clCreateImage2D(context, CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, &imageFormat, outputBufferWidth, outputBufferHeight, 0, outputFloatBuffer, &error); + cl_mem clOutputBuffer = clCreateImage2D(device->getContext(), CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR, &imageFormat, outputBufferWidth, outputBufferHeight, 0, outputFloatBuffer, &error); if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } // STEP 2 @@ -139,7 +140,7 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr clMemToCleanUp->push_back(clOutputBuffer); list *clKernelsToCleanUp = new list(); - this->input->executeOpenCL(context, program, queue, outputBuffer, clOutputBuffer, inputMemoryBuffers, clMemToCleanUp, clKernelsToCleanUp); + this->input->executeOpenCL(device, outputBuffer, clOutputBuffer, inputMemoryBuffers, clMemToCleanUp, clKernelsToCleanUp); // STEP 3 @@ -149,9 +150,9 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr // clFlush(queue); // clFinish(queue); - error = clEnqueueBarrier(queue); + error = clEnqueueBarrier(device->getQueue()); if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - error = clEnqueueReadImage(queue, clOutputBuffer, CL_TRUE, origin, region, 0, 0, outputFloatBuffer, 0, NULL, NULL); + error = clEnqueueReadImage(device->getQueue(), clOutputBuffer, CL_TRUE, origin, region, 0, 0, outputFloatBuffer, 0, NULL, NULL); if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } this->getMemoryProxy()->getBuffer()->copyContentFrom(outputBuffer); diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.h b/source/blender/compositor/operations/COM_WriteBufferOperation.h index 321eed7240a..ccc20584186 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.h +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.h @@ -44,7 +44,7 @@ public: void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer **memoryBuffers); void initExecution(); void deinitExecution(); - void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer **memoryBuffers, MemoryBuffer *outputBuffer); + void executeOpenCLRegion(OpenCLDevice* device, rcti *rect, unsigned int chunkNumber, MemoryBuffer **memoryBuffers, MemoryBuffer *outputBuffer); void readResolutionFromInputSocket(); }; -- cgit v1.2.3