diff --git a/CHANGELOG.md b/CHANGELOG.md index 393f04faf..982c1c511 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 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. diff --git a/rocprim/include/rocprim/device/config_types.hpp b/rocprim/include/rocprim/device/config_types.hpp index 2a7c7115b..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) + + // 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) 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));