Develop stream 2025-01-23 by Snektron · Pull Request #678 · ROCm/rocPRIM

Snektron · 2025-01-23T14:55:31Z

This pull request is based on top of #674, that one should be merged first.

This pull request contains our changes from the past few months. A summary of the most notable changes:

Changes required for updating hipCUB and rocThrust to CCCL 2.6.0.
Add rocPRIM versions of UnrolledThreadLoad and UnrolledThreadCopy
Device histogram workaround for dynamic shared memory not working on Windows
Removed ROCPRIM_TARGET_ARCH, which allowed users to specialize rocPRIM for a particular architecture. This is now mostly resolved by the new ROCPRIM_TARGET_(CDNA3|CDNA2|CDNA1|RDNA3|RDNA2|RNDA1|GCN5) macros, which are automatically determined from the (device) compilation target. This does not affect any downstream user code, but does affect rocThrust a little bit.
Undeprecated and improved rocprim::thread_load and rocprim::thread_store. Most of the inline assembly has been removed, which should make the compiler able to optimize these operations better. These functions allow programmers to load memory more efficiently by providing hints for how memory will be used, for example, rocprim::thread_load<rocprim::load_nontemporal /* previously: rocprim::load_cs */>(ptr) can be used effectively to improve memory throughput on MI300.
Added rocprim::merge_inplace. This algorithm is like rocprim::merge, but (as the name suggests) can operate in-place. While rocprim::merge is usually faster, this algorithm can be used when not enough memory is available for an out-of-place buffer.
Adopt visibility rules to avoid linking problems when linking with separate versions of rocPRIM.
Add 128-bit integer benchmarking and tuning to all algorithms.

spolifroni-amd

Small changes needed to the changelog.

Snektron · 2025-01-23T19:41:37Z

Thanks for the quick review, @spolifroni-amd, your work is appreciated :) I'll update the pull request tomorrow.

Snektron · 2025-01-24T11:38:48Z

Applied the suggestions

umfranzw · 2025-02-07T20:29:23Z

There's a CI test failure for test_device_merge_inplace on gfx942. I'm able to reproduce it, and I'm trying to narrow things down. It looks like the failure happens for the largest test size used in the DeviceMergeInPlace test case:
DeviceMergeInplaceParams<int32_t, random_data_generator<int32_t, 2, 4>, random_data_generator<int32_t, 2, 4>, large_sizes>

Naraenda · 2025-02-11T09:00:57Z

@umfranzw This might be part of the hard memory usage limiter on the CI. IIRC, it will watch for processess that use more than a certain amount of memory and kill those off. If we know the limit we can set the limit defined in ae4709f, better.

umfranzw · 2025-02-11T15:05:31Z

@umfranzw This might be part of the hard memory usage limiter on the CI. IIRC, it will watch for processess that use more than a certain amount of memory and kill those off. If we know the limit we can set the limit defined in ae4709f, better.

@Naraenda I ran some experiments, and it looks like if I allocate 2^35 total bytes, everything runs fine, but if I allocate 2^36 the process gets killed. I believe if you modify your existing check in ae4709f to look at total_bytes instead of size_total, it should work fine.

umfranzw

I think this looks good - I just found a few small things I wanted to mention.

Naraenda · 2025-02-12T09:23:32Z

@Naraenda I ran some experiments, and it looks like if I allocate 2^35 total bytes, everything runs fine, but if I allocate 2^36 the process gets killed. I believe if you modify your existing check in ae4709f to look at total_bytes instead of size_total, it should work fine.

I've updated the commit that introduced this: 46a1f80

umfranzw

I think this looks ready to go.

…3.18 (ROCm#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).

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.

…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.

…get_storage_size as ROCPRIM_HOST_DEVICE."

…support"

…sts to grow too big

NB4444 · 2025-02-14T09:51:51Z

I resolved the merge conflicts.

rocm_install adds all libraries to ${PROJECT_NAME_LOWER}-targets (i.e. rocprim-targets.cmake) unless a custom export file is specified.

* 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. * 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 * Fix static asserts and remove template specialization * Fix tests triggering static asserts * Fix benchmarks triggering static asserts * Set explicit return types * Resolve "Investigate RLE (non trivial runs) failure" * Revert "Fix benchmarks triggering static asserts" This reverts commit a1634a9. * Revert "Fix tests triggering static asserts" This reverts commit 24a36cb. * Revert "Fix static asserts and remove template specialization" This reverts commit b2089ee. * Fix formatting and copyright date * Temp fix: histogram dynamic memory failure on windows * add ROCPRIM_ENABLE_ASSERTS to cmake to enable asserts in release build set ROCPRIM_ENABLE_ASSERTS for CI tests * Add missing docs for `block_run_length_decode` * Resolve "Add documentation and test for out-of-range behaviour on merge path search" add test for out of bounds diagonal test for rocprim::merge_path_search document behaviour for out of bounds diagonal test for rocprim::merge_path_search * Resolve "Replace __launch_bounds__ with ROCPRIM_LAUNCH_BOUNDS in rocPRIM API" * Update rmake.py copyright * fix double buffers in device radix sort * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in tests with device_ptr" * fix device adjacent difference tuning benchmark * Fix warp dpp reduction * Fix static assertion for int128 in ONNX Runtime with GNU extensions * fix GenerateResourceSpec for gfx90a * Resolve "ROCPRIM_TARGET_ARCH deprecation" * Fix warnings when generating sphinx docs * Performance improvements with shuffled warp exchange * log more information during configure * Update documentation generation instructions * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 4 * Resolve "search_n optimization" * CI changes to speedup development * Remove select_arch and select_arch_case * fix missing equality operator and function constness * Print arguments and timing info in `batch_memcpy_func` when `debug_synchronous` is true * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 3 * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 1 * Resolve "Replace hipMalloc, hipMallocHelper, and hipFree in benchmarks and tests with device_ptr" Part 2 * Fix rules overwrite for other then MR pipelines * Resolve "Add stream, memset, array, and unique_ptr support to device_ptr" * Resolve "Reduce amount of test runs based on environment variable." * Resolve "Get clarity on skipping of adjacent difference tests" * Remove windows CI check * Resolve "Add UnrolledThreadLoad and UnrolledThreadCopy functions" * Resolve "General changes CCCL 2.6" * Reduce CI artifact retention time * Resolve "Add thread_load_volatile_pointer to rocprim" * Fix test:install on rocm/dev-ubuntu * Resolve "Nth element - Investigate smaller index sizes where possible" * Enable int128 tests on Windows * Resolve "Benchmark and tuning for 128-bit integers" * Resolve "Verify sorting on device" * Resolve "Expose lookback_scan_state::create and lookback_scan_state::get_storage_size as ROCPRIM_HOST_DEVICE." * add naive in-place api for device merge * To remove error in CI change it to AMDGPU_TARGETS * fix device_ptr assign move and void types * Resolve "Add block- and warp-level inclusive_scan with initial value support" * Fix C++ includes showing up in the docs * Resolve "Performance regressions during tuning: Adjacent find" * CI skip label fix * fix doxygen warning in merge.rst * Resolve "Autotuning for segmented reduce" * Resolve "Update thread_store for parity" * Resolve "Fix compile errors in "rocprim-release-tune" mode" * Documentation to CI upload * Resolve "Remove ASM from thread_load and thread_store" * Add documentation steps to developer guidelines * Resolve "Refactor lookback_scan_state API" * Adopt symbol visibility rules in rocPRIM * remove nm debug messages from cmake * Remove extra dashes in the documentation its \param and \tparam * update changelog * fix(test_device_merge_inplace.cpp): prevent merge inplace tests to tests to grow too big * Fix test_linking: install the test's libraries too * chore(rmake.py): bump copyright * fix(example_type_traits_interface.cpp): explicit return from main * Fix unintentionally exported shared libs of test_linking rocm_install adds all libraries to ${PROJECT_NAME_LOWER}-targets (i.e. rocprim-targets.cmake) unless a custom export file is specified. --------- Co-authored-by: Wayne Franz <wayfranz@amd.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: Jaap Blok <jaap@streamhpc.com> Co-authored-by: Nick Breed <nick@streamhpc.com> Co-authored-by: Bence Parajdi <bence@streamhpc.com> Co-authored-by: Sander Bos <sander@streamhpc.com> Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: Anton Gorenko <anton@streamhpc.com> Co-authored-by: Milo Lurati <milo@streamhpc.com> Co-authored-by: Saiyang Zhang <saiyang@streamhpc.com> Co-authored-by: Enrico Degregori <enrico@streamhpc.com> Co-authored-by: Borys Petrov <borys@streamhpc.com> [ROCm/rocPRIM commit: 678701d]

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

spolifroni-amd suggested changes Jan 23, 2025

View reviewed changes

Snektron force-pushed the develop_stream_20250123 branch from e6cb034 to 5149685 Compare January 24, 2025 12:10

spolifroni-amd approved these changes Jan 27, 2025

View reviewed changes

NB4444 force-pushed the develop_stream_20250123 branch from d9ffa89 to 25c571f Compare February 7, 2025 06:54

umfranzw mentioned this pull request Feb 7, 2025

Fix static assertion for int128 in ONNX Runtime with GNU extensions #683

Merged

umfranzw reviewed Feb 11, 2025

View reviewed changes

Comment thread example/example_type_traits_interface.cpp Resolved

Comment thread rmake.py Resolved

Naraenda force-pushed the develop_stream_20250123 branch from 25c571f to a5361e7 Compare February 12, 2025 09:19

umfranzw approved these changes Feb 13, 2025

View reviewed changes

NB4444 force-pushed the develop_stream_20250123 branch from 281a187 to a5361e7 Compare February 14, 2025 09:37

umfranzw and others added 10 commits February 14, 2025 09:41

Resolve "Add hipGraph compatibility for device run length encode"

4903d4a

Resolve "Investigate RLE (non trivial runs) failure"

29970bf

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

cfeb4a4

…ode"

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

d3aa223

…ded types and third-party's types"

MI200 & MI300 tuning

2e8622f

Optimizations for MI300 unique performance

b60bbab

improve type traits interface docs

84c6d8d

cenxuantian and others added 24 commits February 14, 2025 09:47

Resolve "Expose lookback_scan_state::create and lookback_scan_state::…

d8101b4

…get_storage_size as ROCPRIM_HOST_DEVICE."

add naive in-place api for device merge

b80bba6

To remove error in CI change it to AMDGPU_TARGETS

1189d00

fix device_ptr assign move and void types

e3b35ce

Resolve "Add block- and warp-level inclusive_scan with initial value …

d4a7092

…support"

Fix C++ includes showing up in the docs

ae1ce56

Resolve "Performance regressions during tuning: Adjacent find"

9995fa4

CI skip label fix

42b5518

fix doxygen warning in merge.rst

2764058

Resolve "Autotuning for segmented reduce"

2add4a3

Resolve "Update thread_store for parity"

e0617cb

Resolve "Fix compile errors in "rocprim-release-tune" mode"

0c36f82

Documentation to CI upload

bc6605b

Resolve "Remove ASM from thread_load and thread_store"

bf8985d

Add documentation steps to developer guidelines

2d96209

Resolve "Refactor lookback_scan_state API"

033ced9

Adopt symbol visibility rules in rocPRIM

a7dec4b

remove nm debug messages from cmake

b38f541

Remove extra dashes in the documentation its \param and \tparam

44d19af

update changelog

0b2f462

fix(test_device_merge_inplace.cpp): prevent merge inplace tests to te…

bc8ae44

…sts to grow too big

Fix test_linking: install the test's libraries too

6ba76f6

chore(rmake.py): bump copyright

65c355d

fix(example_type_traits_interface.cpp): explicit return from main

c1e9924

NB4444 force-pushed the develop_stream_20250123 branch from a5361e7 to c1e9924 Compare February 14, 2025 09:50

stanleytsang-amd approved these changes Feb 20, 2025

View reviewed changes

Fix unintentionally exported shared libs of test_linking

72e1ff7

rocm_install adds all libraries to ${PROJECT_NAME_LOWER}-targets (i.e. rocprim-targets.cmake) unless a custom export file is specified.

stanleytsang-amd merged commit 678701d into ROCm:develop Feb 28, 2025

Conversation

Snektron commented Jan 23, 2025

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!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Snektron commented Jan 23, 2025

Uh oh!

Snektron commented Jan 24, 2025

Uh oh!

umfranzw commented Feb 7, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Naraenda commented Feb 11, 2025

Uh oh!

umfranzw commented Feb 11, 2025

Uh oh!

umfranzw left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Naraenda commented Feb 12, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

umfranzw left a comment

Choose a reason for hiding this comment

Uh oh!

NB4444 commented Feb 14, 2025

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

16 participants

umfranzw commented Feb 7, 2025 •

edited

Loading

Naraenda commented Feb 12, 2025 •

edited

Loading