From 6993d3bd2d86b4e5b7273b11929bc885c3669461 Mon Sep 17 00:00:00 2001
From: Andrea Bocci <andrea.bocci@cern.ch>
Date: Thu, 14 Dec 2023 16:23:37 +0200
Subject: [PATCH] Add protections for 0-sized allocations in ROCm 5.2-5.5

---
 include/alpaka/core/ApiHipRt.hpp | 20 ++++++++++++++++++--
 1 file changed, 18 insertions(+), 2 deletions(-)

diff --git a/include/alpaka/core/ApiHipRt.hpp b/include/alpaka/core/ApiHipRt.hpp
index e76294f068fb..905ec58c73a0 100644
--- a/include/alpaka/core/ApiHipRt.hpp
+++ b/include/alpaka/core/ApiHipRt.hpp
@@ -182,8 +182,15 @@ 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
+#    if HIP_VERSION >= 50'600'000
             return ::hipFreeAsync(devPtr, stream);
+#    elif HIP_VERSION >= 50'200'000
+            // before ROCm 5.6.0, hipFreeAsync fails on a null pointer deallocation
+            if (devPtr) {
+              return ::hipFreeAsync(devPtr, stream);
+            } else {
+              return ::hipSuccess;
+            }
 #    else
             // Not implemented.
             return errorUnknown;
@@ -290,8 +297,17 @@ namespace alpaka
             [[maybe_unused]] Stream_t stream)
         {
             // hipMallocAsync is implemented only in ROCm 5.2.0 and later.
-#    if HIP_VERSION >= 50'200'000
+#    if HIP_VERSION >= 50'600'000
             return ::hipMallocAsync(devPtr, size, stream);
+#    elif HIP_VERSION >= 50'200'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;