Go to the documentation of this file.
8 static const char source_map[] = R
"(
12 // memory bank conflict-free address and local buffer size
13 #ifdef LM_NUM_MEM_BANKS
14 #define LM_ADDR(address) (address + ((address) / LM_NUM_MEM_BANKS))
15 #define LM_SIZE(size) (size + (size) / LM_NUM_MEM_BANKS)
18 #define SWAP_KEYS(x, y) \
23 #define SWAP_VALUES(x, y) \
28 // nearest power of two number greater equals n
29 uint ceil_to_pow2(uint n) {
35 // find first element in a sorted array such x <= element
36 uint lower_bound(const uint x,
39 __global const uint* array) {
43 if (array[first + step] < x) {
44 first = first + step + 1;
53 // find first element in a sorted array such x <= element
54 uint lower_bound_local(const uint x,
57 __local const uint* array) {
61 if (array[first + step] < x) {
62 first = first + step + 1;
70 __kernel void map(__global TYPE* restrict g_r,
71 __global const TYPE* restrict g_v,
73 const uint gid = get_global_id(0);
74 const uint gstride = get_global_size(0);
76 for (uint i = gid; i < N; i += gstride) {
77 g_r[i] = OP_UNARY(g_v[i]);