26 #define VATOMIC_INC16(histogram, win_pos) \ 28 atomic_inc(histogram + win_pos.s0); \ 29 atomic_inc(histogram + win_pos.s1); \ 30 atomic_inc(histogram + win_pos.s2); \ 31 atomic_inc(histogram + win_pos.s3); \ 32 atomic_inc(histogram + win_pos.s4); \ 33 atomic_inc(histogram + win_pos.s5); \ 34 atomic_inc(histogram + win_pos.s6); \ 35 atomic_inc(histogram + win_pos.s7); \ 36 atomic_inc(histogram + win_pos.s8); \ 37 atomic_inc(histogram + win_pos.s9); \ 38 atomic_inc(histogram + win_pos.sa); \ 39 atomic_inc(histogram + win_pos.sb); \ 40 atomic_inc(histogram + win_pos.sc); \ 41 atomic_inc(histogram + win_pos.sd); \ 42 atomic_inc(histogram + win_pos.se); \ 43 atomic_inc(histogram + win_pos.sf); \ 69 __local uint *histogram_local,
77 uint local_id_x = get_local_id(0);
81 if(num_bins > local_x_size)
85 histogram_local[i] = 0;
90 if(local_id_x <= num_bins)
92 histogram_local[local_id_x] = 0;
96 uint16 vals = convert_uint16(vload16(0, input_buffer.
ptr));
98 uint16 win_pos =
select(num_bins, ((vals - offset) * num_bins) / range, (vals >= offset && vals < offrange));
100 barrier(CLK_LOCAL_MEM_FENCE);
102 barrier(CLK_LOCAL_MEM_FENCE);
104 if(num_bins > local_x_size)
106 for(
int i = local_id_x; i < num_bins; i +=
local_x_size)
108 atomic_add(histogram + i, histogram_local[i]);
113 if(local_id_x <= num_bins)
115 atomic_add(histogram + local_id_x, histogram_local[local_id_x]);
150 uint val = (uint)(*input_buffer.
ptr);
152 uint win_pos = (val >=
offset) ? (((val - offset) * num_bins) / range) : 0;
154 if(val >= offset && (val < offrange))
156 atomic_inc(histogram + win_pos);
179 __local uint *histogram_local,
184 uint local_index = get_local_id(0);
189 histogram_local[i] = 0;
192 uint16 vals = convert_uint16(vload16(0, input_buffer.
ptr));
194 barrier(CLK_LOCAL_MEM_FENCE);
196 atomic_inc(histogram_local + vals.s0);
197 atomic_inc(histogram_local + vals.s1);
198 atomic_inc(histogram_local + vals.s2);
199 atomic_inc(histogram_local + vals.s3);
200 atomic_inc(histogram_local + vals.s4);
201 atomic_inc(histogram_local + vals.s5);
202 atomic_inc(histogram_local + vals.s6);
203 atomic_inc(histogram_local + vals.s7);
204 atomic_inc(histogram_local + vals.s8);
205 atomic_inc(histogram_local + vals.s9);
206 atomic_inc(histogram_local + vals.sa);
207 atomic_inc(histogram_local + vals.sb);
208 atomic_inc(histogram_local + vals.sc);
209 atomic_inc(histogram_local + vals.sd);
210 atomic_inc(histogram_local + vals.se);
211 atomic_inc(histogram_local + vals.sf);
213 barrier(CLK_LOCAL_MEM_FENCE);
217 atomic_add(histogram + i, histogram_local[i]);
242 atomic_inc(histogram + *input_buffer.
ptr);
__global uchar * offset(const Image *img, int x, int y)
Get the pointer position of a Image.
#define CONVERT_TO_IMAGE_STRUCT(name)
#define VATOMIC_INC16(histogram, win_pos)
#define IMAGE_DECLARATION(name)
__kernel void hist_local_kernel(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __local uint *histogram_local, __global uint *restrict histogram, uint num_bins, uint offset, uint range, uint offrange)
Calculate the histogram of an 8 bit grayscale image.
__kernel void hist_local_kernel_fixed(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __local uint *histogram_local, __global uint *restrict histogram)
Calculate the histogram of an 8 bit grayscale image with bin size of 256 and window size of 1...
__kernel void hist_border_kernel(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __global uint *restrict histogram, uint num_bins, uint offset, uint range, uint offrange)
Calculate the histogram of an 8 bit grayscale image's border.
SimpleTensor< T > range(SimpleTensor< T > &dst, float start, const size_t num_of_elements, float step)
Structure to hold Image information.
__global uchar * ptr
Pointer to the starting postion of the buffer.
SimpleTensor< T > select(const SimpleTensor< uint8_t > &c, const SimpleTensor< T > &x, const SimpleTensor< T > &y)
SimpleTensor< uint32_t > histogram(const SimpleTensor< T > &src, size_t num_bins, int32_t offset, uint32_t range)
constexpr unsigned int local_x_size
__kernel void hist_border_kernel_fixed(__global uchar *input_ptr, uint input_stride_x, uint input_step_x, uint input_stride_y, uint input_step_y, uint input_offset_first_element_in_bytes, __global uint *restrict histogram)
Calculate the histogram of an 8 bit grayscale image with bin size as 256 and window size as 1...