1 #ifndef AMREX_GPU_UTILITY_H_
2 #define AMREX_GPU_UTILITY_H_
3 #include <AMReX_Config.H>
20 #include <curand_kernel.h>
29 #if defined(AMREX_USE_CUDA)
40 #if defined(AMREX_USE_CUDA)
50 CUpointer_attribute attrib = CU_POINTER_ATTRIBUTE_IS_MANAGED;
51 unsigned int is_managed = 0;
52 void* data[] = { (
void*)(&is_managed) };
53 CUresult
r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
54 return r == CUDA_SUCCESS && is_managed;
55 #elif defined(AMREX_USE_SYCL)
56 auto type = sycl::get_pointer_type(p,Device::syclContext());
57 return type == sycl::usm::alloc::shared;
65 #if defined(AMREX_USE_HIP)
66 hipPointerAttribute_t attrib;
67 hipError_t
r = hipPointerGetAttributes(&attrib, p);
68 #if defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR < 6)
69 return r == hipSuccess && attrib.memoryType == hipMemoryTypeDevice;
71 return r == hipSuccess && attrib.type == hipMemoryTypeDevice;
73 #elif defined(AMREX_USE_CUDA)
74 CUpointer_attribute attrib = CU_POINTER_ATTRIBUTE_MEMORY_TYPE;
75 CUmemorytype mem_type =
static_cast<CUmemorytype
>(0);
76 void* data[] = { (
void*)(&mem_type) };
77 CUresult
r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
78 return r == CUDA_SUCCESS && mem_type == CU_MEMORYTYPE_DEVICE;
79 #elif defined(AMREX_USE_SYCL)
80 auto type = sycl::get_pointer_type(p,Device::syclContext());
81 return type == sycl::usm::alloc::device;
89 #if defined(AMREX_USE_HIP)
90 hipPointerAttribute_t attrib;
91 hipError_t
r = hipPointerGetAttributes(&attrib, p);
92 #if defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR < 6)
93 return r == hipSuccess && attrib.memoryType == hipMemoryTypeHost;
95 return r == hipSuccess && attrib.type == hipMemoryTypeHost;
97 #elif defined(AMREX_USE_CUDA)
98 CUpointer_attribute attrib = CU_POINTER_ATTRIBUTE_MEMORY_TYPE;
99 CUmemorytype mem_type =
static_cast<CUmemorytype
>(0);
100 void* data[] = { (
void*)(&mem_type) };
101 CUresult
r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
102 return r == CUDA_SUCCESS && mem_type == CU_MEMORYTYPE_HOST;
103 #elif defined(AMREX_USE_SYCL)
104 auto type = sycl::get_pointer_type(p,Device::syclContext());
105 return type == sycl::usm::alloc::host;
113 #if defined(AMREX_USE_HIP)
117 hipPointerAttribute_t attrib;
118 hipError_t
r = hipPointerGetAttributes(&attrib, p);
119 #if defined(HIP_VERSION_MAJOR) && (HIP_VERSION_MAJOR < 6)
120 return r == hipSuccess &&
121 (attrib.memoryType == hipMemoryTypeHost ||
122 attrib.memoryType == hipMemoryTypeDevice);
124 return r == hipSuccess &&
125 (attrib.type == hipMemoryTypeHost ||
126 attrib.type == hipMemoryTypeDevice);
129 #elif defined(AMREX_USE_CUDA)
130 CUpointer_attribute attrib = CU_POINTER_ATTRIBUTE_MEMORY_TYPE;
131 CUmemorytype mem_type =
static_cast<CUmemorytype
>(0);
132 void* data[] = { (
void*)(&mem_type) };
133 CUresult
r = cuPointerGetAttributes(1, &attrib, data, (CUdeviceptr)p);
134 return r == CUDA_SUCCESS &&
135 (mem_type == CU_MEMORYTYPE_HOST ||
136 mem_type == CU_MEMORYTYPE_DEVICE ||
137 mem_type == CU_MEMORYTYPE_ARRAY ||
138 mem_type == CU_MEMORYTYPE_UNIFIED);
139 #elif defined(AMREX_USE_SYCL)
140 auto type = sycl::get_pointer_type(p,Device::syclContext());
152 #if defined(AMREX_USE_SYCL)
163 #if defined(AMREX_USE_SYCL)
184 StreamIter (
int n,
bool is_thread_safe=
true) noexcept;
198 #if !defined(AMREX_USE_GPU)
205 void init () noexcept;
214 void*
memcpy (
void* dest, const
void* src, std::
size_t count)
216 #if defined(__HIP_DEVICE_COMPILE__) && defined(AMREX_USE_HIP)
226 std::ostream&
operator<< (std::ostream& os,
const dim3& d);
#define AMREX_FORCE_INLINE
Definition: AMReX_Extension.H:119
#define AMREX_IF_ON_DEVICE(CODE)
Definition: AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition: AMReX_GpuQualifiers.H:58
#define AMREX_GPU_HOST_DEVICE
Definition: AMReX_GpuQualifiers.H:20
Definition: AMReX_GpuUtility.H:182
int operator()() const noexcept
Definition: AMReX_GpuUtility.H:194
void operator++() noexcept
Definition: AMReX_GpuUtility.cpp:81
void init() noexcept
Definition: AMReX_GpuUtility.cpp:36
~StreamIter()
Definition: AMReX_GpuUtility.cpp:65
bool isValid() const noexcept
Definition: AMReX_GpuUtility.H:196
void operator=(StreamIter const &)=delete
int m_n
Definition: AMReX_GpuUtility.H:207
bool m_threadsafe
Definition: AMReX_GpuUtility.H:209
StreamIter(StreamIter &&)=delete
int m_i
Definition: AMReX_GpuUtility.H:208
StreamIter(StreamIter const &)=delete
bool m_sync
Definition: AMReX_GpuUtility.H:210
StreamIter(int n, bool is_thread_safe=true) noexcept
Definition: AMReX_GpuUtility.cpp:23
static constexpr Type_t unknown
Definition: AMReX_EB2_Graph.H:41
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool isinf(T m) noexcept
Definition: AMReX_GpuUtility.H:161
bool isManaged(void const *p) noexcept
Definition: AMReX_GpuUtility.H:48
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T LDG(Array4< T > const &a, int i, int j, int k) noexcept
Definition: AMReX_GpuUtility.H:28
bool isGpuPtr(void const *p) noexcept
Definition: AMReX_GpuUtility.H:112
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void * memcpy(void *dest, const void *src, std::size_t count)
Definition: AMReX_GpuUtility.H:214
bool isPinnedPtr(void const *p) noexcept
Definition: AMReX_GpuUtility.H:88
bool inNoSyncRegion() noexcept
Definition: AMReX_GpuControl.H:146
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool isnan(T m) noexcept
Definition: AMReX_GpuUtility.H:150
bool isDevicePtr(void const *p) noexcept
Definition: AMReX_GpuUtility.H:64
Definition: AMReX_Amr.cpp:49
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition: AMReX.H:111
std::ostream & operator<<(std::ostream &os, AmrMesh const &amr_mesh)
Definition: AMReX_AmrMesh.cpp:1236
Definition: AMReX_Array4.H:61
Definition: AMReX_GpuUtility.H:171
StreamItInfo() noexcept
Definition: AMReX_GpuUtility.H:173
StreamItInfo & DisableDeviceSync() noexcept
Definition: AMReX_GpuUtility.H:175
bool device_sync
Definition: AMReX_GpuUtility.H:172