[gegl/opencl-ops: 10/14] Add OpenCL support for reinhard05
- From: Ãyvind KolÃs <ok src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl/opencl-ops: 10/14] Add OpenCL support for reinhard05
- Date: Tue, 20 Mar 2012 13:50:37 +0000 (UTC)
commit 33e15be9b54915b5584f1fa9e597b83254e34934
Author: Victor Oliveira <victormatheus gmail com>
Date: Wed Mar 7 13:42:10 2012 +0800
Add OpenCL support for reinhard05
operations/common/reinhard05.c | 268 +++++++++++++++++++++++++++++++++++-----
1 files changed, 236 insertions(+), 32 deletions(-)
---
diff --git a/operations/common/reinhard05.c b/operations/common/reinhard05.c
index 7b6931e..0d2bb83 100644
--- a/operations/common/reinhard05.c
+++ b/operations/common/reinhard05.c
@@ -115,6 +115,126 @@ reinhard05_stats_finish (stats *s)
s->range = s->max - s->min;
}
+#include "opencl/gegl-cl.h"
+#include "buffer/gegl-buffer-cl-iterator.h"
+
+static const char* kernel_source =
+"__kernel void reinhard05_1 (__global const float4 * pix, \n"
+" __global float4 * pix_out, \n"
+" __global const float * lum, \n"
+" float chrom, \n"
+" float light, \n"
+" float intensity, \n"
+" float contrast, \n"
+" float4 channel_avg, \n"
+" float world_lin_avg) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 pix_v = pix[gid]; \n"
+" float lum_v = lum[gid]; \n"
+" float3 local_; \n"
+" float3 global_; \n"
+" float3 adapt; \n"
+" \n"
+" if (lum_v == 0.0f) return; \n"
+" \n"
+" local_ = chrom * pix_v.xyz + (1.0f - chrom) * lum_v; \n"
+" global_ = chrom * channel_avg.xyz + (1.0f - chrom) * world_lin_avg; \n"
+" adapt = light * local_ + (1.0f - light) * global_; \n"
+" pix_v.xyz /= pix_v.xyz + pow (intensity * adapt, contrast); \n"
+" \n"
+" pix_out[gid] = pix_v; \n"
+"} \n"
+" \n"
+"__kernel void reinhard05_2 (__global float4 * src, \n"
+" __global float4 * dst, \n"
+" float min, \n"
+" float range) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" dst[gid] = (src[gid]-min) / range; \n"
+"} \n";
+
+static gegl_cl_run_data * cl_data = NULL;
+
+static cl_int
+cl_reinhard05_1 (cl_mem in_tex,
+ cl_mem lum_tex,
+ cl_mem out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi,
+ gfloat chrom,
+ gfloat light,
+ gfloat intensity,
+ gfloat contrast,
+ stats world_lin,
+ stats channel[])
+{
+ cl_int cl_err = 0;
+ cl_float4 channel_avg = {channel[0].avg, channel[1].avg, channel[2].avg, 1.0f};
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"reinhard05_1", "reinhard05_2", NULL};
+ cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ }
+
+ if (!cl_data) return 1;
+
+ 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_mem), (void*)&lum_tex);
+
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 3, sizeof(cl_float), (void*)&chrom);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 4, sizeof(cl_float), (void*)&light);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 5, sizeof(cl_float), (void*)&intensity);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 6, sizeof(cl_float), (void*)&contrast);
+
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 7, sizeof(cl_float4), (void*)&channel_avg);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[0], 8, sizeof(cl_float), (void*)&world_lin.avg);
+ if (cl_err != CL_SUCCESS) return cl_err;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 1,
+ NULL, &global_worksize, NULL,
+ 0, NULL, NULL);
+ if (cl_err != CL_SUCCESS) return cl_err;
+ return cl_err;
+}
+
+static cl_int
+cl_reinhard05_2 (cl_mem in_tex,
+ cl_mem out_tex,
+ size_t global_worksize,
+ const GeglRectangle *roi,
+ gfloat min,
+ gfloat range)
+{
+ cl_int cl_err = 0;
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = {"reinhard05_1", "reinhard05_2", NULL};
+ cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ }
+
+ if (!cl_data) return 1;
+
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 0, sizeof(cl_mem), (void*)&in_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 1, sizeof(cl_mem), (void*)&out_tex);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 2, sizeof(cl_float), (void*)&min);
+ cl_err |= gegl_clSetKernelArg(cl_data->kernel[1], 3, sizeof(cl_float), (void*)&range);
+ if (cl_err != CL_SUCCESS) return cl_err;
+
+ cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue (),
+ cl_data->kernel[1], 1,
+ NULL, &global_worksize, NULL,
+ 0, NULL, NULL);
+ if (cl_err != CL_SUCCESS) return cl_err;
+ return cl_err;
+}
+
+
static gboolean
reinhard05_process (GeglOperation *operation,
@@ -201,47 +321,130 @@ reinhard05_process (GeglOperation *operation,
g_return_val_if_fail (contrast >= 0.3 && contrast <= 1.0, FALSE);
- /* Apply the operator */
- for (i = 0; i < result->width * result->height; ++i)
+ if (cl_state.is_accelerated)
{
- gfloat local, global, adapt;
+ const Babl *in_format = gegl_operation_get_format (operation, "input");
+ const Babl *out_format = gegl_operation_get_format (operation, "output");
+ const Babl *lum_format = babl_format("Y float");
- if (lum[i] == 0.0)
- continue;
+ GeglBuffer *pix_out = gegl_buffer_new (result, in_format);
+ gint j, k;
+ cl_int err, cl_err;
- for (c = 0; c < RGB; ++c)
+ {
+ gfloat *pix_map;
+
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (pix_out, result, in_format, GEGL_CL_BUFFER_WRITE);
+ gint read = gegl_buffer_cl_iterator_add (i, input, result, in_format, GEGL_CL_BUFFER_READ);
+ gint lum_ = gegl_buffer_cl_iterator_add (i, input, result, lum_format, GEGL_CL_BUFFER_READ);
+
+ while (gegl_buffer_cl_iterator_next (i, &err))
{
- gfloat *_p = pix + i * pix_stride + c,
- p = *_p;
-
- local = chrom * p +
- chrom_comp * lum[i];
- global = chrom * channel[c].avg +
- chrom_comp * world_lin.avg;
- adapt = light * local +
- light_comp * global;
-
- p /= p + powf (intensity * adapt, contrast);
- *_p = p;
- reinhard05_stats_update (&normalise, p);
+ if (err) return FALSE;
+ for (j=0; j < i->n; j++)
+ {
+
+ cl_err = cl_reinhard05_1(i->tex[read][j], i->tex[lum_][j], i->tex[0][j], i->size[0][j], &i->roi[0][j],
+ chrom, light, intensity, contrast,
+ world_lin, channel);
+ if (cl_err != CL_SUCCESS)
+ {
+ g_warning("[OpenCL] Error in gegl:reinhard05: %s\n", gegl_cl_errstring(cl_err));
+ return FALSE;
+ }
+
+ pix_map = gegl_clEnqueueMapBuffer(gegl_cl_get_command_queue(), i->tex[0][j], CL_TRUE, CL_MAP_READ, 0, i->size[0][j] * babl_format_get_bytes_per_pixel (in_format),
+ 0, NULL, NULL, &cl_err);
+ if (CL_SUCCESS != cl_err) return cl_err;
+
+ for (k = 0; k < i->size[0][j]; ++k)
+ {
+ if(lum[k] == 0.0f)
+ continue;
+ for(c=0; c < RGB; c++)
+ reinhard05_stats_update (&normalise, pix_map[k*pix_stride + c]);
+ }
+
+ cl_err = gegl_clEnqueueUnmapMemObject(gegl_cl_get_command_queue(), i->tex[0][j], pix_map,
+ 0, NULL, NULL);
+ if (CL_SUCCESS != cl_err) return cl_err;
+
+ }
}
- }
+ }
+
+ /* Normalise the pixel values */
+ reinhard05_stats_finish (&normalise);
+ {
+ GeglBufferClIterator *i = gegl_buffer_cl_iterator_new (output, result, out_format, GEGL_CL_BUFFER_WRITE);
+ gint read = gegl_buffer_cl_iterator_add (i, pix_out, result, in_format, GEGL_CL_BUFFER_READ);
+ while (gegl_buffer_cl_iterator_next (i, &err))
+ {
+ if (err) return FALSE;
+ for (j=0; j < i->n; j++)
+ {
- /* Normalise the pixel values */
- reinhard05_stats_finish (&normalise);
+ cl_err = cl_reinhard05_2(i->tex[read][j], i->tex[0][j], i->size[0][j], &i->roi[0][j],
+ normalise.min, normalise.range);
+ if (cl_err != CL_SUCCESS)
+ {
+ g_warning("[OpenCL] Error in gegl:reinhard05: %s\n", gegl_cl_errstring(cl_err));
+ return FALSE;
- for (i = 0; i < result->width * result->height; ++i)
+ }
+
+ }
+ }
+ }
+
+ gegl_buffer_destroy (pix_out);
+ }
+ else
{
- for (c = 0; c < pix_stride; ++c)
+
+ /* Apply the operator */
+ for (i = 0; i < result->width * result->height; ++i)
+ {
+ gfloat local, global, adapt;
+
+ if (lum[i] == 0.0)
+ continue;
+
+ for (c = 0; c < RGB; ++c)
+ {
+ gfloat *_p = pix + i * pix_stride + c,
+ p = *_p;
+
+ local = chrom * p +
+ chrom_comp * lum[i];
+ global = chrom * channel[c].avg +
+ chrom_comp * world_lin.avg;
+ adapt = light * local +
+ light_comp * global;
+
+ p /= p + powf (intensity * adapt, contrast);
+ *_p = p;
+ reinhard05_stats_update (&normalise, p);
+ }
+ }
+
+ /* Normalise the pixel values */
+ reinhard05_stats_finish (&normalise);
+
+ for (i = 0; i < result->width * result->height; ++i)
{
- gfloat *p = pix + i * pix_stride + c;
- *p = (*p - normalise.min) / normalise.range;
+ for (c = 0; c < pix_stride; ++c)
+ {
+ gfloat *p = pix + i * pix_stride + c;
+ *p = (*p - normalise.min) / normalise.range;
+ }
}
+
+ /* Cleanup and set the output */
+ gegl_buffer_set (output, result, babl_format (OUTPUT_FORMAT), pix,
+ GEGL_AUTO_ROWSTRIDE);
}
- /* Cleanup and set the output */
- gegl_buffer_set (output, result, babl_format (OUTPUT_FORMAT), pix,
- GEGL_AUTO_ROWSTRIDE);
g_free (pix);
g_free (lum);
@@ -249,8 +452,7 @@ reinhard05_process (GeglOperation *operation,
}
-/*
- */
+/**/
static void
gegl_chant_class_init (GeglChantClass *klass)
{
@@ -261,6 +463,7 @@ gegl_chant_class_init (GeglChantClass *klass)
filter_class = GEGL_OPERATION_FILTER_CLASS (klass);
filter_class->process = reinhard05_process;
+ operation_class->opencl_support = TRUE;
operation_class->prepare = reinhard05_prepare;
operation_class->get_required_for_output = reinhard05_get_required_for_output;
@@ -270,10 +473,11 @@ gegl_chant_class_init (GeglChantClass *klass)
operation_class->categories = "tonemapping";
operation_class->description =
_("Adapt an image, which may have a high dynamic range, for "
- "presentation using a low dynamic range. This is an efficient "
+ "presentation using a low dynamic range. This is an efficient "
"global operator derived from simple physiological observations, "
"producing luminance within the range 0.0-1.0");
}
#endif
+
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
