[gegl] OpenCL support in some point-filter operators
- From: Ãyvind KolÃs <ok src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] OpenCL support in some point-filter operators
- Date: Tue, 20 Mar 2012 13:52:34 +0000 (UTC)
commit 6307eb054efd201f24f2cebe6e2f404e13a071cd
Author: Victor Oliveira <victormatheus gmail com>
Date: Wed Feb 8 16:46:04 2012 -0200
OpenCL support in some point-filter operators
* color-temperature
* invert
* value-invert
* whitebalance
operations/common/color-temperature.c | 64 ++++++++++++++++++++++++++
operations/common/invert.c | 50 ++++++++++++++++++++
operations/common/value-invert.c | 81 +++++++++++++++++++++++++++++++++
operations/common/whitebalance.c | 72 +++++++++++++++++++++++++++++
4 files changed, 267 insertions(+), 0 deletions(-)
---
diff --git a/operations/common/color-temperature.c b/operations/common/color-temperature.c
index 0015552..aab7d6a 100644
--- a/operations/common/color-temperature.c
+++ b/operations/common/color-temperature.c
@@ -170,6 +170,68 @@ process (GeglOperation *op,
return TRUE;
}
+#include "opencl/gegl-cl.h"
+
+static const char* kernel_source =
+"__kernel void kernel_temp(__global const float4 *in, \n"
+" __global float4 *out, \n"
+" float3 coeff) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v.xyz = in_v.xyz * coeff.xyz; \n"
+" out_v.w = in_v.w; \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 out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi)
+{
+ /* Retrieve a pointer to GeglChantO structure which contains all the
+ * chanted properties
+ */
+
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
+ const gfloat *coeffs = o->chant_data;
+
+ cl_int cl_err = 0;
+
+ if (! coeffs)
+ {
+ coeffs = o->chant_data = preprocess (o);
+ }
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"kernel_temp", 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);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float3), (void*)coeffs);
+ 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)
@@ -188,8 +250,10 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class->prepare = prepare;
point_filter_class->process = process;
+ point_filter_class->cl_process = cl_process;
operation_class->name = "gegl:color-temperature";
+ operation_class->opencl_support = TRUE;
operation_class->categories = "color";
operation_class->description =
_("Allows changing the color temperature of an image.");
diff --git a/operations/common/invert.c b/operations/common/invert.c
index 613cea0..448a45c 100644
--- a/operations/common/invert.c
+++ b/operations/common/invert.c
@@ -60,6 +60,54 @@ 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v.xyz = (1.0f - in_v.xyz); \n"
+" out_v.w = in_v.w; \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 out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi)
+{
+ 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)
{
@@ -70,8 +118,10 @@ 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->name = "gegl:invert";
+ operation_class->opencl_support = TRUE;
operation_class->categories = "color";
operation_class->description =
_("Inverts the components (except alpha), the result is the "
diff --git a/operations/common/value-invert.c b/operations/common/value-invert.c
index f359f44..7a29df6 100644
--- a/operations/common/value-invert.c
+++ b/operations/common/value-invert.c
@@ -115,6 +115,85 @@ 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" \n"
+" float value = fmax (in_v.x, fmax (in_v.y, in_v.z)); \n"
+" float minv = fmin (in_v.x, fmin (in_v.y, in_v.z)); \n"
+" float delta = value - minv; \n"
+" \n"
+" if (value == 0.0f || delta == 0.0f) \n"
+" { \n"
+" out_v.xyz = (float3) (1.0f - value); \n"
+" } \n"
+" else \n"
+" { \n"
+" if (in_v.x == value) \n"
+" { \n"
+" out_v.xzy = (float3) ((1.0f - in_v.x), \n"
+" (1.0f - in_v.x) * in_v.z / value, \n"
+" (1.0f - in_v.x) * in_v.y / value); \n"
+" } \n"
+" else if (in_v.y == value) \n"
+" { \n"
+" out_v.yxz = (float3) ((1.0f - in_v.y), \n"
+" (1.0f - in_v.y) * in_v.x / value, \n"
+" (1.0f - in_v.y) * in_v.z / value); \n"
+" } \n"
+" else \n"
+" { \n"
+" out_v.zyx = (float3) ((1.0f - in_v.z), \n"
+" (1.0f - in_v.z) * in_v.y / value, \n"
+" (1.0f - in_v.z) * in_v.x / value); \n"
+" } \n"
+" } \n"
+" \n"
+" out_v.w = in_v.w; \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 out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi)
+{
+
+ 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)
@@ -126,8 +205,10 @@ 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->name = "gegl:value-invert";
+ operation_class->opencl_support = TRUE;
operation_class->categories = "color";
operation_class->description =
_("Inverts just the value component, the result is the corresponding "
diff --git a/operations/common/whitebalance.c b/operations/common/whitebalance.c
index dcb015a..e494d14 100644
--- a/operations/common/whitebalance.c
+++ b/operations/common/whitebalance.c
@@ -85,6 +85,76 @@ process (GeglOperation *op,
return TRUE;
}
+#include "opencl/gegl-cl.h"
+
+static const char* kernel_source =
+"__kernel void kernel_wb(__global const float4 *in, \n"
+" __global float4 *out, \n"
+" float a_base, \n"
+" float a_scale, \n"
+" float b_base, \n"
+" float b_scale, \n"
+" float saturation) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = (float4) (in_v.x, \n"
+" (in_v.y + in_v.x * a_scale + a_base) * saturation, \n"
+" (in_v.z + in_v.x * b_scale + b_base) * saturation, \n"
+" in_v.w); \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 out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi)
+{
+ /* Retrieve a pointer to GeglChantO structure which contains all the
+ * chanted properties
+ */
+
+ GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
+
+ gfloat a_scale = (o->high_a_delta - o->low_a_delta);
+ gfloat a_base = o->low_a_delta;
+ gfloat b_scale = (o->high_b_delta - o->low_b_delta);
+ gfloat b_base = o->low_b_delta;
+ gfloat saturation = o->saturation;
+
+ cl_int cl_err = 0;
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"kernel_wb", 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);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&a_base);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&a_scale);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&b_base);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_float), (void*)&b_scale);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_float), (void*)&saturation);
+ 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)
@@ -96,9 +166,11 @@ 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->name = "gegl:whitebalance";
+ operation_class->opencl_support = TRUE;
operation_class->categories = "color";
operation_class->description =
_("Allows changing the whitepoint and blackpoint of an image.");
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
