* Added OpenCL kernel for bokeh blur

 * Uncomment COM_OPENCL_ENABLED from COM_defines.h to test

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