spla
Loading...
Searching...
No Matches
auto_map.hpp
Go to the documentation of this file.
1
2// Copyright (c) 2021 - 2023 SparseLinearAlgebra
3// Autogenerated file, do not modify
5
6#pragma once
7
8static 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
29uint 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
36uint 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
54uint 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)";