Commit 6914269c authored by Victor Oliveira's avatar Victor Oliveira Committed by Øyvind "pippin" Kolås

Using opencl buffers instead of Image2D

changing memory flags also.
parent 0afb9cc8
......@@ -20,7 +20,7 @@ typedef struct GeglBufferClIterators
{
/* current region of interest */
gint n;
size_t size [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX][2]; /* length of current data in pixels */
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];
......@@ -42,9 +42,9 @@ typedef struct GeglBufferClIterators
GeglBuffer *buffer [GEGL_CL_BUFFER_MAX_ITERATORS];
/* buffer->format */
cl_image_format buf_cl_format [GEGL_CL_BUFFER_MAX_ITERATORS];
size_t buf_cl_format_size [GEGL_CL_BUFFER_MAX_ITERATORS];
/* format */
cl_image_format op_cl_format [GEGL_CL_BUFFER_MAX_ITERATORS];
size_t op_cl_format_size [GEGL_CL_BUFFER_MAX_ITERATORS];
gegl_cl_color_op conv [GEGL_CL_BUFFER_MAX_ITERATORS];
......@@ -93,8 +93,8 @@ gegl_buffer_cl_iterator_add (GeglBufferClIterator *iterator,
else
i->conv[self] = gegl_cl_color_supported (buffer->format, format);
gegl_cl_color_babl (buffer->format, &i->buf_cl_format[self], NULL);
gegl_cl_color_babl (format, &i->op_cl_format [self], NULL);
gegl_cl_color_babl (buffer->format, NULL, &i->buf_cl_format_size[self]);
gegl_cl_color_babl (format, NULL, &i->op_cl_format_size [self]);
if (self!=0)
{
......@@ -137,8 +137,6 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
gint no, j;
cl_int cl_err = 0;
const size_t origin_zero[3] = {0, 0, 0};
if (i->is_finished)
g_error ("%s called on finished buffer iterator", G_STRFUNC);
if (i->iteration_no == 0)
......@@ -185,38 +183,36 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
for (j=0; j < i->n; j++)
{
gpointer data;
size_t pitch;
const size_t region[3] = {i->roi[no][j].width, i->roi[no][j].height, 1};
/* tile-ize */
if (i->conv[no] == GEGL_CL_COLOR_NOT_SUPPORTED)
{
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_READ,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &cl_err);
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_READ,
0, i->size[no][j] * i->op_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion using BABL */
gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->format[no], data, pitch);
gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->format[no], data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_op[no][j], data,
0, NULL, NULL);
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
else
{
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_READ,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &cl_err);
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_READ,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* 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);
gegl_buffer_set (i->buffer[no], &i->roi[no][j], i->buffer[no]->format, data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
0, NULL, NULL);
if (cl_err != CL_SUCCESS) CL_ERROR;
}
}
......@@ -254,9 +250,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
i->roi_all[i->roi_no+j].width,
i->roi_all[i->roi_no+j].height};
i->roi [no][j] = r;
i->size[no][j][0] = r.width;
i->size[no][j][1] = r.height;
i->size[no][j] = r.width * r.height;
}
if (i->flags[no] == GEGL_CL_BUFFER_READ)
......@@ -264,8 +258,6 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
for (j=0; j < i->n; j++)
{
gpointer data;
size_t pitch;
const size_t region[3] = {i->roi[no][j].width, i->roi[no][j].height, 1};
/* un-tile */
switch (i->conv[no])
......@@ -274,23 +266,21 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->op_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_WRITE,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &cl_err);
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_op[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->op_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* color conversion using BABL */
gegl_buffer_get (i->buffer[no], 1.0, &i->roi[no][j], i->format[no], data, pitch);
gegl_buffer_get (i->buffer[no], 1.0, &i->roi[no][j], i->format[no], data, GEGL_AUTO_ROWSTRIDE);
i->tex[no][j] = i->tex_op[no][j];
......@@ -301,23 +291,21 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->buf_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &cl_err);
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* 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);
gegl_buffer_get (i->buffer[no], 1.0, &i->roi[no][j], i->buffer[no]->format, data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
......@@ -332,32 +320,28 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->buf_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->op_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* pre-pinned memory */
data = gegl_clEnqueueMapImage(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
origin_zero, region, &pitch, NULL,
0, NULL, NULL, &cl_err);
data = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex_buf[no][j], CL_TRUE,
CL_MAP_WRITE,
0, i->size[no][j] * i->buf_cl_format_size [no],
0, NULL, NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
/* 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);
gegl_buffer_get (i->buffer[no], 1.0, &i->roi[no][j], i->buffer[no]->format, data, GEGL_AUTO_ROWSTRIDE);
cl_err = gegl_clEnqueueUnmapMemObject (gegl_cl_get_command_queue(), i->tex_buf[no][j], data,
0, NULL, NULL);
......@@ -369,7 +353,7 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
i->buffer[no]->format, i->format[no]);
if (cl_err == FALSE) CL_ERROR;
i->tex[no][j] = i->tex_buf[no][j];
i->tex[no][j] = i->tex_op[no][j];
break;
}
......@@ -390,12 +374,10 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->op_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_WRITE_ONLY,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
......@@ -407,12 +389,10 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->buf_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_WRITE_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_buf[no][j];
......@@ -424,21 +404,17 @@ gegl_buffer_cl_iterator_next (GeglBufferClIterator *iterator, gboolean *err)
{
g_assert (i->tex_buf[no][j] == NULL);
i->tex_buf[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->buf_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_buf[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_WRITE_ONLY,
i->size[no][j] * i->buf_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
g_assert (i->tex_op[no][j] == NULL);
i->tex_op[no][j] = gegl_clCreateImage2D (gegl_cl_get_context (),
CL_MEM_ALLOC_HOST_PTR | CL_MEM_READ_WRITE,
&i->op_cl_format [no],
i->roi[no][j].width,
i->roi[no][j].height,
0, NULL, &cl_err);
i->tex_op[no][j] = gegl_clCreateBuffer (gegl_cl_get_context (),
CL_MEM_READ_WRITE,
i->size[no][j] * i->op_cl_format_size [no],
NULL, &cl_err);
if (cl_err != CL_SUCCESS) CL_ERROR;
i->tex[no][j] = i->tex_op[no][j];
......
......@@ -16,7 +16,7 @@ enum
typedef struct GeglBufferClIterator
{
gint n;
size_t size [GEGL_CL_BUFFER_MAX_ITERATORS][GEGL_CL_NTEX][2]; /* length of current data in pixels */
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];
} GeglBufferClIterator;
......
......@@ -22,118 +22,115 @@ static const char* kernel_color_source =
" return value / 12.92f; \n"
"} \n"
" \n"
"__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | \n"
" CLK_ADDRESS_NONE | \n"
" CLK_FILTER_NEAREST; \n"
" \n"
"/* RGBA float -> RaGaBaA float */ \n"
"__kernel void non_premultiplied_to_premultiplied (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void non_premultiplied_to_premultiplied (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = in_v * in_v.w; \n"
" out_v.w = in_v.w; \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* RaGaBaA float -> RGBA float */ \n"
"__kernel void premultiplied_to_non_premultiplied (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void premultiplied_to_non_premultiplied (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = (in_v.w > BABL_ALPHA_THRESHOLD)? in_v / in_v.w : (float4)(0.0f); \n"
" out_v.w = in_v.w; \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* RGBA float -> R'G'B'A float */ \n"
"__kernel void rgba2rgba_gamma_2_2 (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgba2rgba_gamma_2_2 (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = (float4)(linear_to_gamma_2_2(in_v.x), \n"
" linear_to_gamma_2_2(in_v.y), \n"
" linear_to_gamma_2_2(in_v.z), \n"
" in_v.w); \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* R'G'B'A float -> RGBA float */ \n"
"__kernel void rgba_gamma_2_22rgba (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgba_gamma_2_22rgba (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = (float4)(gamma_2_2_to_linear(in_v.x), \n"
" gamma_2_2_to_linear(in_v.y), \n"
" gamma_2_2_to_linear(in_v.z), \n"
" in_v.w); \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* RGBA float -> R'aG'aB'aA float */ \n"
"__kernel void rgba2rgba_gamma_2_2_premultiplied (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgba2rgba_gamma_2_2_premultiplied (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = (float4)(linear_to_gamma_2_2(in_v.x) * in_v.w, \n"
" linear_to_gamma_2_2(in_v.y) * in_v.w, \n"
" linear_to_gamma_2_2(in_v.z) * in_v.w, \n"
" in_v.w); \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* R'aG'aB'aA float -> RGBA float */ \n"
"__kernel void rgba_gamma_2_2_premultiplied2rgba (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgba_gamma_2_2_premultiplied2rgba (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v = (in_v.w > BABL_ALPHA_THRESHOLD)? (float4)(linear_to_gamma_2_2(in_v.x) / in_v.w,\n"
" linear_to_gamma_2_2(in_v.y) / in_v.w,\n"
" linear_to_gamma_2_2(in_v.z) / in_v.w,\n"
" in_v.w) : \n"
" (float4)(0.0f); \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* RGBA float -> RGBA u8 */ \n"
"__kernel void rgbaf_to_rgbau8 (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgbaf_to_rgbau8 (__global const float4 * in, \n"
" __global uchar4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" float4 out_v = in_v; \n"
" write_imagef(out, gid, out_v); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v = in_v * 255.0f; \n"
" out[gid] = convert_uchar4_sat_rte(out_v); \n"
"} \n"
" \n"
"/* RGBAu8 -> RGBA float */ \n"
"__kernel void rgbau8_to_rgbaf (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgbau8_to_rgbaf (__global const uchar4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" float4 out_v = in_v; \n"
" write_imagef(out, gid, out_v); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = convert_float4(in[gid]); \n"
" float4 out_v = in_v / 255.0f; \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* RGBA float -> Y'CbCrA float */ \n"
" \n"
"__kernel void rgba_to_ycbcra (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void rgba_to_ycbcra (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" \n"
" float4 rgb = (float4)(linear_to_gamma_2_2(in_v.x), \n"
......@@ -145,16 +142,16 @@ static const char* kernel_color_source =
" -0.168736f * rgb.x - 0.331264f * rgb.y + 0.5f * rgb.z, \n"
" 0.5f * rgb.x - 0.418688f * rgb.y - 0.081312f * rgb.z, \n"
" in_v.w); \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n"
" \n"
"/* Y'CbCrA float -> RGBA float */ \n"
" \n"
"__kernel void ycbcra_to_rgba (__read_only image2d_t in, \n"
" __write_only image2d_t out) \n"
"__kernel void ycbcra_to_rgba (__global const float4 * in, \n"
" __global float4 * out) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" \n"
" float4 rgb = (float4)(1.0f * in_v.x + 0.0f * in_v.y + 1.40200f * in_v.z, \n"
......@@ -166,5 +163,5 @@ static const char* kernel_color_source =
" linear_to_gamma_2_2(rgb.y), \n"
" linear_to_gamma_2_2(rgb.z), \n"
" in_v.w); \n"
" write_imagef(out, gid, out_v); \n"
"} \n";
" out[gid] = out_v; \n"
"} \n";
\ No newline at end of file
......@@ -130,7 +130,7 @@ gegl_cl_color_supported (const Babl *in_format, const Babl *out_format)
#define CL_ERROR {g_printf("[OpenCL] Error in %s:%d@%s - %s\n", __FILE__, __LINE__, __func__, gegl_cl_errstring(errcode)); return FALSE;}
gboolean
gegl_cl_color_conv (cl_mem in_tex, cl_mem out_tex, const size_t size[2],
gegl_cl_color_conv (cl_mem in_tex, cl_mem out_tex, const size_t size,
const Babl *in_format, const Babl *out_format)
{
int errcode;
......@@ -140,13 +140,13 @@ gegl_cl_color_conv (cl_mem in_tex, cl_mem out_tex, const size_t size[2],
if (in_format == out_format)
{
const size_t origin[3] = {0, 0, 0};
const size_t region[3] = {size[0], size[1], 1};
size_t s;
gegl_cl_color_babl (in_format, NULL, &s);
/* just copy in_tex to out_tex */
errcode = gegl_clEnqueueCopyImage (gegl_cl_get_command_queue(),
in_tex, out_tex, origin, origin, region,
0, NULL, NULL);
errcode = gegl_clEnqueueCopyBuffer (gegl_cl_get_command_queue(),
in_tex, out_tex, 0, 0, size * s,
0, NULL, NULL);
if (errcode != CL_SUCCESS) CL_ERROR
errcode = gegl_clEnqueueBarrier(gegl_cl_get_command_queue());
......@@ -163,8 +163,8 @@ gegl_cl_color_conv (cl_mem in_tex, cl_mem out_tex, const size_t size[2],
if (errcode != CL_SUCCESS) CL_ERROR
errcode = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
kernels_color->kernel[k], 2,
NULL, size, NULL,
kernels_color->kernel[k], 1,
NULL, &size, NULL,
0, NULL, NULL);
if (errcode != CL_SUCCESS) CL_ERROR
......
......@@ -17,7 +17,7 @@ gboolean gegl_cl_color_babl (const Babl *buffer_format, cl_image_format *cl_form
gegl_cl_color_op gegl_cl_color_supported (const Babl *in_format, const Babl *out_format);
gboolean gegl_cl_color_conv (cl_mem in_tex, cl_mem aux_tex, const size_t size[2],
gboolean gegl_cl_color_conv (cl_mem in_tex, cl_mem aux_tex, const size_t size,
const Babl *in_format, const Babl *out_format);
#endif
......@@ -62,7 +62,7 @@ struct _GeglOperationPointFilterClass
cl_int (* cl_process) (GeglOperation *self,
cl_mem in_tex,
cl_mem out_tex,
const size_t global_worksize[2],
size_t global_worksize,
const GeglRectangle *roi);
};
......
......@@ -107,20 +107,17 @@ process (GeglOperation *op,
#include "opencl/gegl-cl.h"
static const char* kernel_source =
"__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | \n"
" CLK_ADDRESS_NONE | \n"
" CLK_FILTER_NEAREST; \n"
"__kernel void kernel_bc(__read_only image2d_t in, \n"
" __write_only image2d_t out, \n"
" float brightness, \n"
" float contrast) \n"
"__kernel void kernel_bc(__global const float4 *in, \n"
" __global float4 *out, \n"
" float brightness, \n"
" float contrast) \n"
"{ \n"
" int2 gid = (int2)(get_global_id(0), get_global_id(1)); \n"
" float4 in_v = read_imagef(in, sampler, gid); \n"
" int gid = get_global_id(0); \n"
" float4 in_v = in[gid]; \n"
" float4 out_v; \n"
" out_v.xyz = (in_v.xyz - 0.5f) * contrast + brightness + 0.5f;\n"
" out_v.w = in_v.w; \n"
" write_imagef(out, gid, out_v); \n"
" out[gid] = out_v; \n"
"} \n";
static gegl_cl_run_data *cl_data = NULL;
......@@ -130,7 +127,7 @@ static cl_int
cl_process (GeglOperation *op,
cl_mem in_tex,
cl_mem out_tex,
const size_t global_worksize[2],
size_t global_worksize,
const GeglRectangle *roi)
{
/* Retrieve a pointer to GeglChantO structure which contains all the
......@@ -159,8 +156,8 @@ cl_process (GeglOperation *op,
if (cl_err != CL_SUCCESS) return cl_err;
cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
cl_data->kernel[0], 2,
NULL, global_worksize, NULL,
cl_data->kernel[0], 1,
NULL, &global_worksize, NULL,
0, NULL, NULL);
if (cl_err != CL_SUCCESS) 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