Commit 1f74bccd authored by Carlos Zubieta's avatar Carlos Zubieta Committed by Téo Mazars

Added OpenCL support to color-to-alpha

parent 4f35eba8
__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;
}
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"
;
......@@ -130,6 +130,58 @@ color_to_alpha (const gfloat *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;
else
{
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
......@@ -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",
......
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