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

error handling in gegl-buffer-cl-iterator

parent df54426d
......@@ -14,7 +14,7 @@
#include "gegl-tile-storage.h"
#include "gegl-utils.h"
#define CL_ERROR {g_printf("[OpenCL] Error in %s:%d@%s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(errcode)); g_assert(FALSE);}
#define CL_ERROR {g_printf("[OpenCL] Error in %s:%d@%s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(cl_err)); goto error;}
typedef struct GeglBufferClIterators
{
......@@ -129,11 +129,12 @@ gegl_buffer_cl_iterator_add (GeglBufferClIterator *iterator,
}
gboolean
gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
GeglBufferClIterators *i = (gpointer)iterator;
gboolean result = FALSE;
gint no;
gint no, j;
cl_int cl_err = 0;
const size_t origin_zero[3] = {0, 0, 0};
......@@ -157,17 +158,14 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
}
else
{
gint j;
cl_int errcode = 0;
/* complete pending write work */
for (no=0; no<i->iterators;no++)
{
if (i->flags[no] == GEGL_CL_BUFFER_WRITE)
{
/* Wait Processing */
errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (errcode != CL_SUCCESS) CL_ERROR;
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion in the GPU (output) */
if (i->conv[no] == CL_COLOR_CONVERT)
......@@ -179,17 +177,17 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
&i->buf_cl_format[no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &errcode);
if (errcode != CL_SUCCESS) CL_ERROR;
0, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
errcode = gegl_cl_color_conv (&i->tex[no][j], &i->tex_aux[no][j], i->size[no][j],
cl_err = gegl_cl_color_conv (&i->tex[no][j], &i->tex_aux[no][j], i->size[no][j],
i->format[no], i->buffer[no]->format);
if (errcode != CL_SUCCESS) CL_ERROR;
if (cl_err == FALSE) CL_ERROR;
}
/* Wait Processing */
errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (errcode != CL_SUCCESS) CL_ERROR;
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
/* GPU -> CPU */
for (j=0; j < i->n; j++)
......@@ -201,8 +199,8 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex[no][j], CL_TRUE,
CL_MAP_READ,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &errcode);
if (errcode != CL_SUCCESS) CL_ERROR;
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* tile-ize */
if (i->conv[no] == CL_COLOR_NOT_SUPPORTED)
......@@ -212,16 +210,16 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
/* color conversion has already been performed in the GPU */
gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->buffer[no]->format, data, pitch);
errcode = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex[no][j], data,
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex[no][j], data,
0, NULL, NULL);
if (errcode != CL_SUCCESS) CL_ERROR;
if (cl_err != CL_SUCCESS) CL_ERROR;
}
}
}
/* Run! */
errcode = gegl_clFinish(gegl_cl_get_command_queue());
if (errcode != CL_SUCCESS) CL_ERROR;
cl_err = gegl_clFinish(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
for (no=0; no < i->iterators; no++)
for (j=0; j < i->n; j++)
......@@ -242,9 +240,6 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
/* then we iterate all */
for (no=0; no<i->iterators;no++)
{
int j;
int errcode = 0;
for (j = 0; j < i->n; j++)
{
GeglRectangle r = {i->rect[no].x + i->roi_all[i->roi_no+j].x,
......@@ -273,15 +268,15 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
&i->op_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &errcode);
if (errcode != CL_SUCCESS) CL_ERROR;
0, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex[no][j], CL_TRUE,
CL_MAP_WRITE,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &errcode);
if (errcode != CL_SUCCESS) CL_ERROR;
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* un-tile */
if (i->conv[no] == CL_COLOR_NOT_SUPPORTED)
......@@ -291,13 +286,13 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
/* color conversion will be performed in the GPU later */
gegl_buffer_get (i->buffer[no], 1.0, &i->roi[no][j], i->buffer[no]->format, data, pitch);
errcode = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex[no][j], data,
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex[no][j], data,
0, NULL, NULL);
if (errcode != CL_SUCCESS) CL_ERROR;
if (cl_err != CL_SUCCESS) CL_ERROR;
}
errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (errcode != CL_SUCCESS) CL_ERROR;
cl_err = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion in the GPU (input) */
if (i->conv[no] == CL_COLOR_CONVERT)
......@@ -310,12 +305,12 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
&i->op_cl_format[no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &errcode);
if (errcode != CL_SUCCESS) CL_ERROR;
0, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
errcode = gegl_cl_color_conv (&i->tex[no][j], &i->tex_aux[no][j], i->size[no][j],
cl_err = gegl_cl_color_conv (&i->tex[no][j], &i->tex_aux[no][j], i->size[no][j],
i->buffer[no]->format, i->format[no]);
if (errcode == FALSE) CL_ERROR;
if (cl_err == FALSE) CL_ERROR;
}
}
else if (i->flags[no] == GEGL_CL_BUFFER_WRITE)
......@@ -329,8 +324,8 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
&i->op_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &errcode);
if (errcode != CL_SUCCESS) CL_ERROR;
0, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
}
else
......@@ -370,7 +365,22 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator)
g_slice_free (GeglBufferClIterators, i);
}
*err = FALSE;
return result;
error:
for (no=0; no<i->iterators;no++)
for (j=0; j < i->n; j++)
{
if (i->tex_aux[no][j]) gegl_clReleaseMemObject (i->tex_aux[no][j]);
if (i->tex[no][j]) gegl_clReleaseMemObject (i->tex[no][j]);
i->tex_aux[no][j] = NULL;
i->tex[no][j] = NULL;
}
*err = TRUE;
return FALSE;
}
GeglBufferClIterator *
......
......@@ -27,7 +27,7 @@ gint gegl_buffer_cl_iterator_add (GeglBufferClIterator *iterator,
const Babl *format,
guint flags);
gboolean gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator);
gboolean gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err);
GeglBufferClIterator *gegl_buffer_cl_iterator_new (GeglBuffer *buffer,
const GeglRectangle *roi,
......
......@@ -72,9 +72,6 @@ gegl_operation_point_filter_init (GeglOperationPointFilter *self)
{
}
//#define CL_ERROR {g_assert(0);}
#define CL_ERROR {g_printf("[OpenCL] Error in %s:%d@%s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(errcode)); goto error;}
static gboolean
gegl_operation_point_filter_cl_process (GeglOperation *operation,
GeglBuffer *input,
......@@ -87,7 +84,8 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
GeglOperationPointFilterClass *point_filter_class = GEGL_OPERATION_POINT_FILTER_GET_CLASS (operation);
gint j;
cl_int errcode = 0;
cl_int cl_err = 0;
gboolean err;
/* non-texturizable format! */
if (!gegl_cl_color_babl (in_format, NULL, NULL) ||
......@@ -109,23 +107,25 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
{
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format, GEGL_CL_BUFFER_READ);
while (gegl_buffer_cl_iterator_next (i))
for (j=0; j < i->n; j++)
{
errcode = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
i->size[0][j], &i->roi[0][j]);
if (errcode != CL_SUCCESS) CL_ERROR;
}
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
cl_err = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
i->size[0][j], &i->roi[0][j]);
if (cl_err != CL_SUCCESS)
{
g_warning("[OpenCL] Error in %s [GeglOperationPointFilter] Kernel\n",
GEGL_OPERATION_CLASS (operation)->name);
return FALSE;
}
}
}
}
return TRUE;
error:
return FALSE;
}
#undef CL_ERROR
static gboolean
gegl_operation_point_filter_process (GeglOperation *operation,
GeglBuffer *input,
......
......@@ -59,7 +59,7 @@ struct _GeglOperationPointFilterClass
checkerboard op for
semantics */
gboolean (* cl_process) (GeglOperation *self,
cl_int (* cl_process) (GeglOperation *self,
cl_mem in_tex,
cl_mem out_tex,
const size_t global_worksize[2],
......
......@@ -126,7 +126,7 @@ static const char* kernel_source =
static gegl_cl_run_data *cl_data = NULL;
/* OpenCL processing function */
static gboolean
static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
......@@ -142,7 +142,7 @@ cl_process (GeglOperation *op,
gfloat brightness = o->brightness;
gfloat contrast = o->contrast;
cl_int errcode = 0;
cl_int cl_err = 0;
if (!cl_data)
{
......@@ -152,23 +152,19 @@ cl_process (GeglOperation *op,
if (!cl_data) return 1;
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 0, sizeof(cl_mem), (void*)&in_tex));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 1, sizeof(cl_mem), (void*)&out_tex));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 2, sizeof(cl_float), (void*)&brightness));
CL_SAFE_CALL(errcode = gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&contrast));
CL_SAFE_CALL(errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_worksize, NULL,
0, NULL, NULL) );
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*)&brightness);
cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&contrast);
if (cl_err != CL_SUCCESS) return cl_err;
if (errcode != CL_SUCCESS)
{
g_warning("[OpenCL] Error in Brightness-Constrast Kernel\n");
return errcode;
}
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_worksize, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) return cl_err;
return errcode;
return cl_err;
}
/*
......
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