Develop Stream: fix warp dpp reduction by Beanavil · Pull Request #660 · ROCm/rocPRIM

Beanavil · 2024-12-02T13:34:33Z

Fix intra-warp reduction. By using rotations instead of shifts in warp_reduce_dpp.hpp we make sure that no registers are left with invalid values. We thus substitute the dpp_ctrl values of 0x114 and 0x118 for 0x124 and 0x128, respectively.
Fix inter-warp reduction. After performing the intra-warp reduction, a second round of reduction is done to reduce the warps' results. In this case, we still may have some threads reading invalid values because currently there is divergence in this reduction (only the first warps_no_ threads will participate in the final reduction). So potentially many threads will be disabled and then the dpp instruction will resort to what bound_ctrl dictates just like described before. Thus, to ensure that all threads read and write valid values, we remove this divergence from block_reduce_warp_reduce.hpp.
Additionally, to prevent the same segfaults happening when the BlockSize used is less than the hardware warp size, we added a fallback to the warp_shuffle in that case.

spolifroni-amd · 2024-12-02T15:21:55Z

+
+### Resolved issues
+
+* Fixed the warp-level reduction `rocprim::warp_reduce.reduce` DPP implementation to avoid undefined intermediate values during the reduction.


If this has a github issue number, it can be added here.

There is none afaik, this is for an internal ticket

spolifroni-amd

If there is a github issue number for the resolved issue, adding it to the changelog would be useful for users. Otherwise, all good.

stanleytsang-amd

Hi @Beanavil, can you target this PR to the release-staging/rocm-rel-6.4 branch? Thank you.

Co-authored-by: Robin Voetter <robin@streamhpc.com>

NguyenNhuDi

lgtm

umfranzw

This change looks good to me.

* 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> --------- 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>

* 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> --------- 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> [ROCm/rocPRIM commit: f1d714d]

* 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]

Beanavil requested review from a team, RobsonRLemos, lawruble13, stanleytsang-amd and umfranzw as code owners December 2, 2024 13:34

Beanavil force-pushed the develop_stream_fix_dpp_reduction branch from 30453d9 to 947cbae Compare December 2, 2024 13:38

spolifroni-amd reviewed Dec 2, 2024

View reviewed changes

spolifroni-amd approved these changes Dec 2, 2024

View reviewed changes

stanleytsang-amd requested changes Dec 2, 2024

View reviewed changes

Comment thread CHANGELOG.md Resolved

Beanavil force-pushed the develop_stream_fix_dpp_reduction branch from 947cbae to cdcdbb5 Compare December 3, 2024 11:52

Beanavil changed the base branch from develop to release-staging/rocm-rel-6.4 December 3, 2024 11:53

Beanavil force-pushed the develop_stream_fix_dpp_reduction branch from cdcdbb5 to 03c4b53 Compare December 3, 2024 11:53

Fix warp dpp reduction

03c4b53

Co-authored-by: Robin Voetter <robin@streamhpc.com>

NguyenNhuDi approved these changes Dec 3, 2024

View reviewed changes

umfranzw approved these changes Dec 4, 2024

View reviewed changes

stanleytsang-amd merged commit 277a7cb into ROCm:release-staging/rocm-rel-6.4 Dec 5, 2024

Navigation Menu

Repository navigation

Develop Stream: fix warp dpp reduction - #660

Develop Stream: fix warp dpp reduction#660
stanleytsang-amd merged 1 commit into
ROCm:release-staging/rocm-rel-6.4from
StreamHPC:develop_stream_fix_dpp_reduction

Beanavil commented Dec 2, 2024

Uh oh!

spolifroni-amd Dec 2, 2024

Uh oh!

Beanavil Dec 2, 2024

Uh oh!

spolifroni-amd left a comment

Uh oh!

stanleytsang-amd left a comment

Uh oh!

Uh oh!

NguyenNhuDi left a comment

Uh oh!

umfranzw left a comment

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants


		### Resolved issues

		* Fixed the warp-level reduction `rocprim::warp_reduce.reduce` DPP implementation to avoid undefined intermediate values during the reduction.

Conversation

Beanavil commented Dec 2, 2024

Uh oh!

spolifroni-amd Dec 2, 2024

Choose a reason for hiding this comment

Uh oh!

Beanavil Dec 2, 2024

Choose a reason for hiding this comment

Uh oh!

spolifroni-amd left a comment

Choose a reason for hiding this comment

Uh oh!

stanleytsang-amd left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

NguyenNhuDi left a comment

Choose a reason for hiding this comment

Uh oh!

umfranzw left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants