diff options
| author | Carlos Zubieta <czubieta.dev@gmail.com> | 2013-07-23 18:49:25 -0500 |
|---|---|---|
| committer | Téo Mazars <teo.mazars@ensimag.fr> | 2013-10-31 11:41:43 +0100 |
| commit | 8a6eedf9c5b8676c0c1e2be5ba0b91553dde301f (patch) | |
| tree | 602461635e61116afc6160eb095e6a85e34cfe6b /opencl | |
| parent | 13b09cab60952d79e1c99b27e91f3f0304bd6647 (diff) | |
Added OpenCL support to box-max
Diffstat (limited to 'opencl')
| -rw-r--r-- | opencl/box-max.cl | 49 | ||||
| -rw-r--r-- | opencl/box-max.cl.h | 51 |
2 files changed, 100 insertions, 0 deletions
diff --git a/opencl/box-max.cl b/opencl/box-max.cl new file mode 100644 index 00000000..b780ecea --- /dev/null +++ b/opencl/box-max.cl @@ -0,0 +1,49 @@ +__kernel void kernel_max_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 max; + float4 in_v; + + max = (float4)(-1000000000.0f); + + if (get_global_id(1) < width) + { + for (i=-radius; i <= radius; i++) + { + in_v = in[in_index + i]; + max = max < in_v ? in_v : max; + } + aux[aux_index] = max; + } +} + +__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 max; + float4 aux_v; + + max = (float4)(-1000000000.0f); + + if(get_global_id(1) < width) + { + for (i=-radius; i <= radius; i++) + { + aux_v = aux[aux_index]; + max = max < aux_v ? aux_v : max; + aux_index += width; + } + out[out_index] = max; + } +} diff --git a/opencl/box-max.cl.h b/opencl/box-max.cl.h new file mode 100644 index 00000000..983625dd --- /dev/null +++ b/opencl/box-max.cl.h @@ -0,0 +1,51 @@ +static const char* box_max_cl_source = +"__kernel void kernel_max_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 max; \n" +" float4 in_v; \n" +" \n" +" max = (float4)(-1000000000.0f); \n" +" \n" +" if (get_global_id(1) < width) \n" +" { \n" +" for (i=-radius; i <= radius; i++) \n" +" { \n" +" in_v = in[in_index + i]; \n" +" max = max < in_v ? in_v : max; \n" +" } \n" +" aux[aux_index] = max; \n" +" } \n" +"} \n" +" \n" +"__kernel void kernel_max_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 aux_index = out_index; \n" +" int i; \n" +" float4 max; \n" +" float4 aux_v; \n" +" \n" +" max = (float4)(-1000000000.0f); \n" +" \n" +" if(get_global_id(1) < width) \n" +" { \n" +" for (i=-radius; i <= radius; i++) \n" +" { \n" +" aux_v = aux[aux_index]; \n" +" max = max < aux_v ? aux_v : max; \n" +" aux_index += width; \n" +" } \n" +" out[out_index] = max; \n" +" } \n" +"} \n" +; |