Commit 911d312b authored by Carlos Zubieta's avatar Carlos Zubieta Committed by Téo Mazars

Added OpenCL support to weighted-blend

parent 1f74bccd
__kernel void cl_copy_weigthed_blend(__global const float4 *in,
__global float4 *out)
{
int gid = get_global_id(0);
float4 in_v = in[gid];
out[gid] = in_v;
}
__kernel void cl_weighted_blend(__global const float4 *in,
__global const float4 *aux,
__global float4 *out)
{
int gid = get_global_id(0);
float4 in_v = in[gid];
float4 aux_v = aux[gid];
float4 out_v;
float in_weight;
float aux_weight;
float total_alpha = in_v.w + aux_v.w;
total_alpha = total_alpha == 0 ? 1 : total_alpha;
in_weight = in_v.w / total_alpha;
aux_weight = 1.0f - in_weight;
out_v.xyz = in_weight * in_v.xyz + aux_weight * aux_v.xyz;
out_v.w = total_alpha;
out[gid] = out_v;
}
static const char* weighted_blend_cl_source =
"__kernel void cl_copy_weigthed_blend(__global const float4 *in, \n"
" __global float4 *out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" out[gid] = in_v; \n"
"} \n"
" \n"
"__kernel void cl_weighted_blend(__global const float4 *in, \n"
" __global const float4 *aux, \n"
" __global float4 *out) \n"
"{ \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 aux_v = aux[gid]; \n"
" float4 out_v; \n"
" float in_weight; \n"
" float aux_weight; \n"
" float total_alpha = in_v.w + aux_v.w; \n"
" \n"
" total_alpha = total_alpha == 0 ? 1 : total_alpha; \n"
" \n"
" in_weight = in_v.w / total_alpha; \n"
" aux_weight = 1.0f - in_weight; \n"
" \n"
" out_v.xyz = in_weight * in_v.xyz + aux_weight * aux_v.xyz; \n"
" out_v.w = total_alpha; \n"
" out[gid] = out_v; \n"
"} \n"
;
......@@ -22,8 +22,6 @@
#ifdef GEGL_CHANT_PROPERTIES
gegl_chant_double (value, _("Value"), -G_MAXDOUBLE, G_MAXDOUBLE, 0.0, _("global value used if aux doesn't contain data"))
#else
#define GEGL_CHANT_TYPE_POINT_COMPOSER
......@@ -40,6 +38,68 @@ static void prepare (GeglOperation *operation)
gegl_operation_set_format (operation, "output", format);
}
#include "opencl/gegl-cl.h"
#include "opencl/weighted-blend.cl.h"
static GeglClRunData *cl_data = NULL;
static gboolean
cl_process (GeglOperation *self,
cl_mem in_tex,
cl_mem aux_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
gint cl_err = 0;
if (!cl_data)
{
const char *kernel_name[] = {"cl_copy_weigthed_blend",
"cl_weighted_blend",
NULL};
cl_data = gegl_cl_compile_and_build (weighted_blend_cl_source,
kernel_name);
}
if (!cl_data) return TRUE;
if (!aux_tex)
{
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex);
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;
}
else
{
cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&aux_tex);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_mem), (void*)&out_tex);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[1], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
CL_CHECK;
}
return FALSE;
error:
return TRUE;
}
static gboolean
process (GeglOperation *op,
void *in_buf,
......@@ -58,7 +118,6 @@ process (GeglOperation *op,
{
/* there is no auxilary buffer.
* output the input buffer.
* gfloat value = GEGL_CHANT_PROPERTIES (op)->value;
*/
for (i = 0; i < n_pixels; i++)
{
......@@ -67,6 +126,8 @@ process (GeglOperation *op,
{
out[j] = in[j];
}
in += 4;
out += 4;
}
}
else
......@@ -114,8 +175,10 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class = GEGL_OPERATION_CLASS (klass);
point_composer_class = GEGL_OPERATION_POINT_COMPOSER_CLASS (klass);
point_composer_class->process = process;
operation_class->prepare = prepare;
point_composer_class->process = process;
point_composer_class->cl_process = cl_process;
operation_class->prepare = prepare;
operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name" , "gegl:weighted-blend",
......
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