[Bf-blender-cvs] SVN commit: /data/svn/bf-blender [48678] trunk/blender/source/blender/ compositor: * Added OpenCL implementation of the Defocus node
Jeroen Bakker
j.bakker at atmind.nl
Fri Jul 6 13:31:41 CEST 2012
Revision: 48678
http://projects.blender.org/scm/viewvc.php?view=rev&root=bf-blender&revision=48678
Author: jbakker
Date: 2012-07-06 11:31:40 +0000 (Fri, 06 Jul 2012)
Log Message:
-----------
* Added OpenCL implementation of the Defocus node
* Always disable two phase compositing during rendering
- At Mind -
Modified Paths:
--------------
trunk/blender/source/blender/compositor/intern/COM_compositor.cpp
trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl
trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl.h
trunk/blender/source/blender/compositor/operations/COM_VariableSizeBokehBlurOperation.cpp
trunk/blender/source/blender/compositor/operations/COM_VariableSizeBokehBlurOperation.h
Modified: trunk/blender/source/blender/compositor/intern/COM_compositor.cpp
===================================================================
--- trunk/blender/source/blender/compositor/intern/COM_compositor.cpp 2012-07-06 11:24:43 UTC (rev 48677)
+++ trunk/blender/source/blender/compositor/intern/COM_compositor.cpp 2012-07-06 11:31:40 UTC (rev 48678)
@@ -57,7 +57,7 @@
/* set progress bar to 0% and status to init compositing*/
editingtree->progress(editingtree->prh, 0.0);
- bool twopass = (editingtree->flag&NTREE_TWO_PASS) > 0 || rendering;
+ bool twopass = (editingtree->flag&NTREE_TWO_PASS) > 0 && !rendering;
/* initialize execution system */
if (twopass) {
ExecutionSystem *system = new ExecutionSystem(rd, editingtree, rendering, twopass);
Modified: trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl
===================================================================
--- trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl 2012-07-06 11:24:43 UTC (rev 48677)
+++ trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl 2012-07-06 11:31:40 UTC (rev 48678)
@@ -51,6 +51,68 @@
write_imagef(output, coords, color);
}
+//KERNEL --- DEFOCUS /VARIABLESIZEBOKEHBLUR ---
+__kernel void defocusKernel(__read_only image2d_t inputImage, __read_only image2d_t bokehImage,
+ __read_only image2d_t inputDepth, __read_only image2d_t inputSize,
+ __write_only image2d_t output, int2 offsetInput, int2 offsetOutput,
+ int step, int maxBlur, float threshold, int2 dimension, int2 offset)
+{
+ float4 color = {1.0f, 0.0f, 0.0f, 1.0f};
+ int2 coords = {get_global_id(0), get_global_id(1)};
+ coords += offset;
+ const int2 realCoordinate = coords + offsetOutput;
+
+ float4 readColor;
+ float4 bokeh;
+ float tempSize;
+ float tempDepth;
+ float4 multiplier_accum = {1.0f, 1.0f, 1.0f, 1.0f};
+ float4 color_accum;
+
+ int minx = max(realCoordinate.s0 - maxBlur, 0);
+ int miny = max(realCoordinate.s1 - maxBlur, 0);
+ int maxx = min(realCoordinate.s0 + maxBlur, dimension.s0);
+ int maxy = min(realCoordinate.s1 + maxBlur, dimension.s1);
+
+ {
+ int2 inputCoordinate = realCoordinate - offsetInput;
+ float size = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;
+ float depth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0 + threshold;
+ color_accum = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);
+
+ for (int ny = miny; ny < maxy; ny += step) {
+ for (int nx = minx; nx < maxx; nx += step) {
+ if (nx >= 0 && nx < dimension.s0 && ny >= 0 && ny < dimension.s1) {
+ inputCoordinate.s0 = nx - offsetInput.s0;
+ inputCoordinate.s1 = ny - offsetInput.s1;
+ tempDepth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0;
+ if (tempDepth < depth) {
+ tempSize = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;
+
+ if ((size > threshold && tempSize > threshold) || tempSize <= threshold) {
+ float dx = nx - realCoordinate.s0;
+ float dy = ny - realCoordinate.s1;
+ if (dx != 0 || dy != 0) {
+ if (tempSize >= fabs(dx) && tempSize >= fabs(dy)) {
+ float2 uv = { 256.0f + dx * 256.0f / tempSize, 256.0f + dy * 256.0f / tempSize};
+ bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);
+ readColor = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);
+ color_accum += bokeh*readColor;
+ multiplier_accum += bokeh;
+ }
+ }
+ }
+ }
+ }
+ }
+ }
+ }
+
+ color = color_accum * (1.0f / multiplier_accum);
+ write_imagef(output, coords, color);
+}
+
+
// KERNEL --- DILATE ---
__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,
int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,
Modified: trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl.h
===================================================================
--- trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl.h 2012-07-06 11:24:43 UTC (rev 48677)
+++ trunk/blender/source/blender/compositor/operations/COM_OpenCLKernels.cl.h 2012-07-06 11:31:40 UTC (rev 48678)
@@ -16,7 +16,7 @@
" coords += offset;\n" \
" float tempBoundingBox;\n" \
" float4 color = {0.0f,0.0f,0.0f,0.0f};\n" \
-" float4 multiplier = {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" \
@@ -40,10 +40,10 @@
" 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" \
-" multiplier += bokeh;\n" \
+" multiplyer += bokeh;\n" \
" }\n" \
" }\n" \
-" color /= multiplier;\n" \
+" color /= multiplyer;\n" \
"\n" \
" } else {\n" \
" int2 imageCoordinates = realCoordinate - offsetInput;\n" \
@@ -53,6 +53,68 @@
" write_imagef(output, coords, color);\n" \
"}\n" \
"\n" \
+"//KERNEL --- DEFOCUS /VARIABLESIZEBOKEHBLUR ---\n" \
+"__kernel void defocusKernel(__read_only image2d_t inputImage, __read_only image2d_t bokehImage,\n" \
+" __read_only image2d_t inputDepth, __read_only image2d_t inputSize,\n" \
+" __write_only image2d_t output, int2 offsetInput, int2 offsetOutput,\n" \
+" int step, int maxBlur, float threshold, int2 dimension, int2 offset)\n" \
+"{\n" \
+" float4 color = {1.0f, 0.0f, 0.0f, 1.0f};\n" \
+" int2 coords = {get_global_id(0), get_global_id(1)};\n" \
+" coords += offset;\n" \
+" const int2 realCoordinate = coords + offsetOutput;\n" \
+"\n" \
+" float4 readColor;\n" \
+" float4 bokeh;\n" \
+" float tempSize;\n" \
+" float tempDepth;\n" \
+" float4 multiplier_accum = {1.0f, 1.0f, 1.0f, 1.0f};\n" \
+" float4 color_accum;\n" \
+"\n" \
+" int minx = max(realCoordinate.s0 - maxBlur, 0);\n" \
+" int miny = max(realCoordinate.s1 - maxBlur, 0);\n" \
+" int maxx = min(realCoordinate.s0 + maxBlur, dimension.s0);\n" \
+" int maxy = min(realCoordinate.s1 + maxBlur, dimension.s1);\n" \
+"\n" \
+" {\n" \
+" int2 inputCoordinate = realCoordinate - offsetInput;\n" \
+" float size = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;\n" \
+" float depth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0 + threshold;\n" \
+" color_accum = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);\n" \
+"\n" \
+" for (int ny = miny; ny < maxy; ny += step) {\n" \
+" for (int nx = minx; nx < maxx; nx += step) {\n" \
+" if (nx >= 0 && nx < dimension.s0 && ny >= 0 && ny < dimension.s1) {\n" \
+" inputCoordinate.s0 = nx - offsetInput.s0;\n" \
+" inputCoordinate.s1 = ny - offsetInput.s1;\n" \
+" tempDepth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0;\n" \
+" if (tempDepth < depth) {\n" \
+" tempSize = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;\n" \
+"\n" \
+" if ((size > threshold && tempSize > threshold) || tempSize <= threshold) {\n" \
+" float dx = nx - realCoordinate.s0;\n" \
+" float dy = ny - realCoordinate.s1;\n" \
+" if (dx != 0 || dy != 0) {\n" \
+" if (tempSize >= fabs(dx) && tempSize >= fabs(dy)) {\n" \
+" float2 uv = { 256.0f + dx * 256.0f / tempSize, 256.0f + dy * 256.0f / tempSize};\n" \
+" bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);\n" \
+" readColor = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);\n" \
+" color_accum += bokeh*readColor;\n" \
+" multiplier_accum += bokeh;\n" \
+" }\n" \
+" }\n" \
+" }\n" \
+" }\n" \
+" }\n" \
+" }\n" \
+" }\n" \
+" }\n" \
+"\n" \
+" color = color_accum * (1.0f / multiplier_accum);\n" \
+" write_imagef(output, coords, color);\n" \
+"}\n" \
+"\n" \
+"\n" \
"// KERNEL --- DILATE ---\n" \
"__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \
" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \
@@ -70,9 +132,9 @@
" int2 inputXy;\n" \
"\n" \
" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \
+" const float deltaY = (realCoordinate.y - ny);\n" \
" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \
" const float deltaX = (realCoordinate.x - nx);\n" \
-" const float deltaY = (realCoordinate.y - ny);\n" \
" const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \
" if (measuredDistance <= distanceSquared) {\n" \
" value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \
Modified: trunk/blender/source/blender/compositor/operations/COM_VariableSizeBokehBlurOperation.cpp
===================================================================
--- trunk/blender/source/blender/compositor/operations/COM_VariableSizeBokehBlurOperation.cpp 2012-07-06 11:24:43 UTC (rev 48677)
+++ trunk/blender/source/blender/compositor/operations/COM_VariableSizeBokehBlurOperation.cpp 2012-07-06 11:31:40 UTC (rev 48678)
@@ -22,6 +22,7 @@
#include "COM_VariableSizeBokehBlurOperation.h"
#include "BLI_math.h"
+#include "COM_OpenCLDevice.h"
extern "C" {
#include "RE_pipeline.h"
@@ -38,6 +39,7 @@
#endif
this->addOutputSocket(COM_DT_COLOR);
this->setComplex(true);
+ this->setOpenCL(true);
this->m_inputProgram = NULL;
this->m_inputBokehProgram = NULL;
@@ -128,6 +130,33 @@
}
+static cl_kernel defocusKernel = 0;
+void VariableSizeBokehBlurOperation::executeOpenCL(OpenCLDevice* device,
+ MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer,
+ MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp,
+ list<cl_kernel> *clKernelsToCleanUp)
+{
+ if (!defocusKernel) {
+ defocusKernel = device->COM_clCreateKernel("defocusKernel", NULL);
+ }
+ cl_int step = this->getStep();
+ cl_int maxBlur = this->m_maxBlur;
+ cl_float threshold = this->m_threshold;
+
+ device->COM_clAttachMemoryBufferToKernelParameter(defocusKernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputProgram);
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list