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

alpaka-group · Dec 14, 2023 · 39aa327 · 39aa327
1 parent 76a20bd
commit 39aa327
Showing 1 changed file with 24 additions and 2 deletions.
diff --git a/include/alpaka/core/ApiHipRt.hpp b/include/alpaka/core/ApiHipRt.hpp
@@ -182,8 +182,18 @@ 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 +300,20 @@ 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;