[gegl] opencl: Improving OpenCL ops API
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: Improving OpenCL ops API
- Date: Fri, 20 Apr 2012 16:38:22 +0000 (UTC)
commit 158dbdd9ada4af13b822b3011042cd76a886d8ce
Author: Victor Oliveira <victormatheus gmail com>
Date: Fri Apr 20 13:35:06 2012 -0300
opencl: Improving OpenCL ops API
- Extending simplified API to point composer ops
- Rewriting ops using the simplified API
gegl/operation/gegl-operation-point-composer.c | 73 ++++++----
gegl/operation/gegl-operation-point-filter.c | 12 +--
operations/common/invert.c | 47 +------
operations/common/mono-mixer.c | 177 ++++++------------------
operations/common/opacity.c | 83 +----------
operations/common/over.c | 54 +-------
operations/common/threshold.c | 87 ++-----------
operations/common/value-invert.c | 49 +------
8 files changed, 113 insertions(+), 469 deletions(-)
---
diff --git a/gegl/operation/gegl-operation-point-composer.c b/gegl/operation/gegl-operation-point-composer.c
index 5740d26..6e3b790 100644
--- a/gegl/operation/gegl-operation-point-composer.c
+++ b/gegl/operation/gegl-operation-point-composer.c
@@ -30,6 +30,8 @@
#include "graph/gegl-pad.h"
#include <string.h>
+#include "gegl-buffer-private.h"
+
static gboolean gegl_operation_point_composer_process
(GeglOperation *operation,
GeglBuffer *input,
@@ -150,12 +152,15 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
const Babl *aux_format = gegl_operation_get_format (operation, "aux");
const Babl *out_format = gegl_operation_get_format (operation, "output");
+ GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
GeglOperationPointComposerClass *point_composer_class = GEGL_OPERATION_POINT_COMPOSER_GET_CLASS (operation);
gint j;
cl_int cl_err = 0;
gboolean err;
+ gint foo;
+
/* non-texturizable format! */
if (!gegl_cl_color_babl (in_format, NULL) ||
!gegl_cl_color_babl (aux_format, NULL) ||
@@ -165,45 +170,52 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
return FALSE;
}
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_COMPOSER: %s", operation_class->name);
+
/* Process */
{
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE);
gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
if (aux)
- {
- gint foo = gegl_buffer_cl_iterator_add (i, aux, result, aux_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
+ foo = gegl_buffer_cl_iterator_add (i, aux, result, aux_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
- while (gegl_buffer_cl_iterator_next (i, &err))
+ while (gegl_buffer_cl_iterator_next (i, &err))
+ {
+ if (err) return FALSE;
+ for (j=0; j < i->n; j++)
{
- if (err) return FALSE;
- for (j=0; j < i->n; j++)
+ if (point_composer_class->cl_process)
{
- cl_err = point_composer_class->cl_process(operation, i->tex[read][j], i->tex[foo][j], i->tex[0][j],
- i->size[0][j], &i->roi[0][j], level);
- if (cl_err != CL_SUCCESS)
- {
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s [GeglOperationPointComposer] Kernel",
- GEGL_OPERATION_CLASS (operation)->name);
- return FALSE;
- }
+ cl_err = point_composer_class->cl_process(operation, i->tex[read][j],
+ (aux)? i->tex[foo][j] : NULL,
+ i->tex[0][j], i->size[0][j], &i->roi[0][j], level);
}
- }
- }
- else
- {
- while (gegl_buffer_cl_iterator_next (i, &err))
- {
- if (err) return FALSE;
- for (j=0; j < i->n; j++)
+ else if (operation_class->cl_data)
+ {
+ gint p = 0;
+ gegl_cl_run_data *cl_data = operation_class->cl_data;
+
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read][j]);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (aux)? (void*)&i->tex[foo][j] : NULL);
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[0][j]);
+
+ gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 1,
+ NULL, &i->size[0][j], NULL,
+ 0, NULL, NULL);
+ }
+ else
+ {
+ g_warning ("OpenCL support enabled, but no way to execute");
+ return FALSE;
+ }
+
+ if (cl_err != CL_SUCCESS)
{
- cl_err = point_composer_class->cl_process(operation, i->tex[read][j], NULL, i->tex[0][j],
- i->size[0][j], &i->roi[0][j], level);
- if (cl_err != CL_SUCCESS)
- {
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in %s [GeglOperationPointComposer] Kernel",
- GEGL_OPERATION_CLASS (operation)->name);
- return FALSE;
- }
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
+ return FALSE;
}
}
}
@@ -219,6 +231,7 @@ gegl_operation_point_composer_process (GeglOperation *operation,
const GeglRectangle *result,
gint level)
{
+ GeglOperationClass *operation_class = GEGL_OPERATION_GET_CLASS (operation);
GeglOperationPointComposerClass *point_composer_class = GEGL_OPERATION_POINT_COMPOSER_GET_CLASS (operation);
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *aux_format = gegl_operation_get_format (operation, "aux");
@@ -226,7 +239,7 @@ gegl_operation_point_composer_process (GeglOperation *operation,
if ((result->width > 0) && (result->height > 0))
{
- if (gegl_cl_is_accelerated () && point_composer_class->cl_process)
+ if (gegl_cl_is_accelerated () && (operation_class->cl_data || point_composer_class->cl_process))
{
if (gegl_operation_point_composer_cl_process (operation, input, aux, output, result, level))
return TRUE;
diff --git a/gegl/operation/gegl-operation-point-filter.c b/gegl/operation/gegl-operation-point-filter.c
index edef3a8..d81c4f1 100644
--- a/gegl/operation/gegl-operation-point-filter.c
+++ b/gegl/operation/gegl-operation-point-filter.c
@@ -100,13 +100,7 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
return FALSE;
}
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "BABL formats: (%s,%s) (%s,%s)\n \t Tile Size:(%d, %d)\n",
- babl_get_name(input->format),
- babl_get_name(in_format),
- babl_get_name(out_format),
- babl_get_name(output->format),
- input->tile_storage->tile_width,
- input->tile_storage->tile_height);
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_FILTER: %s", operation_class->name);
/* Process */
{
@@ -145,7 +139,7 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
if (cl_err != CL_SUCCESS)
{
- GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointFilter Kernel: %sS", gegl_cl_errstring(cl_err));
+ GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointFilter Kernel: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
}
@@ -170,7 +164,7 @@ gegl_operation_point_filter_process (GeglOperation *operation,
if ((result->width > 0) && (result->height > 0))
{
- if (gegl_cl_is_accelerated () && operation_class->cl_data)
+ if (gegl_cl_is_accelerated () && (operation_class->cl_data || point_filter_class->cl_process))
{
if (gegl_operation_point_filter_cl_process (operation, input, output, result, level))
return TRUE;
diff --git a/operations/common/invert.c b/operations/common/invert.c
index 4b62697..f5e0afa 100644
--- a/operations/common/invert.c
+++ b/operations/common/invert.c
@@ -61,11 +61,9 @@ process (GeglOperation *op,
return TRUE;
}
-#include "opencl/gegl-cl.h"
-
static const char* kernel_source =
-"__kernel void kernel_inv(__global const float4 *in, \n"
-" __global float4 *out) \n"
+"__kernel void gegl_invert (__global const float4 *in, \n"
+" __global float4 *out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
@@ -75,41 +73,6 @@ static const char* kernel_source =
" out[gid] = out_v; \n"
"} \n";
-static gegl_cl_run_data *cl_data = NULL;
-
-/* OpenCL processing function */
-static cl_int
-cl_process (GeglOperation *op,
- cl_mem in_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *roi,
- int level)
-{
- cl_int cl_err = 0;
-
- if (!cl_data)
- {
- const char *kernel_name[] = {"kernel_inv", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
- }
-
- if (!cl_data) return 1;
-
- 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);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- return cl_err;
-
-}
-
static void
gegl_chant_class_init (GeglChantClass *klass)
{
@@ -120,9 +83,6 @@ 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->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:invert",
@@ -130,7 +90,8 @@ gegl_chant_class_init (GeglChantClass *klass)
"description",
_("Inverts the components (except alpha), the result is the "
"corresponding \"negative\" image."),
- NULL);
+ "cl-source" , kernel_source,
+ NULL);
}
#endif
diff --git a/operations/common/mono-mixer.c b/operations/common/mono-mixer.c
index 75e29ec..53a68ea 100644
--- a/operations/common/mono-mixer.c
+++ b/operations/common/mono-mixer.c
@@ -32,7 +32,7 @@ gegl_chant_double_ui (blue, _("Blue"), -10.0, 10.0, 0.25, -1.0, 1.0, 1.0,
#else
-#define GEGL_CHANT_TYPE_FILTER
+#define GEGL_CHANT_TYPE_POINT_FILTER
#define GEGL_CHANT_C_FILE "mono-mixer.c"
#include "gegl-chant.h"
@@ -40,166 +40,69 @@ gegl_chant_double_ui (blue, _("Blue"), -10.0, 10.0, 0.25, -1.0, 1.0, 1.0,
static void prepare (GeglOperation *operation)
{
/* set the babl format this operation prefers to work on */
- gegl_operation_set_format (operation, "input", babl_format ("RGBA float"));
+ gegl_operation_set_format (operation, "input", babl_format ("RGBA float"));
gegl_operation_set_format (operation, "output", babl_format ("YA float"));
}
-#include "opencl/gegl-cl.h"
-#include "buffer/gegl-buffer-cl-iterator.h"
-
-static const char* kernel_source =
-"__kernel void Mono_mixer_cl(__global const float4 *src_buf, \n"
-" float4 color, \n"
-" __global float2 *dst_buf) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 tmp = src_buf[gid] * color; \n"
-" dst_buf[gid].x = tmp.x + tmp.y + tmp.z; \n"
-" dst_buf[gid].y = tmp.w; \n"
-"} \n";
-
-static gegl_cl_run_data * cl_data = NULL;
-
-static cl_int
-cl_mono_mixer(cl_mem in_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *roi,
- gfloat red,
- gfloat green,
- gfloat blue)
-{
- cl_int cl_err = 0;
- if (!cl_data)
- {
- const char *kernel_name[] = {"Mono_mixer_cl", NULL};
- cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
- }
- if (!cl_data) return 0;
-
- {
- cl_float4 color;
- color.s[0] = red;
- color.s[1] = green;
- color.s[2] = blue;
- color.s[3] = 1.0f;
-
- 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_float4), (void*)&color);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(), cl_data->kernel[0],
- 1, NULL,
- &global_worksize, NULL,
- 0, NULL, NULL);
- }
-
- return cl_err;
-}
-
static gboolean
-cl_process (GeglOperation *operation,
- GeglBuffer *input,
- GeglBuffer *output,
- const GeglRectangle *result)
-{
- const Babl *in_format = gegl_operation_get_format (operation, "input");
- const Babl *out_format = gegl_operation_get_format (operation, "output");
- gint err;
- gint j;
- cl_int cl_err;
-
- GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
-
- GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE, GEGL_ABYSS_NONE);
- gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
- while (gegl_buffer_cl_iterator_next (i, &err))
- {
- if (err) return FALSE;
- for (j=0; j < i->n; j++)
- {
- cl_err = cl_mono_mixer(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], o->red ,o->green , o->blue);
-
- if (cl_err != CL_SUCCESS)
- {
- g_warning("[OpenCL] Error in gegl:mono-mixer: %s", gegl_cl_errstring(cl_err));
- return FALSE;
- }
- }
- }
- return TRUE;
-}
-
-
-static gboolean
-process (GeglOperation *operation,
- GeglBuffer *input,
- GeglBuffer *output,
- const GeglRectangle *result,
+process (GeglOperation *op,
+ void *in_buf,
+ void *out_buf,
+ glong n_pixels,
+ const GeglRectangle *roi,
gint level)
{
- GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
gfloat red = o->red;
gfloat green = o->green;
gfloat blue = o->blue;
- gfloat *in_buf;
- gfloat *out_buf;
-
- if (gegl_cl_is_accelerated ())
- if (cl_process (operation, input, output, result))
- return TRUE;
-
- if ((result->width > 0) && (result->height > 0))
- {
- gint num_pixels = result->width * result->height;
- gint i;
- gfloat *in_pixel, *out_pixel;
-
- in_buf = g_new (gfloat, 4 * num_pixels);
- out_buf = g_new (gfloat, 2 * num_pixels);
-
- gegl_buffer_get (input, result, 1.0, babl_format ("RGBA float"), in_buf, GEGL_AUTO_ROWSTRIDE, GEGL_ABYSS_NONE);
+ gfloat * GEGL_ALIGNED in_pixel;
+ gfloat * GEGL_ALIGNED out_pixel;
+ glong i;
- in_pixel = in_buf;
- out_pixel = out_buf;
- for (i = 0; i < num_pixels; ++i)
- {
- out_pixel[0] = in_pixel[0] * red + in_pixel[1] * green + in_pixel[2] * blue;
- out_pixel[1] = in_pixel[3];
+ in_pixel = in_buf;
+ out_pixel = out_buf;
- in_pixel += 4;
- out_pixel += 2;
- }
-
- gegl_buffer_set (output, result, 0, babl_format ("YA float"), out_buf,
- GEGL_AUTO_ROWSTRIDE);
-
- g_free (in_buf);
- g_free (out_buf);
- }
-
- return TRUE;
+ for (i=0; i<n_pixels; i++)
+ {
+ out_pixel[0] = in_pixel[0] * red + in_pixel[1] * green + in_pixel[2] * blue;
+ out_pixel[1] = in_pixel[3];
+ in_pixel += 4;
+ out_pixel += 2;
+ }
+ return TRUE;
}
+static const char* kernel_source =
+"__kernel void gegl_mono_mixer (__global const float4 *src_buf, \n"
+" __global float2 *dst_buf, \n"
+" float red, \n"
+" float green, \n"
+" float blue) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = src_buf[gid]; \n"
+" dst_buf[gid].x = in_v.x * red + in_v.y * green + in_v.z * blue; \n"
+" dst_buf[gid].y = in_v.w; \n"
+"} \n";
+
static void
gegl_chant_class_init (GeglChantClass *klass)
{
- GeglOperationClass *operation_class;
- GeglOperationFilterClass *filter_class;
+ GeglOperationClass *operation_class;
+ GeglOperationPointFilterClass *point_filter_class;
- operation_class = GEGL_OPERATION_CLASS (klass);
- filter_class = GEGL_OPERATION_FILTER_CLASS (klass);
+ operation_class = GEGL_OPERATION_CLASS (klass);
+ point_filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
- filter_class->process = process;
operation_class->prepare = prepare;
-
- operation_class->opencl_support = TRUE;
+ point_filter_class->process = process;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:mono-mixer",
"categories" , "color",
"description", _("Monochrome channel mixer"),
+ "cl-source" , kernel_source,
NULL);
}
diff --git a/operations/common/opacity.c b/operations/common/opacity.c
index 3b1fee6..76a9508 100644
--- a/operations/common/opacity.c
+++ b/operations/common/opacity.c
@@ -89,89 +89,20 @@ process (GeglOperation *op,
return TRUE;
}
-#include "opencl/gegl-cl.h"
-
static const char* kernel_source =
-"__kernel void kernel_op_3 (__global const float4 *in, \n"
-" __global const float *aux, \n"
-" __global float4 *out, \n"
-" float value) \n"
+"__kernel void gegl_opacity (__global const float4 *in, \n"
+" __global const float *aux, \n"
+" __global float4 *out, \n"
+" float value) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in [gid]; \n"
-" float aux_v = aux[gid]; \n"
+" float aux_v = (aux)? aux[gid] : 1.0f; \n"
" float4 out_v; \n"
" out_v = in_v * aux_v * value; \n"
" out[gid] = out_v; \n"
-"} \n"
-
-"__kernel void kernel_op_2 (__global const float4 *in, \n"
-" __global float4 *out, \n"
-" float value) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 in_v = in [gid]; \n"
-" float4 out_v; \n"
-" out_v = in_v * value; \n"
-" out[gid] = out_v; \n"
"} \n";
-
-static gegl_cl_run_data *cl_data = NULL;
-
-/* OpenCL processing function */
-static cl_int
-cl_process (GeglOperation *op,
- cl_mem in_tex,
- cl_mem aux_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *roi,
- gint level)
-{
- gfloat value = GEGL_CHANT_PROPERTIES (op)->value;
-
- cl_int cl_err = 0;
-
- if (!cl_data)
- {
- const char *kernel_name[] = {"kernel_op_3", "kernel_op_2", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
- }
-
- if (!cl_data) return 1;
-
- if (aux_tex)
- {
- 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*)&aux_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&value);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
- }
- else
- {
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_float), (void*)&value);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[1], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
- }
-
- return cl_err;
-}
-
/* Fast path when opacity is a no-op
*/
static gboolean operation_process (GeglOperation *operation,
@@ -214,9 +145,6 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class->prepare = prepare;
operation_class->process = operation_process;
point_composer_class->process = process;
- point_composer_class->cl_process = cl_process;
-
- operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:opacity",
@@ -224,6 +152,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"description",
_("Weights the opacity of the input both the value of the aux"
" input and the global value property."),
+ "cl-source" , kernel_source,
NULL);
}
diff --git a/operations/common/over.c b/operations/common/over.c
index 1f68058..7cbe0dd 100644
--- a/operations/common/over.c
+++ b/operations/common/over.c
@@ -70,12 +70,10 @@ process (GeglOperation *op,
return TRUE;
}
-#include "opencl/gegl-cl.h"
-
static const char* kernel_source =
-"__kernel void kernel_over(__global const float4 *in, \n"
-" __global const float4 *aux, \n"
-" __global float4 *out) \n"
+"__kernel void svg_src_over (__global const float4 *in, \n"
+" __global const float4 *aux, \n"
+" __global float4 *out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in [gid]; \n"
@@ -86,48 +84,6 @@ static const char* kernel_source =
" out[gid] = out_v; \n"
"} \n";
-static gegl_cl_run_data *cl_data = NULL;
-
-static cl_int
-cl_process (GeglOperation *op,
- cl_mem in_tex,
- cl_mem aux_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *roi,
- gint level)
-{
- /* Retrieve a pointer to GeglChantO structure which contains all the
- * chanted properties
- */
-
- cl_int cl_err = 0;
-
- if (aux_tex == NULL)
- return 0;
-
- if (!cl_data)
- {
- const char *kernel_name[] = {"kernel_over", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
- }
-
- if (!cl_data) return 1;
-
- 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*)&aux_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- return cl_err;
-}
-
/* Fast paths */
static gboolean operation_process (GeglOperation *operation,
GeglOperationContext *context,
@@ -187,9 +143,6 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class->process = operation_process;
point_composer_class->process = process;
- point_composer_class->cl_process = cl_process;
-
- operation_class->opencl_support = TRUE;
operation_class->compat_name = "gegl:over";
@@ -198,6 +151,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"categories" , "compositors:porter-duff",
"description",
_("Porter Duff operation over (d = cA + cB * (1 - aA))"),
+ "cl-source" , kernel_source,
NULL);
}
diff --git a/operations/common/threshold.c b/operations/common/threshold.c
index 6b15b03..74b18f2 100644
--- a/operations/common/threshold.c
+++ b/operations/common/threshold.c
@@ -34,8 +34,8 @@ gegl_chant_double_ui (value, _("Threshold"), -10.0, 10.0, 0.5, 0.0, 1.0, 1.0,
static void prepare (GeglOperation *operation)
{
- gegl_operation_set_format (operation, "input", babl_format ("YA float"));
- gegl_operation_set_format (operation, "aux", babl_format ("Y float"));
+ gegl_operation_set_format (operation, "input", babl_format ("YA float"));
+ gegl_operation_set_format (operation, "aux", babl_format ("Y float"));
gegl_operation_set_format (operation, "output", babl_format ("YA float"));
}
@@ -89,88 +89,21 @@ process (GeglOperation *op,
return TRUE;
}
-#include "opencl/gegl-cl.h"
-
static const char* kernel_source =
-"__kernel void kernel_thr_3 (__global const float2 *in, \n"
-" __global const float *aux, \n"
-" __global float2 *out) \n"
+"__kernel void gegl_threshold (__global const float2 *in, \n"
+" __global const float *aux, \n"
+" __global float2 *out, \n"
+" float value) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float2 in_v = in [gid]; \n"
-" float aux_v = aux[gid]; \n"
+" float aux_v = (aux)? aux[gid] : value; \n"
" float2 out_v; \n"
" out_v.x = (in_v.x > aux_v)? 1.0f : 0.0f; \n"
" out_v.y = in_v.y; \n"
" out[gid] = out_v; \n"
-"} \n"
-
-"__kernel void kernel_thr_2 (__global const float2 *in, \n"
-" __global float2 *out, \n"
-" float value) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float2 in_v = in [gid]; \n"
-" float2 out_v; \n"
-" out_v.x = (in_v.x > value)? 1.0f : 0.0f; \n"
-" out_v.y = in_v.y; \n"
-" out[gid] = out_v; \n"
"} \n";
-static gegl_cl_run_data *cl_data = NULL;
-
-/* OpenCL processing function */
-static cl_int
-cl_process (GeglOperation *op,
- cl_mem in_tex,
- cl_mem aux_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *roi,
- gint level)
-{
- gfloat value = GEGL_CHANT_PROPERTIES (op)->value;
-
- cl_int cl_err = 0;
-
- if (!cl_data)
- {
- const char *kernel_name[] = {"kernel_thr_3", "kernel_thr_2", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
- }
-
- if (!cl_data) return 1;
-
- if (aux_tex)
- {
- 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*)&aux_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_mem), (void*)&out_tex);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
- }
- else
- {
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex);
- cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_float), (void*)&value);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[1], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
- }
-
- return cl_err;
-}
-
static void
gegl_chant_class_init (GeglChantClass *klass)
{
@@ -181,18 +114,16 @@ gegl_chant_class_init (GeglChantClass *klass)
point_composer_class = GEGL_OPERATION_POINT_COMPOSER_CLASS (klass);
point_composer_class->process = process;
- point_composer_class->cl_process = cl_process;
operation_class->prepare = prepare;
- operation_class->opencl_support = TRUE;
-
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:threshold",
"categories" , "color",
"description",
_("Thresholds the image to white/black based on either the global value "
"set in the value property, or per pixel from the aux input."),
- NULL);
+ "cl-source" , kernel_source,
+ NULL);
}
#endif
diff --git a/operations/common/value-invert.c b/operations/common/value-invert.c
index b9b4be5..2403582 100644
--- a/operations/common/value-invert.c
+++ b/operations/common/value-invert.c
@@ -129,11 +129,9 @@ process (GeglOperation *op,
return TRUE;
}
-#include "opencl/gegl-cl.h"
-
static const char* kernel_source =
-"__kernel void kernel_vinv(__global const float4 *in, \n"
-" __global float4 *out) \n"
+"__kernel void gegl_value_invert (__global const float4 *in, \n"
+" __global float4 *out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
@@ -173,43 +171,6 @@ static const char* kernel_source =
" out[gid] = out_v; \n"
"} \n";
-static gegl_cl_run_data *cl_data = NULL;
-
-/* OpenCL processing function */
-static cl_int
-cl_process (GeglOperation *op,
- cl_mem in_tex,
- cl_mem out_tex,
- size_t global_worksize,
- const GeglRectangle *roi,
- int level)
-{
-
- cl_int cl_err = 0;
-
- if (!cl_data)
- {
- const char *kernel_name[] = {"kernel_vinv", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
- }
-
- if (!cl_data) return 1;
-
- 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);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
- cl_data->kernel[0], 1,
- NULL, &global_worksize, NULL,
- 0, NULL, NULL);
- if (cl_err != CL_SUCCESS) return cl_err;
-
- return cl_err;
-}
-
-
-
static void
gegl_chant_class_init (GeglChantClass *klass)
{
@@ -221,9 +182,6 @@ gegl_chant_class_init (GeglChantClass *klass)
point_filter_class->process = process;
operation_class->prepare = prepare;
- point_filter_class->cl_process = cl_process;
-
- operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:value-invert",
@@ -231,7 +189,8 @@ gegl_chant_class_init (GeglChantClass *klass)
"description",
_("Inverts just the value component, the result is the corresponding "
"`inverted' image."),
- NULL);
+ "cl-source" , kernel_source,
+ NULL);
}
#endif
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
