Commit cc7edb66 authored by Victor Oliveira's avatar Victor Oliveira Committed by Øyvind "pippin" Kolås

opencl support for gegl:threshold and YA float color format

parent 4a20977e
......@@ -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";
......@@ -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);
}
......
......@@ -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 "
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment