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 | |
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')
-rw-r--r-- | source/blender/compositor/operations/COM_OpenCLKernels.cl | 66 |
1 files changed, 64 insertions, 2 deletions
diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl b/source/blender/compositor/operations/COM_OpenCLKernels.cl index aeccfcab8b5..e1f175b318a 100644 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl +++ b/source/blender/compositor/operations/COM_OpenCLKernels.cl @@ -6,8 +6,8 @@ const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS __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, +__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage, + __read_only image2d_t bokehImage, __write_only image2d_t output, int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset) { int2 coords = {get_global_id(0), get_global_id(1)}; @@ -50,3 +50,65 @@ __kernel void bokehBlurKernel(__global __read_only image2d_t boundingBox, __glob 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, + int2 offset) +{ + int2 coords = {get_global_id(0), get_global_id(1)}; + coords += offset; + const int2 realCoordinate = coords + offsetOutput; + + const int2 minXY = max(realCoordinate - scope, zero); + const int2 maxXY = min(realCoordinate + scope, dimension); + + float value = 0.0f; + int nx, ny; + int2 inputXy; + + for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) { + for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) { + const float deltaX = (realCoordinate.x - nx); + const float deltaY = (realCoordinate.y - ny); + const float measuredDistance = deltaX*deltaX+deltaY*deltaY; + if (measuredDistance <= distanceSquared) { + value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0); + } + } + } + + float4 color = {value,0.0f,0.0f,0.0f}; + write_imagef(output, coords, color); +} + +// KERNEL --- DILATE --- +__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output, + int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension, + int2 offset) +{ + int2 coords = {get_global_id(0), get_global_id(1)}; + coords += offset; + const int2 realCoordinate = coords + offsetOutput; + + const int2 minXY = max(realCoordinate - scope, zero); + const int2 maxXY = min(realCoordinate + scope, dimension); + + float value = 1.0f; + int nx, ny; + int2 inputXy; + + for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) { + for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) { + const float deltaX = (realCoordinate.x - nx); + const float deltaY = (realCoordinate.y - ny); + const float measuredDistance = deltaX*deltaX+deltaY*deltaY; + if (measuredDistance <= distanceSquared) { + value = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0); + } + } + } + + float4 color = {value,0.0f,0.0f,0.0f}; + write_imagef(output, coords, color); +} |