Commit be8583c3 authored by Daniel Sabo's avatar Daniel Sabo

gegl:checkerboard: Add OpenCL version for float4 formats

This only supports float4 formats for now because OpenCL doesn't
necessarily allow byte addressing. Other formats will fall back
to the CPU version.
parent 4f8cc46f
......@@ -55,7 +55,8 @@ gegl_chant_format (format, _("Babl Format"),
#define GEGL_CHANT_C_FILE "checkerboard.c"
#include "gegl-chant.h"
#include <gegl-utils.h>
#include <gegl-buffer-cl-iterator.h>
#include <gegl-debug.h>
static void
prepare (GeglOperation *operation)
......@@ -74,13 +75,131 @@ get_bounding_box (GeglOperation *operation)
return gegl_rectangle_infinite_plane ();
}
static const char* checkerboard_cl_source =
"static inline int tile_index (int coordinate, int stride) \n"
"{ \n"
" if (coordinate >= 0) \n"
" return coordinate / stride; \n"
" else \n"
" return ((coordinate + 1) / stride) - 1; \n"
"} \n"
" \n"
"__kernel void kernel_checkerboard (__global float4 *out, \n"
" float4 color1, \n"
" float4 color2, \n"
" int square_width, \n"
" int square_height, \n"
" int x_offset, \n"
" int y_offset, \n"
" int roi_x, \n"
" int roi_y, \n"
" int roi_width) \n"
"{ \n"
" int gidx = 0; \n"
" int gidy = get_global_id(0); \n"
" float4 cur_color; \n"
" bool in_color1; \n"
" \n"
" int x = roi_x + gidx - x_offset; \n"
" int y = roi_y + gidy - y_offset; \n"
" \n"
" int tilex = tile_index (x, square_width); \n"
" int tiley = tile_index (y, square_height); \n"
" \n"
" if ((tilex + tiley) % 2 == 0) \n"
" { \n"
" cur_color = color1; \n"
" in_color1 = true; \n"
" } \n"
" else \n"
" { \n"
" cur_color = color2; \n"
" in_color1 = false; \n"
" } \n"
" \n"
" int stripe_end = (tilex + 1) * square_width; \n"
" int stripe_width = stripe_end - x; \n"
" int gidx_max = roi_width; \n"
" \n"
" while (gidx < gidx_max) \n"
" { \n"
" out[gidx++ + gidy * roi_width] = cur_color; \n"
" stripe_width--; \n"
" \n"
" if (stripe_width == 0) \n"
" { \n"
" stripe_width = square_width; \n"
" \n"
" if (in_color1) \n"
" cur_color = color2; \n"
" else \n"
" cur_color = color1; \n"
" in_color1 = !in_color1; \n"
" } \n"
" } \n"
"} \n";
#define TILE_INDEX(coordinate,stride) \
(((coordinate) >= 0)?\
(coordinate) / (stride):\
((((coordinate) + 1) /(stride)) - 1))
static GeglClRunData *cl_data = NULL;
static gboolean
process (GeglOperation *operation,
checkerboard_cl_process (GeglOperation *operation,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *roi,
gint level)
{
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
const Babl *out_format = gegl_operation_get_format (operation, "output");
const size_t gbl_size[1] = {roi->height};
cl_int cl_err = 0;
float color1[4];
float color2[4];
if (!cl_data)
{
const char *kernel_name[] = {"kernel_checkerboard", NULL};
cl_data = gegl_cl_compile_and_build (checkerboard_cl_source, kernel_name);
if (!cl_data)
return TRUE;
}
gegl_color_get_pixel (o->color1, out_format, color1);
gegl_color_get_pixel (o->color2, out_format, color2);
cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
sizeof(cl_mem), &out_tex,
sizeof(color1), &color1,
sizeof(color2), &color2,
sizeof(cl_int), &o->x,
sizeof(cl_int), &o->y,
sizeof(cl_int), &o->x_offset,
sizeof(cl_int), &o->y_offset,
sizeof(cl_int), &roi->x,
sizeof(cl_int), &roi->y,
sizeof(cl_int), &roi->width,
NULL);
CL_CHECK;
cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, gbl_size, NULL,
0, NULL, NULL);
CL_CHECK;
return FALSE;
error:
return TRUE;
}
static gboolean
checkerboard_process (GeglOperation *operation,
void *out_buf,
glong n_pixels,
const GeglRectangle *roi,
......@@ -140,16 +259,66 @@ process (GeglOperation *operation,
return TRUE;
}
static gboolean
operation_source_process (GeglOperation *operation,
GeglBuffer *output,
const GeglRectangle *result,
gint level)
{
const Babl *out_format = gegl_operation_get_format (operation, "output");
if ((result->width > 0) && (result->height > 0))
{
GeglBufferIterator *iter;
if (gegl_operation_use_opencl (operation) &&
babl_format_get_n_components (out_format) == 4 &&
babl_format_get_type (out_format, 0) == babl_type ("float"))
{
GeglBufferClIterator *cl_iter;
gboolean err;
GEGL_NOTE (GEGL_DEBUG_OPENCL, "GEGL_OPERATION_POINT_RENDER: %s", GEGL_OPERATION_GET_CLASS (operation)->name);
cl_iter = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
while (gegl_buffer_cl_iterator_next (cl_iter, &err) && !err)
{
err = checkerboard_cl_process (operation, cl_iter->tex[0], cl_iter->size[0], &cl_iter->roi[0], level);
if (err)
{
gegl_buffer_cl_iterator_stop (cl_iter);
break;
}
}
if (err)
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", GEGL_OPERATION_GET_CLASS (operation)->name);
else
return TRUE;
}
iter = gegl_buffer_iterator_new (output, result, level, out_format, GEGL_BUFFER_WRITE, GEGL_ABYSS_NONE);
while (gegl_buffer_iterator_next (iter))
checkerboard_process (operation, iter->data[0], iter->length, &iter->roi[0], level);
}
return TRUE;
}
static void
gegl_chant_class_init (GeglChantClass *klass)
{
GeglOperationClass *operation_class;
GeglOperationSourceClass *source_class;
GeglOperationPointRenderClass *point_render_class;
operation_class = GEGL_OPERATION_CLASS (klass);
source_class = GEGL_OPERATION_SOURCE_CLASS (klass);
point_render_class = GEGL_OPERATION_POINT_RENDER_CLASS (klass);
point_render_class->process = process;
source_class->process = operation_source_process;
operation_class->get_bounding_box = get_bounding_box;
operation_class->prepare = prepare;
......
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