diff options
Diffstat (limited to 'source/blender/compositor/intern/COM_OpenCLDevice.cpp')
| -rw-r--r-- | source/blender/compositor/intern/COM_OpenCLDevice.cpp | 334 |
1 files changed, 192 insertions, 142 deletions
diff --git a/source/blender/compositor/intern/COM_OpenCLDevice.cpp b/source/blender/compositor/intern/COM_OpenCLDevice.cpp index fd74c7ebfa4..12e072b9b12 100644 --- a/source/blender/compositor/intern/COM_OpenCLDevice.cpp +++ b/source/blender/compositor/intern/COM_OpenCLDevice.cpp @@ -19,200 +19,250 @@ #include "COM_OpenCLDevice.h" #include "COM_WorkScheduler.h" -typedef enum COM_VendorID {NVIDIA = 0x10DE, AMD = 0x1002} COM_VendorID; +typedef enum COM_VendorID { NVIDIA = 0x10DE, AMD = 0x1002 } COM_VendorID; const cl_image_format IMAGE_FORMAT_COLOR = { - CL_RGBA, - CL_FLOAT, + CL_RGBA, + CL_FLOAT, }; const cl_image_format IMAGE_FORMAT_VECTOR = { - CL_RGB, - CL_FLOAT, + CL_RGB, + CL_FLOAT, }; const cl_image_format IMAGE_FORMAT_VALUE = { - CL_R, - CL_FLOAT, + CL_R, + CL_FLOAT, }; -OpenCLDevice::OpenCLDevice(cl_context context, cl_device_id device, cl_program program, cl_int vendorId) +OpenCLDevice::OpenCLDevice(cl_context context, + cl_device_id device, + cl_program program, + cl_int vendorId) { - this->m_device = device; - this->m_context = context; - this->m_program = program; - this->m_queue = NULL; - this->m_vendorID = vendorId; + this->m_device = device; + this->m_context = context; + this->m_program = program; + this->m_queue = NULL; + this->m_vendorID = vendorId; } bool OpenCLDevice::initialize() { - cl_int error; - this->m_queue = clCreateCommandQueue(this->m_context, this->m_device, 0, &error); - return false; + cl_int error; + this->m_queue = clCreateCommandQueue(this->m_context, this->m_device, 0, &error); + return false; } void OpenCLDevice::deinitialize() { - if (this->m_queue) { - clReleaseCommandQueue(this->m_queue); - } + if (this->m_queue) { + clReleaseCommandQueue(this->m_queue); + } } void OpenCLDevice::execute(WorkPackage *work) { - const unsigned int chunkNumber = work->getChunkNumber(); - ExecutionGroup *executionGroup = work->getExecutionGroup(); - rcti rect; + const unsigned int chunkNumber = work->getChunkNumber(); + ExecutionGroup *executionGroup = work->getExecutionGroup(); + rcti rect; - executionGroup->determineChunkRect(&rect, chunkNumber); - MemoryBuffer **inputBuffers = executionGroup->getInputBuffersOpenCL(chunkNumber); - MemoryBuffer *outputBuffer = executionGroup->allocateOutputBuffer(chunkNumber, &rect); + executionGroup->determineChunkRect(&rect, chunkNumber); + MemoryBuffer **inputBuffers = executionGroup->getInputBuffersOpenCL(chunkNumber); + MemoryBuffer *outputBuffer = executionGroup->allocateOutputBuffer(chunkNumber, &rect); - executionGroup->getOutputOperation()->executeOpenCLRegion(this, &rect, - chunkNumber, inputBuffers, outputBuffer); + executionGroup->getOutputOperation()->executeOpenCLRegion( + this, &rect, chunkNumber, inputBuffers, outputBuffer); - delete outputBuffer; + delete outputBuffer; - executionGroup->finalizeChunkExecution(chunkNumber, inputBuffers); + executionGroup->finalizeChunkExecution(chunkNumber, inputBuffers); } -cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex, - list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers, +cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, + int parameterIndex, + int offsetIndex, + list<cl_mem> *cleanup, + MemoryBuffer **inputMemoryBuffers, SocketReader *reader) { - return COM_clAttachMemoryBufferToKernelParameter(kernel, parameterIndex, offsetIndex, cleanup, inputMemoryBuffers, (ReadBufferOperation *)reader); + return COM_clAttachMemoryBufferToKernelParameter(kernel, + parameterIndex, + offsetIndex, + cleanup, + inputMemoryBuffers, + (ReadBufferOperation *)reader); } const cl_image_format *OpenCLDevice::determineImageFormat(MemoryBuffer *memoryBuffer) { - const cl_image_format *imageFormat; - int num_channels = memoryBuffer->get_num_channels(); - if (num_channels == 1) { - imageFormat = &IMAGE_FORMAT_VALUE; - } - else if (num_channels == 3) { - imageFormat = &IMAGE_FORMAT_VECTOR; - } - else { - imageFormat = &IMAGE_FORMAT_COLOR; - } - - return imageFormat; + const cl_image_format *imageFormat; + int num_channels = memoryBuffer->get_num_channels(); + if (num_channels == 1) { + imageFormat = &IMAGE_FORMAT_VALUE; + } + else if (num_channels == 3) { + imageFormat = &IMAGE_FORMAT_VECTOR; + } + else { + imageFormat = &IMAGE_FORMAT_COLOR; + } + + return imageFormat; } -cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, int parameterIndex, int offsetIndex, - list<cl_mem> *cleanup, MemoryBuffer **inputMemoryBuffers, +cl_mem OpenCLDevice::COM_clAttachMemoryBufferToKernelParameter(cl_kernel kernel, + int parameterIndex, + int offsetIndex, + list<cl_mem> *cleanup, + MemoryBuffer **inputMemoryBuffers, ReadBufferOperation *reader) { - cl_int error; - - MemoryBuffer *result = reader->getInputMemoryBuffer(inputMemoryBuffers); - - const cl_image_format *imageFormat = determineImageFormat(result); - - cl_mem clBuffer = clCreateImage2D(this->m_context, CL_MEM_READ_ONLY | CL_MEM_COPY_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; + cl_int error; + + MemoryBuffer *result = reader->getInputMemoryBuffer(inputMemoryBuffers); + + const cl_image_format *imageFormat = determineImageFormat(result); + + cl_mem clBuffer = clCreateImage2D(this->m_context, + CL_MEM_READ_ONLY | CL_MEM_COPY_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) +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)); } - } + 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) +void OpenCLDevice::COM_clAttachSizeToKernelParameter(cl_kernel kernel, + int offsetIndex, + NodeOperation *operation) { - if (offsetIndex != -1) { - cl_int error; - cl_int2 offset = {{(cl_int)operation->getWidth(), (cl_int)operation->getHeight()}}; - - error = clSetKernelArg(kernel, offsetIndex, sizeof(cl_int2), &offset); - if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - } + if (offsetIndex != -1) { + cl_int error; + cl_int2 offset = {{(cl_int)operation->getWidth(), (cl_int)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) +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)); } + 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[] = {(size_t)outputMemoryBuffer->getWidth(), (size_t)outputMemoryBuffer->getHeight()}; - - error = clEnqueueNDRangeKernel(this->m_queue, kernel, 2, NULL, size, 0, 0, 0, NULL); - if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + cl_int error; + const size_t size[] = {(size_t)outputMemoryBuffer->getWidth(), + (size_t)outputMemoryBuffer->getHeight()}; + + error = clEnqueueNDRangeKernel(this->m_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) +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->m_vendorID == NVIDIA) { - localSize = 32; - } - - bool breaked = false; - for (offsety = 0; offsety < height && (!breaked); offsety += localSize) { - offset.s[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.s[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->m_queue, kernel, 2, NULL, size, 0, 0, 0, NULL); - if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - clFlush(this->m_queue); - if (operation->isBreaked()) { - breaked = false; - } - } - } + 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->m_vendorID == NVIDIA) { + localSize = 32; + } + + bool breaked = false; + for (offsety = 0; offsety < height && (!breaked); offsety += localSize) { + offset.s[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.s[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->m_queue, kernel, 2, NULL, size, 0, 0, 0, NULL); + if (error != CL_SUCCESS) { + printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); + } + clFlush(this->m_queue); + if (operation->isBreaked()) { + breaked = false; + } + } + } } -cl_kernel OpenCLDevice::COM_clCreateKernel(const char *kernelname, list<cl_kernel> *clKernelsToCleanUp) +cl_kernel OpenCLDevice::COM_clCreateKernel(const char *kernelname, + list<cl_kernel> *clKernelsToCleanUp) { - cl_int error; - cl_kernel kernel = clCreateKernel(this->m_program, kernelname, &error); - if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - else { - if (clKernelsToCleanUp) clKernelsToCleanUp->push_back(kernel); - } - return kernel; - + cl_int error; + cl_kernel kernel = clCreateKernel(this->m_program, kernelname, &error); + if (error != CL_SUCCESS) { + printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); + } + else { + if (clKernelsToCleanUp) + clKernelsToCleanUp->push_back(kernel); + } + return kernel; } |