Skip to content

Accept hip stream legacy #689

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
umfranzw merged 4 commits into from
Feb 20, 2025
Merged

Accept hip stream legacy #689

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
1effc36
Support getting device ID from stream hipStreamLegacy (#687)
umfranzw Feb 12, 2025
14f6f8d
Add hipStreamLegacy HIP version check
umfranzw Feb 13, 2025
84b35b8
Add get_device_id_from_stream test for hipStreamLegacy
umfranzw Feb 13, 2025
c05b6f1
Update CHANGELOG.md
umfranzw Feb 20, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions CHANGELOG.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Comment thread
umfranzw marked this conversation as resolved.
Resolved

### 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.
Expand Down
10 changes: 9 additions & 1 deletion rocprim/include/rocprim/device/config_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
6 changes: 6 additions & 0 deletions test/rocprim/test_config_dispatch.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand Down