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