[gegl] Add OpenCL support for gegl:vignette
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] Add OpenCL support for gegl:vignette
- Date: Sun, 1 Apr 2012 18:03:36 +0000 (UTC)
commit bc0d2f449e7a7d099408b1dac115a67b61c1db02
Author: Victor Oliveira <victormatheus gmail com>
Date: Mon Mar 19 17:31:51 2012 -0300
Add OpenCL support for gegl:vignette
operations/common/vignette.c | 174 +++++++++++++++++++++++++++++++++++++++++-
1 files changed, 170 insertions(+), 4 deletions(-)
---
diff --git a/operations/common/vignette.c b/operations/common/vignette.c
index 354ef97..a53facc 100644
--- a/operations/common/vignette.c
+++ b/operations/common/vignette.c
@@ -19,6 +19,7 @@
#include "config.h"
#include <glib/gi18n-lib.h>
+
#ifdef GEGL_CHANT_PROPERTIES
gegl_chant_int (shape, _("Shape"), 0, 2, 0, _("Shape to use: 0=circle 1=diamond 2=square"))
gegl_chant_color (color, _("Color"), "black", _("Defaults to 'black', you can use transparency here to erase portions of an image"))
@@ -79,6 +80,170 @@ static float scale_to_aspect (float scale)
}
#endif
+#include "opencl/gegl-cl.h"
+
+static const char* kernel_source =
+"__kernel void vignette_cl (__global const float4 *in, \n"
+" __global float4 *out, \n"
+" float4 color, \n"
+" float scale, \n"
+" float cost, \n"
+" float sint, \n"
+" int roi_x, \n"
+" int roi_y, \n"
+" int midx, \n"
+" int midy, \n"
+" int o_shape, \n"
+" float gamma, \n"
+" float length, \n"
+" float radius0, \n"
+" float rdiff) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int gid = gidx + gidy * get_global_size(0); \n"
+" float strength = 0.0f; \n"
+" float u,v,costy,sinty; \n"
+" int x,y; \n"
+" x = gidx + roi_x; \n"
+" y = gidy + roi_y; \n"
+" sinty = sint * (y-midy) - midx; \n"
+" costy = cost * (y-midy) + midy; \n"
+" \n"
+" u = cost * (x-midx) - sinty; \n"
+" v = sint * (x-midx) + costy; \n"
+" \n"
+" if (length == 0.0f) \n"
+" strength = 0.0f; \n"
+" else \n"
+" { \n"
+" switch (o_shape) \n"
+" { \n"
+" case 0: \n"
+" strength = hypot ((u-midx) / scale, v-midy); \n"
+" break; \n"
+" \n"
+" case 1: \n"
+" strength = fmax (fabs(u-midx)/scale, fabs(v-midy)); \n"
+" break; \n"
+" \n"
+" case 2: \n"
+" strength = fabs (u-midx) / scale + fabs(v-midy); \n"
+" break; \n"
+" } \n"
+" strength /= length; \n"
+" strength = (strength-radius0) / rdiff; \n"
+" } \n"
+" \n"
+" if (strength < 0.0f) strength = 0.0f; \n"
+" if (strength > 1.0f) strength = 1.0f; \n"
+" \n"
+" if (gamma > 0.9999f && gamma < 2.0001f) \n"
+" strength *= strength; \n"
+" else if (gamma != 1.0f) \n"
+" strength = pow(strength, gamma); \n"
+" \n"
+" out[gid] = in[gid]*(1.0f-strength) + color * strength; \n"
+"} \n";
+
+static gegl_cl_run_data * cl_data = NULL;
+
+static cl_int
+cl_process (GeglOperation *operation,
+ cl_mem in_tex,
+ cl_mem out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi,
+ gint level)
+{
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+ gfloat scale;
+ gfloat radius0, radius1;
+ gint roi_x, roi_y,x;
+ gint midx, midy;
+ GeglRectangle *bounds = gegl_operation_source_get_bounding_box (operation, "input");
+
+ gfloat length = hypot (bounds->width, bounds->height)/2;
+ gfloat rdiff;
+ gfloat cost, sint;
+ gfloat color[4];
+
+ scale = bounds->width / (1.0 * bounds->height);
+ scale = scale * (o->proportion) + 1.0 * (1.0-o->proportion);
+ scale *= aspect_to_scale (o->squeeze);
+ length = (bounds->width/2.0);
+
+ if (scale > 1.0)
+ length /= scale;
+
+ gegl_color_get_rgba4f (o->color, color);
+
+ for (x=0; x<3; x++) /* premultiply */
+ color[x] *= color[3];
+
+ radius0 = o->radius * (1.0-o->softness);
+ radius1 = o->radius;
+ rdiff = radius1-radius0;
+ if (fabs (rdiff) < 0.0001)
+ rdiff = 0.0001;
+
+ midx = bounds->x + bounds->width * o->x;
+ midy = bounds->y + bounds->height * o->y;
+
+ roi_x = roi->x;
+ roi_y = roi->y;
+
+ /* constant for all pixels */
+ cost = cos(-o->rotation * (G_PI*2/360.0));
+ sint = sin(-o->rotation * (G_PI*2/360.0));
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"vignette_cl",NULL};
+ cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ }
+ if (!cl_data) return 1;
+
+ {
+ const size_t gbl_size[2] = {roi->width, roi->height};
+
+ gint shape = o->shape;
+ gfloat gamma = o->gamma;
+
+ cl_int cl_err = 0;
+ cl_float4 f_color;
+ f_color.s[0] = color[0];
+ f_color.s[1] = color[1];
+ f_color.s[2] = color[2];
+ f_color.s[3] = color[3];
+
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&scale);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&cost);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_float), (void*)&sint);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_int), (void*)&roi_x);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_int), (void*)&roi_y);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_int), (void*)&midx);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 9, sizeof(cl_int), (void*)&midy);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 10, sizeof(cl_int), (void*)&shape);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 11, sizeof(cl_float), (void*)&gamma);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 12, sizeof(cl_float), (void*)&length);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 13, sizeof(cl_float), (void*)&radius0);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 14, sizeof(cl_float), (void*)&rdiff);
+ if (cl_err != CL_SUCCESS) return cl_err;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 2,
+ NULL, gbl_size, NULL,
+ 0, NULL, NULL);
+ if (cl_err != CL_SUCCESS) return cl_err;
+ }
+
+ return CL_SUCCESS;
+}
+
static gboolean
process (GeglOperation *operation,
void *in_buf,
@@ -94,8 +259,7 @@ process (GeglOperation *operation,
gfloat radius0, radius1;
gint x, y;
gint midx, midy;
- GeglRectangle *bounds = gegl_operation_source_get_bounding_box (
- operation, "input");
+ GeglRectangle *bounds = gegl_operation_source_get_bounding_box (operation, "input");
gfloat length = hypot (bounds->width, bounds->height)/2;
gfloat rdiff;
gfloat cost, sint;
@@ -199,8 +363,6 @@ process (GeglOperation *operation,
return TRUE;
}
-
-
static void
gegl_chant_class_init (GeglChantClass *klass)
{
@@ -211,9 +373,13 @@ gegl_chant_class_init (GeglChantClass *klass)
point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
point_filter_class->process = process;
+ point_filter_class->cl_process = cl_process;
+
operation_class->prepare = prepare;
operation_class->no_cache = TRUE;
+ operation_class->opencl_support = TRUE;
+
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:vignette",
"categories" , "render",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
