diff options
5 files changed, 167 insertions, 5 deletions
diff --git a/source/blender/compositor/nodes/COM_DirectionalBlurNode.cpp b/source/blender/compositor/nodes/COM_DirectionalBlurNode.cpp index dee0e6a88da..85fc63ae8cb 100644 --- a/source/blender/compositor/nodes/COM_DirectionalBlurNode.cpp +++ b/source/blender/compositor/nodes/COM_DirectionalBlurNode.cpp @@ -37,6 +37,7 @@ void DirectionalBlurNode::convertToOperations(ExecutionSystem *graph, Compositor DirectionalBlurOperation *operation = new DirectionalBlurOperation(); operation->setQuality(context->getQuality()); operation->setData(data); + operation->setbNode(this->getbNode()); this->getInputSocket(0)->relinkConnections(operation->getInputSocket(0), 0, graph); this->getOutputSocket(0)->relinkConnections(operation->getOutputSocket()); graph->addOperation(operation); diff --git a/source/blender/compositor/operations/COM_DirectionalBlurOperation.cpp b/source/blender/compositor/operations/COM_DirectionalBlurOperation.cpp index da7336afc07..76ee5b2c2a1 100644 --- a/source/blender/compositor/operations/COM_DirectionalBlurOperation.cpp +++ b/source/blender/compositor/operations/COM_DirectionalBlurOperation.cpp @@ -22,7 +22,7 @@ #include "COM_DirectionalBlurOperation.h" #include "BLI_math.h" - +#include "COM_OpenCLDevice.h" extern "C" { #include "RE_pipeline.h" } @@ -33,6 +33,7 @@ DirectionalBlurOperation::DirectionalBlurOperation() : NodeOperation() this->addOutputSocket(COM_DT_COLOR); this->setComplex(true); + this->setOpenCL(true); this->m_inputProgram = NULL; } @@ -97,9 +98,35 @@ void DirectionalBlurOperation::executePixel(float *color, int x, int y, MemoryBu lsc += this->m_sc; } - mul_v4_v4fl(color, col2, 1.0f / iterations); + mul_v4_v4fl(color, col2, 1.0f / (iterations+1)); } +void DirectionalBlurOperation::executeOpenCL(OpenCLDevice* device, + MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, + MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, + list<cl_kernel> *clKernelsToCleanUp) +{ + cl_kernel directionalBlurKernel = device->COM_clCreateKernel("directionalBlurKernel", NULL); + + cl_int iterations = pow(2.0f, this->m_data->iter); + cl_float2 ltxy = {this->m_tx, this->m_ty}; + cl_float2 centerpix = {this->m_center_x_pix, this->m_center_y_pix}; + cl_float lsc = this->m_sc; + cl_float lrot = this->m_rot; + + device->COM_clAttachMemoryBufferToKernelParameter(directionalBlurKernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputProgram); + device->COM_clAttachOutputMemoryBufferToKernelParameter(directionalBlurKernel, 1, clOutputBuffer); + device->COM_clAttachMemoryBufferOffsetToKernelParameter(directionalBlurKernel, 2, outputMemoryBuffer); + clSetKernelArg(directionalBlurKernel, 3, sizeof(cl_int), &iterations); + clSetKernelArg(directionalBlurKernel, 4, sizeof(cl_float), &lsc); + clSetKernelArg(directionalBlurKernel, 5, sizeof(cl_float), &lrot); + clSetKernelArg(directionalBlurKernel, 6, sizeof(cl_float2), <xy); + clSetKernelArg(directionalBlurKernel, 7, sizeof(cl_float2), ¢erpix); + + device->COM_clEnqueueRange(directionalBlurKernel, outputMemoryBuffer, 8, this); +} + + void DirectionalBlurOperation::deinitExecution() { this->m_inputProgram = NULL; diff --git a/source/blender/compositor/operations/COM_DirectionalBlurOperation.h b/source/blender/compositor/operations/COM_DirectionalBlurOperation.h index 329f855871e..0ab5e9dff2d 100644 --- a/source/blender/compositor/operations/COM_DirectionalBlurOperation.h +++ b/source/blender/compositor/operations/COM_DirectionalBlurOperation.h @@ -55,5 +55,11 @@ public: bool determineDependingAreaOfInterest(rcti *input, ReadBufferOperation *readOperation, rcti *output); void setData(NodeDBlurData *data) { this->m_data = data; } + + void executeOpenCL(OpenCLDevice* device, + MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, + MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, + list<cl_kernel> *clKernelsToCleanUp); + }; #endif diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl b/source/blender/compositor/operations/COM_OpenCLKernels.cl index f9b6d34bfc3..41838e41fba 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl @@ -1,7 +1,30 @@ +/* + * Copyright 2011, Blender Foundation. + * + * This program is free software; you can redistribute it and/or + * modify it under the terms of the GNU General Public License + * as published by the Free Software Foundation; either version 2 + * of the License, or (at your option) any later version. + * + * This program is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with this program; if not, write to the Free Software Foundation, + * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. + * + * Contributor: + * Jeroen Bakker + * Monique Dewanchand + */ + /// This file contains all opencl kernels for node-operation implementations // Global SAMPLERS -const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; +const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST; +const sampler_t SAMPLER_NEAREST_CLAMP = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; __constant const int2 zero = {0,0}; @@ -168,3 +191,44 @@ __kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2 float4 color = {value,0.0f,0.0f,0.0f}; write_imagef(output, coords, color); } + +// KERNEL --- DIRECTIONAL BLUR --- +__kernel void directionalBlurKernel(__read_only image2d_t inputImage, __write_only image2d_t output, + int2 offsetOutput, int iterations, float scale, float rotation, float2 translate, + float2 center, int2 offset) +{ + int2 coords = {get_global_id(0), get_global_id(1)}; + coords += offset; + const int2 realCoordinate = coords + offsetOutput; + + float4 col; + float2 ltxy = translate; + float lsc = scale; + float lrot = rotation; + + col = read_imagef(inputImage, SAMPLER_NEAREST, realCoordinate); + + /* blur the image */ + for (int i = 0; i < iterations; ++i) { + const float cs = cos(lrot), ss = sin(lrot); + const float isc = 1.0f / (1.0f + lsc); + + const float v = isc * (realCoordinate.s1 - center.s1) + ltxy.s1; + const float u = isc * (realCoordinate.s0 - center.s0) + ltxy.s0; + float2 uv = { + cs * u + ss * v + center.s0, + cs * v - ss * u + center.s1 + }; + + col += read_imagef(inputImage, SAMPLER_NEAREST_CLAMP, uv); + + /* double transformations */ + ltxy += translate; + lrot += rotation; + lsc += scale; + } + + col *= (1.0f/(iterations+1)); + + write_imagef(output, coords, col); +} diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h index 2c9ec76501d..d57aa1366de 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h @@ -1,9 +1,32 @@ /* clkernelstoh output of file <COM_OpenCLKernels_cl> */ -const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all opencl kernels for node-operation implementations\n" \ +const char * clkernelstoh_COM_OpenCLKernels_cl = "/*\n" \ +" * Copyright 2011, Blender Foundation.\n" \ +" *\n" \ +" * This program is free software; you can redistribute it and/or\n" \ +" * modify it under the terms of the GNU General Public License\n" \ +" * as published by the Free Software Foundation; either version 2\n" \ +" * of the License, or (at your option) any later version.\n" \ +" *\n" \ +" * This program is distributed in the hope that it will be useful,\n" \ +" * but WITHOUT ANY WARRANTY; without even the implied warranty of\n" \ +" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the\n" \ +" * GNU General Public License for more details.\n" \ +" *\n" \ +" * You should have received a copy of the GNU General Public License\n" \ +" * along with this program; if not, write to the Free Software Foundation,\n" \ +" * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.\n" \ +" *\n" \ +" * Contributor:\n" \ +" * Jeroen Bakker\n" \ +" * Monique Dewanchand\n" \ +" */\n" \ +"\n" \ +"/// 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" \ +"const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n" \ +"const sampler_t SAMPLER_NEAREST_CLAMP = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n" \ "\n" \ "__constant const int2 zero = {0,0};\n" \ "\n" \ @@ -170,4 +193,45 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope " float4 color = {value,0.0f,0.0f,0.0f};\n" \ " write_imagef(output, coords, color);\n" \ "}\n" \ +"\n" \ +"// KERNEL --- DIRECTIONAL BLUR ---\n" \ +"__kernel void directionalBlurKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \ +" int2 offsetOutput, int iterations, float scale, float rotation, float2 translate,\n" \ +" float2 center, int2 offset)\n" \ +"{\n" \ +" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ +" coords += offset;\n" \ +" const int2 realCoordinate = coords + offsetOutput;\n" \ +"\n" \ +" float4 col;\n" \ +" float2 ltxy = translate;\n" \ +" float lsc = scale;\n" \ +" float lrot = rotation;\n" \ +"\n" \ +" col = read_imagef(inputImage, SAMPLER_NEAREST, realCoordinate);\n" \ +"\n" \ +" /* blur the image */\n" \ +" for (int i = 0; i < iterations; ++i) {\n" \ +" const float cs = cos(lrot), ss = sin(lrot);\n" \ +" const float isc = 1.0f / (1.0f + lsc);\n" \ +"\n" \ +" const float v = isc * (realCoordinate.s1 - center.s1) + ltxy.s1;\n" \ +" const float u = isc * (realCoordinate.s0 - center.s0) + ltxy.s0;\n" \ +" float2 uv = {\n" \ +" cs * u + ss * v + center.s0,\n" \ +" cs * v - ss * u + center.s1\n" \ +" };\n" \ +"\n" \ +" col += read_imagef(inputImage, SAMPLER_NEAREST_CLAMP, uv);\n" \ +"\n" \ +" /* double transformations */\n" \ +" ltxy += translate;\n" \ +" lrot += rotation;\n" \ +" lsc += scale;\n" \ +" }\n" \ +"\n" \ +" col *= (1.0f/(iterations+1));\n" \ +"\n" \ +" write_imagef(output, coords, col);\n" \ +"}\n" \ "\0"; |