[gegl/soc-2013-opecl-ops: 1/6] Added OpenCL support to color-to-alpha
- From: Carlos Zubieta <czubieta src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl/soc-2013-opecl-ops: 1/6] Added OpenCL support to color-to-alpha
- Date: Wed, 24 Jul 2013 00:24:25 +0000 (UTC)
commit 81f5de46655b71c773b4447c6396fe61ca12a6ef
Author: Carlos Zubieta <czubieta dev gmail com>
Date: Tue Jul 23 18:08:06 2013 -0500
Added OpenCL support to color-to-alpha
opencl/color-to-alpha.cl | 63 ++++++++++++++++++++++++++++++
opencl/color-to-alpha.cl.h | 65 +++++++++++++++++++++++++++++++
operations/common/color-to-alpha.c | 75 +++++++++++++++++++++++++++++++-----
3 files changed, 193 insertions(+), 10 deletions(-)
---
diff --git a/opencl/color-to-alpha.cl b/opencl/color-to-alpha.cl
new file mode 100644
index 0000000..9ca0894
--- /dev/null
+++ b/opencl/color-to-alpha.cl
@@ -0,0 +1,63 @@
+__kernel void cl_color_to_alpha(__global const float4 *in,
+ __global float4 *out,
+ float4 color)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v = in_v;
+ float4 alpha;
+
+ alpha.w = in_v.w;
+
+ /*First component*/
+ if ( color.x < 0.00001f )
+ alpha.x = in_v.x;
+ else if ( in_v.x > color.x + 0.00001f )
+ alpha.x = (in_v.x - color.x) / (1.0f - color.x);
+ else if ( in_v.x < color.x - 0.00001f )
+ alpha.x = (color.x - in_v.x) / color.x;
+ else
+ alpha.x = 0.0f;
+ /*Second component*/
+ if ( color.y < 0.00001f )
+ alpha.y = in_v.y;
+ else if ( in_v.y > color.y + 0.00001f )
+ alpha.y = (in_v.y - color.y) / (1.0f - color.y);
+ else if ( in_v.y < color.y - 0.00001f )
+ alpha.y = (color.y - in_v.y) / color.y;
+ else
+ alpha.y = 0.0f;
+ /*Third component*/
+ if ( color.z < 0.00001f )
+ alpha.z = in_v.z;
+ else if ( in_v.z > color.z + 0.00001f )
+ alpha.z = (in_v.z - color.z) / (1.0f - color.z);
+ else if ( in_v.z < color.z - 0.00001f )
+ alpha.z = (color.z - in_v.z) / color.z;
+ else
+ alpha.z = 0.0f;
+
+ if (alpha.x > alpha.y)
+ {
+ if (alpha.x > alpha.z)
+ out_v.w = alpha.x;
+ else
+ out_v.w = alpha.z;
+ }
+ else if (alpha.y > alpha.z)
+ {
+ out_v.w = alpha.y;
+ }
+ else
+ {
+ out_v.w = alpha.z;
+ }
+
+ if (out_v.w >= 0.00001f)
+ {
+ out_v.xyz = (out_v.xyz - color.xyz) / out_v.www + color.xyz;
+ out_v.w *= alpha.w;
+ }
+
+ out[gid] = out_v;
+}
diff --git a/opencl/color-to-alpha.cl.h b/opencl/color-to-alpha.cl.h
new file mode 100644
index 0000000..0efc5af
--- /dev/null
+++ b/opencl/color-to-alpha.cl.h
@@ -0,0 +1,65 @@
+static const char* color_to_alpha_cl_source =
+"__kernel void cl_color_to_alpha(__global const float4 *in, \n"
+" __global float4 *out, \n"
+" float4 color) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v = in_v; \n"
+" float4 alpha; \n"
+" \n"
+" alpha.w = in_v.w; \n"
+" \n"
+" /*First component*/ \n"
+" if ( color.x < 0.00001f ) \n"
+" alpha.x = in_v.x; \n"
+" else if ( in_v.x > color.x + 0.00001f ) \n"
+" alpha.x = (in_v.x - color.x) / (1.0f - color.x); \n"
+" else if ( in_v.x < color.x - 0.00001f ) \n"
+" alpha.x = (color.x - in_v.x) / color.x; \n"
+" else \n"
+" alpha.x = 0.0f; \n"
+" /*Second component*/ \n"
+" if ( color.y < 0.00001f ) \n"
+" alpha.y = in_v.y; \n"
+" else if ( in_v.y > color.y + 0.00001f ) \n"
+" alpha.y = (in_v.y - color.y) / (1.0f - color.y); \n"
+" else if ( in_v.y < color.y - 0.00001f ) \n"
+" alpha.y = (color.y - in_v.y) / color.y; \n"
+" else \n"
+" alpha.y = 0.0f; \n"
+" /*Third component*/ \n"
+" if ( color.z < 0.00001f ) \n"
+" alpha.z = in_v.z; \n"
+" else if ( in_v.z > color.z + 0.00001f ) \n"
+" alpha.z = (in_v.z - color.z) / (1.0f - color.z); \n"
+" else if ( in_v.z < color.z - 0.00001f ) \n"
+" alpha.z = (color.z - in_v.z) / color.z; \n"
+" else \n"
+" alpha.z = 0.0f; \n"
+" \n"
+" if (alpha.x > alpha.y) \n"
+" { \n"
+" if (alpha.x > alpha.z) \n"
+" out_v.w = alpha.x; \n"
+" else \n"
+" out_v.w = alpha.z; \n"
+" } \n"
+" else if (alpha.y > alpha.z) \n"
+" { \n"
+" out_v.w = alpha.y; \n"
+" } \n"
+" else \n"
+" { \n"
+" out_v.w = alpha.z; \n"
+" } \n"
+" \n"
+" if (out_v.w >= 0.00001f) \n"
+" { \n"
+" out_v.xyz = (out_v.xyz - color.xyz) / out_v.www + color.xyz; \n"
+" out_v.w *= alpha.w; \n"
+" } \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+;
diff --git a/operations/common/color-to-alpha.c b/operations/common/color-to-alpha.c
index aae21c4..df96b70 100644
--- a/operations/common/color-to-alpha.c
+++ b/operations/common/color-to-alpha.c
@@ -42,9 +42,9 @@ static void
prepare (GeglOperation *operation)
{
gegl_operation_set_format (operation, "input",
- babl_format ("R'G'B'A double"));
+ babl_format ("R'G'B'A float"));
gegl_operation_set_format (operation, "output",
- babl_format ("R'G'B'A double"));
+ babl_format ("R'G'B'A float"));
}
/*
@@ -81,12 +81,12 @@ prepare (GeglOperation *operation)
*/
static void
-color_to_alpha (const gdouble *color,
- const gdouble *src,
- gdouble *dst)
+color_to_alpha (const gfloat *color,
+ const gfloat *src,
+ gfloat *dst)
{
gint i;
- gdouble alpha[4];
+ gfloat alpha[4];
for (i=0; i<4; i++)
dst[i] = src[i];
@@ -130,6 +130,58 @@ color_to_alpha (const gdouble *color,
dst[3] *= alpha[3];
}
+#include "opencl/gegl-cl.h"
+#include "opencl/color-to-alpha.cl.h"
+
+static GeglClRunData * cl_data = NULL;
+
+static gboolean
+cl_process (GeglOperation *operation,
+ cl_mem in,
+ cl_mem out,
+ size_t global_worksize,
+ const GeglRectangle *roi,
+ gint level)
+{
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
+ gfloat color[4];
+ gegl_color_get_pixel (o->color, babl_format ("R'G'B'A float"), color);
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"cl_color_to_alpha",NULL};
+ cl_data = gegl_cl_compile_and_build (color_to_alpha_cl_source, kernel_name);
+ }
+ if (!cl_data) return TRUE;
+
+ {
+ 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);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out);
+ CL_CHECK;
+ cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float4),(void*)&f_color);
+ CL_CHECK;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 1,
+ NULL, &global_worksize, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
+ }
+
+ return FALSE;
+
+error:
+ return TRUE;
+
+}
static gboolean
@@ -141,12 +193,12 @@ process (GeglOperation *operation,
gint level)
{
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
- const Babl *format = babl_format ("R'G'B'A double");
- gdouble color[4];
+ const Babl *format = babl_format ("R'G'B'A float");
+ gfloat color[4];
gint x;
- gdouble *in_buff = in_buf;
- gdouble *out_buff = out_buf;
+ gfloat *in_buff = in_buf;
+ gfloat *out_buff = out_buf;
gegl_color_get_pixel (o->color, format, color);
@@ -192,7 +244,10 @@ gegl_chant_class_init (GeglChantClass *klass)
filter_class = GEGL_OPERATION_POINT_FILTER_CLASS (klass);
filter_class->process = process;
+ filter_class->cl_process = cl_process;
+
operation_class->prepare = prepare;
+ operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name", "gegl:color-to-alpha",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
