[gegl] opencl: splitting opencl kernels in separate files for point ops
- From: Victor Matheus de Araujo Oliveira <vmaolive src gnome org>
- To: commits-list gnome org
- Cc:
- Subject: [gegl] opencl: splitting opencl kernels in separate files for point ops
- Date: Fri, 11 Jan 2013 21:16:53 +0000 (UTC)
commit 785a59f673ead2f4610d2ffb0c34008556aeea63
Author: Victor Oliveira <victormatheus gmail com>
Date: Fri Jan 11 19:15:51 2013 -0200
opencl: splitting opencl kernels in separate files for point ops
opencl/brightness-contrast.cl.h | 2 +-
opencl/invert.cl | 10 ++++++++++
opencl/invert.cl.h | 12 ++++++++++++
opencl/mono-mixer.cl | 11 +++++++++++
opencl/mono-mixer.cl.h | 13 +++++++++++++
opencl/svg-src-over.cl | 12 ++++++++++++
opencl/svg-src-over.cl.h | 14 ++++++++++++++
opencl/threshold.cl | 13 +++++++++++++
opencl/threshold.cl.h | 15 +++++++++++++++
opencl/value-invert.cl | 30 ++++++++++++++++++++++++++++++
opencl/value-invert.cl.h | 32 ++++++++++++++++++++++++++++++++
11 files changed, 163 insertions(+), 1 deletions(-)
---
diff --git a/opencl/brightness-contrast.cl.h b/opencl/brightness-contrast.cl.h
index c21ef7b..cb758f2 100644
--- a/opencl/brightness-contrast.cl.h
+++ b/opencl/brightness-contrast.cl.h
@@ -1,4 +1,4 @@
-static const char* brightnesscontrast_cl_source =
+static const char* brightness_contrast_cl_source =
"__kernel void gegl_brightness_contrast(__global const float4 *in, \n"
" __global float4 *out, \n"
" float contrast, \n"
diff --git a/opencl/invert.cl b/opencl/invert.cl
new file mode 100644
index 0000000..33fdb48
--- /dev/null
+++ b/opencl/invert.cl
@@ -0,0 +1,10 @@
+__kernel void gegl_invert (__global const float4 *in,
+ __global float4 *out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+ out_v.xyz = (1.0f - in_v.xyz);
+ out_v.w = in_v.w;
+ out[gid] = out_v;
+}
diff --git a/opencl/invert.cl.h b/opencl/invert.cl.h
new file mode 100644
index 0000000..05b9d57
--- /dev/null
+++ b/opencl/invert.cl.h
@@ -0,0 +1,12 @@
+static const char* invert_cl_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"
+;
diff --git a/opencl/mono-mixer.cl b/opencl/mono-mixer.cl
new file mode 100644
index 0000000..88b337b
--- /dev/null
+++ b/opencl/mono-mixer.cl
@@ -0,0 +1,11 @@
+__kernel void gegl_mono_mixer (__global const float4 *src_buf,
+ __global float2 *dst_buf,
+ float red,
+ float green,
+ float blue)
+{
+ int gid = get_global_id(0);
+ float4 in_v = src_buf[gid];
+ dst_buf[gid].x = in_v.x * red + in_v.y * green + in_v.z * blue;
+ dst_buf[gid].y = in_v.w;
+}
diff --git a/opencl/mono-mixer.cl.h b/opencl/mono-mixer.cl.h
new file mode 100644
index 0000000..7b2fbc8
--- /dev/null
+++ b/opencl/mono-mixer.cl.h
@@ -0,0 +1,13 @@
+static const char* mono_mixer_cl_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"
+;
diff --git a/opencl/svg-src-over.cl b/opencl/svg-src-over.cl
new file mode 100644
index 0000000..13a1d19
--- /dev/null
+++ b/opencl/svg-src-over.cl
@@ -0,0 +1,12 @@
+__kernel void svg_src_over (__global const float4 *in,
+ __global const float4 *aux,
+ __global float4 *out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in [gid];
+ float4 aux_v = aux[gid];
+ float4 out_v;
+ out_v.xyz = aux_v.xyz + in_v.xyz * (1.0f - aux_v.w);
+ out_v.w = aux_v.w + in_v.w - aux_v.w * in_v.w;
+ out[gid] = out_v;
+}
diff --git a/opencl/svg-src-over.cl.h b/opencl/svg-src-over.cl.h
new file mode 100644
index 0000000..b69eaa1
--- /dev/null
+++ b/opencl/svg-src-over.cl.h
@@ -0,0 +1,14 @@
+static const char* svg_src_over_cl_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"
+;
diff --git a/opencl/threshold.cl b/opencl/threshold.cl
new file mode 100644
index 0000000..a546290
--- /dev/null
+++ b/opencl/threshold.cl
@@ -0,0 +1,13 @@
+__kernel void gegl_threshold (__global const float2 *in,
+ __global const float *aux,
+ __global float2 *out,
+ float value)
+{
+ int gid = get_global_id(0);
+ float2 in_v = in [gid];
+ float aux_v = (aux)? aux[gid] : value;
+ float2 out_v;
+ out_v.x = (in_v.x > aux_v)? 1.0f : 0.0f;
+ out_v.y = in_v.y;
+ out[gid] = out_v;
+}
diff --git a/opencl/threshold.cl.h b/opencl/threshold.cl.h
new file mode 100644
index 0000000..d14148f
--- /dev/null
+++ b/opencl/threshold.cl.h
@@ -0,0 +1,15 @@
+static const char* threshold_cl_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"
+;
diff --git a/opencl/value-invert.cl b/opencl/value-invert.cl
new file mode 100644
index 0000000..de029c6
--- /dev/null
+++ b/opencl/value-invert.cl
@@ -0,0 +1,30 @@
+__kernel void gegl_value_invert (__global const float4 *in,
+ __global float4 *out)
+{
+ int gid = get_global_id(0);
+ float4 in_v = in[gid];
+ float4 out_v;
+
+ float value = fmax (in_v.x, fmax (in_v.y, in_v.z));
+ float minv = fmin (in_v.x, fmin (in_v.y, in_v.z));
+ float delta = value - minv;
+
+ if (value == 0.0f || delta == 0.0f)
+ {
+ out_v = (float4) ((1.0f - value),
+ (1.0f - value),
+ (1.0f - value),
+ in_v.w);
+ }
+ else
+ {
+ out_v = (float4) ((1.0f - value) * in_v.x / value,
+ (1.0f - value) * in_v.y / value,
+ (1.0f - value) * in_v.z / value,
+ in_v.w);
+ }
+
+ out_v.w = in_v.w;
+ out[gid] = out_v;
+}
+
diff --git a/opencl/value-invert.cl.h b/opencl/value-invert.cl.h
new file mode 100644
index 0000000..b9c7b2d
--- /dev/null
+++ b/opencl/value-invert.cl.h
@@ -0,0 +1,32 @@
+static const char* value_invert_cl_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"
+" \n"
+;
[
Date Prev][
Date Next] [
Thread Prev][
Thread Next]
[
Thread Index]
[
Date Index]
[
Author Index]
