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

ROCm 5.3 and later support asynchronous memory operations #2197

Merged
merged 1 commit into from
Dec 19, 2023
Merged

ROCm 5.3 and later support asynchronous memory operations #2197

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
30 changes: 25 additions & 5 deletions include/alpaka/core/ApiHipRt.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -181,9 +181,17 @@ namespace alpaka

static inline Error_t freeAsync([[maybe_unused]] void* devPtr, [[maybe_unused]] Stream_t stream)
{
// hipFreeAsync is implemented only in ROCm 5.2.0 and later.
# if HIP_VERSION >= 50'200'000
return ::hipFreeAsync(devPtr, stream);
// stream-ordered memory operations are fully implemented only in ROCm 5.3.0 and later.
# if HIP_VERSION >= 50'300'000
// hipFreeAsync fails on a null pointer deallocation
if(devPtr)
{
return ::hipFreeAsync(devPtr, stream);
}
else
{
return ::hipSuccess;
}
# else
// Not implemented.
return errorUnknown;
Expand Down Expand Up @@ -289,9 +297,21 @@ namespace alpaka
[[maybe_unused]] size_t size,
[[maybe_unused]] Stream_t stream)
{
// hipMallocAsync is implemented only in ROCm 5.2.0 and later.
# if HIP_VERSION >= 50'200'000
// stream-ordered memory operations are fully implemented only in ROCm 5.3.0 and later.
# if HIP_VERSION >= 50'600'000
return ::hipMallocAsync(devPtr, size, stream);
# elif HIP_VERSION >= 50'300'000
// before ROCm 5.6.0, hipMallocAsync fails for an allocation of 0 bytes
if(size > 0)
{
return ::hipMallocAsync(devPtr, size, stream);
}
else
{
// make sure the pointer can safely be passed to hipFreeAsync
*devPtr = nullptr;
return ::hipSuccess;
}
# else
// Not implemented.
return errorUnknown;
Expand Down
18 changes: 12 additions & 6 deletions include/alpaka/mem/buf/BufUniformCudaHipRt.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -280,8 +280,8 @@ namespace alpaka
# endif
# if defined(ALPAKA_ACC_GPU_HIP_ENABLED)
static_assert(
!std::is_same_v<TApi, ApiHipRt>,
"HIP devices do not support stream-ordered memory buffers.");
std::is_same_v<TApi, ApiHipRt> && TApi::version >= BOOST_VERSION_NUMBER(5, 3, 0),
"Support for stream-ordered memory buffers requires HIP/ROCm 5.3 or higher.");
# endif
static_assert(
TDim::value <= 1,
Expand Down Expand Up @@ -317,16 +317,22 @@ namespace alpaka
}
};

# if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
//! The CUDA/HIP stream-ordered memory allocation capability trait specialization.
template<typename TApi, typename TDim>
struct HasAsyncBufSupport<TDim, DevUniformCudaHipRt<TApi>>
: std::bool_constant<
std::is_same_v<TApi, ApiCudaRt> && TApi::version >= BOOST_VERSION_NUMBER(11, 2, 0)
&& TDim::value <= 1>
TDim::value <= 1
&& (
# if defined(ALPAKA_ACC_GPU_CUDA_ENABLED)
std::is_same_v<TApi, ApiCudaRt> && TApi::version >= BOOST_VERSION_NUMBER(11, 2, 0)
# elif defined(ALPAKA_ACC_GPU_HIP_ENABLED)
std::is_same_v<TApi, ApiHipRt> && TApi::version >= BOOST_VERSION_NUMBER(5, 3, 0)
# else
false
# endif
)>
{
};
# endif

//! The pinned/mapped memory allocation trait specialization for the CUDA/HIP devices.
template<typename TApi, typename TElem, typename TDim, typename TIdx>
Expand Down
Loading