Commit 25fe8d70 authored by Victor Oliveira's avatar Victor Oliveira

Add OpenCL support for gegl:levels

parent 06aed852
......@@ -85,6 +85,82 @@ process (GeglOperation *op,
return TRUE;
}
#include "opencl/gegl-cl.h"
static const char* kernel_source =
"__kernel void kernel_levels(__global const float4 *in, \n"
" __global float4 *out, \n"
" float in_offset, \n"
" float out_offset, \n"
" float scale) \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 - in_offset) * scale + out_offset; \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,
gint level)
{
/* Retrieve a pointer to GeglChantO structure which contains all the
* chanted properties
*/
GeglChantO *o = GEGL_CHANT_PROPERTIES (op);
gfloat in_range;
gfloat out_range;
gfloat in_offset;
gfloat out_offset;
gfloat scale;
cl_int cl_err = 0;
in_offset = o->in_low * 1.0;
out_offset = o->out_low * 1.0;
in_range = o->in_high-o->in_low;
out_range = o->out_high-o->out_low;
if (in_range == 0.0)
in_range = 0.00000001;
scale = out_range/in_range;
if (!cl_data)
{
const char *kernel_name[] = {"kernel_levels", 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*)&in_offset);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&out_offset);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&scale);
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,6 +172,9 @@ 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->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:levels",
......
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