[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [48138] trunk/blender/source/blender/ compositor: Refactoring of tiles opencl implementation:
Monique Dewanchand
m.dewanchand at atmind.nl
Wed Jun 20 22:05:26 CEST 2012
Revision: 48138
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=48138
Author: mdewanchand
Date: 2012-06-20 20:05:21 +0000 (Wed, 20 Jun 2012)
Log Message:
-----------
Refactoring of tiles opencl implementation:
- Moved methods from NodeOperation to OpenCLDevice
- Added check on Nvidia for local size
Modified Paths:
--------------
trunk/blender/source/blender/compositor/intern/COM_Device.h
trunk/blender/source/blender/compositor/intern/COM_Node.h
trunk/blender/source/blender/compositor/intern/COM_NodeOperation.cpp
trunk/blender/source/blender/compositor/intern/COM_NodeOperation.h
trunk/blender/source/blender/compositor/intern/COM_OpenCLDevice.cpp
trunk/blender/source/blender/compositor/intern/COM_OpenCLDevice.h
trunk/blender/source/blender/compositor/intern/COM_WorkPackage.h
trunk/blender/source/blender/compositor/intern/COM_WorkScheduler.cpp
trunk/blender/source/blender/compositor/operations/COM_BokehBlurOperation.cpp
trunk/blender/source/blender/compositor/operations/COM_BokehBlurOperation.h
trunk/blender/source/blender/compositor/operations/COM_DilateErodeOperation.cpp
trunk/blender/source/blender/compositor/operations/COM_DilateErodeOperation.h
trunk/blender/source/blender/compositor/operations/COM_WriteBufferOperation.cpp
trunk/blender/source/blender/compositor/operations/COM_WriteBufferOperation.h
Modified: trunk/blender/source/blender/compositor/intern/COM_Device.h
===================================================================
--- trunk/blender/source/blender/compositor/intern/COM_Device.h 2012-06-20 19:23:12 UTC (rev 48137)
+++ trunk/blender/source/blender/compositor/intern/COM_Device.h 2012-06-20 20:05:21 UTC (rev 48138)
@@ -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.
Modified: trunk/blender/source/blender/compositor/intern/COM_Node.h
===================================================================
--- trunk/blender/source/blender/compositor/intern/COM_Node.h 2012-06-20 19:23:12 UTC (rev 48137)
+++ trunk/blender/source/blender/compositor/intern/COM_Node.h 2012-06-20 20:05:21 UTC (rev 48138)
@@ -29,6 +29,7 @@
#include "COM_CompositorContext.h"
#include "DNA_node_types.h"
#include "BKE_text.h"
+#include "COM_ExecutionSystem.h"
#include <vector>
#include <string>
Modified: trunk/blender/source/blender/compositor/intern/COM_NodeOperation.cpp
===================================================================
--- trunk/blender/source/blender/compositor/intern/COM_NodeOperation.cpp 2012-06-20 19:23:12 UTC (rev 48137)
+++ trunk/blender/source/blender/compositor/intern/COM_NodeOperation.cpp 2012-06-20 20:05:21 UTC (rev 48138)
@@ -140,118 +140,3 @@
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;
-
- 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<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;
-
-}
Modified: trunk/blender/source/blender/compositor/intern/COM_NodeOperation.h
===================================================================
--- trunk/blender/source/blender/compositor/intern/COM_NodeOperation.h 2012-06-20 19:23:12 UTC (rev 48137)
+++ trunk/blender/source/blender/compositor/intern/COM_NodeOperation.h 2012-06-20 20:05:21 UTC (rev 48138)
@@ -22,9 +22,7 @@
#ifndef _COM_NodeOperation_h
#define _COM_NodeOperation_h
-
-class NodeOperation;
-
+class OpenCLDevice;
#include "COM_Node.h"
#include <string>
#include <sstream>
@@ -150,7 +148,7 @@
* @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 @@
* @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<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp) {}
+ virtual void executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp) {}
virtual void deinitExecution();
bool isResolutionSet() {
@@ -272,15 +270,6 @@
* @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<cl_mem> *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<cl_kernel> *clKernelsToCleanUp);
-
};
#endif
Modified: trunk/blender/source/blender/compositor/intern/COM_OpenCLDevice.cpp
===================================================================
--- trunk/blender/source/blender/compositor/intern/COM_OpenCLDevice.cpp 2012-06-20 19:23:12 UTC (rev 48137)
+++ trunk/blender/source/blender/compositor/intern/COM_OpenCLDevice.cpp 2012-06-20 20:05:21 UTC (rev 48138)
@@ -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 @@
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;
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list