New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

OpenCL implementations of noise-hsv, gaussian-blur-selective, motion-blur-circular, diffraction-patterns #72

Merged
merged 4 commits into from May 17, 2017
Jump to file or symbol
Failed to load files and symbols.
+211 −1
Diff settings

Always

Just for now

Viewing a subset of changes. View all

gaussian-blur-selective: Add CL implementation

  • Loading branch information...
ex-rzr committed May 12, 2017
commit f407e25046a017b635354da14bb5f010dd4c7d11
@@ -0,0 +1,46 @@
kernel void
cl_gblur_selective(global const float4 *in,
global const float4 *delta,
global float4 *out,
const float radius,
const float max_delta)
{
const int gidx = get_global_id(0);
const int gidy = get_global_id(1);
const int iradius = (int)radius;
const int dst_width = get_global_size(0);
const int src_width = dst_width + iradius * 2;
const int center_gid1d = (gidy + iradius) * src_width + gidx + iradius;
const float4 center_pix = in[center_gid1d];
const float3 center_delta = delta[center_gid1d].xyz;
float3 accumulated = 0.0f;
float3 count = 0.0f;
for (int v = -iradius; v <= iradius; v++)
{
for (int u = -iradius; u <= iradius; u++)
{
const int i = gidx + iradius + u;
const int j = gidy + iradius + v;
const int gid1d = i + j * src_width;
const float4 src_pix = in[gid1d];
const float3 delta_pix = delta[gid1d].xyz;
const float gaussian_weight = exp(-0.5f * (u * u + v * v) / radius);
const float weight = gaussian_weight * src_pix.w;
const float3 diff = center_delta - delta_pix;
const float3 w = convert_float3 (fabs (diff) <= max_delta);
accumulated += w * weight * src_pix.xyz;
count += w * weight;
}
}
const float3 out_v = select (center_pix.xyz,
accumulated / count,
count != 0.0f);
out[gidx + gidy * dst_width] = (float4)(out_v, center_pix.w);
}
@@ -0,0 +1,48 @@
static const char* gaussian_blur_selective_cl_source =
"kernel void \n"
"cl_gblur_selective(global const float4 *in, \n"
" global const float4 *delta, \n"
" global float4 *out, \n"
" const float radius, \n"
" const float max_delta) \n"
"{ \n"
" const int gidx = get_global_id(0); \n"
" const int gidy = get_global_id(1); \n"
" const int iradius = (int)radius; \n"
" const int dst_width = get_global_size(0); \n"
" const int src_width = dst_width + iradius * 2; \n"
" \n"
" const int center_gid1d = (gidy + iradius) * src_width + gidx + iradius; \n"
" const float4 center_pix = in[center_gid1d]; \n"
" const float3 center_delta = delta[center_gid1d].xyz; \n"
" \n"
" float3 accumulated = 0.0f; \n"
" float3 count = 0.0f; \n"
" \n"
" for (int v = -iradius; v <= iradius; v++) \n"
" { \n"
" for (int u = -iradius; u <= iradius; u++) \n"
" { \n"
" const int i = gidx + iradius + u; \n"
" const int j = gidy + iradius + v; \n"
" const int gid1d = i + j * src_width; \n"
" \n"
" const float4 src_pix = in[gid1d]; \n"
" const float3 delta_pix = delta[gid1d].xyz; \n"
" \n"
" const float gaussian_weight = exp(-0.5f * (u * u + v * v) / radius);\n"
" \n"
" const float weight = gaussian_weight * src_pix.w; \n"
" const float3 diff = center_delta - delta_pix; \n"
" const float3 w = convert_float3 (fabs (diff) <= max_delta); \n"
" accumulated += w * weight * src_pix.xyz; \n"
" count += w * weight; \n"
" } \n"
" } \n"
" \n"
" const float3 out_v = select (center_pix.xyz, \n"
" accumulated / count, \n"
" count != 0.0f); \n"
" out[gidx + gidy * dst_width] = (float4)(out_v, center_pix.w); \n"
"} \n"
;
@@ -218,6 +218,118 @@ gblur_selective (GeglBuffer *input,
return TRUE;
}
#include "opencl/gegl-cl.h"
#include "gegl-buffer-cl-iterator.h"
#include "opencl/gaussian-blur-selective.cl.h"
static GeglClRunData *cl_data = NULL;
static gboolean
cl_gblur_selective (cl_mem in,
cl_mem delta,
cl_mem out,
size_t global_worksize,
const GeglRectangle *roi,
gfloat radius,
gfloat max_delta)
{
cl_int cl_err = 0;
size_t global_ws[2];
if (!cl_data)
{
const char *kernel_name[] = { "cl_gblur_selective", NULL };
cl_data = gegl_cl_compile_and_build (gaussian_blur_selective_cl_source,
kernel_name);
}
if (!cl_data)
return TRUE;
global_ws[0] = roi->width;
global_ws[1] = roi->height;
gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &in,
sizeof(cl_mem), &delta,
sizeof(cl_mem), &out,
sizeof(cl_float), &radius,
sizeof(cl_float), &max_delta,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_ws, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *aux,
GeglBuffer *output,
const GeglRectangle *result)
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *aux_format = gegl_operation_get_format (operation, "aux");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
GeglProperties *o = GEGL_PROPERTIES (operation);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
gint radius = o->blur_radius;
gint read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
radius, radius, radius, radius,
GEGL_ABYSS_CLAMP);
gint delta = !aux ?
read :
gegl_buffer_cl_iterator_add_2 (i,
aux,
result,
aux_format,
GEGL_CL_BUFFER_READ,
radius, radius, radius, radius,
GEGL_ABYSS_CLAMP);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err)
return FALSE;
err = cl_gblur_selective(i->tex[read],
i->tex[delta],
i->tex[0],
i->size[0],
&i->roi[0],
o->blur_radius,
o->max_delta);
if (err)
return FALSE;
}
return TRUE;
}
static gboolean
process (GeglOperation *operation,
GeglBuffer *input,
@@ -232,6 +344,10 @@ process (GeglOperation *operation,
compute = get_required_for_output (operation, "input", result);
if (gegl_operation_use_opencl (operation))
if (cl_process (operation, input, aux, output, result))
return TRUE;
success = gblur_selective (input, &compute,
aux,
output, result,
@@ -251,7 +367,7 @@ gegl_op_class_init (GeglOpClass *klass)
operation_class->prepare = prepare;
operation_class->get_required_for_output = get_required_for_output;
operation_class->get_invalidated_by_change = get_invalidated_by_change;
operation_class->opencl_support = FALSE;
operation_class->opencl_support = TRUE;
composer_class->process = process;