MFEM  v4.5.2
Finite element discretization library
hip.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_HIP_HPP
13 #define MFEM_HIP_HPP
14 
15 #include "../config/config.hpp"
16 #include "error.hpp"
17 
18 // HIP block size used by MFEM.
19 #define MFEM_HIP_BLOCKS 256
20 
21 #ifdef MFEM_USE_HIP
22 #define MFEM_USE_CUDA_OR_HIP
23 #define MFEM_DEVICE __device__
24 #define MFEM_LAMBDA __host__ __device__
25 #define MFEM_HOST_DEVICE __host__ __device__
26 #define MFEM_DEVICE_SYNC MFEM_GPU_CHECK(hipDeviceSynchronize())
27 #define MFEM_STREAM_SYNC MFEM_GPU_CHECK(hipStreamSynchronize(0))
28 // Define a HIP error check macro, MFEM_GPU_CHECK(x), where x returns/is of
29 // type 'hipError_t'. This macro evaluates 'x' and raises an error if the
30 // result is not hipSuccess.
31 #define MFEM_GPU_CHECK(x) \
32  do \
33  { \
34  hipError_t err = (x); \
35  if (err != hipSuccess) \
36  { \
37  mfem_hip_error(err, #x, _MFEM_FUNC_NAME, __FILE__, __LINE__); \
38  } \
39  } \
40  while (0)
41 #endif // MFEM_USE_HIP
42 
43 // Define the MFEM inner threading macros
44 #if defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)
45 #define MFEM_SHARED __shared__
46 #define MFEM_SYNC_THREAD __syncthreads()
47 #define MFEM_BLOCK_ID(k) hipBlockIdx_ ##k
48 #define MFEM_THREAD_ID(k) hipThreadIdx_ ##k
49 #define MFEM_THREAD_SIZE(k) hipBlockDim_ ##k
50 #define MFEM_FOREACH_THREAD(i,k,N) \
51  for(int i=hipThreadIdx_ ##k; i<N; i+=hipBlockDim_ ##k)
52 #endif
53 
54 namespace mfem
55 {
56 
57 #ifdef MFEM_USE_HIP
58 // Function used by the macro MFEM_GPU_CHECK.
59 void mfem_hip_error(hipError_t err, const char *expr, const char *func,
60  const char *file, int line);
61 #endif
62 
63 /// Allocates device memory
64 void* HipMemAlloc(void **d_ptr, size_t bytes);
65 
66 /// Allocates managed device memory
67 void* HipMallocManaged(void **d_ptr, size_t bytes);
68 
69 /// Allocates page-locked (pinned) host memory
70 void* HipMemAllocHostPinned(void **ptr, size_t bytes);
71 
72 /// Frees device memory
73 void* HipMemFree(void *d_ptr);
74 
75 /// Frees page-locked (pinned) host memory and returns destination ptr.
76 void* HipMemFreeHostPinned(void *ptr);
77 
78 /// Copies memory from Host to Device
79 void* HipMemcpyHtoD(void *d_dst, const void *h_src, size_t bytes);
80 
81 /// Copies memory from Host to Device
82 void* HipMemcpyHtoDAsync(void *d_dst, const void *h_src, size_t bytes);
83 
84 /// Copies memory from Device to Device
85 void* HipMemcpyDtoD(void *d_dst, const void *d_src, size_t bytes);
86 
87 /// Copies memory from Device to Device
88 void* HipMemcpyDtoDAsync(void *d_dst, const void *d_src, size_t bytes);
89 
90 /// Copies memory from Device to Host
91 void* HipMemcpyDtoH(void *h_dst, const void *d_src, size_t bytes);
92 
93 /// Copies memory from Device to Host
94 void* HipMemcpyDtoHAsync(void *h_dst, const void *d_src, size_t bytes);
95 
96 /// Check the error code returned by hipGetLastError(), aborting on error.
97 void HipCheckLastError();
98 
99 /// Get the number of HIP devices
100 int HipGetDeviceCount();
101 
102 } // namespace mfem
103 
104 #endif // MFEM_HIP_HPP
void * HipMallocManaged(void **dptr, size_t bytes)
Allocates managed device memory.
Definition: hip.cpp:49
void * HipMemcpyHtoDAsync(void *dst, const void *src, size_t bytes)
Copies memory from Host to Device.
Definition: hip.cpp:124
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 HipCheckLastError()
Check the error code returned by hipGetLastError(), aborting on error.
Definition: hip.cpp:178
int HipGetDeviceCount()
Get the number of HIP devices.
Definition: hip.cpp:185
void * HipMemcpyDtoDAsync(void *dst, const void *src, size_t bytes)
Copies memory from Device to Device.
Definition: hip.cpp:147
void * HipMemFreeHostPinned(void *ptr)
Frees page-locked (pinned) host memory and returns destination ptr.
Definition: hip.cpp:94
void * HipMemcpyDtoHAsync(void *dst, const void *src, size_t bytes)
Copies memory from Device to Host.
Definition: hip.cpp:170
void * HipMemAllocHostPinned(void **ptr, size_t bytes)
Allocates page-locked (pinned) host memory.
Definition: hip.cpp:64
void * HipMemFree(void *dptr)
Frees device memory.
Definition: hip.cpp:79
void mfem_hip_error(hipError_t err, const char *expr, const char *func, const char *file, int line)
Definition: hip.cpp:23
void * HipMemcpyDtoH(void *dst, const void *src, size_t bytes)
Copies memory from Device to Host.
Definition: hip.cpp:155
void * HipMemAlloc(void **dptr, size_t bytes)
Allocates device memory.
Definition: hip.cpp:34
void * HipMemcpyHtoD(void *dst, const void *src, size_t bytes)
Copies memory from Host to Device.
Definition: hip.cpp:109
void * HipMemcpyDtoD(void *dst, const void *src, size_t bytes)
Copies memory from Device to Device.
Definition: hip.cpp:132