Actual source code: cupmatomics.hpp
1: #pragma once
3: /*====================================================================================*/
4: /* Atomic operations on device */
5: /*====================================================================================*/
6: #include <petscdevice_cupm.h>
7: #include <petscsystypes.h>
9: /* In terms of function overloading, long long int is a different type than int64_t, which PetscInt might be defined to.
10: We prefer long long int over PetscInt (int64_t), since CUDA atomics are built around (unsigned) long long int.
11: */
12: typedef long long int llint;
13: typedef unsigned long long int ullint;
15: #if PetscDefined(USING_NVCC)
16: PETSC_PRAGMA_DIAGNOSTIC_IGNORED_BEGIN("-Wunused-function")
17: /*
18: Atomic Insert (exchange) operations
20: CUDA C Programming Guide V10.1 Chapter B.12.1.3:
22: int atomicExch(int* address, int val);
23: unsigned int atomicExch(unsigned int* address, unsigned int val);
24: unsigned long long int atomicExch(unsigned long long int* address, unsigned long long int val);
25: float atomicExch(float* address, float val);
27: reads the 32-bit or 64-bit word old located at the address in global or shared
28: memory and stores val back to memory at the same address. These two operations are
29: performed in one atomic transaction. The function returns old.
31: PETSc notes:
33: It may be useful in PetscSFFetchAndOp with op = MPI_REPLACE.
35: VecScatter with multiple entries scattered to the same location using INSERT_VALUES does not need
36: atomic insertion, since it does not need the old value. A 32-bit or 64-bit store instruction should
37: be atomic itself.
39: With bs>1 and a unit > 64-bits, the current element-wise atomic approach can not guarantee the whole
40: insertion is atomic. Hope no user codes rely on that.
41: */
42: __device__ static double atomicExch(double *address, double val)
43: {
44: return __longlong_as_double(atomicExch((ullint *)address, __double_as_longlong(val)));
45: }
47: __device__ static llint atomicExch(llint *address, llint val)
48: {
49: return (llint)(atomicExch((ullint *)address, (ullint)val));
50: }
52: template <typename Type>
53: struct AtomicInsert {
54: __device__ Type operator()(Type &x, Type y) const { return atomicExch(&x, y); }
55: };
57: #if PetscDefined(HAVE_COMPLEX)
58: #if PetscDefined(USE_REAL_DOUBLE)
59: /* CUDA does not support 128-bit atomics. Users should not insert different 128-bit PetscComplex values to the same location */
60: template <>
61: struct AtomicInsert<PetscComplex> {
62: __device__ PetscComplex operator()(PetscComplex &x, PetscComplex y) const
63: {
64: PetscComplex old, *z = &old;
65: double *xp = (double *)&x, *yp = (double *)&y;
66: AtomicInsert<double> op;
67: z[0] = op(xp[0], yp[0]);
68: z[1] = op(xp[1], yp[1]);
69: return old; /* The returned value may not be atomic. It can be mix of two ops. Caller should discard it. */
70: }
71: };
72: #elif PetscDefined(USE_REAL_SINGLE)
73: template <>
74: struct AtomicInsert<PetscComplex> {
75: __device__ PetscComplex operator()(PetscComplex &x, PetscComplex y) const
76: {
77: double *xp = (double *)&x, *yp = (double *)&y;
78: AtomicInsert<double> op;
79: return op(xp[0], yp[0]);
80: }
81: };
82: #endif
83: #endif
85: /*
86: Atomic add operations
88: CUDA C Programming Guide V10.1 Chapter B.12.1.1:
90: int atomicAdd(int* address, int val);
91: unsigned int atomicAdd(unsigned int* address,unsigned int val);
92: unsigned long long int atomicAdd(unsigned long long int* address,unsigned long long int val);
93: float atomicAdd(float* address, float val);
94: double atomicAdd(double* address, double val);
95: __half2 atomicAdd(__half2 *address, __half2 val);
96: __half atomicAdd(__half *address, __half val);
98: reads the 16-bit, 32-bit or 64-bit word old located at the address in global or shared memory, computes (old + val),
99: and stores the result back to memory at the same address. These three operations are performed in one atomic transaction. The
100: function returns old.
102: The 32-bit floating-point version of atomicAdd() is only supported by devices of compute capability 2.x and higher.
103: The 64-bit floating-point version of atomicAdd() is only supported by devices of compute capability 6.x and higher.
104: The 32-bit __half2 floating-point version of atomicAdd() is only supported by devices of compute capability 6.x and
105: higher. The atomicity of the __half2 add operation is guaranteed separately for each of the two __half elements;
106: the entire __half2 is not guaranteed to be atomic as a single 32-bit access.
107: The 16-bit __half floating-point version of atomicAdd() is only supported by devices of compute capability 7.x and higher.
108: */
109: __device__ static llint atomicAdd(llint *address, llint val)
110: {
111: return (llint)atomicAdd((ullint *)address, (ullint)val);
112: }
114: template <typename Type>
115: struct AtomicAdd {
116: __device__ Type operator()(Type &x, Type y) const { return atomicAdd(&x, y); }
117: };
119: template <>
120: struct AtomicAdd<double> {
121: __device__ double operator()(double &x, double y) const
122: {
123: #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 600)
124: return atomicAdd(&x, y);
125: #else
126: double *address = &x, val = y;
127: ullint *address_as_ull = (ullint *)address;
128: ullint old = *address_as_ull, assumed;
129: do {
130: assumed = old;
131: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val + __longlong_as_double(assumed)));
132: /* Note: uses integer comparison to avoid hang in case of NaN (since NaN !=NaN) */
133: } while (assumed != old);
134: return __longlong_as_double(old);
135: #endif
136: }
137: };
139: template <>
140: struct AtomicAdd<float> {
141: __device__ float operator()(float &x, float y) const
142: {
143: #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 200)
144: return atomicAdd(&x, y);
145: #else
146: float *address = &x, val = y;
147: int *address_as_int = (int *)address;
148: int old = *address_as_int, assumed;
149: do {
150: assumed = old;
151: old = atomicCAS(address_as_int, assumed, __float_as_int(val + __int_as_float(assumed)));
152: /* Note: uses integer comparison to avoid hang in case of NaN (since NaN !=NaN) */
153: } while (assumed != old);
154: return __int_as_float(old);
155: #endif
156: }
157: };
159: #if PetscDefined(HAVE_COMPLEX)
160: template <>
161: struct AtomicAdd<PetscComplex> {
162: __device__ PetscComplex operator()(PetscComplex &x, PetscComplex y) const
163: {
164: PetscComplex old, *z = &old;
165: PetscReal *xp = (PetscReal *)&x, *yp = (PetscReal *)&y;
166: AtomicAdd<PetscReal> op;
167: z[0] = op(xp[0], yp[0]);
168: z[1] = op(xp[1], yp[1]);
169: return old; /* The returned value may not be atomic. It can be mix of two ops. Caller should discard it. */
170: }
171: };
172: #endif
174: /*
175: Atomic Mult operations:
177: CUDA has no atomicMult at all, so we build our own with atomicCAS
178: */
179: #if PetscDefined(USE_REAL_DOUBLE)
180: __device__ static double atomicMult(double *address, double val)
181: {
182: ullint *address_as_ull = (ullint *)(address);
183: ullint old = *address_as_ull, assumed;
184: do {
185: assumed = old;
186: /* Other threads can access and modify value of *address_as_ull after the read above and before the write below */
187: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val * __longlong_as_double(assumed)));
188: } while (assumed != old);
189: return __longlong_as_double(old);
190: }
191: #elif PetscDefined(USE_REAL_SINGLE)
192: __device__ static float atomicMult(float *address, float val)
193: {
194: int *address_as_int = (int *)(address);
195: int old = *address_as_int, assumed;
196: do {
197: assumed = old;
198: old = atomicCAS(address_as_int, assumed, __float_as_int(val * __int_as_float(assumed)));
199: } while (assumed != old);
200: return __int_as_float(old);
201: }
202: #endif
204: __device__ static int atomicMult(int *address, int val)
205: {
206: int *address_as_int = (int *)(address);
207: int old = *address_as_int, assumed;
208: do {
209: assumed = old;
210: old = atomicCAS(address_as_int, assumed, val * assumed);
211: } while (assumed != old);
212: return (int)old;
213: }
215: __device__ static llint atomicMult(llint *address, llint val)
216: {
217: ullint *address_as_ull = (ullint *)(address);
218: ullint old = *address_as_ull, assumed;
219: do {
220: assumed = old;
221: old = atomicCAS(address_as_ull, assumed, (ullint)(val * (llint)assumed));
222: } while (assumed != old);
223: return (llint)old;
224: }
226: template <typename Type>
227: struct AtomicMult {
228: __device__ Type operator()(Type &x, Type y) const { return atomicMult(&x, y); }
229: };
231: /*
232: Atomic Min/Max operations
234: CUDA C Programming Guide V10.1 Chapter B.12.1.4~5:
236: int atomicMin(int* address, int val);
237: unsigned int atomicMin(unsigned int* address,unsigned int val);
238: unsigned long long int atomicMin(unsigned long long int* address,unsigned long long int val);
240: reads the 32-bit or 64-bit word old located at the address in global or shared
241: memory, computes the minimum of old and val, and stores the result back to memory
242: at the same address. These three operations are performed in one atomic transaction.
243: The function returns old.
244: The 64-bit version of atomicMin() is only supported by devices of compute capability 3.5 and higher.
246: atomicMax() is similar.
247: */
249: #if PetscDefined(HAVE_HIPCUDA)
250: #define CUPM_HAVE_HIP_ATOMIC_MINMAX PETSC_PKG_HIP_VERSION_GE(5, 3, 0)
251: #else
252: #define CUPM_HAVE_HIP_ATOMIC_MINMAX 0
253: #endif
255: // HIP's NVIDIA backend provides these floating-point atomics starting in 5.3.
256: #if !CUPM_HAVE_HIP_ATOMIC_MINMAX
257: #if PetscDefined(USE_REAL_DOUBLE)
258: __device__ static double atomicMin(double *address, double val)
259: {
260: ullint *address_as_ull = (ullint *)(address);
261: ullint old = *address_as_ull, assumed;
262: do {
263: assumed = old;
264: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(PetscMin(val, __longlong_as_double(assumed))));
265: } while (assumed != old);
266: return __longlong_as_double(old);
267: }
269: __device__ static double atomicMax(double *address, double val)
270: {
271: ullint *address_as_ull = (ullint *)(address);
272: ullint old = *address_as_ull, assumed;
273: do {
274: assumed = old;
275: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(PetscMax(val, __longlong_as_double(assumed))));
276: } while (assumed != old);
277: return __longlong_as_double(old);
278: }
279: #elif PetscDefined(USE_REAL_SINGLE)
280: __device__ static float atomicMin(float *address, float val)
281: {
282: int *address_as_int = (int *)(address);
283: int old = *address_as_int, assumed;
284: do {
285: assumed = old;
286: old = atomicCAS(address_as_int, assumed, __float_as_int(PetscMin(val, __int_as_float(assumed))));
287: } while (assumed != old);
288: return __int_as_float(old);
289: }
291: __device__ static float atomicMax(float *address, float val)
292: {
293: int *address_as_int = (int *)(address);
294: int old = *address_as_int, assumed;
295: do {
296: assumed = old;
297: old = atomicCAS(address_as_int, assumed, __float_as_int(PetscMax(val, __int_as_float(assumed))));
298: } while (assumed != old);
299: return __int_as_float(old);
300: }
301: #endif
302: #endif
303: #undef CUPM_HAVE_HIP_ATOMIC_MINMAX
305: /*
306: atomicMin/Max(long long *, long long) are not in Nvidia's documentation. But on OLCF Summit we found
307: atomicMin/Max/And/Or/Xor(long long *, long long) in /sw/summit/cuda/10.1.243/include/sm_32_atomic_functions.h.
308: This causes compilation errors with pgi compilers and 64-bit indices:
309: error: function "atomicMin(long long *, long long)" has already been defined
311: So we add extra conditions defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 320)
312: */
313: #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 320)
314: __device__ static llint atomicMin(llint *address, llint val)
315: {
316: ullint *address_as_ull = (ullint *)(address);
317: ullint old = *address_as_ull, assumed;
318: do {
319: assumed = old;
320: old = atomicCAS(address_as_ull, assumed, (ullint)(PetscMin(val, (llint)assumed)));
321: } while (assumed != old);
322: return (llint)old;
323: }
325: __device__ static llint atomicMax(llint *address, llint val)
326: {
327: ullint *address_as_ull = (ullint *)(address);
328: ullint old = *address_as_ull, assumed;
329: do {
330: assumed = old;
331: old = atomicCAS(address_as_ull, assumed, (ullint)(PetscMax(val, (llint)assumed)));
332: } while (assumed != old);
333: return (llint)old;
334: }
335: #endif
337: template <typename Type>
338: struct AtomicMin {
339: __device__ Type operator()(Type &x, Type y) const { return atomicMin(&x, y); }
340: };
341: template <typename Type>
342: struct AtomicMax {
343: __device__ Type operator()(Type &x, Type y) const { return atomicMax(&x, y); }
344: };
346: /*
347: Atomic bitwise operations
349: CUDA C Programming Guide V10.1 Chapter B.12.2.1 ~ B.12.2.3:
351: int atomicAnd(int* address, int val);
352: unsigned int atomicAnd(unsigned int* address,unsigned int val);
353: unsigned long long int atomicAnd(unsigned long long int* address,unsigned long long int val);
355: reads the 32-bit or 64-bit word old located at the address in global or shared
356: memory, computes (old & val), and stores the result back to memory at the same
357: address. These three operations are performed in one atomic transaction.
358: The function returns old.
360: The 64-bit version of atomicAnd() is only supported by devices of compute capability 3.5 and higher.
362: atomicOr() and atomicXor are similar.
363: */
365: #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 320) /* Why 320? see comments at atomicMin() above */
366: __device__ static llint atomicAnd(llint *address, llint val)
367: {
368: ullint *address_as_ull = (ullint *)(address);
369: ullint old = *address_as_ull, assumed;
370: do {
371: assumed = old;
372: old = atomicCAS(address_as_ull, assumed, (ullint)(val & (llint)assumed));
373: } while (assumed != old);
374: return (llint)old;
375: }
376: __device__ static llint atomicOr(llint *address, llint val)
377: {
378: ullint *address_as_ull = (ullint *)(address);
379: ullint old = *address_as_ull, assumed;
380: do {
381: assumed = old;
382: old = atomicCAS(address_as_ull, assumed, (ullint)(val | (llint)assumed));
383: } while (assumed != old);
384: return (llint)old;
385: }
387: __device__ static llint atomicXor(llint *address, llint val)
388: {
389: ullint *address_as_ull = (ullint *)(address);
390: ullint old = *address_as_ull, assumed;
391: do {
392: assumed = old;
393: old = atomicCAS(address_as_ull, assumed, (ullint)(val ^ (llint)assumed));
394: } while (assumed != old);
395: return (llint)old;
396: }
397: #endif
399: template <typename Type>
400: struct AtomicBAND {
401: __device__ Type operator()(Type &x, Type y) const { return atomicAnd(&x, y); }
402: };
403: template <typename Type>
404: struct AtomicBOR {
405: __device__ Type operator()(Type &x, Type y) const { return atomicOr(&x, y); }
406: };
407: template <typename Type>
408: struct AtomicBXOR {
409: __device__ Type operator()(Type &x, Type y) const { return atomicXor(&x, y); }
410: };
412: /*
413: Atomic logical operations:
415: CUDA has no atomic logical operations at all. We support them on integer types.
416: */
418: /* A template without definition makes any instantiation not using given specializations erroneous at compile time,
419: which is what we want since we only support 32-bit and 64-bit integers.
420: */
421: template <typename Type, class Op, int size /* sizeof(Type) */>
422: struct AtomicLogical;
424: template <typename Type, class Op>
425: struct AtomicLogical<Type, Op, 4> {
426: __device__ Type operator()(Type &x, Type y) const
427: {
428: int *address_as_int = (int *)(&x);
429: int old = *address_as_int, assumed;
430: Op op;
431: do {
432: assumed = old;
433: old = atomicCAS(address_as_int, assumed, (int)(op((Type)assumed, y)));
434: } while (assumed != old);
435: return (Type)old;
436: }
437: };
439: template <typename Type, class Op>
440: struct AtomicLogical<Type, Op, 8> {
441: __device__ Type operator()(Type &x, Type y) const
442: {
443: ullint *address_as_ull = (ullint *)(&x);
444: ullint old = *address_as_ull, assumed;
445: Op op;
446: do {
447: assumed = old;
448: old = atomicCAS(address_as_ull, assumed, (ullint)(op((Type)assumed, y)));
449: } while (assumed != old);
450: return (Type)old;
451: }
452: };
454: /* Note land/lor/lxor below are different from LAND etc above. Here we pass arguments by value and return result of ops (not old value) */
455: template <typename Type>
456: struct land {
457: __device__ Type operator()(Type x, Type y) { return x && y; }
458: };
459: template <typename Type>
460: struct lor {
461: __device__ Type operator()(Type x, Type y) { return x || y; }
462: };
463: template <typename Type>
464: struct lxor {
465: __device__ Type operator()(Type x, Type y) { return !x != !y; }
466: };
468: template <typename Type>
469: struct AtomicLAND {
470: __device__ Type operator()(Type &x, Type y) const
471: {
472: AtomicLogical<Type, land<Type>, sizeof(Type)> op;
473: return op(x, y);
474: }
475: };
476: template <typename Type>
477: struct AtomicLOR {
478: __device__ Type operator()(Type &x, Type y) const
479: {
480: AtomicLogical<Type, lor<Type>, sizeof(Type)> op;
481: return op(x, y);
482: }
483: };
484: template <typename Type>
485: struct AtomicLXOR {
486: __device__ Type operator()(Type &x, Type y) const
487: {
488: AtomicLogical<Type, lxor<Type>, sizeof(Type)> op;
489: return op(x, y);
490: }
491: };
492: PETSC_PRAGMA_DIAGNOSTIC_IGNORED_END()
493: #elif PetscDefined(USING_HCC)
495: /*
496: Atomic Insert (exchange) operations
498: See Cuda version
499: */
500: #if PETSC_PKG_HIP_VERSION_LT(4, 4, 0)
501: __device__ static double atomicExch(double *address, double val)
502: {
503: return __longlong_as_double(atomicExch((ullint *)address, __double_as_longlong(val)));
504: }
505: #endif
507: __device__ static inline llint atomicExch(llint *address, llint val)
508: {
509: return (llint)(atomicExch((ullint *)address, (ullint)val));
510: }
512: template <typename Type>
513: struct AtomicInsert {
514: __device__ Type operator()(Type &x, Type y) const { return atomicExch(&x, y); }
515: };
517: #if PetscDefined(HAVE_COMPLEX)
518: #if PetscDefined(USE_REAL_DOUBLE)
519: template <>
520: struct AtomicInsert<PetscComplex> {
521: __device__ PetscComplex operator()(PetscComplex &x, PetscComplex y) const
522: {
523: PetscComplex old, *z = &old;
524: double *xp = (double *)&x, *yp = (double *)&y;
525: AtomicInsert<double> op;
526: z[0] = op(xp[0], yp[0]);
527: z[1] = op(xp[1], yp[1]);
528: return old; /* The returned value may not be atomic. It can be mix of two ops. Caller should discard it. */
529: }
530: };
531: #elif PetscDefined(USE_REAL_SINGLE)
532: template <>
533: struct AtomicInsert<PetscComplex> {
534: __device__ PetscComplex operator()(PetscComplex &x, PetscComplex y) const
535: {
536: double *xp = (double *)&x, *yp = (double *)&y;
537: AtomicInsert<double> op;
538: return op(xp[0], yp[0]);
539: }
540: };
541: #endif
542: #endif
544: /*
545: Atomic add operations
547: */
548: __device__ static inline llint atomicAdd(llint *address, llint val)
549: {
550: return (llint)atomicAdd((ullint *)address, (ullint)val);
551: }
553: template <typename Type>
554: struct AtomicAdd {
555: __device__ Type operator()(Type &x, Type y) const { return atomicAdd(&x, y); }
556: };
558: template <>
559: struct AtomicAdd<double> {
560: __device__ double operator()(double &x, double y) const
561: {
562: /* Cuda version does more checks that may be needed */
563: return atomicAdd(&x, y);
564: }
565: };
567: template <>
568: struct AtomicAdd<float> {
569: __device__ float operator()(float &x, float y) const
570: {
571: /* Cuda version does more checks that may be needed */
572: return atomicAdd(&x, y);
573: }
574: };
576: #if PetscDefined(HAVE_COMPLEX)
577: template <>
578: struct AtomicAdd<PetscComplex> {
579: __device__ PetscComplex operator()(PetscComplex &x, PetscComplex y) const
580: {
581: PetscComplex old, *z = &old;
582: PetscReal *xp = (PetscReal *)&x, *yp = (PetscReal *)&y;
583: AtomicAdd<PetscReal> op;
584: z[0] = op(xp[0], yp[0]);
585: z[1] = op(xp[1], yp[1]);
586: return old; /* The returned value may not be atomic. It can be mix of two ops. Caller should discard it. */
587: }
588: };
589: #endif
591: /*
592: Atomic Mult operations:
594: HIP has no atomicMult at all, so we build our own with atomicCAS
595: */
596: #if PetscDefined(USE_REAL_DOUBLE)
597: __device__ static inline double atomicMult(double *address, double val)
598: {
599: ullint *address_as_ull = (ullint *)(address);
600: ullint old = *address_as_ull, assumed;
601: do {
602: assumed = old;
603: /* Other threads can access and modify value of *address_as_ull after the read above and before the write below */
604: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val * __longlong_as_double(assumed)));
605: } while (assumed != old);
606: return __longlong_as_double(old);
607: }
608: #elif PetscDefined(USE_REAL_SINGLE)
609: __device__ static inline float atomicMult(float *address, float val)
610: {
611: int *address_as_int = (int *)(address);
612: int old = *address_as_int, assumed;
613: do {
614: assumed = old;
615: old = atomicCAS(address_as_int, assumed, __float_as_int(val * __int_as_float(assumed)));
616: } while (assumed != old);
617: return __int_as_float(old);
618: }
619: #endif
621: __device__ static inline int atomicMult(int *address, int val)
622: {
623: int *address_as_int = (int *)(address);
624: int old = *address_as_int, assumed;
625: do {
626: assumed = old;
627: old = atomicCAS(address_as_int, assumed, val * assumed);
628: } while (assumed != old);
629: return (int)old;
630: }
632: __device__ static inline llint atomicMult(llint *address, llint val)
633: {
634: ullint *address_as_ull = (ullint *)(address);
635: ullint old = *address_as_ull, assumed;
636: do {
637: assumed = old;
638: old = atomicCAS(address_as_ull, assumed, (ullint)(val * (llint)assumed));
639: } while (assumed != old);
640: return (llint)old;
641: }
643: template <typename Type>
644: struct AtomicMult {
645: __device__ Type operator()(Type &x, Type y) const { return atomicMult(&x, y); }
646: };
648: /*
649: Atomic Min/Max operations
651: See CUDA version for comments.
652: */
653: #if PETSC_PKG_HIP_VERSION_LT(4, 4, 0)
654: #if PetscDefined(USE_REAL_DOUBLE)
655: __device__ static double atomicMin(double *address, double val)
656: {
657: ullint *address_as_ull = (ullint *)(address);
658: ullint old = *address_as_ull, assumed;
659: do {
660: assumed = old;
661: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(PetscMin(val, __longlong_as_double(assumed))));
662: } while (assumed != old);
663: return __longlong_as_double(old);
664: }
666: __device__ static double atomicMax(double *address, double val)
667: {
668: ullint *address_as_ull = (ullint *)(address);
669: ullint old = *address_as_ull, assumed;
670: do {
671: assumed = old;
672: old = atomicCAS(address_as_ull, assumed, __double_as_longlong(PetscMax(val, __longlong_as_double(assumed))));
673: } while (assumed != old);
674: return __longlong_as_double(old);
675: }
676: #elif PetscDefined(USE_REAL_SINGLE)
677: __device__ static float atomicMin(float *address, float val)
678: {
679: int *address_as_int = (int *)(address);
680: int old = *address_as_int, assumed;
681: do {
682: assumed = old;
683: old = atomicCAS(address_as_int, assumed, __float_as_int(PetscMin(val, __int_as_float(assumed))));
684: } while (assumed != old);
685: return __int_as_float(old);
686: }
688: __device__ static float atomicMax(float *address, float val)
689: {
690: int *address_as_int = (int *)(address);
691: int old = *address_as_int, assumed;
692: do {
693: assumed = old;
694: old = atomicCAS(address_as_int, assumed, __float_as_int(PetscMax(val, __int_as_float(assumed))));
695: } while (assumed != old);
696: return __int_as_float(old);
697: }
698: #endif
699: #endif
701: #if PETSC_PKG_HIP_VERSION_LT(5, 7, 0)
702: __device__ static inline llint atomicMin(llint *address, llint val)
703: {
704: ullint *address_as_ull = (ullint *)(address);
705: ullint old = *address_as_ull, assumed;
706: do {
707: assumed = old;
708: old = atomicCAS(address_as_ull, assumed, (ullint)(PetscMin(val, (llint)assumed)));
709: } while (assumed != old);
710: return (llint)old;
711: }
713: __device__ static inline llint atomicMax(llint *address, llint val)
714: {
715: ullint *address_as_ull = (ullint *)(address);
716: ullint old = *address_as_ull, assumed;
717: do {
718: assumed = old;
719: old = atomicCAS(address_as_ull, assumed, (ullint)(PetscMax(val, (llint)assumed)));
720: } while (assumed != old);
721: return (llint)old;
722: }
723: #endif
725: template <typename Type>
726: struct AtomicMin {
727: __device__ Type operator()(Type &x, Type y) const { return atomicMin(&x, y); }
728: };
729: template <typename Type>
730: struct AtomicMax {
731: __device__ Type operator()(Type &x, Type y) const { return atomicMax(&x, y); }
732: };
734: /*
735: Atomic bitwise operations
736: As of ROCm 3.10, the llint atomicAnd/Or/Xor(llint*, llint) is not supported
737: */
739: __device__ static inline llint atomicAnd(llint *address, llint val)
740: {
741: ullint *address_as_ull = (ullint *)(address);
742: ullint old = *address_as_ull, assumed;
743: do {
744: assumed = old;
745: old = atomicCAS(address_as_ull, assumed, (ullint)(val & (llint)assumed));
746: } while (assumed != old);
747: return (llint)old;
748: }
749: __device__ static inline llint atomicOr(llint *address, llint val)
750: {
751: ullint *address_as_ull = (ullint *)(address);
752: ullint old = *address_as_ull, assumed;
753: do {
754: assumed = old;
755: old = atomicCAS(address_as_ull, assumed, (ullint)(val | (llint)assumed));
756: } while (assumed != old);
757: return (llint)old;
758: }
760: __device__ static inline llint atomicXor(llint *address, llint val)
761: {
762: ullint *address_as_ull = (ullint *)(address);
763: ullint old = *address_as_ull, assumed;
764: do {
765: assumed = old;
766: old = atomicCAS(address_as_ull, assumed, (ullint)(val ^ (llint)assumed));
767: } while (assumed != old);
768: return (llint)old;
769: }
771: template <typename Type>
772: struct AtomicBAND {
773: __device__ Type operator()(Type &x, Type y) const { return atomicAnd(&x, y); }
774: };
775: template <typename Type>
776: struct AtomicBOR {
777: __device__ Type operator()(Type &x, Type y) const { return atomicOr(&x, y); }
778: };
779: template <typename Type>
780: struct AtomicBXOR {
781: __device__ Type operator()(Type &x, Type y) const { return atomicXor(&x, y); }
782: };
784: /*
785: Atomic logical operations:
787: CUDA has no atomic logical operations at all. We support them on integer types.
788: */
790: /* A template without definition makes any instantiation not using given specializations erroneous at compile time,
791: which is what we want since we only support 32-bit and 64-bit integers.
792: */
793: template <typename Type, class Op, int size /* sizeof(Type) */>
794: struct AtomicLogical;
796: template <typename Type, class Op>
797: struct AtomicLogical<Type, Op, 4> {
798: __device__ Type operator()(Type &x, Type y) const
799: {
800: int *address_as_int = (int *)(&x);
801: int old = *address_as_int, assumed;
802: Op op;
803: do {
804: assumed = old;
805: old = atomicCAS(address_as_int, assumed, (int)(op((Type)assumed, y)));
806: } while (assumed != old);
807: return (Type)old;
808: }
809: };
811: template <typename Type, class Op>
812: struct AtomicLogical<Type, Op, 8> {
813: __device__ Type operator()(Type &x, Type y) const
814: {
815: ullint *address_as_ull = (ullint *)(&x);
816: ullint old = *address_as_ull, assumed;
817: Op op;
818: do {
819: assumed = old;
820: old = atomicCAS(address_as_ull, assumed, (ullint)(op((Type)assumed, y)));
821: } while (assumed != old);
822: return (Type)old;
823: }
824: };
826: /* Note land/lor/lxor below are different from LAND etc above. Here we pass arguments by value and return result of ops (not old value) */
827: template <typename Type>
828: struct land {
829: __device__ Type operator()(Type x, Type y) { return x && y; }
830: };
831: template <typename Type>
832: struct lor {
833: __device__ Type operator()(Type x, Type y) { return x || y; }
834: };
835: template <typename Type>
836: struct lxor {
837: __device__ Type operator()(Type x, Type y) { return !x != !y; }
838: };
840: template <typename Type>
841: struct AtomicLAND {
842: __device__ Type operator()(Type &x, Type y) const
843: {
844: AtomicLogical<Type, land<Type>, sizeof(Type)> op;
845: return op(x, y);
846: }
847: };
848: template <typename Type>
849: struct AtomicLOR {
850: __device__ Type operator()(Type &x, Type y) const
851: {
852: AtomicLogical<Type, lor<Type>, sizeof(Type)> op;
853: return op(x, y);
854: }
855: };
856: template <typename Type>
857: struct AtomicLXOR {
858: __device__ Type operator()(Type &x, Type y) const
859: {
860: AtomicLogical<Type, lxor<Type>, sizeof(Type)> op;
861: return op(x, y);
862: }
863: };
864: #endif