Legion Runtime
Loading...
Searching...
No Matches
redop.h
Go to the documentation of this file.
1/* Copyright 2026 Stanford University, NVIDIA Corporation
2 *
3 * Licensed under the Apache License, Version 2.0 (the "License");
4 * you may not use this file except in compliance with the License.
5 * You may obtain a copy of the License at
6 *
7 * http://www.apache.org/licenses/LICENSE-2.0
8 *
9 * Unless required by applicable law or agreed to in writing, software
10 * distributed under the License is distributed on an "AS IS" BASIS,
11 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 * See the License for the specific language governing permissions and
13 * limitations under the License.
14 */
15
16#ifndef __LEGION_REDOP_H__
17#define __LEGION_REDOP_H__
18
24#include <atomic>
25#include <cstring>
26#include <limits>
27
28#include "legion/api/config.h"
29
30#ifdef LEGION_REDOP_COMPLEX
31#ifdef LEGION_REDOP_HALF
32#define COMPLEX_HALF
33#endif
34#include "mathtypes/complex.h"
35#endif
36
37#ifdef LEGION_REDOP_HALF
38#include "mathtypes/half.h"
39#endif
40
41#if defined(__HIP__) && defined(REALM_USE_HIP)
42#include "hip_cuda_compat/hip_cuda.h"
43#include "realm/hip/hiphijack_api.h"
44#endif
45
46namespace Legion {
47
48 template<typename T>
49 class SumReduction {
50 // Empty definition
51 // Specializations provided for each type
52 };
53
54 template<>
55 class SumReduction<bool> {
56 public:
57 typedef bool LHS;
58 typedef bool RHS;
59
60 static constexpr bool identity = false;
61 static constexpr int REDOP_ID = LEGION_REDOP_OR_BOOL;
62
63 template<bool EXCLUSIVE>
64 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
65 template<bool EXCLUSIVE>
66 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
67 };
68
69 template<>
70 class SumReduction<int8_t> {
71 public:
72 typedef int8_t LHS;
73 typedef int8_t RHS;
74
75 static constexpr int8_t identity = 0;
76 static constexpr int REDOP_ID = LEGION_REDOP_SUM_INT8;
77
78 template<bool EXCLUSIVE>
79 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
80 template<bool EXCLUSIVE>
81 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
82 };
83
84 template<>
85 class SumReduction<int16_t> {
86 public:
87 typedef int16_t LHS;
88 typedef int16_t RHS;
89
90 static constexpr int16_t identity = 0;
91 static constexpr int REDOP_ID = LEGION_REDOP_SUM_INT16;
92
93 template<bool EXCLUSIVE>
94 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
95 template<bool EXCLUSIVE>
96 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
97 };
98
99 template<>
100 class SumReduction<int32_t> {
101 public:
102 typedef int32_t LHS;
103 typedef int32_t RHS;
104
105 static constexpr int32_t identity = 0;
106 static constexpr int REDOP_ID = LEGION_REDOP_SUM_INT32;
107
108 template<bool EXCLUSIVE>
109 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
110 template<bool EXCLUSIVE>
111 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
112 };
113
114 template<>
115 class SumReduction<int64_t> {
116 public:
117 typedef int64_t LHS;
118 typedef int64_t RHS;
119
120 static constexpr int64_t identity = 0;
121 static constexpr int REDOP_ID = LEGION_REDOP_SUM_INT64;
122
123 template<bool EXCLUSIVE>
124 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
125 template<bool EXCLUSIVE>
126 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
127 };
128
129 template<>
130 class SumReduction<uint8_t> {
131 public:
132 typedef uint8_t LHS;
133 typedef uint8_t RHS;
134
135 static constexpr uint8_t identity = 0;
136 static constexpr int REDOP_ID = LEGION_REDOP_SUM_UINT8;
137
138 template<bool EXCLUSIVE>
139 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
140 template<bool EXCLUSIVE>
141 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
142 };
143
144 template<>
145 class SumReduction<uint16_t> {
146 public:
147 typedef uint16_t LHS;
148 typedef uint16_t RHS;
149
150 static constexpr uint16_t identity = 0;
151 static constexpr int REDOP_ID = LEGION_REDOP_SUM_UINT16;
152
153 template<bool EXCLUSIVE>
154 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
155 template<bool EXCLUSIVE>
156 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
157 };
158
159 template<>
160 class SumReduction<uint32_t> {
161 public:
162 typedef uint32_t LHS;
163 typedef uint32_t RHS;
164
165 static constexpr uint32_t identity = 0;
166 static constexpr int REDOP_ID = LEGION_REDOP_SUM_UINT32;
167
168 template<bool EXCLUSIVE>
169 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
170 template<bool EXCLUSIVE>
171 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
172 };
173
174 template<>
175 class SumReduction<uint64_t> {
176 public:
177 typedef uint64_t LHS;
178 typedef uint64_t RHS;
179
180 static constexpr uint64_t identity = 0;
181 static constexpr int REDOP_ID = LEGION_REDOP_SUM_UINT64;
182
183 template<bool EXCLUSIVE>
184 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
185 template<bool EXCLUSIVE>
186 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
187 };
188
189#ifdef LEGION_REDOP_HALF
190 template<>
191 class SumReduction<__half> {
192 public:
193 typedef __half LHS;
194 typedef __half RHS;
195
196 static inline const __half identity = __half(0, false /*raw*/);
197 static constexpr int REDOP_ID = LEGION_REDOP_SUM_FLOAT16;
198
199 template<bool EXCLUSIVE>
200 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
201 template<bool EXCLUSIVE>
202 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
203 };
204#endif
205
206 template<>
207 class SumReduction<float> {
208 public:
209 typedef float LHS;
210 typedef float RHS;
211
212 static constexpr float identity = 0.f;
213 static constexpr int REDOP_ID = LEGION_REDOP_SUM_FLOAT32;
214
215 template<bool EXCLUSIVE>
216 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
217 template<bool EXCLUSIVE>
218 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
219 };
220
221 template<>
222 class SumReduction<double> {
223 public:
224 typedef double LHS;
225 typedef double RHS;
226
227 static constexpr double identity = 0.0;
228 static constexpr int REDOP_ID = LEGION_REDOP_SUM_FLOAT64;
229
230 template<bool EXCLUSIVE>
231 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
232 template<bool EXCLUSIVE>
233 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
234 };
235
236#ifdef LEGION_REDOP_COMPLEX
237#ifdef LEGION_REDOP_HALF
238 template<>
239 class SumReduction<complex<__half> > {
240 public:
241 typedef complex<__half> LHS;
242 typedef complex<__half> RHS;
243
244 static inline const complex<__half> identity =
245 complex<__half>(__half(0, false /*raw*/), __half(0, false /*raw*/));
246 static constexpr int REDOP_ID = LEGION_REDOP_SUM_COMPLEX32;
247
248 template<bool EXCLUSIVE>
249 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
250 template<bool EXCLUSIVE>
251 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
252 };
253#endif // LEGION_REDOP_HALF
254
255 template<>
256 class SumReduction<complex<float> > {
257 public:
258 typedef complex<float> LHS;
259 typedef complex<float> RHS;
260
261 static inline const complex<float> identity = complex<float>(0.f, 0.f);
262 static constexpr int REDOP_ID = LEGION_REDOP_SUM_COMPLEX64;
263
264 template<bool EXCLUSIVE>
265 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
266 template<bool EXCLUSIVE>
267 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
268 };
269
270 // WARNING: This operator performs element-wise reductions on real and
271 // imaginary components, and thus has non-linearizable semantics.
272 // Users should be aware of this non-linearizability, which can
273 // lead to inconsistent results.
274 template<>
275 class SumReduction<complex<double> > {
276 public:
277 typedef complex<double> LHS;
278 typedef complex<double> RHS;
279
280 static inline const complex<double> identity = complex<double>(0.0, 0.0);
281 static constexpr int REDOP_ID = LEGION_REDOP_SUM_COMPLEX128;
282
283 template<bool EXCLUSIVE>
284 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
285 template<bool EXCLUSIVE>
286 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
287 };
288#endif // LEGION_REDOP_COMPLEX
289
290 template<typename T>
291 class DiffReduction {
292 // Empty definition
293 // Specializations provided for each type
294 };
295
296 template<>
297 class DiffReduction<int8_t> {
298 public:
299 typedef int8_t LHS;
300 typedef int8_t RHS;
301
302 static constexpr int8_t identity = 0;
303 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_INT8;
304
305 template<bool EXCLUSIVE>
306 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
307 template<bool EXCLUSIVE>
308 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
309 };
310
311 template<>
312 class DiffReduction<int16_t> {
313 public:
314 typedef int16_t LHS;
315 typedef int16_t RHS;
316
317 static constexpr int16_t identity = 0;
318 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_INT16;
319
320 template<bool EXCLUSIVE>
321 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
322 template<bool EXCLUSIVE>
323 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
324 };
325
326 template<>
327 class DiffReduction<int32_t> {
328 public:
329 typedef int32_t LHS;
330 typedef int32_t RHS;
331
332 static constexpr int32_t identity = 0;
333 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_INT32;
334
335 template<bool EXCLUSIVE>
336 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
337 template<bool EXCLUSIVE>
338 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
339 };
340
341 template<>
342 class DiffReduction<int64_t> {
343 public:
344 typedef int64_t LHS;
345 typedef int64_t RHS;
346
347 static constexpr int64_t identity = 0;
348 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_INT64;
349
350 template<bool EXCLUSIVE>
351 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
352 template<bool EXCLUSIVE>
353 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
354 };
355
356 template<>
357 class DiffReduction<uint8_t> {
358 public:
359 typedef uint8_t LHS;
360 typedef uint8_t RHS;
361
362 static constexpr uint8_t identity = 0;
363 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_UINT8;
364
365 template<bool EXCLUSIVE>
366 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
367 template<bool EXCLUSIVE>
368 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
369 };
370
371 template<>
372 class DiffReduction<uint16_t> {
373 public:
374 typedef uint16_t LHS;
375 typedef uint16_t RHS;
376
377 static constexpr uint16_t identity = 0;
378 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_UINT16;
379
380 template<bool EXCLUSIVE>
381 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
382 template<bool EXCLUSIVE>
383 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
384 };
385
386 template<>
387 class DiffReduction<uint32_t> {
388 public:
389 typedef uint32_t LHS;
390 typedef uint32_t RHS;
391
392 static constexpr uint32_t identity = 0;
393 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_UINT32;
394
395 template<bool EXCLUSIVE>
396 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
397 template<bool EXCLUSIVE>
398 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
399 };
400
401 template<>
402 class DiffReduction<uint64_t> {
403 public:
404 typedef uint64_t LHS;
405 typedef uint64_t RHS;
406
407 static constexpr uint64_t identity = 0;
408 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_UINT64;
409
410 template<bool EXCLUSIVE>
411 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
412 template<bool EXCLUSIVE>
413 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
414 };
415
416#ifdef LEGION_REDOP_HALF
417 template<>
418 class DiffReduction<__half> {
419 public:
420 typedef __half LHS;
421 typedef __half RHS;
422
423 static inline const __half identity = __half(0, false /*raw*/);
424 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_FLOAT16;
425
426 template<bool EXCLUSIVE>
427 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
428 template<bool EXCLUSIVE>
429 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
430 };
431#endif
432
433 template<>
434 class DiffReduction<float> {
435 public:
436 typedef float LHS;
437 typedef float RHS;
438
439 static constexpr float identity = 0.f;
440 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_FLOAT32;
441
442 template<bool EXCLUSIVE>
443 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
444 template<bool EXCLUSIVE>
445 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
446 };
447
448 template<>
449 class DiffReduction<double> {
450 public:
451 typedef double LHS;
452 typedef double RHS;
453
454 static constexpr double identity = 0.0;
455 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_FLOAT64;
456
457 template<bool EXCLUSIVE>
458 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
459 template<bool EXCLUSIVE>
460 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
461 };
462
463#ifdef LEGION_REDOP_COMPLEX
464#ifdef LEGION_REDOP_HALF
465 template<>
466 class DiffReduction<complex<__half> > {
467 public:
468 typedef complex<__half> LHS;
469 typedef complex<__half> RHS;
470
471 static inline const complex<__half> identity =
472 complex<__half>(__half(0, false /*raw*/), __half(0, false /*raw*/));
473 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_COMPLEX32;
474
475 template<bool EXCLUSIVE>
476 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
477 template<bool EXCLUSIVE>
478 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
479 };
480#endif // LEGION_REDOP_HALF
481
482 template<>
483 class DiffReduction<complex<float> > {
484 public:
485 typedef complex<float> LHS;
486 typedef complex<float> RHS;
487
488 static inline const complex<float> identity = complex<float>(0.f, 0.f);
489 static constexpr int REDOP_ID = LEGION_REDOP_DIFF_COMPLEX64;
490
491 template<bool EXCLUSIVE>
492 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
493 template<bool EXCLUSIVE>
494 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
495 };
496#endif // LEGION_REDOP_COMPLEX
497
498 template<typename T>
499 class ProdReduction {
500 // Empty definition
501 // Specializations provided for each type
502 };
503
504 template<>
505 class ProdReduction<bool> {
506 public:
507 typedef bool LHS;
508 typedef bool RHS;
509
510 static constexpr bool identity = true;
511 static constexpr int REDOP_ID = LEGION_REDOP_AND_BOOL;
512
513 template<bool EXCLUSIVE>
514 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
515 template<bool EXCLUSIVE>
516 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
517 };
518
519 template<>
520 class ProdReduction<int8_t> {
521 public:
522 typedef int8_t LHS;
523 typedef int8_t RHS;
524
525 static constexpr int8_t identity = 1;
526 static constexpr int REDOP_ID = LEGION_REDOP_PROD_INT8;
527
528 template<bool EXCLUSIVE>
529 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
530 template<bool EXCLUSIVE>
531 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
532 };
533
534 template<>
535 class ProdReduction<int16_t> {
536 public:
537 typedef int16_t LHS;
538 typedef int16_t RHS;
539
540 static constexpr int16_t identity = 1;
541 static constexpr int REDOP_ID = LEGION_REDOP_PROD_INT16;
542
543 template<bool EXCLUSIVE>
544 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
545 template<bool EXCLUSIVE>
546 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
547 };
548
549 template<>
550 class ProdReduction<int32_t> {
551 public:
552 typedef int32_t LHS;
553 typedef int32_t RHS;
554
555 static constexpr int32_t identity = 1;
556 static constexpr int REDOP_ID = LEGION_REDOP_PROD_INT32;
557
558 template<bool EXCLUSIVE>
559 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
560 template<bool EXCLUSIVE>
561 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
562 };
563
564 template<>
565 class ProdReduction<int64_t> {
566 public:
567 typedef int64_t LHS;
568 typedef int64_t RHS;
569
570 static constexpr int64_t identity = 1;
571 static constexpr int REDOP_ID = LEGION_REDOP_PROD_INT64;
572
573 template<bool EXCLUSIVE>
574 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
575 template<bool EXCLUSIVE>
576 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
577 };
578
579 template<>
580 class ProdReduction<uint8_t> {
581 public:
582 typedef uint8_t LHS;
583 typedef uint8_t RHS;
584
585 static constexpr uint8_t identity = 1;
586 static constexpr int REDOP_ID = LEGION_REDOP_PROD_UINT8;
587
588 template<bool EXCLUSIVE>
589 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
590 template<bool EXCLUSIVE>
591 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
592 };
593
594 template<>
595 class ProdReduction<uint16_t> {
596 public:
597 typedef uint16_t LHS;
598 typedef uint16_t RHS;
599
600 static constexpr uint16_t identity = 1;
601 static constexpr int REDOP_ID = LEGION_REDOP_PROD_UINT16;
602
603 template<bool EXCLUSIVE>
604 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
605 template<bool EXCLUSIVE>
606 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
607 };
608
609 template<>
610 class ProdReduction<uint32_t> {
611 public:
612 typedef uint32_t LHS;
613 typedef uint32_t RHS;
614
615 static constexpr uint32_t identity = 1;
616 static constexpr int REDOP_ID = LEGION_REDOP_PROD_UINT32;
617
618 template<bool EXCLUSIVE>
619 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
620 template<bool EXCLUSIVE>
621 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
622 };
623
624 template<>
625 class ProdReduction<uint64_t> {
626 public:
627 typedef uint64_t LHS;
628 typedef uint64_t RHS;
629
630 static constexpr uint64_t identity = 1;
631 static constexpr int REDOP_ID = LEGION_REDOP_PROD_UINT64;
632
633 template<bool EXCLUSIVE>
634 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
635 template<bool EXCLUSIVE>
636 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
637 };
638
639#ifdef LEGION_REDOP_HALF
640 template<>
641 class ProdReduction<__half> {
642 public:
643 typedef __half LHS;
644 typedef __half RHS;
645
646 static inline const __half identity = __half(1, false /*raw*/);
647 static constexpr int REDOP_ID = LEGION_REDOP_PROD_FLOAT16;
648
649 template<bool EXCLUSIVE>
650 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
651 template<bool EXCLUSIVE>
652 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
653 };
654#endif
655
656 template<>
657 class ProdReduction<float> {
658 public:
659 typedef float LHS;
660 typedef float RHS;
661
662 static constexpr float identity = 1.f;
663 static constexpr int REDOP_ID = LEGION_REDOP_PROD_FLOAT32;
664
665 template<bool EXCLUSIVE>
666 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
667 template<bool EXCLUSIVE>
668 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
669 };
670
671 template<>
672 class ProdReduction<double> {
673 public:
674 typedef double LHS;
675 typedef double RHS;
676
677 static constexpr double identity = 1.0;
678 static constexpr int REDOP_ID = LEGION_REDOP_PROD_FLOAT64;
679
680 template<bool EXCLUSIVE>
681 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
682 template<bool EXCLUSIVE>
683 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
684 };
685
686#ifdef LEGION_REDOP_COMPLEX
687#ifdef LEGION_REDOP_HALF
688 template<>
689 class ProdReduction<complex<__half> > {
690 public:
691 typedef complex<__half> LHS;
692 typedef complex<__half> RHS;
693
694 static inline const complex<__half> identity =
695 complex<__half>(__half(1, false /*raw*/), __half(0, false /*raw*/));
696 static constexpr int REDOP_ID = LEGION_REDOP_PROD_COMPLEX32;
697
698 template<bool EXCLUSIVE>
699 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
700 template<bool EXCLUSIVE>
701 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
702 };
703#endif // LEGION_REDOP_HALF
704
705 template<>
706 class ProdReduction<complex<float> > {
707 public:
708 typedef complex<float> LHS;
709 typedef complex<float> RHS;
710
711 static inline const complex<float> identity = complex<float>(1.f, 0.f);
712 static constexpr int REDOP_ID = LEGION_REDOP_PROD_COMPLEX64;
713
714 template<bool EXCLUSIVE>
715 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
716 template<bool EXCLUSIVE>
717 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
718 };
719#endif // LEGION_REDOP_COMPLEX
720
721 template<typename T>
722 class DivReduction {
723 // Empty definition
724 // Specializations provided for each type
725 };
726
727 template<>
728 class DivReduction<int8_t> {
729 public:
730 typedef int8_t LHS;
731 typedef int8_t RHS;
732
733 static constexpr int8_t identity = 1;
734 static constexpr int REDOP_ID = LEGION_REDOP_DIV_INT8;
735
736 template<bool EXCLUSIVE>
737 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
738 template<bool EXCLUSIVE>
739 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
740 };
741
742 template<>
743 class DivReduction<int16_t> {
744 public:
745 typedef int16_t LHS;
746 typedef int16_t RHS;
747
748 static constexpr int16_t identity = 1;
749 static constexpr int REDOP_ID = LEGION_REDOP_DIV_INT16;
750
751 template<bool EXCLUSIVE>
752 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
753 template<bool EXCLUSIVE>
754 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
755 };
756
757 template<>
758 class DivReduction<int32_t> {
759 public:
760 typedef int32_t LHS;
761 typedef int32_t RHS;
762
763 static constexpr int32_t identity = 1;
764 static constexpr int REDOP_ID = LEGION_REDOP_DIV_INT32;
765
766 template<bool EXCLUSIVE>
767 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
768 template<bool EXCLUSIVE>
769 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
770 };
771
772 template<>
773 class DivReduction<int64_t> {
774 public:
775 typedef int64_t LHS;
776 typedef int64_t RHS;
777
778 static constexpr int64_t identity = 1;
779 static constexpr int REDOP_ID = LEGION_REDOP_DIV_INT64;
780
781 template<bool EXCLUSIVE>
782 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
783 template<bool EXCLUSIVE>
784 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
785 };
786
787 template<>
788 class DivReduction<uint8_t> {
789 public:
790 typedef uint8_t LHS;
791 typedef uint8_t RHS;
792
793 static constexpr uint8_t identity = 1;
794 static constexpr int REDOP_ID = LEGION_REDOP_DIV_UINT8;
795
796 template<bool EXCLUSIVE>
797 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
798 template<bool EXCLUSIVE>
799 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
800 };
801
802 template<>
803 class DivReduction<uint16_t> {
804 public:
805 typedef uint16_t LHS;
806 typedef uint16_t RHS;
807
808 static constexpr uint16_t identity = 1;
809 static constexpr int REDOP_ID = LEGION_REDOP_DIV_UINT16;
810
811 template<bool EXCLUSIVE>
812 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
813 template<bool EXCLUSIVE>
814 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
815 };
816
817 template<>
818 class DivReduction<uint32_t> {
819 public:
820 typedef uint32_t LHS;
821 typedef uint32_t RHS;
822
823 static constexpr uint32_t identity = 1;
824 static constexpr int REDOP_ID = LEGION_REDOP_DIV_UINT32;
825
826 template<bool EXCLUSIVE>
827 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
828 template<bool EXCLUSIVE>
829 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
830 };
831
832 template<>
833 class DivReduction<uint64_t> {
834 public:
835 typedef uint64_t LHS;
836 typedef uint64_t RHS;
837
838 static constexpr uint64_t identity = 1;
839 static constexpr int REDOP_ID = LEGION_REDOP_DIV_UINT64;
840
841 template<bool EXCLUSIVE>
842 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
843 template<bool EXCLUSIVE>
844 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
845 };
846
847#ifdef LEGION_REDOP_HALF
848 template<>
849 class DivReduction<__half> {
850 public:
851 typedef __half LHS;
852 typedef __half RHS;
853
854 static inline const __half identity = __half(1, false /*raw*/);
855 static constexpr int REDOP_ID = LEGION_REDOP_DIV_FLOAT16;
856
857 template<bool EXCLUSIVE>
858 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
859 template<bool EXCLUSIVE>
860 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
861 };
862#endif
863
864 template<>
865 class DivReduction<float> {
866 public:
867 typedef float LHS;
868 typedef float RHS;
869
870 static constexpr float identity = 1.f;
871 static constexpr int REDOP_ID = LEGION_REDOP_DIV_FLOAT32;
872
873 template<bool EXCLUSIVE>
874 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
875 template<bool EXCLUSIVE>
876 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
877 };
878
879 template<>
880 class DivReduction<double> {
881 public:
882 typedef double LHS;
883 typedef double RHS;
884
885 static constexpr double identity = 1.0;
886 static constexpr int REDOP_ID = LEGION_REDOP_DIV_FLOAT64;
887
888 template<bool EXCLUSIVE>
889 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
890 template<bool EXCLUSIVE>
891 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
892 };
893
894#ifdef LEGION_REDOP_COMPLEX
895#ifdef LEGION_REDOP_HALF
896 template<>
897 class DivReduction<complex<__half> > {
898 public:
899 typedef complex<__half> LHS;
900 typedef complex<__half> RHS;
901
902 static inline const complex<__half> identity =
903 complex<__half>(__half(1, false /*raw*/), __half(0, false /*raw*/));
904 static constexpr int REDOP_ID = LEGION_REDOP_DIV_COMPLEX32;
905
906 template<bool EXCLUSIVE>
907 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
908 template<bool EXCLUSIVE>
909 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
910 };
911#endif // LEGION_REDOP_HALF
912
913 template<>
914 class DivReduction<complex<float> > {
915 public:
916 typedef complex<float> LHS;
917 typedef complex<float> RHS;
918
919 static inline const complex<float> identity = complex<float>(1.f, 0.f);
920 static constexpr int REDOP_ID = LEGION_REDOP_DIV_COMPLEX64;
921
922 template<bool EXCLUSIVE>
923 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
924 template<bool EXCLUSIVE>
925 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
926 };
927#endif // LEGION_REDOP_COMPLEX
928
929 template<typename T>
930 class MaxReduction {
931 // Empty definition
932 // Specializations provided for each type
933 };
934
935 template<>
936 class MaxReduction<bool> {
937 public:
938 typedef bool LHS;
939 typedef bool RHS;
940
941 static constexpr bool identity = false;
942 static constexpr int REDOP_ID = LEGION_REDOP_MAX_BOOL;
943
944 template<bool EXCLUSIVE>
945 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
946 template<bool EXCLUSIVE>
947 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
948 };
949
950 template<>
951 class MaxReduction<int8_t> {
952 public:
953 typedef int8_t LHS;
954 typedef int8_t RHS;
955
956 static constexpr int8_t identity = std::numeric_limits<int8_t>::min();
957 static constexpr int REDOP_ID = LEGION_REDOP_MAX_INT8;
958
959 template<bool EXCLUSIVE>
960 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
961 template<bool EXCLUSIVE>
962 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
963 };
964
965 template<>
966 class MaxReduction<int16_t> {
967 public:
968 typedef int16_t LHS;
969 typedef int16_t RHS;
970
971 static constexpr int16_t identity = std::numeric_limits<int16_t>::min();
972 static constexpr int REDOP_ID = LEGION_REDOP_MAX_INT16;
973
974 template<bool EXCLUSIVE>
975 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
976 template<bool EXCLUSIVE>
977 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
978 };
979
980 template<>
981 class MaxReduction<int32_t> {
982 public:
983 typedef int32_t LHS;
984 typedef int32_t RHS;
985
986 static constexpr int32_t identity = std::numeric_limits<int32_t>::min();
987 static constexpr int REDOP_ID = LEGION_REDOP_MAX_INT32;
988
989 template<bool EXCLUSIVE>
990 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
991 template<bool EXCLUSIVE>
992 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
993 };
994
995 template<>
996 class MaxReduction<int64_t> {
997 public:
998 typedef int64_t LHS;
999 typedef int64_t RHS;
1000
1001 static constexpr int64_t identity = std::numeric_limits<int64_t>::min();
1002 static constexpr int REDOP_ID = LEGION_REDOP_MAX_INT64;
1003
1004 template<bool EXCLUSIVE>
1005 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1006 template<bool EXCLUSIVE>
1007 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1008 };
1009
1010 template<>
1011 class MaxReduction<uint8_t> {
1012 public:
1013 typedef uint8_t LHS;
1014 typedef uint8_t RHS;
1015
1016 static constexpr uint8_t identity = 0;
1017 static constexpr int REDOP_ID = LEGION_REDOP_MAX_UINT8;
1018
1019 template<bool EXCLUSIVE>
1020 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1021 template<bool EXCLUSIVE>
1022 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1023 };
1024
1025 template<>
1026 class MaxReduction<uint16_t> {
1027 public:
1028 typedef uint16_t LHS;
1029 typedef uint16_t RHS;
1030
1031 static constexpr uint16_t identity = 0;
1032 static constexpr int REDOP_ID = LEGION_REDOP_MAX_UINT16;
1033
1034 template<bool EXCLUSIVE>
1035 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1036 template<bool EXCLUSIVE>
1037 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1038 };
1039
1040 template<>
1041 class MaxReduction<uint32_t> {
1042 public:
1043 typedef uint32_t LHS;
1044 typedef uint32_t RHS;
1045
1046 static constexpr uint32_t identity = 0;
1047 static constexpr int REDOP_ID = LEGION_REDOP_MAX_UINT32;
1048
1049 template<bool EXCLUSIVE>
1050 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1051 template<bool EXCLUSIVE>
1052 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1053 };
1054
1055 template<>
1056 class MaxReduction<uint64_t> {
1057 public:
1058 typedef uint64_t LHS;
1059 typedef uint64_t RHS;
1060
1061 static constexpr uint64_t identity = 0;
1062 static constexpr int REDOP_ID = LEGION_REDOP_MAX_UINT64;
1063
1064 template<bool EXCLUSIVE>
1065 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1066 template<bool EXCLUSIVE>
1067 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1068 };
1069
1070#ifdef LEGION_REDOP_HALF
1071 template<>
1072 class MaxReduction<__half> {
1073 public:
1074 typedef __half LHS;
1075 typedef __half RHS;
1076
1077 static inline const __half identity = __half(-2e10);
1078 static constexpr int REDOP_ID = LEGION_REDOP_MAX_FLOAT16;
1079
1080 template<bool EXCLUSIVE>
1081 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1082 template<bool EXCLUSIVE>
1083 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1084 };
1085#endif
1086
1087 template<>
1088 class MaxReduction<float> {
1089 public:
1090 typedef float LHS;
1091 typedef float RHS;
1092
1093 static constexpr float identity = std::numeric_limits<float>::lowest();
1094 static constexpr int REDOP_ID = LEGION_REDOP_MAX_FLOAT32;
1095
1096 template<bool EXCLUSIVE>
1097 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1098 template<bool EXCLUSIVE>
1099 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1100 };
1101
1102 template<>
1103 class MaxReduction<double> {
1104 public:
1105 typedef double LHS;
1106 typedef double RHS;
1107
1108 static constexpr double identity = std::numeric_limits<double>::lowest();
1109 static constexpr int REDOP_ID = LEGION_REDOP_MAX_FLOAT64;
1110
1111 template<bool EXCLUSIVE>
1112 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1113 template<bool EXCLUSIVE>
1114 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1115 };
1116
1117 template<typename T>
1118 class MinReduction {
1119 // Empty definition
1120 // Specializations provided for each type
1121 };
1122
1123 template<>
1124 class MinReduction<bool> {
1125 public:
1126 typedef bool LHS;
1127 typedef bool RHS;
1128
1129 static constexpr bool identity = true;
1130 static constexpr int REDOP_ID = LEGION_REDOP_MIN_BOOL;
1131
1132 template<bool EXCLUSIVE>
1133 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1134 template<bool EXCLUSIVE>
1135 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1136 };
1137
1138 template<>
1139 class MinReduction<int8_t> {
1140 public:
1141 typedef int8_t LHS;
1142 typedef int8_t RHS;
1143
1144 static constexpr int8_t identity = std::numeric_limits<int8_t>::max();
1145 static constexpr int REDOP_ID = LEGION_REDOP_MIN_INT8;
1146
1147 template<bool EXCLUSIVE>
1148 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1149 template<bool EXCLUSIVE>
1150 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1151 };
1152
1153 template<>
1154 class MinReduction<int16_t> {
1155 public:
1156 typedef int16_t LHS;
1157 typedef int16_t RHS;
1158
1159 static constexpr int16_t identity = std::numeric_limits<int16_t>::max();
1160 static constexpr int REDOP_ID = LEGION_REDOP_MIN_INT16;
1161
1162 template<bool EXCLUSIVE>
1163 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1164 template<bool EXCLUSIVE>
1165 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1166 };
1167
1168 template<>
1169 class MinReduction<int32_t> {
1170 public:
1171 typedef int32_t LHS;
1172 typedef int32_t RHS;
1173
1174 static constexpr int32_t identity = std::numeric_limits<int32_t>::max();
1175 static constexpr int REDOP_ID = LEGION_REDOP_MIN_INT32;
1176
1177 template<bool EXCLUSIVE>
1178 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1179 template<bool EXCLUSIVE>
1180 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1181 };
1182
1183 template<>
1184 class MinReduction<int64_t> {
1185 public:
1186 typedef int64_t LHS;
1187 typedef int64_t RHS;
1188
1189 static constexpr int64_t identity = std::numeric_limits<int64_t>::max();
1190 static constexpr int REDOP_ID = LEGION_REDOP_MIN_INT64;
1191
1192 template<bool EXCLUSIVE>
1193 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1194 template<bool EXCLUSIVE>
1195 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1196 };
1197
1198 template<>
1199 class MinReduction<uint8_t> {
1200 public:
1201 typedef uint8_t LHS;
1202 typedef uint8_t RHS;
1203
1204 static constexpr uint8_t identity = std::numeric_limits<uint8_t>::max();
1205 static constexpr int REDOP_ID = LEGION_REDOP_MIN_UINT8;
1206
1207 template<bool EXCLUSIVE>
1208 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1209 template<bool EXCLUSIVE>
1210 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1211 };
1212
1213 template<>
1214 class MinReduction<uint16_t> {
1215 public:
1216 typedef uint16_t LHS;
1217 typedef uint16_t RHS;
1218
1219 static constexpr uint16_t identity = std::numeric_limits<uint16_t>::max();
1220 static constexpr int REDOP_ID = LEGION_REDOP_MIN_UINT16;
1221
1222 template<bool EXCLUSIVE>
1223 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1224 template<bool EXCLUSIVE>
1225 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1226 };
1227
1228 template<>
1229 class MinReduction<uint32_t> {
1230 public:
1231 typedef uint32_t LHS;
1232 typedef uint32_t RHS;
1233
1234 static constexpr uint32_t identity = std::numeric_limits<uint32_t>::max();
1235 static constexpr int REDOP_ID = LEGION_REDOP_MIN_UINT32;
1236
1237 template<bool EXCLUSIVE>
1238 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1239 template<bool EXCLUSIVE>
1240 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1241 };
1242
1243 template<>
1244 class MinReduction<uint64_t> {
1245 public:
1246 typedef uint64_t LHS;
1247 typedef uint64_t RHS;
1248
1249 static constexpr uint64_t identity = std::numeric_limits<uint64_t>::max();
1250 static constexpr int REDOP_ID = LEGION_REDOP_MIN_UINT64;
1251
1252 template<bool EXCLUSIVE>
1253 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1254 template<bool EXCLUSIVE>
1255 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1256 };
1257
1258#ifdef LEGION_REDOP_HALF
1259 template<>
1260 class MinReduction<__half> {
1261 public:
1262 typedef __half LHS;
1263 typedef __half RHS;
1264
1265 static inline const __half identity = __half(2e10);
1266 static constexpr int REDOP_ID = LEGION_REDOP_MIN_FLOAT16;
1267
1268 template<bool EXCLUSIVE>
1269 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1270 template<bool EXCLUSIVE>
1271 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1272 };
1273#endif
1274
1275 template<>
1276 class MinReduction<float> {
1277 public:
1278 typedef float LHS;
1279 typedef float RHS;
1280
1281 static constexpr float identity = std::numeric_limits<float>::max();
1282 static constexpr int REDOP_ID = LEGION_REDOP_MIN_FLOAT32;
1283
1284 template<bool EXCLUSIVE>
1285 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1286 template<bool EXCLUSIVE>
1287 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1288 };
1289
1290 template<>
1291 class MinReduction<double> {
1292 public:
1293 typedef double LHS;
1294 typedef double RHS;
1295
1296 static constexpr double identity = std::numeric_limits<double>::max();
1297 static constexpr int REDOP_ID = LEGION_REDOP_MIN_FLOAT64;
1298
1299 template<bool EXCLUSIVE>
1300 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1301 template<bool EXCLUSIVE>
1302 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1303 };
1304
1305 template<typename T>
1306 class OrReduction {
1307 // Empty definition
1308 // Specializations provided for each type
1309 };
1310
1311 template<>
1312 class OrReduction<int8_t> {
1313 public:
1314 typedef int8_t LHS;
1315 typedef int8_t RHS;
1316
1317 static constexpr int8_t identity = 0;
1318 static constexpr int REDOP_ID = LEGION_REDOP_OR_INT8;
1319
1320 template<bool EXCLUSIVE>
1321 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1322 template<bool EXCLUSIVE>
1323 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1324 };
1325
1326 template<>
1327 class OrReduction<int16_t> {
1328 public:
1329 typedef int16_t LHS;
1330 typedef int16_t RHS;
1331
1332 static constexpr int16_t identity = 0;
1333 static constexpr int REDOP_ID = LEGION_REDOP_OR_INT16;
1334
1335 template<bool EXCLUSIVE>
1336 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1337 template<bool EXCLUSIVE>
1338 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1339 };
1340
1341 template<>
1342 class OrReduction<int32_t> {
1343 public:
1344 typedef int32_t LHS;
1345 typedef int32_t RHS;
1346
1347 static constexpr int32_t identity = 0;
1348 static constexpr int REDOP_ID = LEGION_REDOP_OR_INT32;
1349
1350 template<bool EXCLUSIVE>
1351 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1352 template<bool EXCLUSIVE>
1353 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1354 };
1355
1356 template<>
1357 class OrReduction<int64_t> {
1358 public:
1359 typedef int64_t LHS;
1360 typedef int64_t RHS;
1361
1362 static constexpr int64_t identity = 0;
1363 static constexpr int REDOP_ID = LEGION_REDOP_OR_INT64;
1364
1365 template<bool EXCLUSIVE>
1366 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1367 template<bool EXCLUSIVE>
1368 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1369 };
1370
1371 template<>
1372 class OrReduction<uint8_t> {
1373 public:
1374 typedef uint8_t LHS;
1375 typedef uint8_t RHS;
1376
1377 static constexpr uint8_t identity = 0;
1378 static constexpr int REDOP_ID = LEGION_REDOP_OR_UINT8;
1379
1380 template<bool EXCLUSIVE>
1381 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1382 template<bool EXCLUSIVE>
1383 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1384 };
1385
1386 template<>
1387 class OrReduction<uint16_t> {
1388 public:
1389 typedef uint16_t LHS;
1390 typedef uint16_t RHS;
1391
1392 static constexpr uint16_t identity = 0;
1393 static constexpr int REDOP_ID = LEGION_REDOP_OR_UINT16;
1394
1395 template<bool EXCLUSIVE>
1396 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1397 template<bool EXCLUSIVE>
1398 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1399 };
1400
1401 template<>
1402 class OrReduction<uint32_t> {
1403 public:
1404 typedef uint32_t LHS;
1405 typedef uint32_t RHS;
1406
1407 static constexpr uint32_t identity = 0;
1408 static constexpr int REDOP_ID = LEGION_REDOP_OR_UINT32;
1409
1410 template<bool EXCLUSIVE>
1411 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1412 template<bool EXCLUSIVE>
1413 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1414 };
1415
1416 template<>
1417 class OrReduction<uint64_t> {
1418 public:
1419 typedef uint64_t LHS;
1420 typedef uint64_t RHS;
1421
1422 static constexpr uint64_t identity = 0;
1423 static constexpr int REDOP_ID = LEGION_REDOP_OR_UINT64;
1424
1425 template<bool EXCLUSIVE>
1426 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1427 template<bool EXCLUSIVE>
1428 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1429 };
1430
1431 template<typename T>
1432 class AndReduction {
1433 // Empty definition
1434 // Specializations provided for each type
1435 };
1436
1437 template<>
1438 class AndReduction<int8_t> {
1439 public:
1440 typedef int8_t LHS;
1441 typedef int8_t RHS;
1442
1443 static constexpr int8_t identity = (int8_t)0xFF;
1444 static constexpr int REDOP_ID = LEGION_REDOP_AND_INT8;
1445
1446 template<bool EXCLUSIVE>
1447 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1448 template<bool EXCLUSIVE>
1449 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1450 };
1451
1452 template<>
1453 class AndReduction<int16_t> {
1454 public:
1455 typedef int16_t LHS;
1456 typedef int16_t RHS;
1457
1458 static constexpr int16_t identity = (int16_t)0xFFFF;
1459 static constexpr int REDOP_ID = LEGION_REDOP_AND_INT16;
1460
1461 template<bool EXCLUSIVE>
1462 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1463 template<bool EXCLUSIVE>
1464 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1465 };
1466
1467 template<>
1468 class AndReduction<int32_t> {
1469 public:
1470 typedef int32_t LHS;
1471 typedef int32_t RHS;
1472
1473 static constexpr int32_t identity = 0xFFFF;
1474 static constexpr int REDOP_ID = LEGION_REDOP_AND_INT32;
1475
1476 template<bool EXCLUSIVE>
1477 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1478 template<bool EXCLUSIVE>
1479 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1480 };
1481
1482 template<>
1483 class AndReduction<int64_t> {
1484 public:
1485 typedef int64_t LHS;
1486 typedef int64_t RHS;
1487
1488 static constexpr int64_t identity = 0xFFFFFFFFL;
1489 static constexpr int REDOP_ID = LEGION_REDOP_AND_INT64;
1490
1491 template<bool EXCLUSIVE>
1492 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1493 template<bool EXCLUSIVE>
1494 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1495 };
1496
1497 template<>
1498 class AndReduction<uint8_t> {
1499 public:
1500 typedef uint8_t LHS;
1501 typedef uint8_t RHS;
1502
1503 static constexpr uint8_t identity = 0xFFU;
1504 static constexpr int REDOP_ID = LEGION_REDOP_AND_UINT8;
1505
1506 template<bool EXCLUSIVE>
1507 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1508 template<bool EXCLUSIVE>
1509 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1510 };
1511
1512 template<>
1513 class AndReduction<uint16_t> {
1514 public:
1515 typedef uint16_t LHS;
1516 typedef uint16_t RHS;
1517
1518 static constexpr uint16_t identity = 0xFFFFU;
1519 static constexpr int REDOP_ID = LEGION_REDOP_AND_UINT16;
1520
1521 template<bool EXCLUSIVE>
1522 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1523 template<bool EXCLUSIVE>
1524 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1525 };
1526
1527 template<>
1528 class AndReduction<uint32_t> {
1529 public:
1530 typedef uint32_t LHS;
1531 typedef uint32_t RHS;
1532
1533 static constexpr uint32_t identity = 0xFFFFFFFFU;
1534 static constexpr int REDOP_ID = LEGION_REDOP_AND_UINT32;
1535
1536 template<bool EXCLUSIVE>
1537 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1538 template<bool EXCLUSIVE>
1539 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1540 };
1541
1542 template<>
1543 class AndReduction<uint64_t> {
1544 public:
1545 typedef uint64_t LHS;
1546 typedef uint64_t RHS;
1547
1548 static constexpr uint64_t identity = 0xFFFFFFFFUL;
1549 static constexpr int REDOP_ID = LEGION_REDOP_AND_UINT64;
1550
1551 template<bool EXCLUSIVE>
1552 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1553 template<bool EXCLUSIVE>
1554 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1555 };
1556
1557 template<typename T>
1558 class XorReduction {
1559 // Empty definition
1560 // Specializations provided for each type
1561 };
1562
1563 template<>
1564 class XorReduction<bool> {
1565 public:
1566 typedef bool LHS;
1567 typedef bool RHS;
1568
1569 static constexpr bool identity = false;
1570 static constexpr int REDOP_ID = LEGION_REDOP_XOR_BOOL;
1571
1572 template<bool EXCLUSIVE>
1573 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1574 template<bool EXCLUSIVE>
1575 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1576 };
1577
1578 template<>
1579 class XorReduction<int8_t> {
1580 public:
1581 typedef int8_t LHS;
1582 typedef int8_t RHS;
1583
1584 static constexpr int8_t identity = 0;
1585 static constexpr int REDOP_ID = LEGION_REDOP_XOR_INT8;
1586
1587 template<bool EXCLUSIVE>
1588 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1589 template<bool EXCLUSIVE>
1590 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1591 };
1592
1593 template<>
1594 class XorReduction<int16_t> {
1595 public:
1596 typedef int16_t LHS;
1597 typedef int16_t RHS;
1598
1599 static constexpr int16_t identity = 0;
1600 static constexpr int REDOP_ID = LEGION_REDOP_XOR_INT16;
1601
1602 template<bool EXCLUSIVE>
1603 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1604 template<bool EXCLUSIVE>
1605 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1606 };
1607
1608 template<>
1609 class XorReduction<int32_t> {
1610 public:
1611 typedef int32_t LHS;
1612 typedef int32_t RHS;
1613
1614 static constexpr int32_t identity = 0;
1615 static constexpr int REDOP_ID = LEGION_REDOP_XOR_INT32;
1616
1617 template<bool EXCLUSIVE>
1618 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1619 template<bool EXCLUSIVE>
1620 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1621 };
1622
1623 template<>
1624 class XorReduction<int64_t> {
1625 public:
1626 typedef int64_t LHS;
1627 typedef int64_t RHS;
1628
1629 static constexpr int64_t identity = 0;
1630 static constexpr int REDOP_ID = LEGION_REDOP_XOR_INT64;
1631
1632 template<bool EXCLUSIVE>
1633 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1634 template<bool EXCLUSIVE>
1635 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1636 };
1637
1638 template<>
1639 class XorReduction<uint8_t> {
1640 public:
1641 typedef uint8_t LHS;
1642 typedef uint8_t RHS;
1643
1644 static constexpr uint8_t identity = 0;
1645 static constexpr int REDOP_ID = LEGION_REDOP_XOR_UINT8;
1646
1647 template<bool EXCLUSIVE>
1648 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1649 template<bool EXCLUSIVE>
1650 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1651 };
1652
1653 template<>
1654 class XorReduction<uint16_t> {
1655 public:
1656 typedef uint16_t LHS;
1657 typedef uint16_t RHS;
1658
1659 static constexpr uint16_t identity = 0;
1660 static constexpr int REDOP_ID = LEGION_REDOP_XOR_UINT16;
1661
1662 template<bool EXCLUSIVE>
1663 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1664 template<bool EXCLUSIVE>
1665 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1666 };
1667
1668 template<>
1669 class XorReduction<uint32_t> {
1670 public:
1671 typedef uint32_t LHS;
1672 typedef uint32_t RHS;
1673
1674 static constexpr uint32_t identity = 0;
1675 static constexpr int REDOP_ID = LEGION_REDOP_XOR_UINT32;
1676
1677 template<bool EXCLUSIVE>
1678 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1679 template<bool EXCLUSIVE>
1680 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1681 };
1682
1683 template<>
1684 class XorReduction<uint64_t> {
1685 public:
1686 typedef uint64_t LHS;
1687 typedef uint64_t RHS;
1688
1689 static constexpr uint64_t identity = 0;
1690 static constexpr int REDOP_ID = LEGION_REDOP_XOR_UINT64;
1691
1692 template<bool EXCLUSIVE>
1693 __LEGION_CUDA_HD__ static void apply(LHS& lhs, RHS rhs);
1694 template<bool EXCLUSIVE>
1695 __LEGION_CUDA_HD__ static void fold(RHS& rhs1, RHS rhs2);
1696 };
1697
1698 // Depending on configuration, we may instantiate these reduction ops in a
1699 // one or more different ways, so define a helper macro that lists all the
1700 // ops to be instantiated
1701 // clang-format off
1702#define LEGION_REDOP_LIST_BASE(__op__) \
1703 /* Sum Reductions */ \
1704 __op__(LEGION_REDOP_SUM_INT8, SumReduction<int8_t>) \
1705 __op__(LEGION_REDOP_SUM_INT16, SumReduction<int16_t>) \
1706 __op__(LEGION_REDOP_SUM_INT32, SumReduction<int32_t>) \
1707 __op__(LEGION_REDOP_SUM_INT64, SumReduction<int64_t>) \
1708 __op__(LEGION_REDOP_SUM_UINT8, SumReduction<uint8_t>) \
1709 __op__(LEGION_REDOP_SUM_UINT16, SumReduction<uint16_t>) \
1710 __op__(LEGION_REDOP_SUM_UINT32, SumReduction<uint32_t>) \
1711 __op__(LEGION_REDOP_SUM_UINT64, SumReduction<uint64_t>) \
1712 __op__(LEGION_REDOP_SUM_FLOAT32, SumReduction<float>) \
1713 __op__(LEGION_REDOP_SUM_FLOAT64, SumReduction<double>) \
1714 /* Difference Reductions */ \
1715 __op__(LEGION_REDOP_DIFF_INT8, DiffReduction<int8_t>) \
1716 __op__(LEGION_REDOP_DIFF_INT16, DiffReduction<int16_t>) \
1717 __op__(LEGION_REDOP_DIFF_INT32, DiffReduction<int32_t>) \
1718 __op__(LEGION_REDOP_DIFF_INT64, DiffReduction<int64_t>) \
1719 __op__(LEGION_REDOP_DIFF_UINT8, DiffReduction<uint8_t>) \
1720 __op__(LEGION_REDOP_DIFF_UINT16, DiffReduction<uint16_t>) \
1721 __op__(LEGION_REDOP_DIFF_UINT32, DiffReduction<uint32_t>) \
1722 __op__(LEGION_REDOP_DIFF_UINT64, DiffReduction<uint64_t>) \
1723 __op__(LEGION_REDOP_DIFF_FLOAT32, DiffReduction<float>) \
1724 __op__(LEGION_REDOP_DIFF_FLOAT64, DiffReduction<double>) \
1725 /* Product Reductions */ \
1726 __op__(LEGION_REDOP_PROD_INT8, ProdReduction<int8_t>) \
1727 __op__(LEGION_REDOP_PROD_INT16, ProdReduction<int16_t>) \
1728 __op__(LEGION_REDOP_PROD_INT32, ProdReduction<int32_t>) \
1729 __op__(LEGION_REDOP_PROD_INT64, ProdReduction<int64_t>) \
1730 __op__(LEGION_REDOP_PROD_UINT8, ProdReduction<uint8_t>) \
1731 __op__(LEGION_REDOP_PROD_UINT16, ProdReduction<uint16_t>) \
1732 __op__(LEGION_REDOP_PROD_UINT32, ProdReduction<uint32_t>) \
1733 __op__(LEGION_REDOP_PROD_UINT64, ProdReduction<uint64_t>) \
1734 __op__(LEGION_REDOP_PROD_FLOAT32, ProdReduction<float>) \
1735 __op__(LEGION_REDOP_PROD_FLOAT64, ProdReduction<double>) \
1736 /* Divide Reductions */ \
1737 __op__(LEGION_REDOP_DIV_INT8, DivReduction<int8_t>) \
1738 __op__(LEGION_REDOP_DIV_INT16, DivReduction<int16_t>) \
1739 __op__(LEGION_REDOP_DIV_INT32, DivReduction<int32_t>) \
1740 __op__(LEGION_REDOP_DIV_INT64, DivReduction<int64_t>) \
1741 __op__(LEGION_REDOP_DIV_UINT8, DivReduction<uint8_t>) \
1742 __op__(LEGION_REDOP_DIV_UINT16, DivReduction<uint16_t>) \
1743 __op__(LEGION_REDOP_DIV_UINT32, DivReduction<uint32_t>) \
1744 __op__(LEGION_REDOP_DIV_UINT64, DivReduction<uint64_t>) \
1745 __op__(LEGION_REDOP_DIV_FLOAT32, DivReduction<float>) \
1746 __op__(LEGION_REDOP_DIV_FLOAT64, DivReduction<double>) \
1747 /* Max Reductions */ \
1748 __op__(LEGION_REDOP_MAX_BOOL, MaxReduction<bool>) \
1749 __op__(LEGION_REDOP_MAX_INT8, MaxReduction<int8_t>) \
1750 __op__(LEGION_REDOP_MAX_INT16, MaxReduction<int16_t>) \
1751 __op__(LEGION_REDOP_MAX_INT32, MaxReduction<int32_t>) \
1752 __op__(LEGION_REDOP_MAX_INT64, MaxReduction<int64_t>) \
1753 __op__(LEGION_REDOP_MAX_UINT8, MaxReduction<uint8_t>) \
1754 __op__(LEGION_REDOP_MAX_UINT16, MaxReduction<uint16_t>) \
1755 __op__(LEGION_REDOP_MAX_UINT32, MaxReduction<uint32_t>) \
1756 __op__(LEGION_REDOP_MAX_UINT64, MaxReduction<uint64_t>) \
1757 __op__(LEGION_REDOP_MAX_FLOAT32, MaxReduction<float>) \
1758 __op__(LEGION_REDOP_MAX_FLOAT64, MaxReduction<double>) \
1759 /* Min Reductions */ \
1760 __op__(LEGION_REDOP_MIN_BOOL, MinReduction<bool>) \
1761 __op__(LEGION_REDOP_MIN_INT8, MinReduction<int8_t>) \
1762 __op__(LEGION_REDOP_MIN_INT16, MinReduction<int16_t>) \
1763 __op__(LEGION_REDOP_MIN_INT32, MinReduction<int32_t>) \
1764 __op__(LEGION_REDOP_MIN_INT64, MinReduction<int64_t>) \
1765 __op__(LEGION_REDOP_MIN_UINT8, MinReduction<uint8_t>) \
1766 __op__(LEGION_REDOP_MIN_UINT16, MinReduction<uint16_t>) \
1767 __op__(LEGION_REDOP_MIN_UINT32, MinReduction<uint32_t>) \
1768 __op__(LEGION_REDOP_MIN_UINT64, MinReduction<uint64_t>) \
1769 __op__(LEGION_REDOP_MIN_FLOAT32, MinReduction<float>) \
1770 __op__(LEGION_REDOP_MIN_FLOAT64, MinReduction<double>) \
1771 /* Bitwise-OR Reductions */ \
1772 __op__(LEGION_REDOP_OR_BOOL, SumReduction<bool>) \
1773 __op__(LEGION_REDOP_OR_INT8, OrReduction<int8_t>) \
1774 __op__(LEGION_REDOP_OR_INT16, OrReduction<int16_t>) \
1775 __op__(LEGION_REDOP_OR_INT32, OrReduction<int32_t>) \
1776 __op__(LEGION_REDOP_OR_INT64, OrReduction<int64_t>) \
1777 __op__(LEGION_REDOP_OR_UINT8, OrReduction<uint8_t>) \
1778 __op__(LEGION_REDOP_OR_UINT16, OrReduction<uint16_t>) \
1779 __op__(LEGION_REDOP_OR_UINT32, OrReduction<uint32_t>) \
1780 __op__(LEGION_REDOP_OR_UINT64, OrReduction<uint64_t>) \
1781 /* Bitwise-AND Reductions */ \
1782 __op__(LEGION_REDOP_AND_BOOL, ProdReduction<bool>) \
1783 __op__(LEGION_REDOP_AND_INT8, AndReduction<int8_t>) \
1784 __op__(LEGION_REDOP_AND_INT16, AndReduction<int16_t>) \
1785 __op__(LEGION_REDOP_AND_INT32, AndReduction<int32_t>) \
1786 __op__(LEGION_REDOP_AND_INT64, AndReduction<int64_t>) \
1787 __op__(LEGION_REDOP_AND_UINT8, AndReduction<uint8_t>) \
1788 __op__(LEGION_REDOP_AND_UINT16, AndReduction<uint16_t>) \
1789 __op__(LEGION_REDOP_AND_UINT32, AndReduction<uint32_t>) \
1790 __op__(LEGION_REDOP_AND_UINT64, AndReduction<uint64_t>) \
1791 /* Bitwise-XOR Reductions */ \
1792 __op__(LEGION_REDOP_XOR_BOOL, XorReduction<bool>) \
1793 __op__(LEGION_REDOP_XOR_INT8, XorReduction<int8_t>) \
1794 __op__(LEGION_REDOP_XOR_INT16, XorReduction<int16_t>) \
1795 __op__(LEGION_REDOP_XOR_INT32, XorReduction<int32_t>) \
1796 __op__(LEGION_REDOP_XOR_INT64, XorReduction<int64_t>) \
1797 __op__(LEGION_REDOP_XOR_UINT8, XorReduction<uint8_t>) \
1798 __op__(LEGION_REDOP_XOR_UINT16, XorReduction<uint16_t>) \
1799 __op__(LEGION_REDOP_XOR_UINT32, XorReduction<uint32_t>) \
1800 __op__(LEGION_REDOP_XOR_UINT64, XorReduction<uint64_t>)
1801
1802#ifdef LEGION_REDOP_HALF
1803 #define LEGION_REDOP_LIST_HALF(__op__) \
1804 __op__(LEGION_REDOP_SUM_FLOAT16, SumReduction<__half>) \
1805 __op__(LEGION_REDOP_DIFF_FLOAT16, DiffReduction<__half>) \
1806 __op__(LEGION_REDOP_PROD_FLOAT16, ProdReduction<__half>) \
1807 __op__(LEGION_REDOP_DIV_FLOAT16, DivReduction<__half>) \
1808 __op__(LEGION_REDOP_MAX_FLOAT16, MaxReduction<__half>) \
1809 __op__(LEGION_REDOP_MIN_FLOAT16, MinReduction<__half>)
1810#else
1811 #define LEGION_REDOP_LIST_HALF(__op__)
1812#endif
1813#ifdef LEGION_REDOP_COMPLEX
1814 #define LEGION_REDOP_LIST_COMPLEX(__op__) \
1815 __op__(LEGION_REDOP_SUM_COMPLEX64, SumReduction<complex<float> >) \
1816 __op__(LEGION_REDOP_DIFF_COMPLEX64, DiffReduction<complex<float> >) \
1817 __op__(LEGION_REDOP_PROD_COMPLEX64, ProdReduction<complex<float> >) \
1818 __op__(LEGION_REDOP_DIV_COMPLEX64, DivReduction<complex<float> >) \
1819 __op__(LEGION_REDOP_SUM_COMPLEX128, SumReduction<complex<double> >)
1820
1821 #ifdef LEGION_REDOP_HALF
1822 #define LEGION_REDOP_LIST_HALF_COMPLEX(__op__) \
1823 __op__(LEGION_REDOP_SUM_COMPLEX32, SumReduction<complex<__half> >) \
1824 __op__(LEGION_REDOP_DIFF_COMPLEX32, DiffReduction<complex<__half> >) \
1825 __op__(LEGION_REDOP_PROD_COMPLEX32, ProdReduction<complex<__half> >) \
1826 __op__(LEGION_REDOP_DIV_COMPLEX32, DivReduction<complex<__half> >)
1827 #else
1828 #define LEGION_REDOP_LIST_HALF_COMPLEX(__op__)
1829 #endif
1830#else
1831 #define LEGION_REDOP_LIST_COMPLEX(__op__)
1832 #define LEGION_REDOP_LIST_HALF_COMPLEX(__op__)
1833#endif
1834
1835#define LEGION_REDOP_LIST(__op__) \
1836 LEGION_REDOP_LIST_BASE(__op__) \
1837 LEGION_REDOP_LIST_HALF(__op__) \
1838 LEGION_REDOP_LIST_COMPLEX(__op__) \
1839 LEGION_REDOP_LIST_HALF_COMPLEX(__op__)
1840
1841 // clang-format on
1842
1843} // namespace Legion
1844
1845#include "legion/api/redop.inl"
1846
1847#endif // __LEGION_REDOP_H__
Legion::AndReduction
Definition redop.h:1432
Legion::DiffReduction
Definition redop.h:291
Legion::DivReduction
Definition redop.h:722
Legion::MaxReduction
Definition redop.h:930
Legion::MinReduction
Definition redop.h:1118
Legion::OrReduction
Definition redop.h:1306
Legion::ProdReduction
Definition redop.h:499
Legion::SumReduction
Definition redop.h:49
Legion::XorReduction
Definition redop.h:1558
Legion
Definition accessors.h:29