Actual source code: petscdevicetypes.h

  1: #pragma once

  3: #include <petscsys.h>

  5: // Some overzealous older gcc versions warn that the comparisons below are always true. Neat
  6: // that it can detect this, but the tautology *is* the point of the static_assert()!
  7: #if defined(__GNUC__) && __GNUC__ >= 6 && !PetscDefined(HAVE_WINDOWS_COMPILERS)
  8:   #define PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING 1
  9: #else
 10:   #define PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING 0
 11: #endif

 13: /* SUBMANSEC = Sys */

 15: /*E
 16:   PetscMemType - Memory type of a pointer

 18:   Level: intermediate

 20:   Note:
 21:   `PETSC_MEMTYPE_KOKKOS` depends on the Kokkos backend configuration

 23:   Developer Notes:
 24:   This enum uses a function (`PetscMemTypeToString()`) to convert to string representation so
 25:   cannot be used in `PetscOptionsEnum()`.

 27:   Encoding of the bitmask in binary\: xxxxyyyz
 28: .vb
 29:  z = 0                - Host memory
 30:  z = 1                - Device memory
 31:  yyy = 000            - CUDA-related memory
 32:  yyy = 001            - HIP-related memory
 33:  yyy = 010            - SYCL-related memory
 34:  xxxxyyy1 = 0000,0001 - CUDA memory
 35:  xxxxyyy1 = 0001,0001 - CUDA NVSHMEM memory
 36:  xxxxyyy1 = 0000,0011 - HIP memory
 37:  xxxxyyy1 = 0000,0101 - SYCL memory
 38: .ve

 40:   Other types of memory, e.g., CUDA managed memory, can be added when needed.

 42: .seealso: `PetscMemTypeToString()`, `VecGetArrayAndMemType()`,
 43: `PetscSFBcastWithMemTypeBegin()`, `PetscSFReduceWithMemTypeBegin()`
 44: E*/
 45: typedef enum {
 46:   PETSC_MEMTYPE_HOST    = 0,
 47:   PETSC_MEMTYPE_DEVICE  = 0x01,
 48:   PETSC_MEMTYPE_CUDA    = 0x01,
 49:   PETSC_MEMTYPE_NVSHMEM = 0x11,
 50:   PETSC_MEMTYPE_HIP     = 0x03,
 51:   PETSC_MEMTYPE_SYCL    = 0x05,
 52: } PetscMemType;
 53: #if PetscDefined(HAVE_CUDA)
 54:   #define PETSC_MEMTYPE_KOKKOS PETSC_MEMTYPE_CUDA
 55: #elif PetscDefined(HAVE_HIP)
 56:   #define PETSC_MEMTYPE_KOKKOS PETSC_MEMTYPE_HIP
 57: #elif PetscDefined(HAVE_SYCL)
 58:   #define PETSC_MEMTYPE_KOKKOS PETSC_MEMTYPE_SYCL
 59: #else
 60:   #define PETSC_MEMTYPE_KOKKOS PETSC_MEMTYPE_HOST
 61: #endif

 63: #define PetscMemTypeHost(m)    (((m) & 0x1) == PETSC_MEMTYPE_HOST)
 64: #define PetscMemTypeDevice(m)  (((m) & 0x1) == PETSC_MEMTYPE_DEVICE)
 65: #define PetscMemTypeCUDA(m)    (((m) & 0xF) == PETSC_MEMTYPE_CUDA)
 66: #define PetscMemTypeHIP(m)     (((m) & 0xF) == PETSC_MEMTYPE_HIP)
 67: #define PetscMemTypeSYCL(m)    (((m) & 0xF) == PETSC_MEMTYPE_SYCL)
 68: #define PetscMemTypeNVSHMEM(m) ((m) == PETSC_MEMTYPE_NVSHMEM)

 70: #if defined(__cplusplus)
 71:   #if PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
 72:     #pragma GCC diagnostic push
 73:     #pragma GCC diagnostic ignored "-Wtautological-compare"
 74:   #endif
 75: static_assert(PetscMemTypeHost(PETSC_MEMTYPE_HOST), "");
 76: static_assert(!PetscMemTypeHost(PETSC_MEMTYPE_DEVICE), "");
 77: static_assert(!PetscMemTypeHost(PETSC_MEMTYPE_CUDA), "");
 78: static_assert(!PetscMemTypeHost(PETSC_MEMTYPE_HIP), "");
 79: static_assert(!PetscMemTypeHost(PETSC_MEMTYPE_SYCL), "");
 80: static_assert(!PetscMemTypeHost(PETSC_MEMTYPE_NVSHMEM), "");

 82: static_assert(!PetscMemTypeDevice(PETSC_MEMTYPE_HOST), "");
 83: static_assert(PetscMemTypeDevice(PETSC_MEMTYPE_DEVICE), "");
 84: static_assert(PetscMemTypeDevice(PETSC_MEMTYPE_CUDA), "");
 85: static_assert(PetscMemTypeDevice(PETSC_MEMTYPE_HIP), "");
 86: static_assert(PetscMemTypeDevice(PETSC_MEMTYPE_SYCL), "");
 87: static_assert(PetscMemTypeDevice(PETSC_MEMTYPE_NVSHMEM), "");

 89: static_assert(PetscMemTypeCUDA(PETSC_MEMTYPE_CUDA), "");
 90: static_assert(PetscMemTypeCUDA(PETSC_MEMTYPE_NVSHMEM), "");
 91:   #if PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
 92:     #pragma GCC diagnostic pop
 93:   #endif
 94: #endif // __cplusplus

 96: PETSC_NODISCARD static inline PETSC_CONSTEXPR_14 const char *PetscMemTypeToString(PetscMemType mtype)
 97: {
 98: #ifdef __cplusplus
 99:   static_assert(PETSC_MEMTYPE_CUDA == PETSC_MEMTYPE_DEVICE, "");
100: #endif
101: #define PETSC_CASE_NAME(v) \
102:   case v: \
103:     return PetscStringize(v)

105:   switch (mtype) {
106:     PETSC_CASE_NAME(PETSC_MEMTYPE_HOST);
107:     /* PETSC_CASE_NAME(PETSC_MEMTYPE_DEVICE); same as PETSC_MEMTYPE_CUDA */
108:     PETSC_CASE_NAME(PETSC_MEMTYPE_CUDA);
109:     PETSC_CASE_NAME(PETSC_MEMTYPE_NVSHMEM);
110:     PETSC_CASE_NAME(PETSC_MEMTYPE_HIP);
111:     PETSC_CASE_NAME(PETSC_MEMTYPE_SYCL);
112:   }
113:   PetscUnreachable();
114:   return "invalid";
115: #undef PETSC_CASE_NAME
116: }

118: #define PETSC_OFFLOAD_VECKOKKOS_DEPRECATED PETSC_OFFLOAD_VECKOKKOS PETSC_DEPRECATED_ENUM(3, 17, 0, "PETSC_OFFLOAD_KOKKOS", )

120: /*E
121:   PetscOffloadMask - indicates which memory (CPU, GPU, or none) contains valid data

123:   Values:
124: + `PETSC_OFFLOAD_UNALLOCATED` - no memory contains valid matrix entries; NEVER used for vectors
125: . `PETSC_OFFLOAD_GPU`         - GPU has valid vector/matrix entries
126: . `PETSC_OFFLOAD_CPU`         - CPU has valid vector/matrix entries
127: . `PETSC_OFFLOAD_BOTH`        - Both GPU and CPU have valid vector/matrix entries and they match
128: - `PETSC_OFFLOAD_KOKKOS`      - Reserved for Kokkos matrix and vector. It means the offload is managed by Kokkos, thus this flag itself cannot tell you where the valid data is.

130:   Level: developer

132:   Developer Note:
133:   This enum uses a function (`PetscOffloadMaskToString()`) to convert to string representation so
134:   cannot be used in `PetscOptionsEnum()`.

136: .seealso: `PetscOffloadMaskToString()`, `PetscOffloadMaskToMemType()`, `PetscOffloadMaskToDeviceCopyMode()`
137: E*/
138: typedef enum {
139:   PETSC_OFFLOAD_UNALLOCATED          = 0x0,
140:   PETSC_OFFLOAD_CPU                  = 0x1,
141:   PETSC_OFFLOAD_GPU                  = 0x2,
142:   PETSC_OFFLOAD_BOTH                 = 0x3,
143:   PETSC_OFFLOAD_VECKOKKOS_DEPRECATED = 0x100,
144:   PETSC_OFFLOAD_KOKKOS               = 0x100
145: } PetscOffloadMask;

147: #define PetscOffloadUnallocated(m) ((m) == PETSC_OFFLOAD_UNALLOCATED)
148: #define PetscOffloadHost(m)        (((m) & PETSC_OFFLOAD_CPU) == PETSC_OFFLOAD_CPU)
149: #define PetscOffloadDevice(m)      (((m) & PETSC_OFFLOAD_GPU) == PETSC_OFFLOAD_GPU)
150: #define PetscOffloadBoth(m)        ((m) == PETSC_OFFLOAD_BOTH)

152: #if defined(__cplusplus)
153:   #if PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
154:     #pragma GCC diagnostic push
155:     #pragma GCC diagnostic ignored "-Wtautological-compare"
156:   #endif
157: static_assert(!PetscOffloadHost(PETSC_OFFLOAD_UNALLOCATED), "");
158: static_assert(PetscOffloadHost(PETSC_OFFLOAD_BOTH), "");
159: static_assert(!PetscOffloadHost(PETSC_OFFLOAD_GPU), "");
160: static_assert(PetscOffloadHost(PETSC_OFFLOAD_BOTH), "");
161: static_assert(!PetscOffloadHost(PETSC_OFFLOAD_KOKKOS), "");

163: static_assert(!PetscOffloadDevice(PETSC_OFFLOAD_UNALLOCATED), "");
164: static_assert(!PetscOffloadDevice(PETSC_OFFLOAD_CPU), "");
165: static_assert(PetscOffloadDevice(PETSC_OFFLOAD_GPU), "");
166: static_assert(PetscOffloadDevice(PETSC_OFFLOAD_BOTH), "");
167: static_assert(!PetscOffloadDevice(PETSC_OFFLOAD_KOKKOS), "");

169: static_assert(PetscOffloadBoth(PETSC_OFFLOAD_BOTH), "");
170: static_assert(!PetscOffloadBoth(PETSC_OFFLOAD_CPU), "");
171: static_assert(!PetscOffloadBoth(PETSC_OFFLOAD_GPU), "");
172: static_assert(!PetscOffloadBoth(PETSC_OFFLOAD_GPU), "");
173: static_assert(!PetscOffloadBoth(PETSC_OFFLOAD_KOKKOS), "");
174:   #if PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
175:     #pragma GCC diagnostic pop
176:   #endif
177: #endif // __cplusplus

179: PETSC_NODISCARD static inline PETSC_CONSTEXPR_14 const char *PetscOffloadMaskToString(PetscOffloadMask mask)
180: {
181: #define PETSC_CASE_RETURN(v) \
182:   case v: \
183:     return PetscStringize(v)

185:   switch (mask) {
186:     PETSC_CASE_RETURN(PETSC_OFFLOAD_UNALLOCATED);
187:     PETSC_CASE_RETURN(PETSC_OFFLOAD_CPU);
188:     PETSC_CASE_RETURN(PETSC_OFFLOAD_GPU);
189:     PETSC_CASE_RETURN(PETSC_OFFLOAD_BOTH);
190:     PETSC_CASE_RETURN(PETSC_OFFLOAD_KOKKOS);
191:   }
192:   PetscUnreachable();
193:   return "invalid";
194: #undef PETSC_CASE_RETURN
195: }

197: PETSC_NODISCARD static inline PETSC_CONSTEXPR_14 PetscMemType PetscOffloadMaskToMemType(PetscOffloadMask mask)
198: {
199:   switch (mask) {
200:   case PETSC_OFFLOAD_UNALLOCATED:
201:   case PETSC_OFFLOAD_CPU:
202:     return PETSC_MEMTYPE_HOST;
203:   case PETSC_OFFLOAD_GPU:
204:   case PETSC_OFFLOAD_BOTH:
205:     return PETSC_MEMTYPE_DEVICE;
206:   case PETSC_OFFLOAD_KOKKOS:
207:     return PETSC_MEMTYPE_KOKKOS;
208:   }
209:   PetscUnreachable();
210:   return PETSC_MEMTYPE_HOST;
211: }

213: /*E
214:   PetscDeviceInitType - Initialization strategy for `PetscDevice`

216:   Values:
217: + `PETSC_DEVICE_INIT_NONE`  - PetscDevice is never initialized
218: . `PETSC_DEVICE_INIT_LAZY`  - PetscDevice is initialized on demand
219: - `PETSC_DEVICE_INIT_EAGER` - PetscDevice is initialized as soon as possible

221:   Level: beginner

223:   Note:
224:   `PETSC_DEVICE_INIT_NONE` implies that any initialization of `PetscDevice` is disallowed and
225:   doing so results in an error. Useful to ensure that no accelerator is used in a program.

227: .seealso: `PetscDevice`, `PetscDeviceType`, `PetscDeviceInitialize()`,
228: `PetscDeviceInitialized()`, `PetscDeviceCreate()`
229: E*/
230: typedef enum {
231:   PETSC_DEVICE_INIT_NONE,
232:   PETSC_DEVICE_INIT_LAZY,
233:   PETSC_DEVICE_INIT_EAGER
234: } PetscDeviceInitType;
235: PETSC_EXTERN const char *const PetscDeviceInitTypes[];

237: /*E
238:   PetscDeviceType - Kind of accelerator device backend

240:   Values:
241: + `PETSC_DEVICE_HOST` - Host, no accelerator backend found
242: . `PETSC_DEVICE_CUDA` - CUDA enabled GPU
243: . `PETSC_DEVICE_HIP`  - ROCM/HIP enabled GPU
244: . `PETSC_DEVICE_SYCL` - SYCL enabled device
245: - `PETSC_DEVICE_MAX`  - Always 1 greater than the largest valid `PetscDeviceType`, invalid type, do not use

247:   Level: beginner

249:   Note:
250:   One can also use the `PETSC_DEVICE_DEFAULT()` routine to get the current default `PetscDeviceType`.

252: .seealso: `PetscDevice`, `PetscDeviceInitType`, `PetscDeviceCreate()`, `PETSC_DEVICE_DEFAULT()`
253: E*/
254: typedef enum {
255:   PETSC_DEVICE_HOST,
256:   PETSC_DEVICE_CUDA,
257:   PETSC_DEVICE_HIP,
258:   PETSC_DEVICE_SYCL,
259:   PETSC_DEVICE_MAX
260: } PetscDeviceType;
261: PETSC_EXTERN const char *const PetscDeviceTypes[];

263: /*E
264:   PetscDeviceAttribute - Attribute detailing a property or feature of a `PetscDevice`

266:   Values:
267: + `PETSC_DEVICE_ATTR_SIZE_T_SHARED_MEM_PER_BLOCK` - The maximum amount of shared memory per block in a device kernel
268: - `PETSC_DEVICE_ATTR_MAX`                         - Invalid attribute, do not use

270:   Level: beginner

272: .seealso: `PetscDevice`, `PetscDeviceGetAttribute()`
273: E*/
274: typedef enum {
275:   PETSC_DEVICE_ATTR_SIZE_T_SHARED_MEM_PER_BLOCK,
276:   PETSC_DEVICE_ATTR_MAX
277: } PetscDeviceAttribute;
278: PETSC_EXTERN const char *const PetscDeviceAttributes[];

280: /*S
281:   PetscDevice - Object to manage an accelerator "device" (usually a GPU)

283:   Level: beginner

285:   Note:
286:   This object is used to house configuration and state of a device, but does not offer any
287:   ability to interact with or drive device computation. This functionality is facilitated
288:   instead by the `PetscDeviceContext` object.

290: .seealso: `PetscDeviceType`, `PetscDeviceInitType`, `PetscDeviceCreate()`,
291: `PetscDeviceConfigure()`, `PetscDeviceDestroy()`, `PetscDeviceContext`,
292: `PetscDeviceContextSetDevice()`, `PetscDeviceContextGetDevice()`, `PetscDeviceGetAttribute()`
293: S*/
294: typedef struct _n_PetscDevice *PetscDevice;

296: /*E
297:   PetscStreamType - indicates how a stream implementation will interact
298:   with other streams and if it blocks the host.

300:   Values:
301: + `PETSC_STREAM_DEFAULT`                  - Same as the default stream in CUDA or HIP. Streams of this type may or may not synchronize implicitly with other streams. It does not block the host.
302: . `PETSC_STREAM_NONBLOCKING`              - Same as the nonblocking stream in CUDA or HIP. Streams of this type is truly asynchronous, and is blocked by nothing. It does not block the host. In CUDA, it is created with cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking).
303: . `PETSC_STREAM_DEFAULT_WITH_BARRIER`     - Same as the default stream in CUDA or HIP. PETSc async functions using this kind of stream will end with a stream synchronization. Stream of this type may or may not synchronize implicitly with other streams.
304: . `PETSC_STREAM_NONBLOCKING_WITH_BARRIER` - Same as the nonblocking stream in CUDA or HIP. PETSc async functions using this kind of stream will end with a stream synchronization. Streams of this type are truly asynchronous and are blocked by nothing.
305: - `PETSC_STREAM_MAX`                - Always 1 greater than the largest `PetscStreamType`, do not use

307:   Level: intermediate

309:   Note:
310:   The default stream, also known as the NULL stream or stream 0, can have two different behaviors: legacy behavior and per-thread behavior.
311:   The behavior is determined at compile time. By default, the legacy default stream is used.
312:   The legacy default stream implicitly synchronizes with per-thread default streams.
313:   The per-thread default stream, like nonblocking streams, does not synchronizes with other per-thread streams, but synchronize with the default stream.
314:   The per-thread default stream may be useful for running kernels launched from different threads concurrently on the same GPU when the Multi-Process Service is not available.
315:   To use the per-thread default stream, one can enable it by using the nvcc option "--default-stream per-thread" or the hipcc option "-fgpu-default-stream=per-thread", depending on the backend used.

317: .seealso: `PetscDeviceContextSetStreamType()`, `PetscDeviceContextGetStreamType()`
318: E*/
319: typedef enum {
320:   PETSC_STREAM_DEFAULT,
321:   PETSC_STREAM_NONBLOCKING,
322:   PETSC_STREAM_DEFAULT_WITH_BARRIER,
323:   PETSC_STREAM_NONBLOCKING_WITH_BARRIER,
324:   PETSC_STREAM_MAX
325: } PetscStreamType;
326: PETSC_EXTERN const char *const PetscStreamTypes[];

328: /*E
329:   PetscDeviceContextJoinMode - Describes the type of join operation to perform in
330:   `PetscDeviceContextJoin()`

332:   Values:
333: + `PETSC_DEVICE_CONTEXT_JOIN_DESTROY` - Destroy all incoming sub-contexts after join.
334: . `PETSC_DEVICE_CONTEXT_JOIN_SYNC`    - Synchronize incoming sub-contexts after join.
335: - `PETSC_DEVICE_CONTEXT_JOIN_NO_SYNC` - Do not synchronize incoming sub-contexts after join.

337:   Level: beginner

339: .seealso: `PetscDeviceContext`, `PetscDeviceContextFork()`, `PetscDeviceContextJoin()`
340: E*/
341: typedef enum {
342:   PETSC_DEVICE_CONTEXT_JOIN_DESTROY,
343:   PETSC_DEVICE_CONTEXT_JOIN_SYNC,
344:   PETSC_DEVICE_CONTEXT_JOIN_NO_SYNC
345: } PetscDeviceContextJoinMode;
346: PETSC_EXTERN const char *const PetscDeviceContextJoinModes[];

348: /*S
349:   PetscDeviceContext - Container to manage stream dependencies and the various solver handles
350:   for asynchronous device compute.

352:   Level: beginner

354: .seealso: `PetscDevice`, `PetscDeviceContextCreate()`, `PetscDeviceContextSetDevice()`,
355: `PetscDeviceContextDestroy()`, `PetscDeviceContextFork()`, `PetscDeviceContextJoin()`
356: S*/
357: typedef struct _p_PetscDeviceContext *PetscDeviceContext;

359: /*E
360:   PetscDeviceCopyMode - Describes the copy direction of a device-aware `memcpy`

362:   Values:
363: + `PETSC_DEVICE_COPY_HTOH` - Copy from host memory to host memory
364: . `PETSC_DEVICE_COPY_DTOH` - Copy from device memory to host memory
365: . `PETSC_DEVICE_COPY_HTOD` - Copy from host memory to device memory
366: . `PETSC_DEVICE_COPY_DTOD` - Copy from device memory to device memory
367: - `PETSC_DEVICE_COPY_AUTO` - Infer the copy direction from the pointers

369:   Level: beginner

371: .seealso: `PetscDeviceArrayCopy()`, `PetscDeviceMemcpy()`
372: E*/
373: typedef enum {
374:   PETSC_DEVICE_COPY_HTOH,
375:   PETSC_DEVICE_COPY_DTOH,
376:   PETSC_DEVICE_COPY_HTOD,
377:   PETSC_DEVICE_COPY_DTOD,
378:   PETSC_DEVICE_COPY_AUTO,
379: } PetscDeviceCopyMode;
380: PETSC_EXTERN const char *const PetscDeviceCopyModes[];

382: PETSC_NODISCARD static inline PetscDeviceCopyMode PetscOffloadMaskToDeviceCopyMode(PetscOffloadMask dest, PetscOffloadMask src)
383: {
384:   PetscDeviceCopyMode mode;

386:   PetscFunctionBegin;
387:   PetscAssertAbort(dest != PETSC_OFFLOAD_UNALLOCATED, PETSC_COMM_SELF, PETSC_ERR_ARG_WRONG, "Cannot copy to unallocated");
388:   PetscAssertAbort(src != PETSC_OFFLOAD_UNALLOCATED, PETSC_COMM_SELF, PETSC_ERR_ARG_WRONG, "Cannot copy from unallocated");

390:   if (PetscOffloadDevice(dest)) {
391:     mode = PetscOffloadHost(src) ? PETSC_DEVICE_COPY_HTOD : PETSC_DEVICE_COPY_DTOD;
392:   } else {
393:     mode = PetscOffloadHost(src) ? PETSC_DEVICE_COPY_HTOH : PETSC_DEVICE_COPY_DTOH;
394:   }
395:   PetscFunctionReturn(mode);
396: }

398: PETSC_NODISCARD static inline PETSC_CONSTEXPR_14 PetscDeviceCopyMode PetscMemTypeToDeviceCopyMode(PetscMemType dest, PetscMemType src)
399: {
400:   if (PetscMemTypeHost(dest)) {
401:     return PetscMemTypeHost(src) ? PETSC_DEVICE_COPY_HTOH : PETSC_DEVICE_COPY_DTOH;
402:   } else {
403:     return PetscMemTypeDevice(src) ? PETSC_DEVICE_COPY_DTOD : PETSC_DEVICE_COPY_HTOD;
404:   }
405: }

407: /*E
408:   PetscMemoryAccessMode - Describes the intended usage of a memory region

410:   Values:
411: + `PETSC_MEMORY_ACCESS_READ`       - Read only
412: . `PETSC_MEMORY_ACCESS_WRITE`      - Write only
413: - `PETSC_MEMORY_ACCESS_READ_WRITE` - Read and write

415:   Level: beginner

417:   Notes:
418:   This `enum` is a bitmask with the following encoding (assuming 2 bit)\:

420: .vb
421:   PETSC_MEMORY_ACCESS_READ       = 0b01
422:   PETSC_MEMORY_ACCESS_WRITE      = 0b10
423:   PETSC_MEMORY_ACCESS_READ_WRITE = 0b11

425:   // consequently
426:   PETSC_MEMORY_ACCESS_READ | PETSC_MEMORY_ACCESS_WRITE = PETSC_MEMORY_ACCESS_READ_WRITE
427: .ve

429:   The following convenience macros are also provided\:

431: + `PetscMemoryAccessRead(mode)` - `true` if `mode` is any kind of read, `false` otherwise
432: - `PetscMemoryAccessWrite(mode)` - `true` if `mode` is any kind of write, `false` otherwise

434:   Developer Note:
435:   This enum uses a function (`PetscMemoryAccessModeToString()`) to convert values to string
436:   representation, so cannot be used in `PetscOptionsEnum()`.

438: .seealso: `PetscMemoryAccessModeToString()`, `PetscDevice`, `PetscDeviceContext`
439: E*/
440: typedef enum {
441:   PETSC_MEMORY_ACCESS_READ       = 0x1, // 01
442:   PETSC_MEMORY_ACCESS_WRITE      = 0x2, // 10
443:   PETSC_MEMORY_ACCESS_READ_WRITE = 0x3, // 11
444: } PetscMemoryAccessMode;

446: #define PetscMemoryAccessRead(m)  (((m) & PETSC_MEMORY_ACCESS_READ) == PETSC_MEMORY_ACCESS_READ)
447: #define PetscMemoryAccessWrite(m) (((m) & PETSC_MEMORY_ACCESS_WRITE) == PETSC_MEMORY_ACCESS_WRITE)

449: #if defined(__cplusplus)
450:   #if PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
451:     #pragma GCC diagnostic push
452:     #pragma GCC diagnostic ignored "-Wtautological-compare"
453:   #endif
454: static_assert(PetscMemoryAccessRead(PETSC_MEMORY_ACCESS_READ), "");
455: static_assert(PetscMemoryAccessRead(PETSC_MEMORY_ACCESS_READ_WRITE), "");
456: static_assert(!PetscMemoryAccessRead(PETSC_MEMORY_ACCESS_WRITE), "");
457: static_assert(PetscMemoryAccessWrite(PETSC_MEMORY_ACCESS_WRITE), "");
458: static_assert(PetscMemoryAccessWrite(PETSC_MEMORY_ACCESS_READ_WRITE), "");
459: static_assert(!PetscMemoryAccessWrite(PETSC_MEMORY_ACCESS_READ), "");
460: static_assert((PETSC_MEMORY_ACCESS_READ | PETSC_MEMORY_ACCESS_WRITE) == PETSC_MEMORY_ACCESS_READ_WRITE, "");
461:   #if PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
462:     #pragma GCC diagnostic pop
463:   #endif
464: #endif

466: PETSC_NODISCARD static inline PETSC_CONSTEXPR_14 const char *PetscMemoryAccessModeToString(PetscMemoryAccessMode mode)
467: {
468: #define PETSC_CASE_RETURN(v) \
469:   case v: \
470:     return PetscStringize(v)

472:   switch (mode) {
473:     PETSC_CASE_RETURN(PETSC_MEMORY_ACCESS_READ);
474:     PETSC_CASE_RETURN(PETSC_MEMORY_ACCESS_WRITE);
475:     PETSC_CASE_RETURN(PETSC_MEMORY_ACCESS_READ_WRITE);
476:   }
477:   PetscUnreachable();
478:   return "invalid";
479: #undef PETSC_CASE_RETURN
480: }

482: #undef PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING