Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

use HIP stream-ordered allocator #3980

Merged
merged 6 commits into from
Jun 9, 2024
Merged

use HIP stream-ordered allocator #3980

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
f2c1084
initial version
BenWibking Jun 9, 2024
2a0f68c
fix macro usage
BenWibking Jun 9, 2024
34e69f3
use uint64_t for HIP
BenWibking Jun 9, 2024
a258a09
fix typo
BenWibking Jun 9, 2024
cad7364
cudaDevAttribute -> hipDeviceAttribute
BenWibking Jun 9, 2024
401fde4
fix [[no_discard]] error
BenWibking Jun 9, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 4 additions & 0 deletions Src/Base/AMReX_GpuControl.H
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,6 +11,10 @@
#define AMREX_CUDA_GE_11_2 1
#endif

#if defined(AMREX_USE_HIP) || defined(AMREX_CUDA_GE_11_2)
#define AMREX_GPU_STREAM_ALLOC_SUPPORT 1
#endif

#if defined(AMREX_USE_HIP)
#define AMREX_HIP_OR_CUDA(a,b) a
#elif defined(AMREX_USE_CUDA)
Expand Down
6 changes: 5 additions & 1 deletion Src/Base/AMReX_GpuDevice.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -386,13 +386,17 @@ Device::initialize_gpu ()
AMREX_HIP_SAFE_CALL(hipStreamCreate(&gpu_stream_pool[i]));
}

#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
AMREX_HIP_SAFE_CALL(hipDeviceGetAttribute(&memory_pools_supported, hipDeviceAttributeMemoryPoolsSupported, device_id));
#endif

#elif defined(AMREX_USE_CUDA)
AMREX_CUDA_SAFE_CALL(cudaGetDeviceProperties(&device_prop, device_id));

AMREX_ALWAYS_ASSERT_WITH_MESSAGE(device_prop.major >= 4 || (device_prop.major == 3 && device_prop.minor >= 5),
"Compute capability must be >= 3.5");

#ifdef AMREX_CUDA_GE_11_2
#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
cudaDeviceGetAttribute(&memory_pools_supported, cudaDevAttrMemoryPoolsSupported, device_id);
#endif

Expand Down
12 changes: 9 additions & 3 deletions Src/Base/AMReX_PArena.H
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,10 +42,16 @@ public:
[[nodiscard]] bool isStreamOrderedArena () const final { return true; }
#endif

#ifdef AMREX_CUDA_GE_11_2
#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
private:
cudaMemPool_t m_pool;
cuuint64_t m_old_release_threshold;
AMREX_HIP_OR_CUDA(
hipMemPool_t m_pool;,
cudaMemPool_t m_pool;
)
AMREX_HIP_OR_CUDA(
uint64_t m_old_release_threshold;,
cuuint64_t m_old_release_threshold;
)
#endif
};

Expand Down
59 changes: 41 additions & 18 deletions Src/Base/AMReX_PArena.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,24 +11,41 @@ namespace amrex {

PArena::PArena (Long release_threshold)
{
#ifdef AMREX_CUDA_GE_11_2
#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
if (Gpu::Device::memoryPoolsSupported()) {
AMREX_CUDA_SAFE_CALL(cudaDeviceGetMemPool(&m_pool, Gpu::Device::deviceId()));
AMREX_CUDA_SAFE_CALL(cudaMemPoolGetAttribute(m_pool, cudaMemPoolAttrReleaseThreshold,
&m_old_release_threshold));
cuuint64_t value = release_threshold;
AMREX_CUDA_SAFE_CALL(cudaMemPoolSetAttribute(m_pool, cudaMemPoolAttrReleaseThreshold, &value));
AMREX_HIP_OR_CUDA(
AMREX_HIP_SAFE_CALL(hipDeviceGetMemPool(&m_pool, Gpu::Device::deviceId()));,
AMREX_CUDA_SAFE_CALL(cudaDeviceGetMemPool(&m_pool, Gpu::Device::deviceId()));
)
AMREX_HIP_OR_CUDA(
AMREX_HIP_SAFE_CALL(hipMemPoolGetAttribute(m_pool, hipMemPoolAttrReleaseThreshold,
&m_old_release_threshold));,
AMREX_CUDA_SAFE_CALL(cudaMemPoolGetAttribute(m_pool, cudaMemPoolAttrReleaseThreshold,
&m_old_release_threshold));
)
AMREX_HIP_OR_CUDA(
uint64_t value = release_threshold;,
cuuint64_t value = release_threshold;
)
AMREX_HIP_OR_CUDA(
AMREX_HIP_SAFE_CALL(hipMemPoolSetAttribute(m_pool, hipMemPoolAttrReleaseThreshold, &value));,
AMREX_CUDA_SAFE_CALL(cudaMemPoolSetAttribute(m_pool, cudaMemPoolAttrReleaseThreshold, &value));
)
}
#endif
amrex::ignore_unused(release_threshold);
}

PArena::~PArena () // NOLINT(modernize-use-equals-default)
{
#ifdef AMREX_CUDA_GE_11_2
#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
if (Gpu::Device::memoryPoolsSupported()) {
AMREX_CUDA_SAFE_CALL(cudaMemPoolSetAttribute(m_pool, cudaMemPoolAttrReleaseThreshold,
&m_old_release_threshold));
AMREX_HIP_OR_CUDA(
AMREX_HIP_SAFE_CALL(hipMemPoolSetAttribute(m_pool, hipMemPoolAttrReleaseThreshold,
&m_old_release_threshold));,
AMREX_CUDA_SAFE_CALL(cudaMemPoolSetAttribute(m_pool, cudaMemPoolAttrReleaseThreshold,
&m_old_release_threshold));
)
}
#endif
}
Expand All @@ -38,10 +55,13 @@ PArena::alloc (std::size_t nbytes)
{
#if defined(AMREX_USE_GPU)

#ifdef AMREX_CUDA_GE_11_2
#if defined(AMREX_GPU_STREAM_ALLOC_SUPPORT)
if (Gpu::Device::memoryPoolsSupported()) {
void* p;
AMREX_CUDA_SAFE_CALL(cudaMallocAsync(&p, nbytes, m_pool, Gpu::gpuStream()));
AMREX_HIP_OR_CUDA(
AMREX_HIP_SAFE_CALL(hipMallocAsync(&p, nbytes, m_pool, Gpu::gpuStream()));,
AMREX_CUDA_SAFE_CALL(cudaMallocAsync(&p, nbytes, m_pool, Gpu::gpuStream()));
)
return p;
} else
#endif
Expand Down Expand Up @@ -71,9 +91,12 @@ PArena::free (void* p)

#if defined(AMREX_USE_GPU)

#ifdef AMREX_CUDA_GE_11_2
#if defined (AMREX_GPU_STREAM_ALLOC_SUPPORT)
if (Gpu::Device::memoryPoolsSupported()) {
AMREX_CUDA_SAFE_CALL(cudaFreeAsync(p, Gpu::gpuStream()));
AMREX_HIP_OR_CUDA(
AMREX_HIP_SAFE_CALL(hipFreeAsync(p, Gpu::gpuStream()));,
AMREX_CUDA_SAFE_CALL(cudaFreeAsync(p, Gpu::gpuStream()));
)
} else
#endif
{
Expand All @@ -99,7 +122,7 @@ bool
PArena::isDeviceAccessible () const
{
#ifdef AMREX_USE_GPU
#ifdef AMREX_CUDA_GE_11_2
#ifdef AMREX_GPU_STREAM_ALLOC_SUPPORT
if (Gpu::Device::memoryPoolsSupported()) {
return true;
} else
Expand All @@ -116,7 +139,7 @@ bool
PArena::isHostAccessible () const
{
#ifdef AMREX_USE_GPU
#ifdef AMREX_CUDA_GE_11_2
#if defined (AMREX_GPU_STREAM_ALLOC_SUPPORT)
if (Gpu::Device::memoryPoolsSupported()) {
return false; // cudaMallocAsync allocates device memory
} else
Expand All @@ -133,7 +156,7 @@ bool
PArena::isManaged () const
{
#ifdef AMREX_USE_GPU
#ifdef AMREX_CUDA_GE_11_2
#if defined(AMREX_GPU_STREAM_ALLOC_SUPPORT)
if (Gpu::Device::memoryPoolsSupported()) {
return false; // cudaMallocAsync allocates device memory
} else
Expand All @@ -150,7 +173,7 @@ bool
PArena::isDevice () const
{
#ifdef AMREX_USE_GPU
#ifdef AMREX_CUDA_GE_11_2
#if defined (AMREX_GPU_STREAM_ALLOC_SUPPORT)
if (Gpu::Device::memoryPoolsSupported()) {
return true; // cudaMallocAsync allocates device memory
} else
Expand All @@ -167,7 +190,7 @@ bool
PArena::isPinned () const
{
#ifdef AMREX_USE_GPU
#ifdef AMREX_CUDA_GE_11_2
#if defined (AMREX_GPU_STREAM_ALLOC_SUPPORT)
if (Gpu::Device::memoryPoolsSupported()) {
return false; // cudaMallocAsync allocates device memory
} else
Expand Down
Loading