MI300 radix sort, run length encode, and reduce by key optimizations by Snektron · Pull Request #674 · ROCm/rocPRIM

Snektron · 2025-01-17T17:23:29Z

This changes a number of things to improve performance of the "unique" operation - which is typically implemented using radix sort and run length encode. Summary of changes:

hipGraph support for run length encode
Autotuning for run length encode
General tuning for MI200 & MI300
Targeted improvements for lookback scan & onesweep

Additionally included is some supporting code related to type traits. The next pull request will add (autotune) benchmarking for 128-bit types, but this pull request already contains the tuning results.

I also added in the commit from Wayne to fix the LINK_LANGUAGE issue. Maybe its better to add that in a separate pull request though.

spolifroni-amd

The documentation needs some work. If the code needs to go in asap, an issue can be opened for the doc work.

Snektron · 2025-01-20T17:46:55Z

@spolifroni-amd thanks, I've updated the documentation with your suggestions in verbatim.

spolifroni-amd

Found a small typo. It's an easy fix and I'll approve so no one is waiting on me.

spolifroni-amd · 2025-01-20T18:05:48Z

@spolifroni-amd thanks, I've updated the documentation with your suggestions in verbatim.

I found a small typo. Fix that and it's all good! I approved in advance so you don't have to wait.

Snektron · 2025-01-20T18:08:01Z

All good, I've already updated it

stanleytsang-amd · 2025-01-22T06:30:46Z

@Snektron Our existing math-ci (which for some reason, does not seem to be showing on GitHub this time), is showing the 'hip/hip_runtime.h' file not found error. I am unable to reproduce it with the same conditions so far. Still looking into it.

Snektron · 2025-01-22T11:12:53Z

Is it running the most recent version? The issues related to that seem to be fixed on the public CI at least.

Snektron · 2025-01-23T13:59:15Z

Rebased Wayne's changes into this branch

stanleytsang-amd · 2025-01-24T16:51:46Z

@Snektron thanks for rebasing Wayne's changes, looks like CI is building properly now. There is a single test failure, block_histogram, but that's a known failure due to compiler bug.
Azure DevOps CI reports more failures with block_adjacent_difference and block_discontinuity. I don't see any clear pattern or commonality with the failures, and the ROCm stack used in DevOps CI is quite bleeding edge (ie. they are not best-known configuration builds) so I don't know how much weight we want to give those failures at this time.

That being said, the Pytorch team has now expressed a desire to get these changes in for 6.4 if possible. Could you retarget this PR to the release-staging/rocm-rel-6.4 branch?

…ode"

…ded types and third-party's types"

A missing call to __syncthreads() caused this algorithm to occasionally produce an invalid index and cause a crash. Instead of just adding that call, we can remove the use of ordered block id by the regular block id, which is already ordered per HSA guarantees. There still remains an issue with a hang, but this seems a runtime issue and will be investigated separately.

Snektron · 2025-01-24T17:25:25Z

Done.

Azure DevOps CI reports more failures with block_adjacent_difference and block_discontinuity. I don't see any clear pattern or commonality with the failures, and the ROCm stack used in DevOps CI is quite bleeding edge (ie. they are not best-known configuration builds) so I don't know how much weight we want to give those failures at this time.

I think we may be able to fix at least one of them. There seems to be something wrong with the setup of the discontinuity tests, but I haven't looked into it too deeply.

stanleytsang-amd · 2025-01-24T17:40:20Z

Done.

Azure DevOps CI reports more failures with block_adjacent_difference and block_discontinuity. I don't see any clear pattern or commonality with the failures, and the ROCm stack used in DevOps CI is quite bleeding edge (ie. they are not best-known configuration builds) so I don't know how much weight we want to give those failures at this time.

I think we may be able to fix at least one of them. There seems to be something wrong with the setup of the discontinuity tests, but I haven't looked into it too deeply.

I see. If you can or want to spare a couple of days to look at it, I think we have time. Let me check and see when the release-staging branch closes for 6.4...after it closes we're no longer able to freely merge in PRs without paperwork

Snektron · 2025-01-30T10:29:43Z

I've managed to workaround the failure in the block discontinuity test... As written in the comment, there is a miscompilation in LLVM. It should be fixed in LLVM's master branch, but that commit has not yet made its way to the ROCm fork. See llvm/llvm-project#121745.

Snektron · 2025-01-31T14:44:53Z

Checked a preview version of the compiler, the issue seems to be no longer there.

stanleytsang-amd · 2025-01-31T19:23:46Z


 Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projects/rocPRIM/en/latest/](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/).

+## (Unreleased) rocPRIM 3.4.x for ROCm 6.4.x


This will be part of the 6.4.0 release, so you can add these entries to the 6.4.0 section.

Im away this week. @Naraenda?

Added the entries to 6.4.0.

umfranzw

I think the code changes look good. Provided we've addressed all the test failures that are within our control (unrelated to compiler bugs), I think this looks ready to go.

stanleytsang-amd · 2025-02-03T18:52:06Z

@umfranzw 2 out of the 3 test failures are accounted for - block_histogram due to compiler issue, block_discontinuity should/will be fixed in an upcoming ROCm 6.4 compiler build. The third block_adjacent_difference I am unable to repro in any BKC build so far and might just be due to using a staging build of the compiler - I think we can see if CQE can reproduce the failure and we'll go from there.

* Add atomic_add overload for double (#652) The rocprim::histogram_even algorithm uses rocprim::atomic_add. There are atomic_add overloads for all of the non-half-type primitive data types we support, with the exception of double. Because of this, if you called rocprim::histogram_even and used double as the counter type, you'd get a compilation error. This changes just adds an atomic_add overload that accepts double, and tacks on a few unit test cases that exercise histogram_even with float and double counters. * Make temp_storage functions __host__ and __device__ (#658) * Enable CMake HIP language (#655) * Restart on HIP language support * linking changes * rearrange --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> * Fix for rocprim::batch_copy when using iterators (#662) * Fix for using iterators in batch_copy * Add indirect_iterator test to batch_memcpy * Add specific iterator test based on reproducer, with iterators in iterators. * Fix warp dpp reduction (#660) Co-authored-by: Robin Voetter <robin@streamhpc.com> * Workaround for shared memory write ordering compiler issue (#672) This change implements a workaround to fix block_histogram while we wait for a compiler fix. The block_histogram_sort algorithm contains a loop that initializes values in a shared memory struct. Then it runs the flag_heads() algorithm using that struct. It appear that in some cases, the optimizer is reordering things such that shared memory writes in flag_heads() run before the loop initialization code completes. Changing the loop condition and removing the inner "if" statement seems to prevent the optimizer from doing this. From an algorithmic standpoint, this should be functionally equivalent to the way the code was before the change. * Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <3.18 (#667) Recently, we added CMake HIP Language support. That change introduced generator expressions that use the LINK_LANGUAGE keyword. This keyword is only supported in CMake 3.18+. The minimum CMake version we require is 3.16. As a result, if you built with CMake 3.16 - 3.18, you'd get a build failure. This change adds checks around all the spots where the LINK_LANGUAGE keyword is used to avoid this problem. It also adds a description of the CMake HIP Language support to the changelog (previously overlooked). * Fix Windows CMake link issue (#675) A recent cmake change added a condition check that avoids evaluating a LINK_LANGUAGE generator expression if HIP_USECXX is OFF (since LINK_LANGUAGE is unsupported in CMake < 3.18). However, it didn't account for the fact that the LINK_LANGUAGE expression is embedded in a conditional generator expression. This change adds an "else" block that ensures the "false" side of the conditional generator expression is still evaluated when the LINK_LANGUAGE code is skipped. This fixes some build issues on Windows. Also updates the readme to include recent changes to the way projects must link to rocPRIM. * MI300 radix sort, run length encode, and reduce by key optimizations (#674) * Resolve "Add hipGraph compatibility for device run length encode" * Resolve "Investigate RLE (non trivial runs) failure" * Resolve "Config tuning and dynamic dispatch for device run length encode" * Resolve "Implement type traits interface for rocPRIM's compiler-extended types and third-party's types" * remove ordered block id from run length encode A missing call to __syncthreads() caused this algorithm to occasionally produce an invalid index and cause a crash. Instead of just adding that call, we can remove the use of ordered block id by the regular block id, which is already ordered per HSA guarantees. There still remains an issue with a hang, but this seems a runtime issue and will be investigated separately. * MI200 & MI300 tuning * Optimizations for MI300 unique performance * improve type traits interface docs * Move changlog items to 6.4.0 --------- Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com> Co-authored-by: Yung-sheng Tu <yung-sheng@streamhpc.com> Co-authored-by: Cenxuan Tian <cenxuan@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> * Add tuned configs for gfx1200 and gfx1201 (#681) Add tuned configs for gfx1200 and gfx1201 for some algorithms. A few algorithms still require some additional work to get tuning working. These will be added in a follow-up commit. * Fix static assertion for int128 in ONNX Runtime with GNU extensions (#683) Co-authored-by: Anton Gorenko <anton@streamhpc.com> * Add remaining gfx1201 tuning configs (#684) This change adds the remaining tuning configs for gfx1201. * Accept hip stream legacy (#689) * 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. * 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. * 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. * Update CHANGELOG.md Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> --------- Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> * Workaround memory allocation issues on Windows for large size test cases (#690) Radix sort and merge sort test cases that use the largest data sizes hang or cause an SEH exception on Windows when using the latest version of the HIP SDK. Unfortunately, the HIP SDK issue interferes with our current strategies for skipping tests that require too much memory for a device (querying available device memory or examining the return value from hipMalloc), making them unreliable. While we wait for a fix, this change introduces a workaround by reducing the maximum size these test cases can use on Windows. * used 678 version of test_device_run_length_encode * got rid of duplicated code * removed ; on if statement --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Beatriz Navidad Vilches <61422851+Beanavil@users.noreply.github.com> Co-authored-by: Robin Voetter <robin@streamhpc.com> Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com> Co-authored-by: Yung-sheng Tu <yung-sheng@streamhpc.com> Co-authored-by: Cenxuan Tian <cenxuan@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>

* Add atomic_add overload for double (#652) The rocprim::histogram_even algorithm uses rocprim::atomic_add. There are atomic_add overloads for all of the non-half-type primitive data types we support, with the exception of double. Because of this, if you called rocprim::histogram_even and used double as the counter type, you'd get a compilation error. This changes just adds an atomic_add overload that accepts double, and tacks on a few unit test cases that exercise histogram_even with float and double counters. * Make temp_storage functions __host__ and __device__ (#658) * Enable CMake HIP language (#655) * Restart on HIP language support * linking changes * rearrange --------- Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> * Fix for rocprim::batch_copy when using iterators (#662) * Fix for using iterators in batch_copy * Add indirect_iterator test to batch_memcpy * Add specific iterator test based on reproducer, with iterators in iterators. * Fix warp dpp reduction (#660) Co-authored-by: Robin Voetter <robin@streamhpc.com> * Workaround for shared memory write ordering compiler issue (#672) This change implements a workaround to fix block_histogram while we wait for a compiler fix. The block_histogram_sort algorithm contains a loop that initializes values in a shared memory struct. Then it runs the flag_heads() algorithm using that struct. It appear that in some cases, the optimizer is reordering things such that shared memory writes in flag_heads() run before the loop initialization code completes. Changing the loop condition and removing the inner "if" statement seems to prevent the optimizer from doing this. From an algorithmic standpoint, this should be functionally equivalent to the way the code was before the change. * Avoid evaluating CMake LINK_LANGUAGE generator expressions in CMake <3.18 (#667) Recently, we added CMake HIP Language support. That change introduced generator expressions that use the LINK_LANGUAGE keyword. This keyword is only supported in CMake 3.18+. The minimum CMake version we require is 3.16. As a result, if you built with CMake 3.16 - 3.18, you'd get a build failure. This change adds checks around all the spots where the LINK_LANGUAGE keyword is used to avoid this problem. It also adds a description of the CMake HIP Language support to the changelog (previously overlooked). * Fix Windows CMake link issue (#675) A recent cmake change added a condition check that avoids evaluating a LINK_LANGUAGE generator expression if HIP_USECXX is OFF (since LINK_LANGUAGE is unsupported in CMake < 3.18). However, it didn't account for the fact that the LINK_LANGUAGE expression is embedded in a conditional generator expression. This change adds an "else" block that ensures the "false" side of the conditional generator expression is still evaluated when the LINK_LANGUAGE code is skipped. This fixes some build issues on Windows. Also updates the readme to include recent changes to the way projects must link to rocPRIM. * MI300 radix sort, run length encode, and reduce by key optimizations (#674) * Resolve "Add hipGraph compatibility for device run length encode" * Resolve "Investigate RLE (non trivial runs) failure" * Resolve "Config tuning and dynamic dispatch for device run length encode" * Resolve "Implement type traits interface for rocPRIM's compiler-extended types and third-party's types" * remove ordered block id from run length encode A missing call to __syncthreads() caused this algorithm to occasionally produce an invalid index and cause a crash. Instead of just adding that call, we can remove the use of ordered block id by the regular block id, which is already ordered per HSA guarantees. There still remains an issue with a hang, but this seems a runtime issue and will be investigated separately. * MI200 & MI300 tuning * Optimizations for MI300 unique performance * improve type traits interface docs * Move changlog items to 6.4.0 --------- Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com> Co-authored-by: Yung-sheng Tu <yung-sheng@streamhpc.com> Co-authored-by: Cenxuan Tian <cenxuan@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> * Add tuned configs for gfx1200 and gfx1201 (#681) Add tuned configs for gfx1200 and gfx1201 for some algorithms. A few algorithms still require some additional work to get tuning working. These will be added in a follow-up commit. * Fix static assertion for int128 in ONNX Runtime with GNU extensions (#683) Co-authored-by: Anton Gorenko <anton@streamhpc.com> * Add remaining gfx1201 tuning configs (#684) This change adds the remaining tuning configs for gfx1201. * Accept hip stream legacy (#689) * 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. * 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. * 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. * Update CHANGELOG.md Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> --------- Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> * Workaround memory allocation issues on Windows for large size test cases (#690) Radix sort and merge sort test cases that use the largest data sizes hang or cause an SEH exception on Windows when using the latest version of the HIP SDK. Unfortunately, the HIP SDK issue interferes with our current strategies for skipping tests that require too much memory for a device (querying available device memory or examining the return value from hipMalloc), making them unreliable. While we wait for a fix, this change introduces a workaround by reducing the maximum size these test cases can use on Windows. * used 678 version of test_device_run_length_encode * got rid of duplicated code * removed ; on if statement --------- Co-authored-by: Wayne Franz <wayfranz@amd.com> Co-authored-by: Nick Breed <78807921+NB4444@users.noreply.github.com> Co-authored-by: Lauren Wrubleski <Lauren.Wrubleski@amd.com> Co-authored-by: Beatriz Navidad Vilches <61422851+Beanavil@users.noreply.github.com> Co-authored-by: Robin Voetter <robin@streamhpc.com> Co-authored-by: Beatriz Navidad Vilches <beatriz@streamhpc.com> Co-authored-by: Yung-sheng Tu <yung-sheng@streamhpc.com> Co-authored-by: Cenxuan Tian <cenxuan@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Jeffrey Novotny <jnovotny@amd.com> [ROCm/rocPRIM commit: 07fda10]

Snektron requested review from a team, RobsonRLemos, lawruble13, stanleytsang-amd and umfranzw as code owners January 17, 2025 17:23

spolifroni-amd suggested changes Jan 17, 2025

View reviewed changes

Comment thread docs/concepts/type_traits.rst Resolved

Comment thread docs/concepts/type_traits.rst Resolved

Comment thread docs/concepts/type_traits.rst Resolved

Comment thread docs/concepts/type_traits.rst Resolved

Comment thread docs/concepts/type_traits.rst Resolved

spolifroni-amd reviewed Jan 20, 2025

View reviewed changes

Comment thread docs/concepts/type_traits.rst Resolved

spolifroni-amd approved these changes Jan 20, 2025

View reviewed changes

Snektron force-pushed the mi300-unique-upstreaming branch from 803a434 to 01adbf7 Compare January 20, 2025 18:07

Snektron force-pushed the mi300-unique-upstreaming branch from 01adbf7 to ae1b2e3 Compare January 23, 2025 13:58

Snektron mentioned this pull request Jan 23, 2025

Develop stream 2025-01-23 #678

Merged

samjwu added the ci:testonly label Jan 23, 2025

Beanavil and others added 8 commits January 24, 2025 18:22

Resolve "Add hipGraph compatibility for device run length encode"

689c730

Resolve "Investigate RLE (non trivial runs) failure"

d30c3ee

Resolve "Config tuning and dynamic dispatch for device run length enc…

a6af335

…ode"

Resolve "Implement type traits interface for rocPRIM's compiler-exten…

dda045a

…ded types and third-party's types"

MI200 & MI300 tuning

2d5707b

Optimizations for MI300 unique performance

fa09122

improve type traits interface docs

81bf9cf

Snektron force-pushed the mi300-unique-upstreaming branch from ae1b2e3 to 81bf9cf Compare January 24, 2025 17:22

Snektron changed the base branch from develop to release-staging/rocm-rel-6.4 January 24, 2025 17:23

Snektron force-pushed the mi300-unique-upstreaming branch from 465d993 to 81bf9cf Compare January 31, 2025 14:44

stanleytsang-amd requested changes Jan 31, 2025

View reviewed changes

Move changlog items to 6.4.0

594012d

umfranzw approved these changes Feb 3, 2025

View reviewed changes

stanleytsang-amd merged commit fdd9777 into ROCm:release-staging/rocm-rel-6.4 Feb 3, 2025


		Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projects/rocPRIM/en/latest/](https://rocm.docs.amd.com/projects/rocPRIM/en/latest/).

		## (Unreleased) rocPRIM 3.4.x for ROCm 6.4.x

Conversation

Snektron commented Jan 17, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

spolifroni-amd left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Snektron commented Jan 20, 2025

Uh oh!

Uh oh!

spolifroni-amd left a comment

Choose a reason for hiding this comment

Uh oh!

spolifroni-amd commented Jan 20, 2025

Uh oh!

Snektron commented Jan 20, 2025

Uh oh!

stanleytsang-amd commented Jan 22, 2025

Uh oh!

Snektron commented Jan 22, 2025

Uh oh!

Snektron commented Jan 23, 2025

Uh oh!

stanleytsang-amd commented Jan 24, 2025

Uh oh!

Snektron commented Jan 24, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

stanleytsang-amd commented Jan 24, 2025

Uh oh!

Snektron commented Jan 30, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Snektron commented Jan 31, 2025

Uh oh!

stanleytsang-amd Jan 31, 2025

Choose a reason for hiding this comment

Uh oh!

Snektron Feb 2, 2025

Choose a reason for hiding this comment

Uh oh!

NB4444 Feb 3, 2025

Choose a reason for hiding this comment

Uh oh!

umfranzw left a comment

Choose a reason for hiding this comment

Uh oh!

stanleytsang-amd commented Feb 3, 2025

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

9 participants

Snektron commented Jan 17, 2025 •

edited

Loading

Snektron commented Jan 24, 2025 •

edited

Loading

Snektron commented Jan 30, 2025 •

edited

Loading