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 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 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 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 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 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 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