New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

OpenCL implementations of noise-hsv, gaussian-blur-selective, motion-blur-circular, diffraction-patterns #72

Merged
merged 4 commits into from May 17, 2017
Jump to file or symbol
Failed to load files and symbols.
+298 −12
Diff settings

Always

Just for now

Viewing a subset of changes. View all
Prev

diffraction-patterns: Add CL implementation

  • Loading branch information...
ex-rzr committed May 16, 2017
commit 855cc073091762cabd67ba513ca7cc1b3b13dfe1
@@ -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;
}
@@ -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"
;
@@ -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",