1 files changed, 108 insertions, 0 deletions
diff --git a/source/blender/compositor/intern/COM_NodeOperation.cpp b/source/blender/compositor/intern/COM_NodeOperation.cpp
index fae652e39d7..650e4af5ae0 100644
--- a/source/blender/compositor/intern/COM_NodeOperation.cpp
+++ b/source/blender/compositor/intern/COM_NodeOperation.cpp
@@ -124,3 +124,111 @@ bool NodeOperation::determineDependingAreaOfInterest(rcti * input, ReadBufferOpe
 		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;
+	
+	for (offsety = 0 ; offsety < height; offsety+=localSize) {
+		offset[1] = offsety;
+		if (offsety+localSize < height) {
+			size[1] = localSize;
+		} else {
+			size[1] = height - offsety;
+		}
+		for (offsetx = 0 ; offsetx < width ; 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);
+		}
+	}
+}
+
+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;
+	
+}