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 --- .../compositor/intern/COM_NodeOperation.cpp | 115 --------------------- 1 file changed, 115 deletions(-) (limited to 'source/blender/compositor/intern/COM_NodeOperation.cpp') 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; - -} -- cgit v1.2.3