summaryrefslogtreecommitdiff
path: root/opencl/box-min.cl
diff options
context:
space:
mode:
authorCarlos Zubieta <czubieta.dev@gmail.com>2013-07-23 18:47:59 -0500
committerTéo Mazars <teo.mazars@ensimag.fr>2013-10-31 11:41:43 +0100
commit13b09cab60952d79e1c99b27e91f3f0304bd6647 (patch)
tree14f492e3a8059e79ce7daa0970074bab737a55fa /opencl/box-min.cl
parent6aea9f790ef65ede7733e07113de9d21c2b78726 (diff)
Added OpenCL support to box-min
Diffstat (limited to 'opencl/box-min.cl')
-rw-r--r--opencl/box-min.cl49
1 files changed, 49 insertions, 0 deletions
diff --git a/opencl/box-min.cl b/opencl/box-min.cl
new file mode 100644
index 00000000..f0978de6
--- /dev/null
+++ b/opencl/box-min.cl
@@ -0,0 +1,49 @@
+__kernel void kernel_min_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 min;
+ float4 in_v;
+
+ min = (float4)(1000000000.0f);
+
+ if (get_global_id(1) < width)
+ {
+ for (i=-radius; i <= radius; i++)
+ {
+ in_v = in[in_index + i];
+ min = min > in_v ? in_v : min;
+ }
+ aux[aux_index] = min;
+ }
+}
+
+__kernel void kernel_max_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 aux_index = out_index;
+ int i;
+ float4 min;
+ float4 aux_v;
+
+ min = (float4)(1000000000.0f);
+
+ if(get_global_id(1) < width)
+ {
+ for (i=-radius; i <= radius; i++)
+ {
+ aux_v = aux[aux_index];
+ min = min > aux_v ? aux_v : min;
+ aux_index += width;
+ }
+ out[out_index] = min;
+ }
+}