[gegl] diffraction-patterns: Add CL implementation
- From: Øyvind Kolås <ok src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] diffraction-patterns: Add CL implementation
- Date: Mon, 22 May 2017 21:18:02 +0000 (UTC)
commit 35acf74b47167d8e84fd9fda4db569e367262844
Author: Anton Gorenko <anton streamcomputing eu>
Date: Tue May 16 15:46:09 2017 +0600
diffraction-patterns: Add CL implementation
opencl/diffraction-patterns.cl | 76 +++++++++++++++
opencl/diffraction-patterns.cl.h | 78 +++++++++++++++
operations/common/diffraction-patterns.c | 156 +++++++++++++++++++++++++++---
3 files changed, 298 insertions(+), 12 deletions(-)
---
diff --git a/opencl/diffraction-patterns.cl b/opencl/diffraction-patterns.cl
new file mode 100644
index 0000000..95ba5a3
--- /dev/null
+++ b/opencl/diffraction-patterns.cl
@@ -0,0 +1,76 @@
+float3
+diff_intensity (const float x,
+ const float y,
+ const float3 lam,
+ const float polarization,
+ const float scattering,
+ const int iterations,
+ const float weird_factor)
+{
+ float3 cxy = (float3)0.0f;
+ float3 sxy = (float3)0.0f;
+
+ for (int i = 0; i <= iterations; i++)
+ {
+ const float a = -M_PI_F + i * (2.0f * M_PI_F / iterations);
+ float sina, cosa;
+ sina = sincos (a, &cosa);
+
+ const float v0 = cosa;
+ const float v1 = 0.75f * sina;
+ const float v2 = 0.5f * (4.0f * cosa * cosa + sina * sina);
+ const float3 p = 4.0f * lam * (v0 * x + v1 * y - v2);
+
+ float3 sinp, cosp;
+ sinp = sincos (p, &cosp);
+ cxy += cosp;
+ sxy += sinp;
+ }
+
+ cxy *= weird_factor;
+ sxy *= weird_factor;
+
+ const float polpi2 = polarization * (M_PI_F / 2.0f);
+ float sinpolpi2, cospolpi2;
+ sinpolpi2 = sincos (polpi2, &cospolpi2);
+
+ return scattering * ((cospolpi2 + sinpolpi2) * cxy * cxy +
+ (cospolpi2 - sinpolpi2) * sxy * sxy);
+}
+
+__kernel void
+cl_diffraction_patterns (__global float *out,
+ const int offset_x,
+ const int offset_y,
+ const int width,
+ const int height,
+ const float3 sedges,
+ const float3 contours,
+ const float3 frequency,
+ const float brightness,
+ const float polarization,
+ const float scattering,
+ const int iterations,
+ const float weird_factor)
+{
+ const int gidx = get_global_id (0);
+ const int gidy = get_global_id (1);
+
+ const int x = gidx + offset_x;
+ const int y = gidy + offset_y;
+
+ const float dh = +10.0f / (width - 1);
+ const float dv = -10.0f / (height - 1);
+
+ const float px = dh * x - 5.0f;
+ const float py = dv * y + 5.0f;
+
+ const float3 di = diff_intensity (px, py, frequency, polarization, scattering,
+ iterations, weird_factor);
+ const float3 out_v = fabs (sedges * sin (contours * atan (brightness * di)));
+
+ const int gid = gidy * get_global_size (0) + gidx;
+ out[gid * 3 + 0] = out_v.x;
+ out[gid * 3 + 1] = out_v.y;
+ out[gid * 3 + 2] = out_v.z;
+}
diff --git a/opencl/diffraction-patterns.cl.h b/opencl/diffraction-patterns.cl.h
new file mode 100644
index 0000000..d1ddb9a
--- /dev/null
+++ b/opencl/diffraction-patterns.cl.h
@@ -0,0 +1,78 @@
+static const char* diffraction_patterns_cl_source =
+"float3 \n"
+"diff_intensity (const float x, \n"
+" const float y, \n"
+" const float3 lam, \n"
+" const float polarization, \n"
+" const float scattering, \n"
+" const int iterations, \n"
+" const float weird_factor) \n"
+"{ \n"
+" float3 cxy = (float3)0.0f; \n"
+" float3 sxy = (float3)0.0f; \n"
+" \n"
+" for (int i = 0; i <= iterations; i++) \n"
+" { \n"
+" const float a = -M_PI_F + i * (2.0f * M_PI_F / iterations); \n"
+" float sina, cosa; \n"
+" sina = sincos (a, &cosa); \n"
+" \n"
+" const float v0 = cosa; \n"
+" const float v1 = 0.75f * sina; \n"
+" const float v2 = 0.5f * (4.0f * cosa * cosa + sina * sina); \n"
+" const float3 p = 4.0f * lam * (v0 * x + v1 * y - v2); \n"
+" \n"
+" float3 sinp, cosp; \n"
+" sinp = sincos (p, &cosp); \n"
+" cxy += cosp; \n"
+" sxy += sinp; \n"
+" } \n"
+" \n"
+" cxy *= weird_factor; \n"
+" sxy *= weird_factor; \n"
+" \n"
+" const float polpi2 = polarization * (M_PI_F / 2.0f); \n"
+" float sinpolpi2, cospolpi2; \n"
+" sinpolpi2 = sincos (polpi2, &cospolpi2); \n"
+" \n"
+" return scattering * ((cospolpi2 + sinpolpi2) * cxy * cxy + \n"
+" (cospolpi2 - sinpolpi2) * sxy * sxy); \n"
+"} \n"
+" \n"
+"__kernel void \n"
+"cl_diffraction_patterns (__global float *out, \n"
+" const int offset_x, \n"
+" const int offset_y, \n"
+" const int width, \n"
+" const int height, \n"
+" const float3 sedges, \n"
+" const float3 contours, \n"
+" const float3 frequency, \n"
+" const float brightness, \n"
+" const float polarization, \n"
+" const float scattering, \n"
+" const int iterations, \n"
+" const float weird_factor) \n"
+"{ \n"
+" const int gidx = get_global_id (0); \n"
+" const int gidy = get_global_id (1); \n"
+" \n"
+" const int x = gidx + offset_x; \n"
+" const int y = gidy + offset_y; \n"
+" \n"
+" const float dh = +10.0f / (width - 1); \n"
+" const float dv = -10.0f / (height - 1); \n"
+" \n"
+" const float px = dh * x - 5.0f; \n"
+" const float py = dv * y + 5.0f; \n"
+" \n"
+" const float3 di = diff_intensity (px, py, frequency, polarization, scattering,\n"
+" iterations, weird_factor); \n"
+" const float3 out_v = fabs (sedges * sin (contours * atan (brightness * di)));\n"
+" \n"
+" const int gid = gidy * get_global_size (0) + gidx; \n"
+" out[gid * 3 + 0] = out_v.x; \n"
+" out[gid * 3 + 1] = out_v.y; \n"
+" out[gid * 3 + 2] = out_v.z; \n"
+"} \n"
+;
diff --git a/operations/common/diffraction-patterns.c b/operations/common/diffraction-patterns.c
index 5a523f3..351f426 100644
--- a/operations/common/diffraction-patterns.c
+++ b/operations/common/diffraction-patterns.c
@@ -98,6 +98,12 @@ property_int (height, _("Height"), 200)
#define GEGL_OP_C_SOURCE diffraction-patterns.c
#include "gegl-op.h"
+#include <gegl-buffer-cl-iterator.h>
+#include <gegl-debug.h>
+
+#include "opencl/diffraction-patterns.cl.h"
+
+static GeglClRunData *cl_data = NULL;
#define ITERATIONS 100
#define WEIRD_FACTOR 0.04
@@ -193,11 +199,90 @@ get_bounding_box (GeglOperation *operation)
}
static gboolean
-process (GeglOperation *operation,
- void *out_buf,
- glong n_pixels,
- const GeglRectangle *roi,
- gint level)
+cl_process (GeglOperation *operation,
+ cl_mem out_tex,
+ const GeglRectangle *roi)
+{
+ GeglProperties *o = GEGL_PROPERTIES (operation);
+ const size_t gbl_size[] = { roi->width, roi->height };
+ cl_int cl_err = 0;
+ cl_int offset_x;
+ cl_int offset_y;
+ cl_int width;
+ cl_int height;
+ cl_float3 sedges;
+ cl_float3 contours;
+ cl_float3 frequency;
+ cl_float brightness;
+ cl_float polarization;
+ cl_float scattering;
+ cl_int iterations;
+ cl_float weird_factor;
+
+ if (!cl_data)
+ {
+ const char *kernel_name[] = { "cl_diffraction_patterns", NULL };
+ cl_data = gegl_cl_compile_and_build (diffraction_patterns_cl_source,
+ kernel_name);
+
+ if (!cl_data)
+ return TRUE;
+ }
+
+ offset_x = roi->x;
+ offset_y = roi->y;
+ width = o->width;
+ height = o->height;
+ sedges.s[0] = o->red_sedges;
+ sedges.s[1] = o->green_sedges;
+ sedges.s[2] = o->blue_sedges;
+ contours.s[0] = o->red_contours;
+ contours.s[1] = o->green_contours;
+ contours.s[2] = o->blue_contours;
+ frequency.s[0] = o->red_frequency;
+ frequency.s[1] = o->green_frequency;
+ frequency.s[2] = o->blue_frequency;
+ brightness = o->brightness;
+ polarization = o->polarization;
+ scattering = o->scattering;
+ iterations = ITERATIONS;
+ weird_factor = WEIRD_FACTOR;
+ cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
+ sizeof(cl_mem), &out_tex,
+ sizeof(cl_int), &offset_x,
+ sizeof(cl_int), &offset_y,
+ sizeof(cl_int), &width,
+ sizeof(cl_int), &height,
+ sizeof(cl_float3), &sedges,
+ sizeof(cl_float3), &contours,
+ sizeof(cl_float3), &frequency,
+ sizeof(cl_float), &brightness,
+ sizeof(cl_float), &polarization,
+ sizeof(cl_float), &scattering,
+ sizeof(cl_int), &iterations,
+ sizeof(cl_float), &weird_factor,
+ NULL);
+ CL_CHECK;
+
+ cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
+ cl_data->kernel[0], 2,
+ NULL, gbl_size, NULL,
+ 0, NULL, NULL);
+ CL_CHECK;
+
+ cl_err = gegl_clFinish (gegl_cl_get_command_queue ());
+ CL_CHECK;
+
+ return FALSE;
+
+error:
+ return TRUE;
+}
+
+static gboolean
+c_process (GeglOperation *operation,
+ void *out_buf,
+ const GeglRectangle *roi)
{
GeglProperties *o = GEGL_PROPERTIES (operation);
gfloat *out;
@@ -230,22 +315,69 @@ process (GeglOperation *operation,
return TRUE;
}
+static gboolean
+process (GeglOperation *operation,
+ GeglBuffer *out_buf,
+ const GeglRectangle *roi,
+ gint level)
+{
+ GeglBufferIterator *iter;
+ const Babl *out_format = gegl_operation_get_format (operation,
+ "output");
+
+ if (gegl_operation_use_opencl (operation))
+ {
+ 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 (out_buf, roi, out_format,
+ GEGL_CL_BUFFER_WRITE);
+
+ while (gegl_buffer_cl_iterator_next (cl_iter, &err) && !err)
+ {
+ err = cl_process (operation, cl_iter->tex[0], cl_iter->roi);
+
+ 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 (out_buf, roi, level, out_format,
+ GEGL_ACCESS_WRITE, GEGL_ABYSS_NONE);
+
+ while (gegl_buffer_iterator_next (iter))
+ c_process (operation, iter->data[0], &iter->roi[0]);
+
+ return TRUE;
+}
+
static void
gegl_op_class_init (GeglOpClass *klass)
{
- GeglOperationClass *operation_class;
- GeglOperationPointRenderClass *point_render_class;
+ GeglOperationClass *operation_class;
+ GeglOperationSourceClass *source_class;
init_luts();
- operation_class = GEGL_OPERATION_CLASS (klass);
- point_render_class = GEGL_OPERATION_POINT_RENDER_CLASS (klass);
-
- point_render_class->process = process;
+ operation_class = GEGL_OPERATION_CLASS (klass);
+ source_class = GEGL_OPERATION_SOURCE_CLASS (klass);
+ source_class->process = process;
operation_class->get_bounding_box = get_bounding_box;
operation_class->prepare = prepare;
- operation_class->opencl_support = FALSE;
+ operation_class->opencl_support = TRUE;
gegl_operation_class_set_keys (operation_class,
"name", "gegl:diffraction-patterns",
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
