1 #pragma once
2
3 #include <petscsys.h> /*I <petscdevicetypes.h> I*/
4
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
12
13 /* SUBMANSEC = Sys */
14
15 /*E
16 PetscMemType - Memory type of a pointer
17
18 Level: intermediate
19
20 Note:
21 `PETSC_MEMTYPE_KOKKOS` depends on the Kokkos backend configuration
22
23 Developer Notes:
24 This enum uses a function (`PetscMemTypeToString()`) to convert to string representation so
25 cannot be used in `PetscOptionsEnum()`.
26
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
39
40 Other types of memory, e.g., CUDA managed memory, can be added when needed.
41
42 .seealso: `PetscMemTypeToString()`, `VecGetArrayAndMemType()`,
43 `PetscSFBcastWithMemTypeBegin()`, `PetscSFReduceWithMemTypeBegin()`
44 E*/
45 typedef enum {
46 PETSC_MEMTYPE_HOST = 0,
47 PETSC_MEMTYPE_DEVICE = 1, /* 0x01 */
48 PETSC_MEMTYPE_CUDA = 1, /* 0x01 */
49 PETSC_MEMTYPE_NVSHMEM = 17, /* 0x11 */
50 PETSC_MEMTYPE_HIP = 3, /* 0x03 */
51 PETSC_MEMTYPE_SYCL = 5 /* 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
62
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)
69
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), "");
81
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), "");
88
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
95
PetscMemTypeToString(PetscMemType mtype)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)
104
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 }
117
118 #define PETSC_OFFLOAD_VECKOKKOS_DEPRECATED PETSC_OFFLOAD_VECKOKKOS PETSC_DEPRECATED_ENUM(3, 17, 0, "PETSC_OFFLOAD_KOKKOS", )
119
120 /*E
121 PetscOffloadMask - indicates which memory (CPU, GPU, or none) contains valid data
122
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.
129
130 Level: developer
131
132 Developer Note:
133 This enum uses a function (`PetscOffloadMaskToString()`) to convert to string representation so
134 cannot be used in `PetscOptionsEnum()`.
135
136 .seealso: `PetscOffloadMaskToString()`, `PetscOffloadMaskToMemType()`, `PetscOffloadMaskToDeviceCopyMode()`
137 E*/
138 typedef enum {
139 PETSC_OFFLOAD_UNALLOCATED = 0, /* 0x0 */
140 PETSC_OFFLOAD_CPU = 1, /* 0x1 */
141 PETSC_OFFLOAD_GPU = 2, /* 0x2 */
142 PETSC_OFFLOAD_BOTH = 3, /* 0x3 */
143 PETSC_OFFLOAD_VECKOKKOS_DEPRECATED = 256, /* 0x100 */
144 PETSC_OFFLOAD_KOKKOS = 256 /* 0x100 */
145 } PetscOffloadMask;
146
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)
151
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), "");
162
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), "");
168
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
178
PetscOffloadMaskToString(PetscOffloadMask mask)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)
184
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 }
196
PetscOffloadMaskToMemType(PetscOffloadMask mask)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 }
212
213 /*E
214 PetscDeviceInitType - Initialization strategy for `PetscDevice`
215
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
220
221 Level: beginner
222
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.
226
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[];
236
237 /*E
238 PetscDeviceType - Kind of accelerator device backend
239
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
246
247 Level: beginner
248
249 Note:
250 One can also use the `PETSC_DEVICE_DEFAULT()` routine to get the current default `PetscDeviceType`.
251
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[];
262
263 /*E
264 PetscDeviceAttribute - Attribute detailing a property or feature of a `PetscDevice`
265
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
269
270 Level: beginner
271
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[];
279
280 /*S
281 PetscDevice - Object to manage an accelerator "device" (usually a GPU)
282
283 Level: beginner
284
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.
289
290 .seealso: `PetscDeviceType`, `PetscDeviceInitType`, `PetscDeviceCreate()`,
291 `PetscDeviceConfigure()`, `PetscDeviceDestroy()`, `PetscDeviceContext`,
292 `PetscDeviceContextSetDevice()`, `PetscDeviceContextGetDevice()`, `PetscDeviceGetAttribute()`
293 S*/
294 typedef struct _n_PetscDevice *PetscDevice;
295
296 /*E
297 PetscStreamType - indicates how a stream implementation will interact
298 with other streams and if it blocks the host.
299
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
306
307 Level: intermediate
308
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.
316
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[];
327
328 /*E
329 PetscDeviceContextJoinMode - Describes the type of join operation to perform in
330 `PetscDeviceContextJoin()`
331
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.
336
337 Level: beginner
338
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[];
347
348 /*S
349 PetscDeviceContext - Container to manage stream dependencies and the various solver handles
350 for asynchronous device compute.
351
352 Level: beginner
353
354 .seealso: `PetscDevice`, `PetscDeviceContextCreate()`, `PetscDeviceContextSetDevice()`,
355 `PetscDeviceContextDestroy()`, `PetscDeviceContextFork()`, `PetscDeviceContextJoin()`
356 S*/
357 typedef struct _p_PetscDeviceContext *PetscDeviceContext;
358
359 /*E
360 PetscDeviceCopyMode - Describes the copy direction of a device-aware `memcpy`
361
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
368
369 Level: beginner
370
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[];
381
PetscOffloadMaskToDeviceCopyMode(PetscOffloadMask dest,PetscOffloadMask src)382 PETSC_NODISCARD static inline PetscDeviceCopyMode PetscOffloadMaskToDeviceCopyMode(PetscOffloadMask dest, PetscOffloadMask src)
383 {
384 PetscDeviceCopyMode mode;
385
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");
389
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 }
397
PetscMemTypeToDeviceCopyMode(PetscMemType dest,PetscMemType src)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 }
406
407 /*E
408 PetscMemoryAccessMode - Describes the intended usage of a memory region
409
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
414
415 Level: beginner
416
417 Notes:
418 This `enum` is a bitmask with the following encoding (assuming 2 bit)\:
419
420 .vb
421 PETSC_MEMORY_ACCESS_READ = 0b01
422 PETSC_MEMORY_ACCESS_WRITE = 0b10
423 PETSC_MEMORY_ACCESS_READ_WRITE = 0b11
424
425 // consequently
426 PETSC_MEMORY_ACCESS_READ | PETSC_MEMORY_ACCESS_WRITE = PETSC_MEMORY_ACCESS_READ_WRITE
427 .ve
428
429 The following convenience macros are also provided\:
430
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
433
434 Developer Note:
435 This enum uses a function (`PetscMemoryAccessModeToString()`) to convert values to string
436 representation, so cannot be used in `PetscOptionsEnum()`.
437
438 .seealso: `PetscMemoryAccessModeToString()`, `PetscDevice`, `PetscDeviceContext`
439 E*/
440 typedef enum {
441 PETSC_MEMORY_ACCESS_READ = 1, /* 01 */
442 PETSC_MEMORY_ACCESS_WRITE = 2, /* 10 */
443 PETSC_MEMORY_ACCESS_READ_WRITE = 3 /* 11 */
444 } PetscMemoryAccessMode;
445
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)
448
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
465
PetscMemoryAccessModeToString(PetscMemoryAccessMode mode)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)
471
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 }
481
482 #undef PETSC_SHOULD_SILENCE_GCC_TAUTOLOGICAL_COMPARE_WARNING
483