diff options
author | Jeroen Bakker <j.bakker@atmind.nl> | 2012-06-08 13:17:07 +0400 |
---|---|---|
committer | Jeroen Bakker <j.bakker@atmind.nl> | 2012-06-08 13:17:07 +0400 |
commit | de7fe937ff24121ce8c66af902639cd96244a55f (patch) | |
tree | 625821742e2a136e10fe1cd76a81303766b8633f /source/blender/compositor/operations | |
parent | 95641388471d178552fea26bb477c13536bd58ef (diff) |
* Added OpenCL kernel for bokeh blur
* Uncomment COM_OPENCL_ENABLED from COM_defines.h to test
Diffstat (limited to 'source/blender/compositor/operations')
15 files changed, 257 insertions, 49 deletions
diff --git a/source/blender/compositor/operations/COM_BokehBlurOperation.cpp b/source/blender/compositor/operations/COM_BokehBlurOperation.cpp index 1050fc57194..c48f3b0a291 100644 --- a/source/blender/compositor/operations/COM_BokehBlurOperation.cpp +++ b/source/blender/compositor/operations/COM_BokehBlurOperation.cpp @@ -34,8 +34,9 @@ BokehBlurOperation::BokehBlurOperation() : NodeOperation() this->addInputSocket(COM_DT_VALUE); this->addOutputSocket(COM_DT_COLOR); this->setComplex(true); + this->setOpenCL(true); - this->size = .01; + this->size = 1.0f; this->inputProgram = NULL; this->inputBokehProgram = NULL; @@ -90,7 +91,7 @@ void BokehBlurOperation::executePixel(float *color, int x, int y, MemoryBuffer * int bufferwidth = inputBuffer->getWidth(); int bufferstartx = inputBuffer->getRect()->xmin; int bufferstarty = inputBuffer->getRect()->ymin; - int pixelSize = this->size*this->getWidth(); + int pixelSize = this->size*this->getWidth()/100.0f; int miny = y - pixelSize; int maxy = y + pixelSize; @@ -142,10 +143,10 @@ bool BokehBlurOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe rcti newInput; rcti bokehInput; - newInput.xmax = input->xmax + (size*this->getWidth()); - newInput.xmin = input->xmin - (size*this->getWidth()); - newInput.ymax = input->ymax + (size*this->getWidth()); - newInput.ymin = input->ymin - (size*this->getWidth()); + newInput.xmax = input->xmax + (size*this->getWidth()/100.0f); + newInput.xmin = input->xmin - (size*this->getWidth()/100.0f); + newInput.ymax = input->ymax + (size*this->getWidth()/100.0f); + newInput.ymin = input->ymin - (size*this->getWidth()/100.0f); NodeOperation *operation = getInputOperation(1); bokehInput.xmax = operation->getWidth(); @@ -165,3 +166,27 @@ bool BokehBlurOperation::determineDependingAreaOfInterest(rcti *input, ReadBuffe } return false; } + +static cl_kernel kernel = 0; +void BokehBlurOperation::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) +{ + if (!kernel) { + kernel = COM_clCreateKernel(program, "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); + clSetKernelArg(kernel, 6, sizeof(cl_int), &radius); + clSetKernelArg(kernel, 7, sizeof(cl_int), &step); + COM_clAttachSizeToKernelParameter(kernel, 8); + + COM_clEnqueueRange(queue, kernel, outputMemoryBuffer, 9); +} diff --git a/source/blender/compositor/operations/COM_BokehBlurOperation.h b/source/blender/compositor/operations/COM_BokehBlurOperation.h index ce14faa8596..3cdd995b1df 100644 --- a/source/blender/compositor/operations/COM_BokehBlurOperation.h +++ b/source/blender/compositor/operations/COM_BokehBlurOperation.h @@ -56,5 +56,7 @@ public: bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); 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<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp); }; #endif diff --git a/source/blender/compositor/operations/COM_ColorCurveOperation.cpp b/source/blender/compositor/operations/COM_ColorCurveOperation.cpp index 8aee54013b1..069bbde5e73 100644 --- a/source/blender/compositor/operations/COM_ColorCurveOperation.cpp +++ b/source/blender/compositor/operations/COM_ColorCurveOperation.cpp @@ -28,6 +28,7 @@ extern "C" { #include "BKE_colortools.h" #ifdef __cplusplus } +#include "MEM_guardedalloc.h" #endif ColorCurveOperation::ColorCurveOperation(): CurveBaseOperation() @@ -59,6 +60,9 @@ void ColorCurveOperation::initExecution() void ColorCurveOperation::executePixel(float *color, float x, float y, PixelSampler sampler, MemoryBuffer *inputBuffers[]) { + CurveMapping* cumap = this->curveMapping; + CurveMapping* workingCopy = (CurveMapping*)MEM_dupallocN(cumap); + float black[4]; float white[4]; float fac[4]; @@ -67,13 +71,13 @@ void ColorCurveOperation::executePixel(float *color, float x, float y, PixelSamp this->inputBlackProgram->read(black, x, y, sampler, inputBuffers); this->inputWhiteProgram->read(white, x, y, sampler, inputBuffers); - curvemapping_set_black_white(this->curveMapping, black, white); + curvemapping_set_black_white(workingCopy, black, white); this->inputFacProgram->read(fac, x, y, sampler, inputBuffers); this->inputImageProgram->read(image, x, y, sampler, inputBuffers); if (fac[0] >= 1.0) - curvemapping_evaluate_premulRGBF(this->curveMapping, color, image); + curvemapping_evaluate_premulRGBF(workingCopy, color, image); else if (*fac<=0.0) { color[0] = image[0]; color[1] = image[1]; @@ -81,12 +85,13 @@ void ColorCurveOperation::executePixel(float *color, float x, float y, PixelSamp } else { float col[4], mfac = 1.0f-*fac; - curvemapping_evaluate_premulRGBF(this->curveMapping, col, image); + curvemapping_evaluate_premulRGBF(workingCopy, col, image); color[0] = mfac*image[0] + *fac*col[0]; color[1] = mfac*image[1] + *fac*col[1]; color[2] = mfac*image[2] + *fac*col[2]; } color[3] = image[3]; + MEM_freeN(workingCopy); } void ColorCurveOperation::deinitExecution() @@ -97,3 +102,61 @@ void ColorCurveOperation::deinitExecution() this->inputWhiteProgram = NULL; curvemapping_premultiply(this->curveMapping, 1); } + + +// Constant level curve mapping + +ConstantLevelColorCurveOperation::ConstantLevelColorCurveOperation(): CurveBaseOperation() +{ + this->addInputSocket(COM_DT_VALUE); + this->addInputSocket(COM_DT_COLOR); + this->addOutputSocket(COM_DT_COLOR); + + this->inputFacProgram = NULL; + this->inputImageProgram = NULL; + + this->setResolutionInputSocketIndex(1); +} +void ConstantLevelColorCurveOperation::initExecution() +{ + CurveBaseOperation::initExecution(); + this->inputFacProgram = this->getInputSocketReader(0); + this->inputImageProgram = this->getInputSocketReader(1); + + curvemapping_premultiply(this->curveMapping, 0); + + curvemapping_set_black_white(this->curveMapping, this->black, this->white); +} + +void ConstantLevelColorCurveOperation::executePixel(float *color, float x, float y, PixelSampler sampler, MemoryBuffer *inputBuffers[]) +{ + float fac[4]; + float image[4]; + + + this->inputFacProgram->read(fac, x, y, sampler, inputBuffers); + this->inputImageProgram->read(image, x, y, sampler, inputBuffers); + + if (fac[0] >= 1.0) + curvemapping_evaluate_premulRGBF(this->curveMapping, color, image); + else if (*fac<=0.0) { + color[0] = image[0]; + color[1] = image[1]; + color[2] = image[2]; + } + else { + float col[4], mfac = 1.0f-*fac; + curvemapping_evaluate_premulRGBF(this->curveMapping, col, image); + color[0] = mfac*image[0] + *fac*col[0]; + color[1] = mfac*image[1] + *fac*col[1]; + color[2] = mfac*image[2] + *fac*col[2]; + } + color[3] = image[3]; +} + +void ConstantLevelColorCurveOperation::deinitExecution() +{ + this->inputFacProgram = NULL; + this->inputImageProgram = NULL; + curvemapping_premultiply(this->curveMapping, 1); +} diff --git a/source/blender/compositor/operations/COM_ColorCurveOperation.h b/source/blender/compositor/operations/COM_ColorCurveOperation.h index 15f9fd25e81..6ce5befb14a 100644 --- a/source/blender/compositor/operations/COM_ColorCurveOperation.h +++ b/source/blender/compositor/operations/COM_ColorCurveOperation.h @@ -53,4 +53,37 @@ public: */ void deinitExecution(); }; + +class ConstantLevelColorCurveOperation : public CurveBaseOperation { +private: + /** + * Cached reference to the inputProgram + */ + SocketReader * inputFacProgram; + SocketReader * inputImageProgram; + float black[3]; + float white[3]; + +public: + ConstantLevelColorCurveOperation(); + + /** + * the inner loop of this program + */ + void executePixel(float *color, float x, float y, PixelSampler sampler, MemoryBuffer *inputBuffers[]); + + /** + * Initialize the execution + */ + void initExecution(); + + /** + * Deinitialize the execution + */ + void deinitExecution(); + + void setBlackLevel(float black[3]) {this->black[0] =black[0];this->black[1] =black[1];this->black[2] =black[2]; } + void setWhiteLevel(float white[3]) {this->white[0] =white[0];this->white[1] =white[1];this->white[2] =white[2]; } +}; + #endif diff --git a/source/blender/compositor/operations/COM_CompositorOperation.h b/source/blender/compositor/operations/COM_CompositorOperation.h index 41d43f896bb..13cb4f28324 100644 --- a/source/blender/compositor/operations/COM_CompositorOperation.h +++ b/source/blender/compositor/operations/COM_CompositorOperation.h @@ -63,7 +63,7 @@ public: bool isOutputOperation(bool rendering) const {return true;} void initExecution(); void deinitExecution(); - const int getRenderPriority() const {return 7;} + const CompositorPriority getRenderPriority() const {return COM_PRIORITY_MEDIUM;} void determineResolution(unsigned int resolution[], unsigned int preferredResolution[]); }; #endif diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl b/source/blender/compositor/operations/COM_OpenCLKernels.cl index 40932e54bc7..aeccfcab8b5 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl @@ -1,10 +1,52 @@ /// This file contains all opencl kernels for node-operation implementations -__kernel void testKernel(__global __write_only image2d_t output) +// Global SAMPLERS +const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; + +__constant const int2 zero = {0,0}; + +// KERNEL --- BOKEH BLUR --- +__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage, + __global __read_only image2d_t bokehImage, __global __write_only image2d_t output, + int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset) { - int x = get_global_id(0); - int y = get_global_id(1); - int2 coords = {x, y}; - float4 color = {0.0f, 1.0f, 0.0f, 1.0f}; + int2 coords = {get_global_id(0), get_global_id(1)}; + coords += offset; + float tempBoundingBox; + float4 color = {0.0f,0.0f,0.0f,0.0f}; + float4 multiplyer = {0.0f,0.0f,0.0f,0.0f}; + float4 bokeh; + const float radius2 = radius*2.0f; + const int2 realCoordinate = coords + offsetOutput; + + tempBoundingBox = read_imagef(boundingBox, SAMPLER_NEAREST, coords).s0; + + if (tempBoundingBox > 0.0f) { + const int2 bokehImageDim = get_image_dim(bokehImage); + const int2 bokehImageCenter = bokehImageDim/2; + const int2 minXY = max(realCoordinate - radius, zero); + const int2 maxXY = min(realCoordinate + radius, dimension); + int nx, ny; + + float2 uv; + int2 inputXy; + + for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny +=step, inputXy.y+=step) { + uv.y = ((realCoordinate.y-ny)/radius2)*bokehImageDim.y+bokehImageCenter.y; + + for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx +=step, inputXy.x+=step) { + uv.x = ((realCoordinate.x-nx)/radius2)*bokehImageDim.x+bokehImageCenter.x; + bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv); + color += bokeh * read_imagef(inputImage, SAMPLER_NEAREST, inputXy); + multiplyer += bokeh; + } + } + color /= multiplyer; + + } else { + int2 imageCoordinates = realCoordinate - offsetInput; + color = read_imagef(inputImage, SAMPLER_NEAREST, imageCoordinates); + } + write_imagef(output, coords, color); } diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp b/source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp deleted file mode 100644 index 1024d460044..00000000000 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl.cpp +++ /dev/null @@ -1,15 +0,0 @@ -/// @todo: this source needs to be generated from COM_OpenCLKernels.cl. -/// not implemented yet. new data to h - -const char *sourcecode = "/// This file contains all opencl kernels for node-operation implementations \n" \ -"\n" \ -"__kernel void testKernel(__global __write_only image2d_t output)\n" \ -"{\n" \ -" int x = get_global_id(0);\n" \ -" int y = get_global_id(1);\n" \ -" int2 coords = {x, y}; \n" \ -" float4 color = {0.0f, 1.0f, 0.0f, 1.0f};\n" \ -" write_imagef(output, coords, color);\n" \ -"}\n" \ -"\0\n"; - diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h new file mode 100644 index 00000000000..3cf33c75272 --- /dev/null +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h @@ -0,0 +1,55 @@ +/* clkernelstoh output of file <COM_OpenCLKernels_cl> */ + +const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all opencl kernels for node-operation implementations\n" \ +"\n" \ +"// Global SAMPLERS\n" \ +"const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n" \ +"\n" \ +"__constant const int2 zero = {0,0};\n" \ +"\n" \ +"// KERNEL --- BOKEH BLUR ---\n" \ +"__kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __global __read_only image2d_t inputImage,\n" \ +" __global __read_only image2d_t bokehImage, __global __write_only image2d_t output,\n" \ +" int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset)\n" \ +"{\n" \ +" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ +" coords += offset;\n" \ +" float tempBoundingBox;\n" \ +" float4 color = {0.0f,0.0f,0.0f,0.0f};\n" \ +" float4 multiplyer = {0.0f,0.0f,0.0f,0.0f};\n" \ +" float4 bokeh;\n" \ +" const float radius2 = radius*2.0f;\n" \ +" const int2 realCoordinate = coords + offsetOutput;\n" \ +"\n" \ +" tempBoundingBox = read_imagef(boundingBox, SAMPLER_NEAREST, coords).s0;\n" \ +"\n" \ +" if (tempBoundingBox > 0.0f) {\n" \ +" const int2 bokehImageDim = get_image_dim(bokehImage);\n" \ +" const int2 bokehImageCenter = bokehImageDim/2;\n" \ +" const int2 minXY = max(realCoordinate - radius, zero);;\n" \ +" const int2 maxXY = min(realCoordinate + radius, dimension);;\n" \ +" int nx, ny;\n" \ +"\n" \ +" float2 uv;\n" \ +" int2 inputXy;\n" \ +"\n" \ +" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny +=step, inputXy.y+=step) {\n" \ +" uv.y = ((realCoordinate.y-ny)/radius2)*bokehImageDim.y+bokehImageCenter.y;\n" \ +"\n" \ +" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx +=step, inputXy.x+=step) {\n" \ +" uv.x = ((realCoordinate.x-nx)/radius2)*bokehImageDim.x+bokehImageCenter.x;\n" \ +" bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);\n" \ +" color += bokeh * read_imagef(inputImage, SAMPLER_NEAREST, inputXy);\n" \ +" multiplyer += bokeh;\n" \ +" }\n" \ +" }\n" \ +" color /= multiplyer;\n" \ +"\n" \ +" } else {\n" \ +" int2 imageCoordinates = realCoordinate - offsetInput;\n" \ +" color = read_imagef(inputImage, SAMPLER_NEAREST, imageCoordinates);\n" \ +" }\n" \ +"\n" \ +" write_imagef(output, coords, color);\n" \ +"}\n" \ +"\0"; diff --git a/source/blender/compositor/operations/COM_OutputFileOperation.h b/source/blender/compositor/operations/COM_OutputFileOperation.h index 0e37432ca5b..9b9fb023467 100644 --- a/source/blender/compositor/operations/COM_OutputFileOperation.h +++ b/source/blender/compositor/operations/COM_OutputFileOperation.h @@ -49,7 +49,7 @@ public: bool isOutputOperation(bool rendering) const {return true;} void initExecution(); void deinitExecution(); - const int getRenderPriority() const {return 7;} + const CompositorPriority getRenderPriority() const {return COM_PRIORITY_LOW;} }; /* extra info for OpenEXR layers */ @@ -83,7 +83,7 @@ public: bool isOutputOperation(bool rendering) const {return true;} void initExecution(); void deinitExecution(); - const int getRenderPriority() const {return 7;} + const CompositorPriority getRenderPriority() const {return COM_PRIORITY_LOW;} }; #endif diff --git a/source/blender/compositor/operations/COM_PreviewOperation.cpp b/source/blender/compositor/operations/COM_PreviewOperation.cpp index a7b6fc93b25..4975f13a285 100644 --- a/source/blender/compositor/operations/COM_PreviewOperation.cpp +++ b/source/blender/compositor/operations/COM_PreviewOperation.cpp @@ -46,7 +46,6 @@ PreviewOperation::PreviewOperation() : NodeOperation() this->input = NULL; this->divider = 1.0f; this->node = NULL; - this->priority = 0; } void PreviewOperation::initExecution() @@ -129,7 +128,7 @@ void PreviewOperation::determineResolution(unsigned int resolution[], unsigned i resolution[1] = height; } -const int PreviewOperation::getRenderPriority() const +const CompositorPriority PreviewOperation::getRenderPriority() const { - return this->priority; + return COM_PRIORITY_LOW; } diff --git a/source/blender/compositor/operations/COM_PreviewOperation.h b/source/blender/compositor/operations/COM_PreviewOperation.h index 8450b7fc556..3d1cd38a5ea 100644 --- a/source/blender/compositor/operations/COM_PreviewOperation.h +++ b/source/blender/compositor/operations/COM_PreviewOperation.h @@ -37,20 +37,18 @@ protected: const bNodeTree *tree; SocketReader *input; float divider; - int priority; public: PreviewOperation(); bool isOutputOperation(bool rendering) const {return true;} void initExecution(); void deinitExecution(); - const int getRenderPriority() const; + const CompositorPriority getRenderPriority() const; void executeRegion(rcti *rect, unsigned int tileNumber, MemoryBuffer **memoryBuffers); void determineResolution(unsigned int resolution[], unsigned int preferredResolution[]); void setbNode(bNode *node) { this->node = node;} void setbNodeTree(const bNodeTree *tree) { this->tree = tree;} bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); - void setPriority(int priority) { this->priority = priority; } }; #endif diff --git a/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp b/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp index 71fc53e8f8e..809c688195f 100644 --- a/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp +++ b/source/blender/compositor/operations/COM_ViewerBaseOperation.cpp @@ -88,12 +88,12 @@ void ViewerBaseOperation::deinitExecution() this->outputBuffer = NULL; } -const int ViewerBaseOperation::getRenderPriority() const +const CompositorPriority ViewerBaseOperation::getRenderPriority() const { if (this->isActiveViewerOutput()) { - return 8; + return COM_PRIORITY_HIGH; } else { - return 0; + return COM_PRIORITY_LOW; } } diff --git a/source/blender/compositor/operations/COM_ViewerBaseOperation.h b/source/blender/compositor/operations/COM_ViewerBaseOperation.h index f5f30809f65..51fa8cecc0d 100644 --- a/source/blender/compositor/operations/COM_ViewerBaseOperation.h +++ b/source/blender/compositor/operations/COM_ViewerBaseOperation.h @@ -56,7 +56,7 @@ public: float getCenterX() { return this->centerX; } float getCenterY() { return this->centerY; } OrderOfChunks getChunkOrder() { return this->chunkOrder; } - const int getRenderPriority() const; + const CompositorPriority getRenderPriority() const; void setColorManagement(bool doColorManagement) {this->doColorManagement = doColorManagement;} void setColorPredivide(bool doColorPredivide) {this->doColorPredivide = doColorPredivide;} bool isViewerOperation() {return true;} diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp index 498add2fc87..222b879645c 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.cpp +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.cpp @@ -111,10 +111,9 @@ 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) +void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** inputMemoryBuffers, MemoryBuffer* outputBuffer) { - MemoryBuffer *outputMemoryBuffer = this->getMemoryProxy()->getBuffer();// @todo wrong implementation needs revision - float *outputFloatBuffer = outputMemoryBuffer->getBuffer(); + float *outputFloatBuffer = outputBuffer->getBuffer(); cl_int error; /* * 1. create cl_mem from outputbuffer @@ -125,8 +124,8 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr * note: list of cl_mem will be filled by 2, and needs to be cleaned up by 4 */ // STEP 1 - const unsigned int outputBufferWidth = outputMemoryBuffer->getWidth(); - const unsigned int outputBufferHeight = outputMemoryBuffer->getHeight(); + const unsigned int outputBufferWidth = outputBuffer->getWidth(); + const unsigned int outputBufferHeight = outputBuffer->getHeight(); const cl_image_format imageFormat = { CL_RGBA, @@ -141,19 +140,26 @@ void WriteBufferOperation::executeOpenCLRegion(cl_context context, cl_program pr clMemToCleanUp->push_back(clOutputBuffer); list<cl_kernel> *clKernelsToCleanUp = new list<cl_kernel>(); - this->input->executeOpenCL(context, program, queue, outputMemoryBuffer, clOutputBuffer, inputMemoryBuffers, clMemToCleanUp, clKernelsToCleanUp); + this->input->executeOpenCL(context, program, queue, outputBuffer, clOutputBuffer, inputMemoryBuffers, clMemToCleanUp, clKernelsToCleanUp); // STEP 3 size_t origin[3] = {0,0,0}; size_t region[3] = {outputBufferWidth,outputBufferHeight,1}; +// clFlush(queue); +// clFinish(queue); + error = clEnqueueBarrier(queue); 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); if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } + + this->getMemoryProxy()->getBuffer()->copyContentFrom(outputBuffer); + // STEP 4 + while (clMemToCleanUp->size()>0) { cl_mem mem = clMemToCleanUp->front(); error = clReleaseMemObject(mem); diff --git a/source/blender/compositor/operations/COM_WriteBufferOperation.h b/source/blender/compositor/operations/COM_WriteBufferOperation.h index b17122d68f0..6f2c49c49bd 100644 --- a/source/blender/compositor/operations/COM_WriteBufferOperation.h +++ b/source/blender/compositor/operations/COM_WriteBufferOperation.h @@ -46,7 +46,7 @@ public: void initExecution(); void deinitExecution(); void setbNodeTree(const bNodeTree *tree) {this->tree = tree;} - void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers); + void executeOpenCLRegion(cl_context context, cl_program program, cl_command_queue queue, rcti *rect, unsigned int chunkNumber, MemoryBuffer** memoryBuffers, MemoryBuffer* outputBuffer); }; #endif |