commit 0e708f3252e6a1849570ded42ab96854306f300b
Author: Victor Oliveira <victormatheus gmail com>
Date: Fri Jan 11 19:56:54 2013 -0200
opencl: more filters updated and forgotten Makefile
opencl/Makefile.am | 1 +
opencl/bilateral-filter.cl | 47 +++++++++
opencl/bilateral-filter.cl.h | 49 +++++++++
opencl/box-blur.cl | 45 ++++++++
opencl/box-blur.cl.h | 47 +++++++++
opencl/c2g.cl | 157 ++++++++++++++++++++++++++++
opencl/c2g.cl.h | 159 +++++++++++++++++++++++++++++
opencl/cltostring.pl | 3 +-
opencl/color-temperature.cl | 12 ++
opencl/color-temperature.cl.h | 14 +++
opencl/gaussian-blur.cl | 52 ++++++++++
opencl/gaussian-blur.cl.h | 54 ++++++++++
operations/common/bilateral-filter.c | 51 +---------
operations/common/box-blur.c | 49 +---------
operations/common/brightness-contrast.c | 2 +-
operations/common/c2g.c | 169 +------------------------------
operations/common/color-temperature.c | 19 +---
operations/common/gaussian-blur.c | 58 +----------
operations/common/invert.c | 14 +--
operations/common/mono-mixer.c | 15 +---
operations/common/over.c | 16 +---
operations/common/threshold.c | 17 +---
operations/common/value-invert.c | 33 +------
23 files changed, 664 insertions(+), 419 deletions(-)
---
diff --git a/opencl/Makefile.am b/opencl/Makefile.am
new file mode 100644
index 0000000..c27f635
--- /dev/null
+++ b/opencl/Makefile.am
@@ -0,0 +1 @@
+EXTRA_DIST = $(wildcard $(srcdir)/*.cl) + $(wildcard $(srcdir)/*.pl)
diff --git a/opencl/bilateral-filter.cl b/opencl/bilateral-filter.cl
new file mode 100644
index 0000000..0205b98
--- /dev/null
+++ b/opencl/bilateral-filter.cl
@@ -0,0 +1,47 @@
+#define POW2(a) ((a) * (a))
+kernel void bilateral_filter(global float4 *in,
+ global float4 *out,
+ const float radius,
+ const float preserve)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+ int n_radius = ceil(radius);
+ int dst_width = get_global_size(0);
+ int src_width = dst_width + n_radius * 2;
+
+ int u, v, i, j;
+ float4 center_pix =
+ in[(gidy + n_radius) * src_width + gidx + n_radius];
+ float4 accumulated = 0.0f;
+ float4 tempf = 0.0f;
+ float count = 0.0f;
+ float diff_map, gaussian_weight, weight;
+
+ for (v = -n_radius;v <= n_radius; ++v)
+ {
+ for (u = -n_radius;u <= n_radius; ++u)
+ {
+ i = gidx + n_radius + u;
+ j = gidy + n_radius + v;
+
+ int gid1d = i + j * src_width;
+ tempf = in[gid1d];
+
+ diff_map = exp (
+ - ( POW2(center_pix.x - tempf.x)
+ + POW2(center_pix.y - tempf.y)
+ + POW2(center_pix.z - tempf.z))
+ * preserve);
+
+ gaussian_weight =
+ exp( - 0.5f * (POW2(u) + POW2(v)) / radius);
+
+ weight = diff_map * gaussian_weight;
+
+ accumulated += tempf * weight;
+ count += weight;
+ }
+ }
+ out[gidx + gidy * dst_width] = accumulated / count;
+}
diff --git a/opencl/bilateral-filter.cl.h b/opencl/bilateral-filter.cl.h
new file mode 100644
index 0000000..e92e27e
--- /dev/null
+++ b/opencl/bilateral-filter.cl.h
@@ -0,0 +1,49 @@
+static const char* bilateral_filter_cl_source =
+"#define POW2(a) ((a) * (a)) \n"
+"kernel void bilateral_filter(global float4 *in, \n"
+" global float4 *out, \n"
+" const float radius, \n"
+" const float preserve) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int n_radius = ceil(radius); \n"
+" int dst_width = get_global_size(0); \n"
+" int src_width = dst_width + n_radius * 2; \n"
+" \n"
+" int u, v, i, j; \n"
+" float4 center_pix = \n"
+" in[(gidy + n_radius) * src_width + gidx + n_radius]; \n"
+" float4 accumulated = 0.0f; \n"
+" float4 tempf = 0.0f; \n"
+" float count = 0.0f; \n"
+" float diff_map, gaussian_weight, weight; \n"
+" \n"
+" for (v = -n_radius;v <= n_radius; ++v) \n"
+" { \n"
+" for (u = -n_radius;u <= n_radius; ++u) \n"
+" { \n"
+" i = gidx + n_radius + u; \n"
+" j = gidy + n_radius + v; \n"
+" \n"
+" int gid1d = i + j * src_width; \n"
+" tempf = in[gid1d]; \n"
+" \n"
+" diff_map = exp ( \n"
+" - ( POW2(center_pix.x - tempf.x) \n"
+" + POW2(center_pix.y - tempf.y) \n"
+" + POW2(center_pix.z - tempf.z)) \n"
+" * preserve); \n"
+" \n"
+" gaussian_weight = \n"
+" exp( - 0.5f * (POW2(u) + POW2(v)) / radius); \n"
+" \n"
+" weight = diff_map * gaussian_weight; \n"
+" \n"
+" accumulated += tempf * weight; \n"
+" count += weight; \n"
+" } \n"
+" } \n"
+" out[gidx + gidy * dst_width] = accumulated / count; \n"
+"} \n"
+;
diff --git a/opencl/box-blur.cl b/opencl/box-blur.cl
new file mode 100644
index 0000000..f99856e
--- /dev/null
+++ b/opencl/box-blur.cl
@@ -0,0 +1,45 @@
+__kernel void kernel_blur_hor (__global const float4 *in,
+ __global float4 *aux,
+ int width, int radius)
+{
+ const int in_index = get_global_id(0) * (width + 2 * radius)
+ + (radius + get_global_id (1));
+
+ const int aux_index = get_global_id(0) * width + get_global_id (1);
+ int i;
+ float4 mean;
+
+ mean = (float4)(0.0f);
+
+ if (get_global_id(1) < width)
+ {
+ for (i=-radius; i <= radius; i++)
+ {
+ mean += in[in_index + i];
+ }
+ aux[aux_index] = mean / (2 * radius + 1);
+ }
+}
+
+__kernel void kernel_blur_ver (__global const float4 *aux,
+ __global float4 *out,
+ int width, int radius)
+{
+
+ const int out_index = get_global_id(0) * width + get_global_id (1);
+ int i;
+ float4 mean;
+
+ mean = (float4)(0.0f);
+ int aux_index = get_global_id(0) * width + get_global_id (1);
+
+ if(get_global_id(1) < width)
+ {
+ for (i=-radius; i <= radius; i++)
+ {
+ mean += aux[aux_index];
+ aux_index += width;
+ }
+ out[out_index] = mean / (2 * radius + 1);
+ }
+}
diff --git a/opencl/box-blur.cl.h b/opencl/box-blur.cl.h
new file mode 100644
index 0000000..3d168fb
--- /dev/null
+++ b/opencl/box-blur.cl.h
@@ -0,0 +1,47 @@
+static const char* box_blur_cl_source =
+"__kernel void kernel_blur_hor (__global const float4 *in, \n"
+" __global float4 *aux, \n"
+" int width, int radius) \n"
+"{ \n"
+" const int in_index = get_global_id(0) * (width + 2 * radius) \n"
+" + (radius + get_global_id (1)); \n"
+" \n"
+" const int aux_index = get_global_id(0) * width + get_global_id (1); \n"
+" int i; \n"
+" float4 mean; \n"
+" \n"
+" mean = (float4)(0.0f); \n"
+" \n"
+" if (get_global_id(1) < width) \n"
+" { \n"
+" for (i=-radius; i <= radius; i++) \n"
+" { \n"
+" mean += in[in_index + i]; \n"
+" } \n"
+" aux[aux_index] = mean / (2 * radius + 1); \n"
+" } \n"
+"} \n"
+" \n"
+"__kernel void kernel_blur_ver (__global const float4 *aux, \n"
+" __global float4 *out, \n"
+" int width, int radius) \n"
+"{ \n"
+" \n"
+" const int out_index = get_global_id(0) * width + get_global_id (1); \n"
+" int i; \n"
+" float4 mean; \n"
+" \n"
+" mean = (float4)(0.0f); \n"
+" int aux_index = get_global_id(0) * width + get_global_id (1); \n"
+" \n"
+" if(get_global_id(1) < width) \n"
+" { \n"
+" for (i=-radius; i <= radius; i++) \n"
+" { \n"
+" mean += aux[aux_index]; \n"
+" aux_index += width; \n"
+" } \n"
+" out[out_index] = mean / (2 * radius + 1); \n"
+" } \n"
+"} \n"
+;
diff --git a/opencl/c2g.cl b/opencl/c2g.cl
new file mode 100644
index 0000000..6442eeb
--- /dev/null
+++ b/opencl/c2g.cl
@@ -0,0 +1,157 @@
+#define ANGLE_PRIME 95273
+#define RADIUS_PRIME 29537
+
+void sample_min_max(const __global float4 *src_buf,
+ int src_width,
+ int src_height,
+ const __global float *radiuses,
+ const __global float *lut_cos,
+ const __global float *lut_sin,
+ int x,
+ int y,
+ int radius,
+ int samples,
+ float4 *min,
+ float4 *max,
+ int j,
+ int iterations)
+{
+ float4 best_min;
+ float4 best_max;
+ float4 center_pix = *(src_buf + src_width * y + x);
+ int i;
+
+ best_min = center_pix;
+ best_max = center_pix;
+
+ int angle_no = (src_width * y + x) * (iterations) *
+ samples + j * samples;
+ int radius_no = angle_no;
+ angle_no %= ANGLE_PRIME;
+ radius_no %= RADIUS_PRIME;
+ for(i=0; i<samples; i++)
+ {
+ int angle;
+ float rmag;
+ /* if we've sampled outside the valid image
+ area, we grab another sample instead, this
+ should potentially work better than mirroring
+ or extending the image */
+
+ angle = angle_no++;
+ rmag = radiuses[radius_no++] * radius;
+
+ if( angle_no >= ANGLE_PRIME)
+ angle_no = 0;
+ if( radius_no >= RADIUS_PRIME)
+ radius_no = 0;
+
+ int u = x + rmag * lut_cos[angle];
+ int v = y + rmag * lut_sin[angle];
+
+ if(u>=src_width || u <0 || v>=src_height || v<0)
+ {
+ //--i;
+ continue;
+ }
+ float4 pixel = *(src_buf + (src_width * v + u));
+ if(pixel.w<=0.0f)
+ {
+ //--i;
+ continue;
+ }
+
+ best_min = pixel < best_min ? pixel : best_min;
+ best_max = pixel > best_max ? pixel : best_max;
+ }
+
+ (*min).xyz = best_min.xyz;
+ (*max).xyz = best_max.xyz;
+}
+
+void compute_envelopes(const __global float4 *src_buf,
+ int src_width,
+ int src_height,
+ const __global float *radiuses,
+ const __global float *lut_cos,
+ const __global float *lut_sin,
+ int x,
+ int y,
+ int radius,
+ int samples,
+ int iterations,
+ float4 *min_envelope,
+ float4 *max_envelope)
+{
+ float4 range_sum = 0;
+ float4 relative_brightness_sum = 0;
+ float4 pixel = *(src_buf + src_width * y + x);
+
+ int i;
+ for(i =0; i<iterations; i++)
+ {
+ float4 min,max;
+ float4 range, relative_brightness;
+
+ sample_min_max(src_buf, src_width, src_height,
+ radiuses, lut_cos, lut_sin, x, y,
+ radius,samples,&min,&max,i,iterations);
+ range = max - min;
+ relative_brightness = range <= 0.0f ?
+ 0.5f : (pixel - min) / range;
+ relative_brightness_sum += relative_brightness;
+ range_sum += range;
+ }
+
+ float4 relative_brightness = relative_brightness_sum / iterations;
+ float4 range = range_sum / iterations;
+
+ if(max_envelope)
+ *max_envelope = pixel + (1.0f - relative_brightness) * range;
+
+ if(min_envelope)
+ *min_envelope = pixel - relative_brightness * range;
+}
+
+__kernel void c2g(const __global float4 *src_buf,
+ int src_width,
+ int src_height,
+ const __global float *radiuses,
+ const __global float *lut_cos,
+ const __global float *lut_sin,
+ __global float2 *dst_buf,
+ int radius,
+ int samples,
+ int iterations)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+
+ int x = gidx + radius;
+ int y = gidy + radius;
+
+ int src_offset = (src_width * y + x);
+ int dst_offset = gidx + get_global_size(0) * gidy;
+ float4 min,max;
+
+ compute_envelopes(src_buf, src_width, src_height,
+ radiuses, lut_cos, lut_sin, x, y,
+ radius, samples, iterations, &min, &max);
+
+ float4 pixel = *(src_buf + src_offset);
+
+ float nominator=0, denominator=0;
+ float4 t1 = (pixel - min) * (pixel - min);
+ float4 t2 = (pixel - max) * (pixel - max);
+
+ nominator = t1.x + t1.y + t1.z;
+ denominator = t2.x + t2.y + t2.z;
+
+ nominator = sqrt(nominator);
+ denominator = sqrt(denominator);
+ denominator+= nominator + denominator;
+
+ dst_buf[dst_offset].x = (denominator > 0.000f)
+ ? (nominator / denominator) : 0.5f;
+ dst_buf[dst_offset].y = src_buf[src_offset].w;
+}
diff --git a/opencl/c2g.cl.h b/opencl/c2g.cl.h
new file mode 100644
index 0000000..3831faf
--- /dev/null
+++ b/opencl/c2g.cl.h
@@ -0,0 +1,159 @@
+static const char* c2g_cl_source =
+"#define ANGLE_PRIME 95273 \n"
+"#define RADIUS_PRIME 29537 \n"
+" \n"
+"void sample_min_max(const __global float4 *src_buf, \n"
+" int src_width, \n"
+" int src_height, \n"
+" const __global float *radiuses, \n"
+" const __global float *lut_cos, \n"
+" const __global float *lut_sin, \n"
+" int x, \n"
+" int y, \n"
+" int radius, \n"
+" int samples, \n"
+" float4 *min, \n"
+" float4 *max, \n"
+" int j, \n"
+" int iterations) \n"
+"{ \n"
+" float4 best_min; \n"
+" float4 best_max; \n"
+" float4 center_pix = *(src_buf + src_width * y + x); \n"
+" int i; \n"
+" \n"
+" best_min = center_pix; \n"
+" best_max = center_pix; \n"
+" \n"
+" int angle_no = (src_width * y + x) * (iterations) * \n"
+" samples + j * samples; \n"
+" int radius_no = angle_no; \n"
+" angle_no %= ANGLE_PRIME; \n"
+" radius_no %= RADIUS_PRIME; \n"
+" for(i=0; i<samples; i++) \n"
+" { \n"
+" int angle; \n"
+" float rmag; \n"
+" /* if we've sampled outside the valid image \n"
+" area, we grab another sample instead, this \n"
+" should potentially work better than mirroring \n"
+" or extending the image */ \n"
+" \n"
+" angle = angle_no++; \n"
+" rmag = radiuses[radius_no++] * radius; \n"
+" \n"
+" if( angle_no >= ANGLE_PRIME) \n"
+" angle_no = 0; \n"
+" if( radius_no >= RADIUS_PRIME) \n"
+" radius_no = 0; \n"
+" \n"
+" int u = x + rmag * lut_cos[angle]; \n"
+" int v = y + rmag * lut_sin[angle]; \n"
+" \n"
+" if(u>=src_width || u <0 || v>=src_height || v<0) \n"
+" { \n"
+" //--i; \n"
+" continue; \n"
+" } \n"
+" float4 pixel = *(src_buf + (src_width * v + u)); \n"
+" if(pixel.w<=0.0f) \n"
+" { \n"
+" //--i; \n"
+" continue; \n"
+" } \n"
+" \n"
+" best_min = pixel < best_min ? pixel : best_min; \n"
+" best_max = pixel > best_max ? pixel : best_max; \n"
+" } \n"
+" \n"
+" (*min).xyz = best_min.xyz; \n"
+" (*max).xyz = best_max.xyz; \n"
+"} \n"
+" \n"
+"void compute_envelopes(const __global float4 *src_buf, \n"
+" int src_width, \n"
+" int src_height, \n"
+" const __global float *radiuses, \n"
+" const __global float *lut_cos, \n"
+" const __global float *lut_sin, \n"
+" int x, \n"
+" int y, \n"
+" int radius, \n"
+" int samples, \n"
+" int iterations, \n"
+" float4 *min_envelope, \n"
+" float4 *max_envelope) \n"
+"{ \n"
+" float4 range_sum = 0; \n"
+" float4 relative_brightness_sum = 0; \n"
+" float4 pixel = *(src_buf + src_width * y + x); \n"
+" \n"
+" int i; \n"
+" for(i =0; i<iterations; i++) \n"
+" { \n"
+" float4 min,max; \n"
+" float4 range, relative_brightness; \n"
+" \n"
+" sample_min_max(src_buf, src_width, src_height, \n"
+" radiuses, lut_cos, lut_sin, x, y, \n"
+" radius,samples,&min,&max,i,iterations); \n"
+" range = max - min; \n"
+" relative_brightness = range <= 0.0f ? \n"
+" 0.5f : (pixel - min) / range; \n"
+" relative_brightness_sum += relative_brightness; \n"
+" range_sum += range; \n"
+" } \n"
+" \n"
+" float4 relative_brightness = relative_brightness_sum / iterations; \n"
+" float4 range = range_sum / iterations; \n"
+" \n"
+" if(max_envelope) \n"
+" *max_envelope = pixel + (1.0f - relative_brightness) * range; \n"
+" \n"
+" if(min_envelope) \n"
+" *min_envelope = pixel - relative_brightness * range; \n"
+"} \n"
+" \n"
+"__kernel void c2g(const __global float4 *src_buf, \n"
+" int src_width, \n"
+" int src_height, \n"
+" const __global float *radiuses, \n"
+" const __global float *lut_cos, \n"
+" const __global float *lut_sin, \n"
+" __global float2 *dst_buf, \n"
+" int radius, \n"
+" int samples, \n"
+" int iterations) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" \n"
+" int x = gidx + radius; \n"
+" int y = gidy + radius; \n"
+" \n"
+" int src_offset = (src_width * y + x); \n"
+" int dst_offset = gidx + get_global_size(0) * gidy; \n"
+" float4 min,max; \n"
+" \n"
+" compute_envelopes(src_buf, src_width, src_height, \n"
+" radiuses, lut_cos, lut_sin, x, y, \n"
+" radius, samples, iterations, &min, &max); \n"
+" \n"
+" float4 pixel = *(src_buf + src_offset); \n"
+" \n"
+" float nominator=0, denominator=0; \n"
+" float4 t1 = (pixel - min) * (pixel - min); \n"
+" float4 t2 = (pixel - max) * (pixel - max); \n"
+" \n"
+" nominator = t1.x + t1.y + t1.z; \n"
+" denominator = t2.x + t2.y + t2.z; \n"
+" \n"
+" nominator = sqrt(nominator); \n"
+" denominator = sqrt(denominator); \n"
+" denominator+= nominator + denominator; \n"
+" \n"
+" dst_buf[dst_offset].x = (denominator > 0.000f) \n"
+" ? (nominator / denominator) : 0.5f; \n"
+" dst_buf[dst_offset].y = src_buf[src_offset].w; \n"
+"} \n"
+;
diff --git a/opencl/cltostring.pl b/opencl/cltostring.pl
index 11f1fc0..8e6ddc5 100755
--- a/opencl/cltostring.pl
+++ b/opencl/cltostring.pl
@@ -12,7 +12,8 @@ open(OUTPUT_FILE, ">${path}.h");
($name,$dir,$ext) = fileparse($path,'\..*');
-$name =~ s/[^a-zA-Z0-9]*//g;
+$name =~ s/-/_/g;
+$name =~ s/:/_/g;
print OUTPUT_FILE "static const char* ${name}_cl_source =\n";
diff --git a/opencl/color-temperature.cl b/opencl/color-temperature.cl
new file mode 100644
index 0000000..7e974dc
--- /dev/null
+++ b/opencl/color-temperature.cl
@@ -0,0 +1,12 @@
+__kernel void gegl_color_temperature(__global const float4 *in,
+ __global float4 *out,
+ float coeff1,
+ float coeff2,
+ float coeff3)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v = in_v * (float4) (coeff1, coeff2, coeff3, 1.0f);
+ out[gid] = out_v;
+}
diff --git a/opencl/color-temperature.cl.h b/opencl/color-temperature.cl.h
new file mode 100644
index 0000000..12a86b3
--- /dev/null
+++ b/opencl/color-temperature.cl.h
@@ -0,0 +1,14 @@
+static const char* color_temperature_cl_source =
+"__kernel void gegl_color_temperature(__global const float4 *in, \n"
+" __global float4 *out, \n"
+" float coeff1, \n"
+" float coeff2, \n"
+" float coeff3) \n"
+"{ \n"
+" int gid = get_global_id(0); \n"
+" float4 in_v = in[gid]; \n"
+" float4 out_v; \n"
+" out_v = in_v * (float4) (coeff1, coeff2, coeff3, 1.0f); \n"
+" out[gid] = out_v; \n"
+"} \n"
+;
diff --git a/opencl/gaussian-blur.cl b/opencl/gaussian-blur.cl
new file mode 100644
index 0000000..03b2502
--- /dev/null
+++ b/opencl/gaussian-blur.cl
@@ -0,0 +1,52 @@
+float4 fir_get_mean_component_1D(const global float4 *buf,
+ int offset,
+ const int delta_offset,
+ constant float *cmatrix,
+ const int matrix_length)
+{
+ float4 acc = 0.0f;
+ int i;
+
+ for(i=0; i<matrix_length; i++)
+ {
+ acc += buf[offset] * cmatrix[i];
+ offset += delta_offset;
+ }
+ return acc;
+}
+
+__kernel void fir_ver_blur(const global float4 *src_buf,
+ const int src_width,
+ global float4 *dst_buf,
+ constant float *cmatrix,
+ const int matrix_length,
+ const int yoff)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+ int gid = gidx + gidy * get_global_size(0);
+
+ int radius = matrix_length / 2;
+ int src_offset = gidx + (gidy - radius + yoff) * src_width;
+
+ dst_buf[gid] = fir_get_mean_component_1D(
+ src_buf, src_offset, src_width, cmatrix, matrix_length);
+}
+
+__kernel void fir_hor_blur(const global float4 *src_buf,
+ const int src_width,
+ global float4 *dst_buf,
+ constant float *cmatrix,
+ const int matrix_length,
+ const int yoff)
+{
+ int gidx = get_global_id(0);
+ int gidy = get_global_id(1);
+ int gid = gidx + gidy * get_global_size(0);
+
+ int radius = matrix_length / 2;
+ int src_offset = gidy * src_width + (gidx - radius + yoff);
+
+ dst_buf[gid] = fir_get_mean_component_1D(
+ src_buf, src_offset, 1, cmatrix, matrix_length);
+}
diff --git a/opencl/gaussian-blur.cl.h b/opencl/gaussian-blur.cl.h
new file mode 100644
index 0000000..acf7f9d
--- /dev/null
+++ b/opencl/gaussian-blur.cl.h
@@ -0,0 +1,54 @@
+static const char* gaussian_blur_cl_source =
+"float4 fir_get_mean_component_1D(const global float4 *buf, \n"
+" int offset, \n"
+" const int delta_offset, \n"
+" constant float *cmatrix, \n"
+" const int matrix_length) \n"
+"{ \n"
+" float4 acc = 0.0f; \n"
+" int i; \n"
+" \n"
+" for(i=0; i<matrix_length; i++) \n"
+" { \n"
+" acc += buf[offset] * cmatrix[i]; \n"
+" offset += delta_offset; \n"
+" } \n"
+" return acc; \n"
+"} \n"
+" \n"
+"__kernel void fir_ver_blur(const global float4 *src_buf, \n"
+" const int src_width, \n"
+" global float4 *dst_buf, \n"
+" constant float *cmatrix, \n"
+" const int matrix_length, \n"
+" const int yoff) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int gid = gidx + gidy * get_global_size(0); \n"
+" \n"
+" int radius = matrix_length / 2; \n"
+" int src_offset = gidx + (gidy - radius + yoff) * src_width; \n"
+" \n"
+" dst_buf[gid] = fir_get_mean_component_1D( \n"
+" src_buf, src_offset, src_width, cmatrix, matrix_length); \n"
+"} \n"
+" \n"
+"__kernel void fir_hor_blur(const global float4 *src_buf, \n"
+" const int src_width, \n"
+" global float4 *dst_buf, \n"
+" constant float *cmatrix, \n"
+" const int matrix_length, \n"
+" const int yoff) \n"
+"{ \n"
+" int gidx = get_global_id(0); \n"
+" int gidy = get_global_id(1); \n"
+" int gid = gidx + gidy * get_global_size(0); \n"
+" \n"
+" int radius = matrix_length / 2; \n"
+" int src_offset = gidy * src_width + (gidx - radius + yoff); \n"
+" \n"
+" dst_buf[gid] = fir_get_mean_component_1D( \n"
+" src_buf, src_offset, 1, cmatrix, matrix_length); \n"
+"} \n"
+;
diff --git a/operations/common/bilateral-filter.c b/operations/common/bilateral-filter.c
index 6a0eb25..97cafd4 100644
--- a/operations/common/bilateral-filter.c
+++ b/operations/common/bilateral-filter.c
@@ -60,54 +60,7 @@ static void prepare (GeglOperation *operation)
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"
-static const char* kernel_source =
-"#define POW2(a) ((a) * (a)) \n"
-"kernel void bilateral_filter(global float4 *in, \n"
-" global float4 *out, \n"
-" const float radius, \n"
-" const float preserve) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" int n_radius = ceil(radius); \n"
-" int dst_width = get_global_size(0); \n"
-" int src_width = dst_width + n_radius * 2; \n"
-" \n"
-" int u, v, i, j; \n"
-" float4 center_pix = \n"
-" in[(gidy + n_radius) * src_width + gidx + n_radius]; \n"
-" float4 accumulated = 0.0f; \n"
-" float4 tempf = 0.0f; \n"
-" float count = 0.0f; \n"
-" float diff_map, gaussian_weight, weight; \n"
-" \n"
-" for (v = -n_radius;v <= n_radius; ++v) \n"
-" { \n"
-" for (u = -n_radius;u <= n_radius; ++u) \n"
-" { \n"
-" i = gidx + n_radius + u; \n"
-" j = gidy + n_radius + v; \n"
-" \n"
-" int gid1d = i + j * src_width; \n"
-" tempf = in[gid1d]; \n"
-" \n"
-" diff_map = exp ( \n"
-" - ( POW2(center_pix.x - tempf.x) \n"
-" + POW2(center_pix.y - tempf.y) \n"
-" + POW2(center_pix.z - tempf.z)) \n"
-" * preserve); \n"
-" \n"
-" gaussian_weight = \n"
-" exp( - 0.5f * (POW2(u) + POW2(v)) / radius); \n"
-" \n"
-" weight = diff_map * gaussian_weight; \n"
-" \n"
-" accumulated += tempf * weight; \n"
-" count += weight; \n"
-" } \n"
-" } \n"
-" out[gidx + gidy * dst_width] = accumulated / count; \n"
-"} \n";
+#include "opencl/bilateral-filter.cl.h"
static GeglClRunData *cl_data = NULL;
@@ -125,7 +78,7 @@ cl_bilateral_filter (cl_mem in_tex,
if (!cl_data)
{
const char *kernel_name[] = {"bilateral_filter", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ cl_data = gegl_cl_compile_and_build (bilateral_filter_cl_source, kernel_name);
}
if (!cl_data) return 1;
diff --git a/operations/common/box-blur.c b/operations/common/box-blur.c
index 13328a9..760c3a7 100644
--- a/operations/common/box-blur.c
+++ b/operations/common/box-blur.c
@@ -180,52 +180,7 @@ static void prepare (GeglOperation *operation)
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"
-static const char* kernel_source =
-"__kernel void kernel_blur_hor (__global const float4 *in, \n"
-" __global float4 *aux, \n"
-" int width, int radius) \n"
-"{ \n"
-" const int in_index = get_global_id(0) * (width + 2 * radius) \n"
-" + (radius + get_global_id (1)); \n"
-" \n"
-" const int aux_index = get_global_id(0) * width + get_global_id (1); \n"
-" int i; \n"
-" float4 mean; \n"
-" \n"
-" mean = (float4)(0.0f); \n"
-" \n"
-" if (get_global_id(1) < width) \n"
-" { \n"
-" for (i=-radius; i <= radius; i++) \n"
-" { \n"
-" mean += in[in_index + i]; \n"
-" } \n"
-" aux[aux_index] = mean / (2 * radius + 1); \n"
-" } \n"
-"} \n"
-
-"__kernel void kernel_blur_ver (__global const float4 *aux, \n"
-" __global float4 *out, \n"
-" int width, int radius) \n"
-"{ \n"
-" \n"
-" const int out_index = get_global_id(0) * width + get_global_id (1); \n"
-" int i; \n"
-" float4 mean; \n"
-" \n"
-" mean = (float4)(0.0f); \n"
-" int aux_index = get_global_id(0) * width + get_global_id (1); \n"
-" \n"
-" if(get_global_id(1) < width) \n"
-" { \n"
-" for (i=-radius; i <= radius; i++) \n"
-" { \n"
-" mean += aux[aux_index]; \n"
-" aux_index += width; \n"
-" } \n"
-" out[out_index] = mean / (2 * radius + 1); \n"
-" } \n"
-"} \n";
+#include "opencl/box-blur.cl.h"
static GeglClRunData *cl_data = NULL;
@@ -244,7 +199,7 @@ cl_box_blur (cl_mem in_tex,
if (!cl_data)
{
const char *kernel_name[] = {"kernel_blur_hor", "kernel_blur_ver", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name);
}
if (!cl_data) return 1;
diff --git a/operations/common/brightness-contrast.c b/operations/common/brightness-contrast.c
index 22ab173..d24e69a 100644
--- a/operations/common/brightness-contrast.c
+++ b/operations/common/brightness-contrast.c
@@ -145,7 +145,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"name", "gegl:brightness-contrast",
"categories", "color",
"description", _("Changes the light level and contrast."),
- "cl-source" , brightnesscontrast_cl_source,
+ "cl-source" , brightness_contrast_cl_source,
"reference-composition", composition,
NULL);
}
diff --git a/operations/common/c2g.c b/operations/common/c2g.c
index 4e692fb..742cf07 100644
--- a/operations/common/c2g.c
+++ b/operations/common/c2g.c
@@ -152,168 +152,7 @@ get_bounding_box (GeglOperation *operation)
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"
-static const char* kernel_source =
-"#define TRUE true \n"
-" \n"
-"#define FALSE false \n"
-"#define ANGLE_PRIME 95273 \n"
-"#define RADIUS_PRIME 29537 \n"
-" \n"
-"void sample_min_max(const __global float4 *src_buf, \n"
-" int src_width, \n"
-" int src_height, \n"
-" const __global float *radiuses, \n"
-" const __global float *lut_cos, \n"
-" const __global float *lut_sin, \n"
-" int x, \n"
-" int y, \n"
-" int radius, \n"
-" int samples, \n"
-" float4 *min, \n"
-" float4 *max, \n"
-" int j, \n"
-" int iterations) \n"
-"{ \n"
-" float4 best_min; \n"
-" float4 best_max; \n"
-" float4 center_pix = *(src_buf + src_width * y + x); \n"
-" int i; \n"
-" \n"
-" best_min = center_pix; \n"
-" best_max = center_pix; \n"
-" \n"
-" int angle_no = (src_width * y + x) * (iterations) * \n"
-" samples + j * samples; \n"
-" int radius_no = angle_no; \n"
-" angle_no %= ANGLE_PRIME; \n"
-" radius_no %= RADIUS_PRIME; \n"
-" for(i=0; i<samples; i++) \n"
-" { \n"
-" int angle; \n"
-" float rmag; \n"
-" /* if we've sampled outside the valid image \n"
-" area, we grab another sample instead, this \n"
-" should potentially work better than mirroring \n"
-" or extending the image */ \n"
-" \n"
-" angle = angle_no++; \n"
-" rmag = radiuses[radius_no++] * radius; \n"
-" \n"
-" if( angle_no >= ANGLE_PRIME) \n"
-" angle_no = 0; \n"
-" if( radius_no >= RADIUS_PRIME) \n"
-" radius_no = 0; \n"
-" \n"
-" int u = x + rmag * lut_cos[angle]; \n"
-" int v = y + rmag * lut_sin[angle]; \n"
-" \n"
-" if(u>=src_width || u <0 || v>=src_height || v<0) \n"
-" { \n"
-" //--i; \n"
-" continue; \n"
-" } \n"
-" float4 pixel = *(src_buf + (src_width * v + u)); \n"
-" if(pixel.w<=0.0f) \n"
-" { \n"
-" //--i; \n"
-" continue; \n"
-" } \n"
-" \n"
-" best_min = pixel < best_min ? pixel : best_min; \n"
-" best_max = pixel > best_max ? pixel : best_max; \n"
-" } \n"
-" \n"
-" (*min).xyz = best_min.xyz; \n"
-" (*max).xyz = best_max.xyz; \n"
-"} \n"
-" \n"
-"void compute_envelopes_CL(const __global float4 *src_buf, \n"
-" int src_width, \n"
-" int src_height, \n"
-" const __global float *radiuses, \n"
-" const __global float *lut_cos, \n"
-" const __global float *lut_sin, \n"
-" int x, \n"
-" int y, \n"
-" int radius, \n"
-" int samples, \n"
-" int iterations, \n"
-" float4 *min_envelope, \n"
-" float4 *max_envelope) \n"
-"{ \n"
-" float4 range_sum = 0; \n"
-" float4 relative_brightness_sum = 0; \n"
-" float4 pixel = *(src_buf + src_width * y + x); \n"
-" \n"
-" int i; \n"
-" for(i =0; i<iterations; i++) \n"
-" { \n"
-" float4 min,max; \n"
-" float4 range, relative_brightness; \n"
-" \n"
-" sample_min_max(src_buf, src_width, src_height, \n"
-" radiuses, lut_cos, lut_sin, x, y, \n"
-" radius,samples,&min,&max,i,iterations); \n"
-" range = max - min; \n"
-" relative_brightness = range <= 0.0f ? \n"
-" 0.5f : (pixel - min) / range; \n"
-" relative_brightness_sum += relative_brightness; \n"
-" range_sum += range; \n"
-" } \n"
-" \n"
-" float4 relative_brightness = relative_brightness_sum / iterations;\n"
-" float4 range = range_sum / iterations; \n"
-" \n"
-" if(max_envelope) \n"
-" *max_envelope = pixel + (1.0f - relative_brightness) * range; \n"
-" \n"
-" if(min_envelope) \n"
-" *min_envelope = pixel - relative_brightness * range; \n"
-"} \n"
-" \n"
-"__kernel void C2g_CL(const __global float4 *src_buf, \n"
-" int src_width, \n"
-" int src_height, \n"
-" const __global float *radiuses, \n"
-" const __global float *lut_cos, \n"
-" const __global float *lut_sin, \n"
-" __global float2 *dst_buf, \n"
-" int radius, \n"
-" int samples, \n"
-" int iterations) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" \n"
-" int x = gidx + radius; \n"
-" int y = gidy + radius; \n"
-" \n"
-" int src_offset = (src_width * y + x); \n"
-" int dst_offset = gidx + get_global_size(0) * gidy; \n"
-" float4 min,max; \n"
-" \n"
-" compute_envelopes_CL(src_buf, src_width, src_height, \n"
-" radiuses, lut_cos, lut_sin, x, y, \n"
-" radius, samples, iterations, &min, &max); \n"
-" \n"
-" float4 pixel = *(src_buf + src_offset); \n"
-" \n"
-" float nominator=0, denominator=0; \n"
-" float4 t1 = (pixel - min) * (pixel - min); \n"
-" float4 t2 = (pixel - max) * (pixel - max); \n"
-" \n"
-" nominator = t1.x + t1.y + t1.z; \n"
-" denominator = t2.x + t2.y + t2.z; \n"
-" \n"
-" nominator = sqrt(nominator); \n"
-" denominator = sqrt(denominator); \n"
-" denominator+= nominator + denominator; \n"
-" \n"
-" dst_buf[dst_offset].x = (denominator > 0.000f) \n"
-" ? (nominator / denominator) : 0.5f; \n"
-" dst_buf[dst_offset].y = src_buf[src_offset].w; \n"
-"} \n"
-" \n";
+#include "opencl/c2g.cl.h"
static GeglClRunData *cl_data = NULL;
@@ -334,8 +173,8 @@ cl_c2g (cl_mem in_tex,
if (!cl_data)
{
- const char *kernel_name[] ={"C2g_CL", NULL};
- cl_data = gegl_cl_compile_and_build(kernel_source, kernel_name);
+ const char *kernel_name[] ={"c2g", NULL};
+ cl_data = gegl_cl_compile_and_build(c2g_cl_source, kernel_name);
}
if (!cl_data) return 0;
@@ -406,7 +245,7 @@ cl_c2g (cl_mem in_tex,
}
static gboolean
-cl_process (GeglOperation *operation,
+cl_process (GeglOperation *operation,
GeglBuffer *input,
GeglBuffer *output,
const GeglRectangle *result)
diff --git a/operations/common/color-temperature.c b/operations/common/color-temperature.c
index bb134fd..a177380 100644
--- a/operations/common/color-temperature.c
+++ b/operations/common/color-temperature.c
@@ -172,20 +172,7 @@ process (GeglOperation *op,
}
#include "opencl/gegl-cl.h"
-
-static const char* kernel_source =
-"__kernel void kernel_temp(__global const float4 *in, \n"
-" __global float4 *out, \n"
-" float coeff1, \n"
-" float coeff2, \n"
-" float coeff3) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 in_v = in[gid]; \n"
-" float4 out_v; \n"
-" out_v = in_v * (float4) (coeff1, coeff2, coeff3, 1.0f); \n"
-" out[gid] = out_v; \n"
-"} \n";
+#include "opencl/color-temperature.cl.h"
static GeglClRunData *cl_data = NULL;
@@ -214,8 +201,8 @@ cl_process (GeglOperation *op,
if (!cl_data)
{
- const char *kernel_name[] = {"kernel_temp", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ const char *kernel_name[] = {"gegl_color_temperature", NULL};
+ cl_data = gegl_cl_compile_and_build (color_temperature_cl_source, kernel_name);
}
if (!cl_data) return 1;
diff --git a/operations/common/gaussian-blur.c b/operations/common/gaussian-blur.c
index 4a749c2..20f8723 100644
--- a/operations/common/gaussian-blur.c
+++ b/operations/common/gaussian-blur.c
@@ -418,59 +418,7 @@ static void prepare (GeglOperation *operation)
#include "opencl/gegl-cl.h"
#include "buffer/gegl-buffer-cl-iterator.h"
-static const char* kernel_source =
-"float4 fir_get_mean_component_1D_CL(const global float4 *buf, \n"
-" int offset, \n"
-" const int delta_offset, \n"
-" constant float *cmatrix, \n"
-" const int matrix_length) \n"
-"{ \n"
-" float4 acc = 0.0f; \n"
-" int i; \n"
-" \n"
-" for(i=0; i<matrix_length; i++) \n"
-" { \n"
-" acc += buf[offset] * cmatrix[i]; \n"
-" offset += delta_offset; \n"
-" } \n"
-" return acc; \n"
-"} \n"
-" \n"
-"__kernel void fir_ver_blur_CL(const global float4 *src_buf, \n"
-" const int src_width, \n"
-" global float4 *dst_buf, \n"
-" constant float *cmatrix, \n"
-" const int matrix_length, \n"
-" const int yoff) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" int gid = gidx + gidy * get_global_size(0); \n"
-" \n"
-" int radius = matrix_length / 2; \n"
-" int src_offset = gidx + (gidy - radius + yoff) * src_width; \n"
-" \n"
-" dst_buf[gid] = fir_get_mean_component_1D_CL( \n"
-" src_buf, src_offset, src_width, cmatrix, matrix_length); \n"
-"} \n"
-" \n"
-"__kernel void fir_hor_blur_CL(const global float4 *src_buf, \n"
-" const int src_width, \n"
-" global float4 *dst_buf, \n"
-" constant float *cmatrix, \n"
-" const int matrix_length, \n"
-" const int yoff) \n"
-"{ \n"
-" int gidx = get_global_id(0); \n"
-" int gidy = get_global_id(1); \n"
-" int gid = gidx + gidy * get_global_size(0); \n"
-" \n"
-" int radius = matrix_length / 2; \n"
-" int src_offset = gidy * src_width + (gidx - radius + yoff); \n"
-" \n"
-" dst_buf[gid] = fir_get_mean_component_1D_CL( \n"
-" src_buf, src_offset, 1, cmatrix, matrix_length); \n"
-"} \n";
+#include "opencl/gaussian-blur.cl.h"
static GeglClRunData *cl_data = NULL;
@@ -497,8 +445,8 @@ cl_gaussian_blur (cl_mem in_tex,
if (!cl_data)
{
- const char *kernel_name[] = {"fir_ver_blur_CL", "fir_hor_blur_CL", NULL};
- cl_data = gegl_cl_compile_and_build (kernel_source, kernel_name);
+ const char *kernel_name[] = {"fir_ver_blur", "fir_hor_blur", NULL};
+ cl_data = gegl_cl_compile_and_build (gaussian_blur_cl_source, kernel_name);
}
if (!cl_data) return 1;
diff --git a/operations/common/invert.c b/operations/common/invert.c
index cd49db4..d8627ff 100644
--- a/operations/common/invert.c
+++ b/operations/common/invert.c
@@ -56,17 +56,7 @@ process (GeglOperation *op,
return TRUE;
}
-static const char* kernel_source =
-"__kernel void gegl_invert (__global const float4 *in, \n"
-" __global float4 *out) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 in_v = in[gid]; \n"
-" float4 out_v; \n"
-" out_v.xyz = (1.0f - in_v.xyz); \n"
-" out_v.w = in_v.w; \n"
-" out[gid] = out_v; \n"
-"} \n";
+#include "opencl/invert.cl.h"
static void
gegl_chant_class_init (GeglChantClass *klass)
@@ -85,7 +75,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"description",
_("Inverts the components (except alpha), the result is the "
"corresponding \"negative\" image."),
- "cl-source" , kernel_source,
+ "cl-source" , invert_cl_source,
NULL);
}
diff --git a/operations/common/mono-mixer.c b/operations/common/mono-mixer.c
index 53a68ea..c289e27 100644
--- a/operations/common/mono-mixer.c
+++ b/operations/common/mono-mixer.c
@@ -73,18 +73,7 @@ process (GeglOperation *op,
return TRUE;
}
-static const char* kernel_source =
-"__kernel void gegl_mono_mixer (__global const float4 *src_buf, \n"
-" __global float2 *dst_buf, \n"
-" float red, \n"
-" float green, \n"
-" float blue) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 in_v = src_buf[gid]; \n"
-" dst_buf[gid].x = in_v.x * red + in_v.y * green + in_v.z * blue; \n"
-" dst_buf[gid].y = in_v.w; \n"
-"} \n";
+#include "opencl/mono-mixer.cl.h"
static void
gegl_chant_class_init (GeglChantClass *klass)
@@ -102,7 +91,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"name" , "gegl:mono-mixer",
"categories" , "color",
"description", _("Monochrome channel mixer"),
- "cl-source" , kernel_source,
+ "cl-source" , mono_mixer_cl_source,
NULL);
}
diff --git a/operations/common/over.c b/operations/common/over.c
index 7cbe0dd..826c283 100644
--- a/operations/common/over.c
+++ b/operations/common/over.c
@@ -70,19 +70,7 @@ process (GeglOperation *op,
return TRUE;
}
-static const char* kernel_source =
-"__kernel void svg_src_over (__global const float4 *in, \n"
-" __global const float4 *aux, \n"
-" __global float4 *out) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 in_v = in [gid]; \n"
-" float4 aux_v = aux[gid]; \n"
-" float4 out_v; \n"
-" out_v.xyz = aux_v.xyz + in_v.xyz * (1.0f - aux_v.w); \n"
-" out_v.w = aux_v.w + in_v.w - aux_v.w * in_v.w; \n"
-" out[gid] = out_v; \n"
-"} \n";
+#include "opencl/svg-src-over.cl.h"
/* Fast paths */
static gboolean operation_process (GeglOperation *operation,
@@ -151,7 +139,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"categories" , "compositors:porter-duff",
"description",
_("Porter Duff operation over (d = cA + cB * (1 - aA))"),
- "cl-source" , kernel_source,
+ "cl-source" , svg_src_over_cl_source,
NULL);
}
diff --git a/operations/common/threshold.c b/operations/common/threshold.c
index 74b18f2..6c06654 100644
--- a/operations/common/threshold.c
+++ b/operations/common/threshold.c
@@ -89,20 +89,7 @@ process (GeglOperation *op,
return TRUE;
}
-static const char* kernel_source =
-"__kernel void gegl_threshold (__global const float2 *in, \n"
-" __global const float *aux, \n"
-" __global float2 *out, \n"
-" float value) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float2 in_v = in [gid]; \n"
-" float aux_v = (aux)? aux[gid] : value; \n"
-" float2 out_v; \n"
-" out_v.x = (in_v.x > aux_v)? 1.0f : 0.0f; \n"
-" out_v.y = in_v.y; \n"
-" out[gid] = out_v; \n"
-"} \n";
+#include "opencl/threshold.cl.h"
static void
gegl_chant_class_init (GeglChantClass *klass)
@@ -122,7 +109,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"description",
_("Thresholds the image to white/black based on either the global value "
"set in the value property, or per pixel from the aux input."),
- "cl-source" , kernel_source,
+ "cl-source" , threshold_cl_source,
NULL);
}
diff --git a/operations/common/value-invert.c b/operations/common/value-invert.c
index 5ab1191..26b4a21 100644
--- a/operations/common/value-invert.c
+++ b/operations/common/value-invert.c
@@ -129,36 +129,7 @@ process (GeglOperation *op,
return TRUE;
}
-static const char* kernel_source =
-"__kernel void gegl_value_invert (__global const float4 *in, \n"
-" __global float4 *out) \n"
-"{ \n"
-" int gid = get_global_id(0); \n"
-" float4 in_v = in[gid]; \n"
-" float4 out_v; \n"
-" \n"
-" float value = fmax (in_v.x, fmax (in_v.y, in_v.z)); \n"
-" float minv = fmin (in_v.x, fmin (in_v.y, in_v.z)); \n"
-" float delta = value - minv; \n"
-" \n"
-" if (value == 0.0f || delta == 0.0f) \n"
-" { \n"
-" out_v = (float4) ((1.0f - value), \n"
-" (1.0f - value), \n"
-" (1.0f - value), \n"
-" in_v.w); \n"
-" } \n"
-" else \n"
-" { \n"
-" out_v = (float4) ((1.0f - value) * in_v.x / value, \n"
-" (1.0f - value) * in_v.y / value, \n"
-" (1.0f - value) * in_v.z / value, \n"
-" in_v.w); \n"
-" } \n"
-" \n"
-" out_v.w = in_v.w; \n"
-" out[gid] = out_v; \n"
-"} \n";
+#include "opencl/value-invert.cl.h"
static void
gegl_chant_class_init (GeglChantClass *klass)
@@ -178,7 +149,7 @@ gegl_chant_class_init (GeglChantClass *klass)
"description",
_("Inverts just the value component, the result is the corresponding "
"`inverted' image."),
- "cl-source" , kernel_source,
+ "cl-source" , value_invert_cl_source,
NULL);
}