[gegl] opencl support for gegl:threshold and YA float color format
- From: Ãyvind KolÃs <ok src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl support for gegl:threshold and YA float color format
- Date: Tue, 20 Mar 2012 13:53:40 +0000 (UTC)
commit cc7edb662fa817f6cd034a7297d64fd111c01a7f
Author: Victor Oliveira <victormatheus gmail com>
Date: Fri Feb 24 17:21:03 2012 -0200
opencl support for gegl:threshold and YA float color format
gegl/opencl/gegl-cl-color-kernel.h | 67 +++++++++++++++++++++++++++++
gegl/opencl/gegl-cl-color.c | 22 +++++++++-
operations/common/threshold.c | 83 ++++++++++++++++++++++++++++++++++++
3 files changed, 171 insertions(+), 1 deletions(-)
---
diff --git a/gegl/opencl/gegl-cl-color-kernel.h b/gegl/opencl/gegl-cl-color-kernel.h
index 0d161e3..6a3589e 100644
--- a/gegl/opencl/gegl-cl-color-kernel.h
+++ b/gegl/opencl/gegl-cl-color-kernel.h
@@ -262,6 +262,7 @@ static const char* kernel_color_source =
/* -- Y u8 -- */
+/* Y u8 -> Y float */
"__kernel void yu8_to_yf (__global const uchar * in, \n"
" __global float * out) \n"
"{ \n"
@@ -270,4 +271,70 @@ static const char* kernel_color_source =
" float out_v; \n"
" out_v = in_v; \n"
" out[gid] = out_v; \n"
+"} \n"
+
+/* -- YA float -- */
+
+" /* source: http://www.poynton.com/ColorFAQ.html */ \n"
+" #define RGB_LUMINANCE_RED (0.212671f) \n"
+" #define RGB_LUMINANCE_GREEN (0.715160f) \n"
+" #define RGB_LUMINANCE_BLUE (0.072169f) \n"
+
+/* RGBA float -> YA float */
+"__kernel void rgbaf_to_yaf (__global const float4 * in, \n"
+" __global float2 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float2 out_v; \n"
+" \n"
+" float luminance = in_v.x * RGB_LUMINANCE_RED + \n"
+" in_v.y * RGB_LUMINANCE_GREEN + \n"
+" in_v.z * RGB_LUMINANCE_BLUE; \n"
+" \n"
+" out_v.x = luminance; \n"
+" out_v.y = in_v.w; \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+
+/* YA float -> RGBA float */
+"__kernel void yaf_to_rgbaf (__global const float2 * in, \n"
+" __global float4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in[gid]; \n"
+" float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y); \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+
+
+/* RGBA u8 -> YA float */
+"__kernel void rgbau8_to_yaf (__global const uchar4 * in, \n"
+" __global float2 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = convert_float4(in[gid]) / 255.0f; \n"
+" float2 out_v; \n"
+" \n"
+" float luminance = in_v.x * RGB_LUMINANCE_RED + \n"
+" in_v.y * RGB_LUMINANCE_GREEN + \n"
+" in_v.z * RGB_LUMINANCE_BLUE; \n"
+" \n"
+" out_v.x = luminance; \n"
+" out_v.y = in_v.w; \n"
+" \n"
+" out[gid] = out_v; \n"
+"} \n"
+
+/* YA float -> RGBA u8 */
+"__kernel void yaf_to_rgbau8 (__global const float2 * in, \n"
+" __global uchar4 * out) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in[gid]; \n"
+" float4 out_v = (float4) (in_v.x, in_v.x, in_v.x, in_v.y); \n"
+" \n"
+" out[gid] = convert_uchar4_sat_rte(255.0f * out_v); \n"
"} \n";
diff --git a/gegl/opencl/gegl-cl-color.c b/gegl/opencl/gegl-cl-color.c
index 83a7d20..62b33d1 100644
--- a/gegl/opencl/gegl-cl-color.c
+++ b/gegl/opencl/gegl-cl-color.c
@@ -6,7 +6,7 @@
static gegl_cl_run_data *kernels_color = NULL;
-#define CL_FORMAT_N 8
+#define CL_FORMAT_N 9
static const Babl *format[CL_FORMAT_N];
@@ -34,6 +34,11 @@ CL_RGBU8_TO_RGBAF = 14,
CL_RGBAF_TO_RGBU8 = 15,
CL_YU8_TO_YF = 16,
+
+CL_RGBAF_TO_YAF = 17,
+CL_YAF_TO_RGBAF = 18,
+CL_RGBAU8_TO_YAF = 19,
+CL_YAF_TO_RGBAU8 = 20,
};
void
@@ -62,6 +67,11 @@ gegl_cl_color_compile_kernels(void)
"yu8_to_yf", /* 16 */
+ "rgbaf_to_yaf", /* 17 */
+ "yaf_to_rgbaf", /* 18 */
+ "rgbau8_to_yaf", /* 19 */
+ "yaf_to_rgbau8", /* 20 */
+
NULL};
format[0] = babl_format ("RGBA u8");
@@ -72,6 +82,7 @@ gegl_cl_color_compile_kernels(void)
format[5] = babl_format ("RGB u8");
format[6] = babl_format ("Y u8");
format[7] = babl_format ("Y float");
+ format[8] = babl_format ("YA float");
kernels_color = gegl_cl_compile_and_build (kernel_color_source, kernel_name);
}
@@ -89,6 +100,7 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
else if (out_format == babl_format ("R'G'B'A float")) kernel = CL_RGBAF_TO_RGBA_GAMMA_F;
else if (out_format == babl_format ("Y'CbCrA float")) kernel = CL_RGBAF_TO_YCBCRAF;
else if (out_format == babl_format ("RGB u8")) kernel = CL_RGBAF_TO_RGBU8;
+ else if (out_format == babl_format ("YA float")) kernel = CL_RGBAF_TO_YAF;
}
else if (in_format == babl_format ("RGBA u8"))
{
@@ -96,6 +108,7 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
else if (out_format == babl_format ("RaGaBaA float")) kernel = CL_RGBAU8_TO_RAGABAF;
else if (out_format == babl_format ("R'G'B'A float")) kernel = CL_RGBAU8_TO_RGBA_GAMMA_F;
else if (out_format == babl_format ("Y'CbCrA float")) kernel = CL_RGBAU8_TO_YCBCRAF;
+ else if (out_format == babl_format ("YA float")) kernel = CL_RGBAU8_TO_YAF;
}
else if (in_format == babl_format ("RaGaBaA float"))
{
@@ -120,6 +133,11 @@ choose_kernel (const Babl *in_format, const Babl *out_format)
{
if (out_format == babl_format ("Y float")) kernel = CL_YU8_TO_YF;
}
+ else if (in_format == babl_format ("YA float"))
+ {
+ if (out_format == babl_format ("RGBA float")) kernel = CL_YAF_TO_RGBAF;
+ else if (out_format == babl_format ("RGBA u8")) kernel = CL_YAF_TO_RGBAU8;
+ }
return kernel;
}
@@ -146,6 +164,8 @@ gegl_cl_color_babl (const Babl *buffer_format, size_t *bytes)
*bytes = sizeof (cl_uchar);
else if (buffer_format == babl_format ("Y float"))
*bytes = sizeof (cl_float);
+ else if (buffer_format == babl_format ("YA float"))
+ *bytes = sizeof (cl_float2);
else
*bytes = sizeof (cl_float4);
}
diff --git a/operations/common/threshold.c b/operations/common/threshold.c
index cef4057..f546c4d 100644
--- a/operations/common/threshold.c
+++ b/operations/common/threshold.c
@@ -88,6 +88,87 @@ 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"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float2 in_v = in [gid]; \n"
+" float aux_v = aux[gid]; \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)
+{
+ 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)
{
@@ -98,9 +179,11 @@ 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->name = "gegl:threshold";
+ operation_class->opencl_support = TRUE;
operation_class->categories = "color";
operation_class->description =
_("Thresholds the image to white/black based on either the global value "
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
