Commit 33bb648b authored by Victor Oliveira's avatar Victor Oliveira

Simplying opencl buffer iterators

opencl buffer iterators now iterate over just
one region at a time, instead of possibly many.

This change is because the overhead of many
clFinish calls is not that great and it was
already happening in many places because of the
gpu caching code.
parent e4fd7c63
This diff is collapsed.
......@@ -22,7 +22,6 @@
#include "gegl-buffer.h"
#include "opencl/gegl-cl.h"
#define GEGL_CL_NTEX 16
#define GEGL_CL_BUFFER_MAX_ITERATORS 6
enum
......@@ -34,10 +33,9 @@ enum
typedef struct GeglBufferClIterator
{
gint n;
size_t size [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX]; /* length of current data in pixels */
cl_mem tex [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
GeglRectangle roi [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX];
size_t size [GEGL_CL_BUFFER_MAX_ITERATORS]; /* length of current data in pixels */
cl_mem tex [GEGL_CL_BUFFER_MAX_ITERATORS];
GeglRectangle roi [GEGL_CL_BUFFER_MAX_ITERATORS];
} GeglBufferClIterator;
gint gegl_buffer_cl_iterator_add (GeglBufferClIterator *iterator,
......
......@@ -51,20 +51,13 @@
{ cl_err = gegl_clReleaseMemObject(obj); \
CL_CHECK; }
#define GEGL_CL_BUFFER_ITERATE_START(I, J, ERR) \
while (gegl_buffer_cl_iterator_next (I, & ERR)) \
#define GEGL_CL_BUFFER_ITERATE_START(I, ERR) \
while (gegl_buffer_cl_iterator_next (I, & ERR)) \
{ \
if (ERR) return FALSE; \
for (J=0; J < I ->n; J++) \
{
#define GEGL_CL_BUFFER_ITERATE_END(ERR) \
if (ERR) \
{ \
g_warning("[OpenCL] Error"); \
return FALSE; \
} \
} \
if (ERR) return FALSE; \
}
......
......@@ -172,19 +172,20 @@ gegl_operation_point_composer_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, GEGL_ABYSS_NONE);
if (aux)
foo = gegl_buffer_cl_iterator_add (i, aux, result, aux_format, GEGL_CL_BUFFER_READ, GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
if (point_composer_class->cl_process)
{
err = point_composer_class->cl_process(operation, i->tex[read][j],
(aux)? i->tex[foo][j] : NULL,
i->tex[0][j], i->size[0][j], &i->roi[0][j], level);
err = point_composer_class->cl_process(operation, i->tex[read],
(aux)? i->tex[foo] : NULL,
i->tex[0], i->size[0], &i->roi[0], level);
if (err)
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
......@@ -196,16 +197,16 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
gint p = 0;
GeglClRunData *cl_data = operation_class->cl_data;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read][j]);
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read]);
CL_CHECK;
if (aux)
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[foo][j]);
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[foo]);
else
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), NULL);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[0][j]);
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[0]);
CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
......@@ -213,7 +214,7 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &i->size[0][j], NULL,
NULL, &i->size[0], NULL,
0, NULL, NULL);
CL_CHECK;
}
......@@ -229,7 +230,7 @@ gegl_operation_point_composer_cl_process (GeglOperation *operation,
return TRUE;
error:
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring(cl_err));
return FALSE;
}
......
......@@ -106,18 +106,20 @@ 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, GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
if (point_filter_class->cl_process)
{
err = point_filter_class->cl_process(operation, i->tex[read][j], i->tex[0][j],
i->size[0][j], &i->roi[0][j], level);
err = point_filter_class->cl_process(operation, i->tex[read], i->tex[0],
i->size[0], &i->roi[0], level);
if (err)
{
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", operation_class->name);
gegl_cl_disable();
return FALSE;
}
}
......@@ -126,9 +128,9 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
gint p = 0;
GeglClRunData *cl_data = operation_class->cl_data;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read][j]);
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[read]);
CL_CHECK;
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[ 0 ][j]);
cl_err = gegl_clSetKernelArg(cl_data->kernel[0], p++, sizeof(cl_mem), (void*)&i->tex[ 0 ]);
CL_CHECK;
gegl_operation_cl_set_kernel_args (operation, cl_data->kernel[0], &p, &cl_err);
......@@ -136,7 +138,7 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 1,
NULL, &i->size[0][j], NULL,
NULL, &i->size[0], NULL,
0, NULL, NULL);
CL_CHECK;
}
......@@ -152,7 +154,8 @@ gegl_operation_point_filter_cl_process (GeglOperation *operation,
return TRUE;
error:
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error in GeglOperationPointComposer Kernel: %s", gegl_cl_errstring(cl_err));
GEGL_NOTE (GEGL_DEBUG_OPENCL, "Error: %s", gegl_cl_errstring(cl_err));
gegl_cl_disable();
return FALSE;
}
......
#define LAPLACE_RADIUS 1
void minmax(float x1, float x2, float x3,
float x4, float x5,
float *min_result,
float *max_result)
{
float min1, min2, max1, max2;
if (x1 > x2)
{
max1 = x1;
min1 = x2;
}
else
{
max1 = x2;
min1 = x1;
}
if (x3 > x4)
{
max2 = x3;
min2 = x4;
}
else
{
max2 = x4;
min2 = x3;
}
if (min1 < min2)
*min_result = fmin(min1, x5);
else
*min_result = fmin(min2, x5);
if (max1 > max2)
*max_result = fmax(max1, x5);
else
*max_result = fmax(max2, x5);
}
kernel void pre_edgelaplace (global float4 *in,
global float4 *out)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
int src_width = get_global_size(0) + LAPLACE_RADIUS * 2;
int src_height = get_global_size(1);
int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;
int gid1d = i + j * src_width;
float pix_fl[4] = {
in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y,
in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w
};
float pix_fm[4] = {
in[gid1d - src_width].x, in[gid1d - src_width].y,
in[gid1d - src_width].z, in[gid1d - src_width].w
};
float pix_fr[4] = {
in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y,
in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w
};
float pix_ml[4] = {
in[gid1d - 1 ].x, in[gid1d - 1 ].y,
in[gid1d - 1 ].z, in[gid1d - 1 ].w
};
float pix_mm[4] = {
in[gid1d ].x, in[gid1d ].y,
in[gid1d ].z, in[gid1d ].w
};
float pix_mr[4] = {
in[gid1d + 1 ].x, in[gid1d + 1 ].y,
in[gid1d + 1 ].z, in[gid1d + 1 ].w
};
float pix_bl[4] = {
in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y,
in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w
};
float pix_bm[4] = {
in[gid1d + src_width].x, in[gid1d + src_width].y,
in[gid1d + src_width].z, in[gid1d + src_width].w
};
float pix_br[4] = {
in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y,
in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w
};
int c;
float minval, maxval;
float gradient[4];
for (c = 0;c < 3; ++c)
{
minmax(pix_fm[c], pix_bm[c], pix_ml[c], pix_mr[c],
pix_mm[c], &minval, &maxval);
gradient[c] = 0.5f *
fmax((maxval - pix_mm[c]),(pix_mm[c] - minval));
gradient[c] =
(pix_fl[c] + pix_fm[c] + pix_fr[c] +
pix_ml[c] + pix_mr[c] + pix_bl[c] +
pix_bm[c] + pix_br[c] - 8.0f * pix_mm[c]) >
0.0f ? gradient[c] : -1.0f * gradient[c];
}
gradient[3] = pix_mm[3];
out[gid1d] = (float4)
(gradient[0], gradient[1], gradient[2], gradient[3]);
}
kernel void knl_edgelaplace (global float4 *in,
global float4 *out)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
int src_width = get_global_size(0) + LAPLACE_RADIUS * 2;
int src_height = get_global_size(1);
int i = gidx + LAPLACE_RADIUS, j = gidy + LAPLACE_RADIUS;
int gid1d = i + j * src_width;
float pix_fl[4] = {
in[gid1d - 1 - src_width].x, in[gid1d - 1 - src_width].y,
in[gid1d - 1 - src_width].z, in[gid1d - 1 - src_width].w
};
float pix_fm[4] = {
in[gid1d - src_width].x, in[gid1d - src_width].y,
in[gid1d - src_width].z, in[gid1d - src_width].w
};
float pix_fr[4] = {
in[gid1d + 1 - src_width].x, in[gid1d + 1 - src_width].y,
in[gid1d + 1 - src_width].z, in[gid1d + 1 - src_width].w
};
float pix_ml[4] = {
in[gid1d - 1 ].x, in[gid1d - 1 ].y,
in[gid1d - 1 ].z, in[gid1d - 1 ].w
};
float pix_mm[4] = {
in[gid1d ].x, in[gid1d ].y,
in[gid1d ].z, in[gid1d ].w
};
float pix_mr[4] = {
in[gid1d + 1 ].x, in[gid1d + 1 ].y,
in[gid1d + 1 ].z, in[gid1d + 1 ].w
};
float pix_bl[4] = {
in[gid1d - 1 + src_width].x, in[gid1d - 1 + src_width].y,
in[gid1d - 1 + src_width].z, in[gid1d - 1 + src_width].w
};
float pix_bm[4] = {
in[gid1d + src_width].x, in[gid1d + src_width].y,
in[gid1d + src_width].z, in[gid1d + src_width].w
};
float pix_br[4] = {
in[gid1d + 1 + src_width].x, in[gid1d + 1 + src_width].y,
in[gid1d + 1 + src_width].z, in[gid1d + 1 + src_width].w
};
int c;
float value[4];
for (c = 0;c < 3; ++c)
{
float current = pix_mm[c];
current =
((current > 0.0f) &&
(pix_fl[c] < 0.0f || pix_fm[c] < 0.0f ||
pix_fr[c] < 0.0f || pix_ml[c] < 0.0f ||
pix_mr[c] < 0.0f || pix_bl[c] < 0.0f ||
pix_bm[c] < 0.0f || pix_br[c] < 0.0f )
) ? current : 0.0f;
value[c] = current;
}
value[3] = pix_mm[3];
out[gidx + gidy * get_global_size(0)] = (float4)
(value[0], value[1], value[2], value[3]);
}
#define SOBEL_RADIUS 1
kernel void kernel_edgesobel(global float4 *in,
global float4 *out,
const int horizontal,
const int vertical,
const int keep_signal,
const int has_alpha)
{
int gidx = get_global_id(0);
int gidy = get_global_id(1);
float4 hor_grad = 0.0f;
float4 ver_grad = 0.0f;
float4 gradient = 0.0f;
int dst_width = get_global_size(0);
int src_width = dst_width + SOBEL_RADIUS * 2;
int i = gidx + SOBEL_RADIUS, j = gidy + SOBEL_RADIUS;
int gid1d = i + j * src_width;
float4 pix_fl = in[gid1d - 1 - src_width];
float4 pix_fm = in[gid1d - src_width];
float4 pix_fr = in[gid1d + 1 - src_width];
float4 pix_ml = in[gid1d - 1 ];
float4 pix_mm = in[gid1d ];
float4 pix_mr = in[gid1d + 1 ];
float4 pix_bl = in[gid1d - 1 + src_width];
float4 pix_bm = in[gid1d + src_width];
float4 pix_br = in[gid1d + 1 + src_width];
if (horizontal)
{
hor_grad +=
- 1.0f * pix_fl + 1.0f * pix_fr
- 2.0f * pix_ml + 2.0f * pix_mr
- 1.0f * pix_bl + 1.0f * pix_br;
}
if (vertical)
{
ver_grad +=
- 1.0f * pix_fl - 2.0f * pix_fm
- 1.0f * pix_fr + 1.0f * pix_bl
+ 2.0f * pix_bm + 1.0f * pix_br;
}
if (horizontal && vertical)
{
gradient = sqrt(
hor_grad * hor_grad +
ver_grad * ver_grad) / 1.41f;
}
else
{
if (keep_signal)
gradient = hor_grad + ver_grad;
else
gradient = fabs(hor_grad + ver_grad);
}
if (has_alpha)
{
gradient.w = pix_mm.w;
}
else
{
gradient.w = 1.0f;
}
out[gidx + gidy * dst_width] = gradient;
}
......@@ -430,18 +430,25 @@ bilateral_cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
gint j;
cl_int cl_err;
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, GEGL_ABYSS_NONE);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
GEGL_CL_BUFFER_ITERATE_START(i, j, err)
gint read = gegl_buffer_cl_iterator_add (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
GEGL_ABYSS_NONE);
GEGL_CL_BUFFER_ITERATE_START(i, err)
{
err = cl_bilateral(i->tex[read][j],
i->tex[0][j],
&i->roi[0][j],
&i->roi[read][j],
err = cl_bilateral(i->tex[read],
i->tex[0],
&i->roi[0],
&i->roi[read],
s_sigma,
r_sigma);
}
......
......@@ -76,10 +76,10 @@ cl_bilateral_filter (cl_mem in_tex,
size_t global_ws[2];
if (!cl_data)
{
const char *kernel_name[] = {"bilateral_filter", NULL};
cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name);
}
{
const char *kernel_name[] = {"bilateral_filter", NULL};
cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name);
}
if (!cl_data) return TRUE;
global_ws[0] = roi->width;
......@@ -115,26 +115,38 @@ cl_process (GeglOperation *operation,
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
gint j;
cl_int cl_err;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
op_area->left,
op_area->right,
op_area->top,
op_area->bottom,
GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
err = cl_bilateral_filter(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], ceil(o->blur_radius), o->edge_preservation);
if (err)
{
g_warning("[OpenCL] Error in gegl:bilateral-filter");
return FALSE;
}
}
err = cl_bilateral_filter(i->tex[read],
i->tex[0],
i->size[0],
&i->roi[0],
ceil(o->blur_radius),
o->edge_preservation);
if (err) return FALSE;
}
return TRUE;
......
......@@ -256,29 +256,53 @@ cl_process (GeglOperation *operation,
{
const Babl *in_format = gegl_operation_get_format (operation, "input");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
gint j;
cl_int cl_err;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ, op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
gint aux = gegl_buffer_cl_iterator_add_2 (i, NULL, result, in_format, GEGL_CL_BUFFER_AUX, 0, 0, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
op_area->left,
op_area->right,
op_area->top,
op_area->bottom,
GEGL_ABYSS_NONE);
gint aux = gegl_buffer_cl_iterator_add_2 (i,
NULL,
result,
in_format,
GEGL_CL_BUFFER_AUX,
0,
0,
op_area->top,
op_area->bottom,
GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
err = cl_box_blur(i->tex[read][j], i->tex[aux][j], i->tex[0][j], i->size[0][j], &i->roi[0][j], ceil (o->radius));
if (err)
{
g_warning("[OpenCL] Error in gegl:box-blur");
return FALSE;
}
}
err = cl_box_blur(i->tex[read],
i->tex[aux],
i->tex[0],
i->size[0],
&i->roi[0],
ceil (o->radius));
if (err) return FALSE;
}
return TRUE;
}
......@@ -297,8 +321,12 @@ process (GeglOperation *operation,
op_area = GEGL_OPERATION_AREA_FILTER (operation);
if (gegl_cl_is_accelerated ())
if (cl_process (operation, input, output, result))
return TRUE;
{
if (cl_process (operation, input, output, result))
return TRUE;
else
gegl_cl_disable();
}
rect = *result;
tmprect = *result;
......
......@@ -167,14 +167,14 @@ static GeglClRunData *cl_data = NULL;
static gboolean
cl_c2g (cl_mem in_tex,
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *src_roi,
const GeglRectangle *roi,
gint radius,
gint samples,
gint iterations,
gdouble rgamma)
cl_mem out_tex,
size_t global_worksize,
const GeglRectangle *src_roi,
const GeglRectangle *roi,
gint radius,
gint samples,
gint iterations,
gdouble rgamma)
{
cl_int cl_err = 0;
cl_mem cl_lut_cos, cl_lut_sin, cl_radiuses;
......@@ -254,36 +254,51 @@ error:
}
static gboolean
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result)
cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result)
{
const Babl *in_format = babl_format("RGBA float");
const Babl *out_format = gegl_operation_get_format (operation, "output");
gint err;
cl_int cl_err;
gint j;
GeglOperationAreaFilter *op_area = GEGL_OPERATION_AREA_FILTER (operation);
GeglChantO *o = GEGL_CHANT_PROPERTIES (operation);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,result, out_format, GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i, input, result, in_format, GEGL_CL_BUFFER_READ,
op_area->left, op_area->right, op_area->top, op_area->bottom, GEGL_ABYSS_NONE);
GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output,
result,
out_format,
GEGL_CL_BUFFER_WRITE);
gint read = gegl_buffer_cl_iterator_add_2 (i,
input,
result,
in_format,
GEGL_CL_BUFFER_READ,
op_area->left,
op_area->right,
op_area->top,
op_area->bottom,
GEGL_ABYSS_NONE);
while (gegl_buffer_cl_iterator_next (i, &err))
{
if (err) return FALSE;
for (j=0; j < i->n; j++)
{
err = cl_c2g(i->tex[read][j], i->tex[0][j],i->size[0][j], &i->roi[read][j], &i->roi[0][j], o->radius, o->samples, o->iterations, RGAMMA);
if (err)
{
g_warning("[OpenCL] Error in gegl:c2g");
return FALSE;
}
}
err = cl_c2g(i->tex[read],
i->tex[0],
i->size[0],
&i->roi[read],
&i->roi[0],
o->radius,
o->samples,
o->iterations,
RGAMMA);
if (err) return FALSE;
}
return TRUE;
}
......
......@@ -239,191 +239,11 @@ edge_laplace (GeglBuffer *src,
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"