Refactoring of tiles opencl implementation:
[blender.git] / source / blender / compositor / intern / COM_OpenCLDevice.cpp
index 9d0058040985fff4f7fde724d52a3ca24bed00a9..c9d27b8543c3a32895331af8eda0a5beaebd3a9c 100644 (file)
 #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<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(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<cl_kernel> *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;
+
+}