From 1effc36786d886b6fd8949b2b4feff56a264da24 Mon Sep 17 00:00:00 2001 From: Wayne Franz Date: Wed, 12 Feb 2025 18:31:18 -0500 Subject: [PATCH 1/4] Support getting device ID from stream hipStreamLegacy (#687) rocprim::get_device_from_stream must retreive the device ID in different ways depending on whether a user-created stream or a built-in stream type is being used. This change modifies the check for built-in stream types so that it includes hipStreamLegacy. This prevents the function from calling hipGetStreamDeviceId(hipStreamLegacy), which currently causes a segmentation fault. --- rocprim/include/rocprim/device/config_types.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/rocprim/include/rocprim/device/config_types.hpp b/rocprim/include/rocprim/device/config_types.hpp index 2a7c7115b..e7af2f99e 100644 --- a/rocprim/include/rocprim/device/config_types.hpp +++ b/rocprim/include/rocprim/device/config_types.hpp @@ -350,7 +350,7 @@ inline hipError_t get_device_arch(int device_id, target_arch& arch) inline hipError_t get_device_from_stream(const hipStream_t stream, int& device_id) { static constexpr hipStream_t default_stream = 0; - if(stream == default_stream || stream == hipStreamPerThread) + if(stream == default_stream || stream == hipStreamPerThread || stream == hipStreamLegacy) { const hipError_t result = hipGetDevice(&device_id); if(result != hipSuccess) From 14f6f8d734eeb7c8c052580ec9165be5e901ec86 Mon Sep 17 00:00:00 2001 From: Wayne Franz Date: Thu, 13 Feb 2025 14:56:59 -0500 Subject: [PATCH 2/4] Add hipStreamLegacy HIP version check The hipStreamLegacy stream handle was added in HIP 6.1.0. Add a check to ensure that we don't reference it if rocPRIM is being built using an older version of HIP. --- rocprim/include/rocprim/device/config_types.hpp | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/rocprim/include/rocprim/device/config_types.hpp b/rocprim/include/rocprim/device/config_types.hpp index e7af2f99e..b1bfa55c7 100644 --- a/rocprim/include/rocprim/device/config_types.hpp +++ b/rocprim/include/rocprim/device/config_types.hpp @@ -350,7 +350,15 @@ inline hipError_t get_device_arch(int device_id, target_arch& arch) inline hipError_t get_device_from_stream(const hipStream_t stream, int& device_id) { static constexpr hipStream_t default_stream = 0; - if(stream == default_stream || stream == hipStreamPerThread || stream == hipStreamLegacy) + + // hipStreamLegacy is supported in HIP >= 6.1.0 +#if (HIP_VERSION_MAJOR >= 6 && HIP_VERSION_MINOR >= 1) + const bool is_legacy_stream = (stream == hipStreamLegacy); +#else + const bool is_legacy_stream = false; +#endif + + if (stream == default_stream || stream == hipStreamPerThread || is_legacy_stream); { const hipError_t result = hipGetDevice(&device_id); if(result != hipSuccess) From 84b35b88393ddbacdb276652277f6cc8f33afd11 Mon Sep 17 00:00:00 2001 From: Wayne Franz Date: Thu, 13 Feb 2025 17:02:23 -0500 Subject: [PATCH 3/4] Add get_device_id_from_stream test for hipStreamLegacy Add a test case to ensure that passing hipStreamLegacy to rocprim::get_device_id_from_stream works correctly. Updated the changelog with a note to reflect the fact that hipStreamLegacy can now be used without issue. --- CHANGELOG.md | 1 + test/rocprim/test_config_dispatch.cpp | 6 ++++++ 2 files changed, 7 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 393f04faf..dc6c1dd13 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -44,6 +44,7 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec * Fixed incorrect 128-bit signed and unsigned integers type traits. * Fixed compilation issue when `rocprim::radix_key_codec<...>` is specialized with a 128-bit integer. * Fixed the warp-level reduction `rocprim::warp_reduce.reduce` DPP implementation to avoid undefined intermediate values during the reduction. +* Fixed an issue that was causing a segmentation fault when hipStreamLegacy was passed to some API functions. ### Upcoming changes * Using the initialisation constructor of `rocprim::reverse_iterator` will throw a deprecation warning. It will be marked as explicit in the next major release. diff --git a/test/rocprim/test_config_dispatch.cpp b/test/rocprim/test_config_dispatch.cpp index 62ebda06b..f7ae1813c 100644 --- a/test/rocprim/test_config_dispatch.cpp +++ b/test/rocprim/test_config_dispatch.cpp @@ -98,6 +98,12 @@ TEST(RocprimConfigDispatchTests, DeviceIdFromStream) HIP_CHECK(get_device_from_stream(hipStreamPerThread, result)); ASSERT_EQ(result, device_id); + // hipStreamLegacy support was added in ROCm 6.1.0 +#if (HIP_VERSION_MAJOR >= 6 && HIP_VERSION_MINOR >= 1) + HIP_CHECK(get_device_from_stream(hipStreamLegacy, result)); + ASSERT_EQ(result, device_id); +#endif + hipStream_t stream; HIP_CHECK(hipStreamCreate(&stream)); HIP_CHECK(get_device_from_stream(stream, result)); From c05b6f148ca95147cbad31e124545cb29d223c5f Mon Sep 17 00:00:00 2001 From: Wayne Franz Date: Thu, 20 Feb 2025 11:20:14 -0500 Subject: [PATCH 4/4] Update CHANGELOG.md Co-authored-by: Jeffrey Novotny --- CHANGELOG.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index dc6c1dd13..982c1c511 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -44,7 +44,7 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec * Fixed incorrect 128-bit signed and unsigned integers type traits. * Fixed compilation issue when `rocprim::radix_key_codec<...>` is specialized with a 128-bit integer. * Fixed the warp-level reduction `rocprim::warp_reduce.reduce` DPP implementation to avoid undefined intermediate values during the reduction. -* Fixed an issue that was causing a segmentation fault when hipStreamLegacy was passed to some API functions. +* Fixed an issue that caused a segmentation fault when `hipStreamLegacy` was passed to some API functions. ### Upcoming changes * Using the initialisation constructor of `rocprim::reverse_iterator` will throw a deprecation warning. It will be marked as explicit in the next major release.