From a1309ae7d8b23846c4a3e7c7a55e854dfb3d6a78 Mon Sep 17 00:00:00 2001 From: Wayne Franz Date: Thu, 9 Jul 2026 15:39:41 -0400 Subject: [PATCH] Avoid zero-sized allocation when HMM is enabled While hipMalloc allows you to allocate buffers of size 0, hipMallocManaged does not. This was causing HipcubDeviceSegmentedRadixSort's SortKeysEmptyData test to fail. This change avoids the zero-sized allocation in the test, and adds a warning message that's printed to stderr whenever we attempt to allocate zero bytes using the hipMallocHelper function. --- projects/hipcub/test/hipcub/common_test_header.hpp | 2 ++ .../test_hipcub_device_segmented_radix_sort.hpp | 11 ++++++++--- 2 files changed, 10 insertions(+), 3 deletions(-) diff --git a/projects/hipcub/test/hipcub/common_test_header.hpp b/projects/hipcub/test/hipcub/common_test_header.hpp index 093bdafea540..737c29142cda 100755 --- a/projects/hipcub/test/hipcub/common_test_header.hpp +++ b/projects/hipcub/test/hipcub/common_test_header.hpp @@ -160,6 +160,8 @@ hipError_t hipMallocHelper(T** devPtr, size_t size) { if (use_hmm()) { + if (size == 0) + std::cerr << "Warning: attempting to allocate a buffer of size 0 with hipMallocManaged." << std::endl; return hipMallocManaged((void**)devPtr, size); } else diff --git a/projects/hipcub/test/hipcub/test_hipcub_device_segmented_radix_sort.hpp b/projects/hipcub/test/hipcub/test_hipcub_device_segmented_radix_sort.hpp index d32259dc4215..470410994aed 100644 --- a/projects/hipcub/test/hipcub/test_hipcub_device_segmented_radix_sort.hpp +++ b/projects/hipcub/test/hipcub/test_hipcub_device_segmented_radix_sort.hpp @@ -266,8 +266,12 @@ inline void sort_keys_empty_data() offsets[0] = 0; offsets[1] = 0; - key_type* d_keys; - HIP_CHECK(test_common_utils::hipMallocHelper(&d_keys, size * sizeof(key_type))); + key_type* d_keys = nullptr; + // hipMallocManaged will not allocate buffers of size 0. + // If HMM is enabled and size is 0, leave d_keys set to nullptr. + if (!(size == 0 && test_common_utils::use_hmm())) + HIP_CHECK(test_common_utils::hipMallocHelper(&d_keys, size * sizeof(key_type))); + HIP_CHECK(hipMemcpy(d_keys, keys_input.data(), size * sizeof(key_type), @@ -337,8 +341,9 @@ inline void sort_keys_empty_data() hipMemcpyDeviceToHost)); HIP_CHECK(hipFree(d_temporary_storage)); - HIP_CHECK(hipFree(d_keys)); HIP_CHECK(hipFree(d_offsets)); + if (d_keys) + HIP_CHECK(hipFree(d_keys)); // Output should not have changed ASSERT_NO_FATAL_FAILURE(test_utils::assert_eq(keys_output, keys_input));