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
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
|
#pragma once
#include <iostream>
#include <thrust/execution_policy.h>
template <typename T, unsigned int N>
struct FixedVector
{
T data[N];
__host__ __device__
FixedVector()
{
for(unsigned int i = 0; i < N; i++)
data[i] = T();
}
__host__ __device__
FixedVector(T init)
{
for(unsigned int i = 0; i < N; i++)
data[i] = init;
}
__host__ __device__
FixedVector operator+(const FixedVector& bs) const
{
FixedVector output;
for(unsigned int i = 0; i < N; i++)
output.data[i] = data[i] + bs.data[i];
return output;
}
__host__ __device__
bool operator<(const FixedVector& bs) const
{
for(unsigned int i = 0; i < N; i++)
{
if(data[i] < bs.data[i])
return true;
else if(bs.data[i] < data[i])
return false;
}
return false;
}
__host__ __device__
bool operator==(const FixedVector& bs) const
{
for(unsigned int i = 0; i < N; i++)
{
if(!(data[i] == bs.data[i]))
return false;
}
return true;
}
};
template<typename Key, typename Value>
struct key_value
{
typedef Key key_type;
typedef Value value_type;
__host__ __device__
key_value(void)
: key(), value()
{}
__host__ __device__
key_value(key_type k, value_type v)
: key(k), value(v)
{}
__host__ __device__
bool operator<(const key_value &rhs) const
{
return key < rhs.key;
}
__host__ __device__
bool operator>(const key_value &rhs) const
{
return key > rhs.key;
}
__host__ __device__
bool operator==(const key_value &rhs) const
{
return key == rhs.key && value == rhs.value;
}
__host__ __device__
bool operator!=(const key_value &rhs) const
{
return !operator==(rhs);
}
friend std::ostream &operator<<(std::ostream &os, const key_value &kv)
{
return os << "(" << kv.key << ", " << kv.value << ")";
}
key_type key;
value_type value;
};
struct user_swappable
{
inline __host__ __device__
user_swappable(bool swapped = false)
: was_swapped(swapped)
{}
bool was_swapped;
};
inline __host__ __device__
bool operator==(const user_swappable &x, const user_swappable &y)
{
return x.was_swapped == y.was_swapped;
}
inline __host__ __device__
void swap(user_swappable &x, user_swappable &y)
{
x.was_swapped = true;
y.was_swapped = false;
}
class my_system : public thrust::device_execution_policy<my_system>
{
public:
my_system(int)
: correctly_dispatched(false),
num_copies(0)
{}
my_system(const my_system &other)
: correctly_dispatched(false),
num_copies(other.num_copies + 1)
{}
void validate_dispatch()
{
correctly_dispatched = (num_copies == 0);
}
bool is_valid()
{
return correctly_dispatched;
}
private:
bool correctly_dispatched;
// count the number of copies so that we can validate
// that dispatch does not introduce any
unsigned int num_copies;
// disallow default construction
my_system();
};
struct my_tag : thrust::device_execution_policy<my_tag> {};
namespace unittest
{
using thrust::detail::int8_t;
using thrust::detail::int16_t;
using thrust::detail::int32_t;
using thrust::detail::int64_t;
using thrust::detail::uint8_t;
using thrust::detail::uint16_t;
using thrust::detail::uint32_t;
using thrust::detail::uint64_t;
}
|
