opencl/box-blur.cl.h


1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124

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 / (float)(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"
"  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 / (float)(2 * radius + 1);                        \n"
"    }                                                                         \n"
"}                                                                             \n"
"                                                                              \n"
"__kernel                                                                      \n"
"__attribute__((reqd_work_group_size(256,1,1)))                                \n"
"void kernel_box_blur_fast (__global const float4 *in,                         \n"
"                           __global       float4 *out,                        \n"
"                           __local        float4 *column_sum,                 \n"
"                                    const int     width,                      \n"
"                                    const int     height,                     \n"
"                                    const int     radius,                     \n"
"                                    const int     size)                       \n"
"{                                                                             \n"
"  const int local_id0 = get_local_id(0);                                      \n"
"  const int twice_radius = 2 * radius;                                        \n"
"  const int in_width = twice_radius + width;                                  \n"
"  const int in_height = twice_radius + height;                                \n"
"  const float4 area = (float4)( (twice_radius + 1) * (twice_radius + 1) );    \n"
"  int column_index_start,column_index_end;                                    \n"
"  int y = get_global_id(1) * size;                                            \n"
"  const int out_x = get_group_id(0)                                           \n"
"                    * ( get_local_size(0) - twice_radius ) + local_id0 - radius;\n"
"  const int in_x = out_x + radius;                                            \n"
"  int tmp_size = size;                                                        \n"
"  int tmp_index = 0;                                                          \n"
"  float4 tmp_sum = (float4)0.0f;                                              \n"
"  float4 total_sum = (float4)0.0f;                                            \n"
"                                                                              \n"
"  if (in_x < in_width)                                                        \n"
"    {                                                                         \n"
"      column_index_start = y;                                                 \n"
"      column_index_end = y + twice_radius;                                    \n"
"                                                                              \n"
"      for (int i = 0; i < twice_radius + 1; ++i)                              \n"
"        {                                                                     \n"
"          tmp_sum += in[(y + i) * in_width + in_x];                           \n"
"        }                                                                     \n"
"                                                                              \n"
"      column_sum[local_id0] = tmp_sum;                                        \n"
"    }                                                                         \n"
"                                                                              \n"
"  barrier(CLK_LOCAL_MEM_FENCE);                                               \n"
"                                                                              \n"
"  while (1)                                                                   \n"
"    {                                                                         \n"
"      if (out_x < width)                                                      \n"
"        {                                                                     \n"
"          if (local_id0 >= radius && local_id0 < get_local_size(0) - radius)  \n"
"            {                                                                 \n"
"                total_sum = (float4)0.0f;                                     \n"
"                                                                              \n"
"                for (int i = 0; i < twice_radius + 1; ++i)                    \n"
"                  {                                                           \n"
"                    total_sum += column_sum[local_id0 - radius + i];          \n"
"                  }                                                           \n"
"                                                                              \n"
"                out[y * width + out_x] = total_sum / area;                    \n"
"            }                                                                 \n"
"        }                                                                     \n"
"                                                                              \n"
"    if (--tmp_size == 0 || y == height - 1)                                   \n"
"      break;                                                                  \n"
"                                                                              \n"
"    barrier(CLK_LOCAL_MEM_FENCE);                                             \n"
"                                                                              \n"
"    ++y;                                                                      \n"
"                                                                              \n"
"    if (in_x < in_width)                                                      \n"
"      {                                                                       \n"
"        tmp_sum = column_sum[local_id0];                                      \n"
"        tmp_sum -= in[(column_index_start)   * in_width + in_x];              \n"
"        tmp_sum += in[(column_index_end + 1) * in_width + in_x];              \n"
"        ++column_index_start;                                                 \n"
"        ++column_index_end;                                                   \n"
"        column_sum[local_id0] = tmp_sum;                                      \n"
"      }                                                                       \n"
"                                                                              \n"
"    barrier(CLK_LOCAL_MEM_FENCE);                                             \n"
"  }                                                                           \n"
"}                                                                             \n"
;