MFEM  v4.2.0
Finite element discretization library
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Pages
hip.hpp
Go to the documentation of this file.
1 // Copyright (c) 2010-2020, 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_DEVICE __device__
23 #define MFEM_LAMBDA __host__ __device__
24 #define MFEM_HOST_DEVICE __host__ __device__
25 #define MFEM_DEVICE_SYNC MFEM_GPU_CHECK(hipDeviceSynchronize())
26 #define MFEM_STREAM_SYNC MFEM_GPU_CHECK(hipStreamSynchronize(0))
27 // Define a HIP error check macro, MFEM_GPU_CHECK(x), where x returns/is of
28 // type 'hipError_t'. This macro evaluates 'x' and raises an error if the
29 // result is not hipSuccess.
30 #define MFEM_GPU_CHECK(x) \
31  do \
32  { \
33  hipError_t err = (x); \
34  if (err != hipSuccess) \
35  { \
36  mfem_hip_error(err, #x, _MFEM_FUNC_NAME, __FILE__, __LINE__); \
37  } \
38  } \
39  while (0)
40 #endif // MFEM_USE_HIP
41 
42 // Define the MFEM inner threading macros
43 #if defined(MFEM_USE_HIP) && defined(__HIP_DEVICE_COMPILE__)
44 #define MFEM_SHARED __shared__
45 #define MFEM_SYNC_THREAD __syncthreads()
46 #define MFEM_THREAD_ID(k) hipThreadIdx_ ##k
47 #define MFEM_THREAD_SIZE(k) hipBlockDim_ ##k
48 #define MFEM_FOREACH_THREAD(i,k,N) \
49  for(int i=hipThreadIdx_ ##k; i<N; i+=hipBlockDim_ ##k)
50 #endif
51 
52 namespace mfem
53 {
54 
55 #ifdef MFEM_USE_HIP
56 // Function used by the macro MFEM_GPU_CHECK.
57 void mfem_hip_error(hipError_t err, const char *expr, const char *func,
58  const char *file, int line);
59 #endif
60 
61 /// Allocates device memory
62 void* HipMemAlloc(void **d_ptr, size_t bytes);
63 
64 /// Allocates managed device memory
65 void* HipMallocManaged(void **d_ptr, size_t bytes);
66 
67 /// Frees device memory
68 void* HipMemFree(void *d_ptr);
69 
70 /// Copies memory from Host to Device
71 void* HipMemcpyHtoD(void *d_dst, const void *h_src, size_t bytes);
72 
73 /// Copies memory from Host to Device
74 void* HipMemcpyHtoDAsync(void *d_dst, const void *h_src, size_t bytes);
75 
76 /// Copies memory from Device to Device
77 void* HipMemcpyDtoD(void *d_dst, const void *d_src, size_t bytes);
78 
79 /// Copies memory from Device to Device
80 void* HipMemcpyDtoDAsync(void *d_dst, const void *d_src, size_t bytes);
81 
82 /// Copies memory from Device to Host
83 void* HipMemcpyDtoH(void *h_dst, const void *d_src, size_t bytes);
84 
85 /// Copies memory from Device to Host
86 void* HipMemcpyDtoHAsync(void *h_dst, const void *d_src, size_t bytes);
87 
88 /// Check the error code returned by hipGetLastError(), aborting on error.
89 void HipCheckLastError();
90 
91 /// Get the number of HIP devices
92 int HipGetDeviceCount();
93 
94 } // namespace mfem
95 
96 #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:94
void HipCheckLastError()
Check the error code returned by hipGetLastError(), aborting on error.
Definition: hip.cpp:148
int HipGetDeviceCount()
Get the number of HIP devices.
Definition: hip.cpp:155
void * HipMemcpyDtoDAsync(void *dst, const void *src, size_t bytes)
Copies memory from Device to Device.
Definition: hip.cpp:117
void * HipMemcpyDtoHAsync(void *dst, const void *src, size_t bytes)
Copies memory from Device to Host.
Definition: hip.cpp:140
void * HipMemFree(void *dptr)
Frees device memory.
Definition: hip.cpp:64
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:125
void * HipMemAlloc(void **dptr, size_t bytes)
Allocates device memory.
Definition: hip.cpp:34
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 * HipMemcpyHtoD(void *dst, const void *src, size_t bytes)
Copies memory from Host to Device.
Definition: hip.cpp:79
void * HipMemcpyDtoD(void *dst, const void *src, size_t bytes)
Copies memory from Device to Device.
Definition: hip.cpp:102