MFEM  v4.5.2
Finite element discretization library
cuda.hpp
Go to the documentation of this file.
1 // Copyright (c) 2010-2023, Lawrence Livermore National Security, LLC. Produced
2 // at the Lawrence Livermore National Laboratory. All Rights reserved. See files
3 // LICENSE and NOTICE for details. LLNL-CODE-806117.
4 //
5 // This file is part of the MFEM library. For more information and source code
6 // availability visit https://mfem.org.
7 //
8 // MFEM is free software; you can redistribute it and/or modify it under the
9 // terms of the BSD-3 license. We welcome feedback and contributions, see file
10 // CONTRIBUTING.md for details.
11 
12 #ifndef MFEM_CUDA_HPP
13 #define MFEM_CUDA_HPP
14 
15 #include "../config/config.hpp"
16 #include "error.hpp"
17 
18 // CUDA block size used by MFEM.
19 #define MFEM_CUDA_BLOCKS 256
20 
21 #ifdef MFEM_USE_CUDA
22 #define MFEM_USE_CUDA_OR_HIP
23 #define MFEM_DEVICE __device__
24 #define MFEM_LAMBDA __host__
25 #define MFEM_HOST_DEVICE __host__ __device__
26 #define MFEM_DEVICE_SYNC MFEM_GPU_CHECK(cudaDeviceSynchronize())
27 #define MFEM_STREAM_SYNC MFEM_GPU_CHECK(cudaStreamSynchronize(0))
28 // Define a CUDA error check macro, MFEM_GPU_CHECK(x), where x returns/is of
29 // type 'cudaError_t'. This macro evaluates 'x' and raises an error if the
30 // result is not cudaSuccess.
31 #define MFEM_GPU_CHECK(x) \
32  do \
33  { \
34  cudaError_t err = (x); \
35  if (err != cudaSuccess) \
36  { \
37  mfem_cuda_error(err, #x, _MFEM_FUNC_NAME, __FILE__, __LINE__); \
38  } \
39  } \
40  while (0)
41 #endif // MFEM_USE_CUDA
42 
43 // Define the MFEM inner threading macros
44 #if defined(MFEM_USE_CUDA) && defined(__CUDA_ARCH__)
45 #define MFEM_SHARED __shared__
46 #define MFEM_SYNC_THREAD __syncthreads()
47 #define MFEM_BLOCK_ID(k) blockIdx.k
48 #define MFEM_THREAD_ID(k) threadIdx.k
49 #define MFEM_THREAD_SIZE(k) blockDim.k
50 #define MFEM_FOREACH_THREAD(i,k,N) for(int i=threadIdx.k; i<N; i+=blockDim.k)
51 #endif
52 
53 namespace mfem
54 {
55 
56 #ifdef MFEM_USE_CUDA
57 // Function used by the macro MFEM_GPU_CHECK.
58 void mfem_cuda_error(cudaError_t err, const char *expr, const char *func,
59  const char *file, int line);
60 #endif
61 
62 /// Allocates device memory and returns destination ptr.
63 void* CuMemAlloc(void **d_ptr, size_t bytes);
64 
65 /// Allocates managed device memory
66 void* CuMallocManaged(void **d_ptr, size_t bytes);
67 
68 /// Allocates page-locked (pinned) host memory
69 void* CuMemAllocHostPinned(void **ptr, size_t bytes);
70 
71 /// Frees device memory and returns destination ptr.
72 void* CuMemFree(void *d_ptr);
73 
74 /// Frees page-locked (pinned) host memory and returns destination ptr.
75 void* CuMemFreeHostPinned(void *ptr);
76 
77 /// Copies memory from Host to Device and returns destination ptr.
78 void* CuMemcpyHtoD(void *d_dst, const void *h_src, size_t bytes);
79 
80 /// Copies memory from Host to Device and returns destination ptr.
81 void* CuMemcpyHtoDAsync(void *d_dst, const void *h_src, size_t bytes);
82 
83 /// Copies memory from Device to Device
84 void* CuMemcpyDtoD(void *d_dst, const void *d_src, size_t bytes);
85 
86 /// Copies memory from Device to Device
87 void* CuMemcpyDtoDAsync(void *d_dst, const void *d_src, size_t bytes);
88 
89 /// Copies memory from Device to Host
90 void* CuMemcpyDtoH(void *h_dst, const void *d_src, size_t bytes);
91 
92 /// Copies memory from Device to Host
93 void* CuMemcpyDtoHAsync(void *h_dst, const void *d_src, size_t bytes);
94 
95 /// Check the error code returned by cudaGetLastError(), aborting on error.
96 void CuCheckLastError();
97 
98 /// Get the number of CUDA devices
99 int CuGetDeviceCount();
100 
101 } // namespace mfem
102 
103 #endif // MFEM_CUDA_HPP
void * CuMemcpyHtoD(void *dst, const void *src, size_t bytes)
Copies memory from Host to Device and returns destination ptr.
Definition: cuda.cpp:109
void * CuMemFree(void *dptr)
Frees device memory and returns destination ptr.
Definition: cuda.cpp:79
OutStream err(std::cerr)
Global stream used by the library for standard error output. Initially it uses the same std::streambu...
Definition: globals.hpp:71
void * CuMemFreeHostPinned(void *ptr)
Frees page-locked (pinned) host memory and returns destination ptr.
Definition: cuda.cpp:94
void CuCheckLastError()
Check the error code returned by cudaGetLastError(), aborting on error.
Definition: cuda.cpp:178
int CuGetDeviceCount()
Get the number of CUDA devices.
Definition: cuda.cpp:185
void * CuMallocManaged(void **dptr, size_t bytes)
Allocates managed device memory.
Definition: cuda.cpp:49
void mfem_cuda_error(cudaError_t err, const char *expr, const char *func, const char *file, int line)
Definition: cuda.cpp:23
void * CuMemcpyDtoD(void *dst, const void *src, size_t bytes)
Copies memory from Device to Device.
Definition: cuda.cpp:132
void * CuMemcpyDtoDAsync(void *dst, const void *src, size_t bytes)
Copies memory from Device to Device.
Definition: cuda.cpp:147
void * CuMemcpyDtoHAsync(void *dst, const void *src, size_t bytes)
Copies memory from Device to Host.
Definition: cuda.cpp:170
void * CuMemcpyHtoDAsync(void *dst, const void *src, size_t bytes)
Copies memory from Host to Device and returns destination ptr.
Definition: cuda.cpp:124
void * CuMemAllocHostPinned(void **ptr, size_t bytes)
Allocates page-locked (pinned) host memory.
Definition: cuda.cpp:64
void * CuMemAlloc(void **dptr, size_t bytes)
Allocates device memory and returns destination ptr.
Definition: cuda.cpp:34
void * CuMemcpyDtoH(void *dst, const void *src, size_t bytes)
Copies memory from Device to Host.
Definition: cuda.cpp:155