• Home
  • Features
  • Pricing
  • Docs
  • Announcements
  • Sign In

celerity / celerity-runtime / 11143617170

02 Oct 2024 12:18PM UTC coverage: 95.009% (-0.08%) from 95.087%
11143617170

Pull #288

github

web-flow
Merge 13df376db into b1cb3bbf5
Pull Request #288: Disambiguate `parallel_for(Integral, ...)`

3013 of 3414 branches covered (88.25%)

Branch coverage included in aggregate %.

13 of 13 new or added lines in 2 files covered. (100.0%)

7 existing lines in 2 files now uncovered.

6620 of 6725 relevant lines covered (98.44%)

1493366.96 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

97.25
/include/ranges.h
1
#pragma once
2

3
#include "sycl_wrappers.h"
4
#include "workaround.h"
5

6
namespace celerity {
7

8
// clang-format off
9
template <int Dims = 1> class range;
10
template <int Dims = 1> class id;
11
template <int Dims = 1> class nd_range;
12
template <int Dims = 1> struct subrange;
13
template <int Dims = 1> struct chunk;
14
// clang-format on
15

16
} // namespace celerity
17

18
namespace celerity::detail {
19

20
struct make_from_t {
21
} inline static constexpr make_from;
22

23
// We need a specialized storage type for coordinates to avoid generating a `size_t values[0]` array which clang interprets as dynamically-sized.
24
// By specializing on the Interface type, id<> and range<> become distinct types "all the way down", so both an id<0> and a range<0> can be included as struct
25
// members with [[no_unique_address]] within another struct and to actually overlap. This is required to ensure that 0-dimensional accessors are pointer-sized,
26
// and would otherwise be prohibited by strict-aliasing rules (because two identical pointers with the same type must point to the same object).
27
template <typename Interface, int Dims>
28
struct coordinate_storage {
29
        constexpr size_t operator[](int dimension) const {
1,346,222,299✔
30
                CELERITY_DETAIL_ASSERT_ON_HOST(dimension < Dims);
1,346,222,299✔
31
                return values[dimension];
1,346,222,299✔
32
        }
33

34
        constexpr size_t& operator[](int dimension) {
426,020,513✔
35
                CELERITY_DETAIL_ASSERT_ON_HOST(dimension < Dims);
426,020,513✔
36
                return values[dimension];
426,020,513✔
37
        }
38

39
        size_t values[Dims] = {};
40
};
41

42
template <typename Interface>
43
struct coordinate_storage<Interface, 0> {
44
        constexpr size_t operator[](int /* dimension */) const { return 0; }
×
45
        // This is UB, but also unreachable. Unfortunately we can't call __builtin_unreachable from a constexpr function.
46
        constexpr size_t& operator[](int /* dimension */) { return *static_cast<size_t*>(static_cast<void*>(this)); }
×
47
};
48

49
// We implement range and id from scratch to allow zero-dimensional structures.
50
template <typename Interface, int Dims>
51
class coordinate {
52
  public:
53
        constexpr static int dimensions = Dims;
54

55
        coordinate() = default;
43,978,513✔
56

57
        template <typename InterfaceIn, int DimsIn>
58
        constexpr coordinate(const make_from_t /* tag */, const coordinate<InterfaceIn, DimsIn>& other, const size_t default_value) {
45,009,093✔
59
                for(int d = 0; d < Dims; ++d) {
137,108,491✔
60
                        (*this)[d] = d < DimsIn ? other[d] : default_value;
92,099,734✔
61
                }
62
        }
45,008,757✔
63

64
        template <typename InterfaceIn>
65
        constexpr coordinate(const make_from_t /* tag */, const coordinate<InterfaceIn, Dims>& other) : coordinate(make_from, other, 0) {}
25,556,496✔
66

67
        template <typename... Values, typename = std::enable_if_t<sizeof...(Values) + 1 == Dims && (... && std::is_convertible_v<Values, size_t>)>>
68
        constexpr coordinate(const size_t dim_0, const Values... dim_n) : m_values{{dim_0, static_cast<size_t>(dim_n)...}} {}
33,813,538✔
69

70
        constexpr size_t get(int dimension) const { return m_values[dimension]; }
10✔
71
        constexpr size_t& operator[](int dimension) { return m_values[dimension]; }
426,026,315✔
72
        constexpr size_t operator[](int dimension) const { return m_values[dimension]; }
762,051,293✔
73

74
        friend constexpr bool operator==(const Interface& lhs, const Interface& rhs) {
29,389,076✔
75
                bool equal = true;
29,389,076✔
76
                for(int d = 0; d < Dims; ++d) {
88,390,267✔
77
                        equal &= lhs[d] == rhs[d];
59,001,191✔
78
                }
79
                return equal;
29,389,076✔
80
        }
81

82
        friend constexpr bool operator!=(const Interface& lhs, const Interface& rhs) { return !(lhs == rhs); }
1,116✔
83

84
#define CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(op)                                                                                         \
85
        friend constexpr Interface operator op(const Interface& lhs, const Interface& rhs) {                                                                       \
86
                Interface result;                                                                                                                                      \
87
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
88
                        result[d] = lhs.m_values[d] op rhs.m_values[d];                                                                                                    \
89
                }                                                                                                                                                      \
90
                return result;                                                                                                                                         \
91
        }                                                                                                                                                          \
92
        friend constexpr Interface operator op(const Interface& lhs, const size_t& rhs) {                                                                          \
93
                Interface result;                                                                                                                                      \
94
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
95
                        result[d] = lhs.m_values[d] op rhs;                                                                                                                \
96
                }                                                                                                                                                      \
97
                return result;                                                                                                                                         \
98
        }
99

100
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(+)
272,769,613✔
101
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(-)
106,017,187✔
102
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(*)
78✔
103
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(/)
99,994✔
104
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(%)
103,155✔
105
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(<<)
78✔
106
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(>>)
78✔
107
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(&)
78✔
108
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(|)
78✔
109
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(^)
26✔
110
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(&&)
898✔
111
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(||)
52!
112
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(<)
100,771,592✔
113
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(>)
78✔
114
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(<=)
69,598✔
115
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR(>=)
100,772,435✔
116

117
#undef CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_LHS_OPERATOR
118

119
#define CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(op)                                                                                      \
120
        friend constexpr Interface& operator op(Interface& lhs, const Interface& rhs) {                                                                            \
121
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
122
                        lhs.m_values[d] op rhs.m_values[d];                                                                                                                \
123
                }                                                                                                                                                      \
124
                return lhs;                                                                                                                                            \
125
        }                                                                                                                                                          \
126
        friend constexpr Interface& operator op(Interface& lhs, const size_t& rhs) {                                                                               \
127
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
128
                        lhs.m_values[d] op rhs;                                                                                                                            \
129
                }                                                                                                                                                      \
130
                return lhs;                                                                                                                                            \
131
        }
132

133
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(+=)
640✔
134
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(-=)
640✔
135
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(*=)
20✔
136
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(/=)
20✔
137
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(%=)
20✔
138
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(<<=)
20✔
139
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(>>=)
20✔
140
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(&=)
20✔
141
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(|=)
20✔
142
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR(^=)
20✔
143

144
#undef CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_INPLACE_LHS_OPERATOR
145

146
#define CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(op)                                                                                         \
147
        friend constexpr Interface operator op(const size_t& lhs, const Interface& rhs) {                                                                          \
148
                Interface result;                                                                                                                                      \
149
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
150
                        result[d] = lhs op rhs.m_values[d];                                                                                                                \
151
                }                                                                                                                                                      \
152
                return result;                                                                                                                                         \
153
        }
154

155
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(+)
26✔
156
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(-)
26✔
157
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(*)
35✔
158
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(/)
26✔
159
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(%)
26✔
160
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(<<)
26✔
161
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(>>)
26✔
162
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(&)
26✔
163
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(|)
26✔
164
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(^)
165
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(&&)
26!
166
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(||)
26!
167
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(<)
56✔
168
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(>)
26✔
169
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(<=)
26✔
170
        CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR(>=)
56✔
171

172
#undef CELERITY_DETAIL_DEFINE_COORDINATE_BINARY_COPY_RHS_OPERATOR
173

174
#define CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_COPY_OPERATOR(op)                                                                                              \
175
        friend constexpr Interface operator op(const Interface& rhs) {                                                                                             \
176
                Interface result;                                                                                                                                      \
177
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
178
                        result[d] = op rhs[d];                                                                                                                             \
179
                }                                                                                                                                                      \
180
                return result;                                                                                                                                         \
181
        }
182

183
        CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_COPY_OPERATOR(+)
26✔
184
        CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_COPY_OPERATOR(-)
26✔
185

186
#undef CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_COPY_OPERATOR
187

188
#define CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_PREFIX_OPERATOR(op)                                                                                            \
189
        friend constexpr Interface& operator op(Interface& rhs) {                                                                                                  \
190
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
191
                        op rhs[d];                                                                                                                                         \
192
                }                                                                                                                                                      \
193
                return rhs;                                                                                                                                            \
194
        }
195

196
        CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_PREFIX_OPERATOR(++)
20✔
197
        CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_PREFIX_OPERATOR(--)
20✔
198

199
#undef CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_PREFIX_OPERATOR
200

201
#define CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_POSTFIX_OPERATOR(op)                                                                                           \
202
        friend constexpr Interface operator op(Interface& lhs, int) {                                                                                              \
203
                Interface result = lhs;                                                                                                                                \
204
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
205
                        lhs[d] op;                                                                                                                                         \
206
                }                                                                                                                                                      \
207
                return result;                                                                                                                                         \
208
        }
209

210
        CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_POSTFIX_OPERATOR(++)
26✔
211
        CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_POSTFIX_OPERATOR(--)
26✔
212

213
#undef CELERITY_DETAIL_DEFINE_COORDINATE_UNARY_POSTFIX_OPERATOR
214

215
  private:
216
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS coordinate_storage<Interface, Dims> m_values;
217
};
218

219
template <int DimsOut, typename InterfaceIn>
220
range<DimsOut> range_cast(const InterfaceIn& in);
221

222
template <int DimsOut, typename InterfaceIn>
223
id<DimsOut> id_cast(const InterfaceIn& in);
224

225
struct zeros_t {
226
} inline static constexpr zeros;
227
struct ones_t {
228
} inline static constexpr ones;
229

230
}; // namespace celerity::detail
231

232
namespace celerity {
233

234
template <int Dims>
235
class range : public detail::coordinate<range<Dims>, Dims> {
236
  private:
237
        using coordinate = detail::coordinate<range<Dims>, Dims>;
238

239
  public:
240
        constexpr range() noexcept = default;
16✔
241

242
        template <typename... Values, typename = std::enable_if_t<sizeof...(Values) + 1 == Dims>>
243
        constexpr range(const size_t dim_0, const Values... dim_n) : coordinate(dim_0, dim_n...) {}
18,901✔
244

245
        constexpr range(const detail::zeros_t /* tag */) {}
65,175✔
246

247
        constexpr range(const detail::ones_t /* tag */) {
5,291✔
248
                for(int d = 0; d < Dims; ++d) {
20,295✔
249
                        (*this)[d] = 1;
15,004✔
250
                }
251
        }
5,291✔
252

253
        template <int D = Dims, typename = std::enable_if_t<D >= 1 && D <= 3>>
254
        range(const sycl::range<Dims>& sycl_range) {
49,824✔
255
                for(int d = 0; d < Dims; ++d) {
196,848✔
256
                        (*this)[d] = sycl_range[d];
147,024✔
257
                }
258
        }
49,824✔
259

260
        constexpr size_t size() const {
81,478✔
261
                size_t s = 1;
81,478✔
262
                for(int d = 0; d < Dims; ++d) {
297,950✔
263
                        s *= (*this)[d];
216,472✔
264
                }
265
                return s;
81,478✔
266
        }
267

268
        template <int D = Dims, typename = std::enable_if_t<D >= 1 && D <= 3>>
269
        operator sycl::range<Dims>() const {
388✔
270
                if constexpr(Dims == 1) {
271
                        return {(*this)[0]};
211✔
272
                } else if constexpr(Dims == 2) {
273
                        return {(*this)[0], (*this)[1]};
142✔
274
                } else {
275
                        return {(*this)[0], (*this)[1], (*this)[2]};
35✔
276
                }
277
        }
278

279
  private:
280
        template <int DimsOut, typename InterfaceIn>
281
        friend range<DimsOut> detail::range_cast(const InterfaceIn& in);
282

283
        template <typename InterfaceIn, int DimsIn>
284
        constexpr range(const detail::make_from_t /* tag */, const detail::coordinate<InterfaceIn, DimsIn>& in)
1,476,193✔
285
            : coordinate(detail::make_from, in, /* default_value= */ 1) {}
1,476,193✔
286
};
287

288
range()->range<0>;
289
range(size_t)->range<1>;
290
range(size_t, size_t)->range<2>;
291
range(size_t, size_t, size_t)->range<3>;
292

293
template <int Dims>
294
class id : public detail::coordinate<id<Dims>, Dims> {
295
  private:
296
        using coordinate = detail::coordinate<id<Dims>, Dims>;
297

298
  public:
299
        constexpr id() noexcept = default;
27,014,150✔
300

301
        template <typename... Values, typename = std::enable_if_t<sizeof...(Values) + 1 == Dims>>
302
        constexpr id(const size_t dim_0, const Values... dim_n) : coordinate(dim_0, dim_n...) {}
33,794,633✔
303

304
        constexpr id(const range<Dims>& range) : coordinate(detail::make_from, range) {}
25,556,505✔
305

306
        constexpr id(const detail::zeros_t /* tag */) {}
3,799✔
307

308
        constexpr id(const detail::ones_t /* tag */) {
140✔
309
                for(int d = 0; d < Dims; ++d) {
554✔
310
                        (*this)[d] = 1;
414✔
311
                }
312
        }
140✔
313

314
        template <int D = Dims, typename = std::enable_if_t<D >= 1 && D <= 3>>
315
        id(const sycl::id<Dims>& sycl_id) {
16,904,718✔
316
                for(int d = 0; d < Dims; ++d) {
50,790,075✔
317
                        (*this)[d] = sycl_id[d];
33,885,357✔
318
                }
319
        }
16,904,718✔
320

321
        template <int D = Dims, typename = std::enable_if_t<D >= 1 && D <= 3>>
322
        operator sycl::id<Dims>() const {
132✔
323
                if constexpr(Dims == 1) {
324
                        return {(*this)[0]};
132✔
325
                } else if constexpr(Dims == 2) {
326
                        return {(*this)[0], (*this)[1]};
×
327
                } else {
UNCOV
328
                        return {(*this)[0], (*this)[1], (*this)[2]};
×
329
                }
330
        }
331

332
  private:
333
        template <int DimsOut, typename InterfaceIn>
334
        friend id<DimsOut> detail::id_cast(const InterfaceIn& in);
335

336
        template <typename InterfaceIn, int DimsIn>
337
        constexpr id(const detail::make_from_t /* tag */, const detail::coordinate<InterfaceIn, DimsIn>& in)
17,976,411✔
338
            : coordinate(detail::make_from, in, /* default_value= */ 0) {}
17,976,411✔
339
};
340

341
id()->id<0>;
342
id(size_t)->id<1>;
343
id(size_t, size_t)->id<2>;
344
id(size_t, size_t, size_t)->id<3>;
345

346
// Celerity nd_range does not deprecate kernel offsets since we can support them at no additional cost in the distributed model.
347
template <int Dims>
348
class nd_range {
349
  public:
350
        template <int D = Dims, typename = std::enable_if_t<D == 0>>
351
        constexpr nd_range() noexcept {}
1✔
352

353
#pragma GCC diagnostic push
354
#pragma GCC diagnostic ignored "-Wdeprecated-declarations"
355
        nd_range(const sycl::nd_range<Dims>& s_range)
356
            : m_global_range(s_range.get_global_range()), m_local_range(s_range.get_local_range()), m_offset(s_range.get_offset()) {}
357
#pragma GCC diagnostic pop
358

359
        nd_range(const range<Dims>& global_range, const range<Dims>& local_range, const id<Dims>& offset = {})
54✔
360
            : m_global_range(global_range), m_local_range(local_range), m_offset(offset) {
54✔
361
#ifndef __SYCL_DEVICE_ONLY__
362
                for(int d = 0; d < Dims; ++d) {
146✔
363
                        if(local_range[d] == 0 || global_range[d] % local_range[d] != 0) { throw std::invalid_argument("global_range is not divisible by local_range"); }
98✔
364
                }
365
#endif
366
        }
48✔
367

368
        operator sycl::nd_range<Dims>() const { return sycl::nd_range<Dims>{m_global_range, m_local_range, m_offset}; }
369

370
        const range<Dims>& get_global_range() const { return m_global_range; }
49✔
371
        const range<Dims>& get_local_range() const { return m_local_range; }
49✔
372
        const range<Dims>& get_group_range() const { return m_global_range / m_local_range; }
373
        const id<Dims>& get_offset() const { return m_offset; }
49✔
374

375
        friend bool operator==(const nd_range& lhs, const nd_range& rhs) {
376
                return lhs.m_global_range == rhs.m_global_range && lhs.m_local_range == rhs.m_local_range && lhs.m_offset == rhs.m_offset;
377
        }
378

379
        friend bool operator!=(const nd_range& lhs, const nd_range& rhs) { return !(lhs == rhs); }
380

381
  private:
382
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS range<Dims> m_global_range;
383
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS range<Dims> m_local_range;
384
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS id<Dims> m_offset;
385
};
386

387
// Non-templated deduction guides allow construction of nd_range from range initializer lists like so: nd_range{{1, 2}, {3, 4}}
388
nd_range(range<1> global_range, range<1> local_range, id<1> offset)->nd_range<1>;
389
nd_range(range<1> global_range, range<1> local_range)->nd_range<1>;
390
nd_range(range<2> global_range, range<2> local_range, id<2> offset)->nd_range<2>;
391
nd_range(range<2> global_range, range<2> local_range)->nd_range<2>;
392
nd_range(range<3> global_range, range<3> local_range, id<3> offset)->nd_range<3>;
393
nd_range(range<3> global_range, range<3> local_range)->nd_range<3>;
394

395
} // namespace celerity
396

397
namespace celerity {
398
namespace detail {
399

400
        template <int TargetDims, typename Target, int SubscriptDim = 0>
401
        class subscript_proxy;
402

403
        template <int TargetDims, typename Target, int SubscriptDim>
404
        inline decltype(auto) subscript(Target& tgt, id<TargetDims> id, const size_t index) {
916✔
405
                static_assert(SubscriptDim < TargetDims);
406
                id[SubscriptDim] = index;
916✔
407
                if constexpr(SubscriptDim == TargetDims - 1) {
408
                        return tgt[std::as_const(id)];
625✔
409
                } else {
410
                        return subscript_proxy<TargetDims, Target, SubscriptDim + 1>{tgt, id};
291✔
411
                }
412
        }
413

414
        template <int TargetDims, typename Target>
415
        inline decltype(auto) subscript(Target& tgt, const size_t index) {
625✔
416
                return subscript<TargetDims, Target, 0>(tgt, id<TargetDims>{}, index);
1,102✔
417
        }
418

419
        template <int TargetDims, typename Target, int SubscriptDim>
420
        class subscript_proxy {
421
          public:
422
                subscript_proxy(Target& tgt, const id<TargetDims> id) : m_tgt(tgt), m_id(id) {}
291✔
423

424
                inline decltype(auto) operator[](const size_t index) const { //
291✔
425
                        return subscript<TargetDims, Target, SubscriptDim>(m_tgt, m_id, index);
291✔
426
                }
427

428
          private:
429
                Target& m_tgt;
430
                id<TargetDims> m_id{};
431
        };
432

433
        inline size_t get_linear_index(const range<0>& /* range */, const id<0>& /* index */) { return 0; }
105✔
434

435
        inline size_t get_linear_index(const range<1>& range, const id<1>& index) { return index[0]; }
4,942✔
436

437
        inline size_t get_linear_index(const range<2>& range, const id<2>& index) { return index[0] * range[1] + index[1]; }
50,355,911✔
438

439
        inline size_t get_linear_index(const range<3>& range, const id<3>& index) { return index[0] * range[1] * range[2] + index[1] * range[2] + index[2]; }
1,214,886✔
440

441
#define CELERITY_DETAIL_MAKE_COMPONENT_WISE_FN(name, coord, op)                                                                                                \
442
        template <int Dims>                                                                                                                                        \
443
        coord<Dims> name(const coord<Dims>& a, const coord<Dims>& b) {                                                                                             \
444
                auto result = a;                                                                                                                                       \
445
                for(int d = 0; d < Dims; ++d) {                                                                                                                        \
446
                        result[d] = op(result[d], b[d]);                                                                                                                   \
447
                }                                                                                                                                                      \
448
                return result;                                                                                                                                         \
449
        }
450

451
        CELERITY_DETAIL_MAKE_COMPONENT_WISE_FN(range_min, range, std::min)
452
        CELERITY_DETAIL_MAKE_COMPONENT_WISE_FN(range_max, range, std::max)
453
        CELERITY_DETAIL_MAKE_COMPONENT_WISE_FN(id_min, id, std::min)
1,299,658✔
454
        CELERITY_DETAIL_MAKE_COMPONENT_WISE_FN(id_max, id, std::max)
1,299,656✔
455

456
#undef CELERITY_DETAIL_MAKE_COMPONENT_WISE_FN
457

458
        template <typename Interface, int Dims>
459
        bool all_true(const coordinate<Interface, Dims>& bools) {
50,403,417✔
460
                for(int d = 0; d < Dims; ++d) {
151,274,825✔
461
                        CELERITY_DETAIL_ASSERT_ON_HOST(bools[d] == 0 || bools[d] == 1);
100,871,712✔
462
                        if(bools[d] == 0) return false;
100,871,712✔
463
                }
464
                return true;
50,403,113✔
465
        }
466

467
} // namespace detail
468

469
template <int Dims>
470
struct chunk {
471
        static constexpr int dimensions = Dims;
472

473
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS id<Dims> offset;
474
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS celerity::range<Dims> range = detail::zeros;
475
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS celerity::range<Dims> global_size = detail::zeros;
476

477
        constexpr chunk() = default;
4✔
478

479
        constexpr chunk(const id<Dims>& offset, const celerity::range<Dims>& range, const celerity::range<Dims>& global_size)
114,289✔
480
            : offset(offset), range(range), global_size(global_size) {}
114,289✔
481

482
        friend bool operator==(const chunk& lhs, const chunk& rhs) {
483
                return lhs.offset == rhs.offset && lhs.range == rhs.range && lhs.global_size == rhs.global_size;
484
        }
485
        friend bool operator!=(const chunk& lhs, const chunk& rhs) { return !operator==(lhs, rhs); }
486
};
487

488
template <int Dims>
489
struct subrange {
490
        static constexpr int dimensions = Dims;
491

492
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS id<Dims> offset;
493
        CELERITY_DETAIL_NO_UNIQUE_ADDRESS celerity::range<Dims> range = detail::zeros;
494

495
        constexpr subrange() = default;
28,940✔
496
        constexpr subrange(const id<Dims>& offset, const celerity::range<Dims>& range) : offset(offset), range(range) {}
375,380✔
497
        constexpr subrange(const chunk<Dims>& other) : offset(other.offset), range(other.range) {}
40,092✔
498

499
        friend bool operator==(const subrange& lhs, const subrange& rhs) { return lhs.offset == rhs.offset && lhs.range == rhs.range; }
1,126!
500
        friend bool operator!=(const subrange& lhs, const subrange& rhs) { return !operator==(lhs, rhs); }
501
};
502

503
} // namespace celerity
504

505
namespace celerity::detail {
506

507
/// Returns the smallest dimensionality that the range can be `range_cast` to.
508
template <int Dims>
509
int get_effective_dims(const range<Dims>& range) {
737,357✔
510
        for(int dims = Dims; dims > 0; --dims) {
1,647,523✔
511
                if(range[dims - 1] > 1) return dims;
1,317,546✔
512
        }
513
        return 0;
329,977✔
514
}
515

516
template <int DimsOut, typename InterfaceIn>
517
range<DimsOut> range_cast(const InterfaceIn& in) {
1,476,214✔
518
        CELERITY_DETAIL_ASSERT_ON_HOST(get_effective_dims(in) <= DimsOut);
1,476,214✔
519
        return range<DimsOut>(make_from, in);
1,476,196✔
520
}
521

522
/// Returns the smallest dimensionality that the id can be `id_cast` to.
523
template <int Dims>
524
int get_effective_dims(const id<Dims>& id) {
19,117,717✔
525
        for(int dims = Dims; dims > 0; --dims) {
20,292,988✔
526
                if(id[dims - 1] > 0) { return dims; }
19,957,975✔
527
        }
528
        return 0;
335,013✔
529
}
530

531
template <int DimsOut, typename InterfaceIn>
532
id<DimsOut> id_cast(const InterfaceIn& in) {
17,976,422✔
533
        CELERITY_DETAIL_ASSERT_ON_HOST(get_effective_dims(in) <= DimsOut);
17,976,422✔
534
        return id<DimsOut>(make_from, in);
17,976,412✔
535
}
536

537
/// Returns the smallest dimensionality that the chunk can be `chunk_cast` to.
538
template <int Dims>
539
int get_effective_dims(const chunk<Dims>& ck) {
46,777✔
540
        return std::max({get_effective_dims(ck.offset), get_effective_dims(ck.range), get_effective_dims(ck.global_size)});
46,777✔
541
}
542

543
/// Returns the smallest dimensionality that the subrange can be `subrange_cast` to.
544
template <int Dims>
545
int get_effective_dims(const subrange<Dims>& sr) {
205,519✔
546
        return std::max(get_effective_dims(sr.offset), get_effective_dims(sr.range));
205,519✔
547
}
548

549
template <int Dims, int OtherDims>
550
chunk<Dims> chunk_cast(const chunk<OtherDims>& other) {
46,777✔
551
        CELERITY_DETAIL_ASSERT_ON_HOST(get_effective_dims(other) <= Dims);
46,777✔
552
        return chunk{detail::id_cast<Dims>(other.offset), detail::range_cast<Dims>(other.range), detail::range_cast<Dims>(other.global_size)};
46,776✔
553
}
554

555
template <int Dims, int OtherDims>
556
subrange<Dims> subrange_cast(const subrange<OtherDims>& other) {
206,424✔
557
        CELERITY_DETAIL_ASSERT_ON_HOST(get_effective_dims(other) <= Dims);
206,424✔
558
        return subrange{detail::id_cast<Dims>(other.offset), detail::range_cast<Dims>(other.range)};
206,439✔
559
}
560

561
} // namespace celerity::detail
STATUS · Troubleshooting · Open an Issue · Sales · Support · CAREERS · ENTERPRISE · START FREE · SCHEDULE DEMO
ANNOUNCEMENTS · TWITTER · TOS & SLA · Supported CI Services · What's a CI service? · Automated Testing

© 2026 Coveralls, Inc