diff options
| author | Jeroen Bakker <j.bakker@atmind.nl> | 2012-06-13 16:34:56 +0400 |
|---|---|---|
| committer | Jeroen Bakker <j.bakker@atmind.nl> | 2012-06-13 16:34:56 +0400 |
| commit | be1b5f82cee09041fdee355697841ee92b31ef70 (patch) | |
| tree | 7032f52aaffb862c228d92a476dc9a0e00261ef1 /source/blender/compositor/operations/COM_OpenCLKernels.cl.h | |
| parent | 4ba456d1754c29b488b8304c8546af45078e8536 (diff) | |
* optimized threading
* break out with glare node
* Added OpenCL kernels compatible with AMD still need some testing.
Diffstat (limited to 'source/blender/compositor/operations/COM_OpenCLKernels.cl.h')
| -rw-r--r-- | source/blender/compositor/operations/COM_OpenCLKernels.cl.h | 70 |
1 files changed, 66 insertions, 4 deletions
diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h index 3cf33c75272..ef8668f6f21 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h @@ -8,8 +8,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope "__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" \ +"__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage,\n" \ +" __read_only image2d_t bokehImage, __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" \ @@ -26,8 +26,8 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope " 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" \ +" const int2 minXY = max(realCoordinate - radius, zero);\n" \ +" const int2 maxXY = min(realCoordinate + radius, dimension);\n" \ " int nx, ny;\n" \ "\n" \ " float2 uv;\n" \ @@ -52,4 +52,66 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope "\n" \ " write_imagef(output, coords, color);\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" \ +" 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" \ +" const int2 minXY = max(realCoordinate - scope, zero);\n" \ +" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \ +"\n" \ +" float value = 0.0f;\n" \ +" int nx, ny;\n" \ +" int2 inputXy;\n" \ +"\n" \ +" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\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" \ +" }\n" \ +" }\n" \ +" }\n" \ +"\n" \ +" float4 color = {value,0.0f,0.0f,0.0f};\n" \ +" write_imagef(output, coords, color);\n" \ +"}\n" \ +"\n" \ +"// KERNEL --- DILATE ---\n" \ +"__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \ +" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \ +" 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" \ +" const int2 minXY = max(realCoordinate - scope, zero);\n" \ +" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \ +"\n" \ +" float value = 1.0f;\n" \ +" int nx, ny;\n" \ +" int2 inputXy;\n" \ +"\n" \ +" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\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 = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \ +" }\n" \ +" }\n" \ +" }\n" \ +"\n" \ +" float4 color = {value,0.0f,0.0f,0.0f};\n" \ +" write_imagef(output, coords, color);\n" \ +"}\n" \ "\0"; |