spla
auto_map.hpp
Go to the documentation of this file.
1 // Copyright (c) 2021 - 2023 SparseLinearAlgebra
3 // Autogenerated file, do not modify
5 
6 #pragma once
7 
8 static const char source_map[] = R"(
9 
10 
11 
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)
16 #endif
17 
18 #define SWAP_KEYS(x, y) \
19  uint tmp1 = x; \
20  x = y; \
21  y = tmp1;
22 
23 #define SWAP_VALUES(x, y) \
24  TYPE tmp2 = x; \
25  x = y; \
26  y = tmp2;
27 
28 // nearest power of two number greater equals n
29 uint ceil_to_pow2(uint n) {
30  uint r = 1;
31  while (r < n) r *= 2;
32  return r;
33 }
34 
35 // find first element in a sorted array such x <= element
36 uint lower_bound(const uint x,
37  uint first,
38  uint size,
39  __global const uint* array) {
40  while (size > 0) {
41  int step = size / 2;
42 
43  if (array[first + step] < x) {
44  first = first + step + 1;
45  size -= step + 1;
46  } else {
47  size = step;
48  }
49  }
50  return first;
51 }
52 
53 // find first element in a sorted array such x <= element
54 uint lower_bound_local(const uint x,
55  uint first,
56  uint size,
57  __local const uint* array) {
58  while (size > 0) {
59  int step = size / 2;
60 
61  if (array[first + step] < x) {
62  first = first + step + 1;
63  size -= step + 1;
64  } else {
65  size = step;
66  }
67  }
68  return first;
69 }
70 __kernel void map(__global TYPE* restrict g_r,
71  __global const TYPE* restrict g_v,
72  const uint N) {
73  const uint gid = get_global_id(0);
74  const uint gstride = get_global_size(0);
75 
76  for (uint i = gid; i < N; i += gstride) {
77  g_r[i] = OP_UNARY(g_v[i]);
78  }
79 }
80 )";