Add protections for 0-sized allocations in ROCm 5.2-5.5

alpaka-group · Dec 14, 2023 · 6993d3b · 6993d3b
1 parent 76a20bd
commit 6993d3b
Showing 1 changed file with 18 additions and 2 deletions.
diff --git 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;