Develop stream 2025 05 06 by NB4444 · Pull Request #731 · ROCm/rocPRIM

NB4444 · 2025-05-06T13:53:42Z

The changes we made to the benchmarks:

Applied new benchmark abstraction to increase maintainability.
Introduced device_ptr to the benchmarks.
Output more data in JSON output of the benchmarks.

Deprecation and removal related to major release:

Remove [[deprecated]]float_bit_mask and all uses of it from the type trait interface, a472e80.
Remove short_radix_bits in segmented_radix_sort_config_params, 38ac5a5.
Remove already deprecated functions, 4d265b7.
Move rocprim::detail::radix_key_codec_base into traits system, 7c664b3.

CCCL 2.7 parity changes:

Add virtual shared memory fallback to device_merge, ca3b0cd.
Add device-level inclusive_scan with initial value support, cff2e16.
Change default scan accumulator type to be in line with (hip) CUB, 5f7accb.
Match CUB's behavior in rocPRIM for device merge, 4239ef5.
Create tests for rocPRIM's bit_cast, 668f913.

Performance optimizations:

Implement tuning for rocprim::search_n, 79b4655.
Make use of vectorized load in rocprim::transform, a755431.

SPIR-V support changes:

ROCPRIM_WAVEFRONT_SIZE' from architecture defines, 68354b7.
Implement mechanism for wavefront size-based dispatching, ff1b0c5.
Fix various compile issues when targeting spir-v, 53a1bea.
SPIR-V: warp sort, 1ed975a.
Disable dispatching with macro for usage with spir-v, 960461d.
Additional fix for warp_sort for SPIR-v, c50cd7a.
fix(intrinsics/atomics.hpp): fix atomics when compiler to spirv, 527c24c.
fix: improve compatibility with spir-v target in algorithms using 'lane_mask_type', 5815656.
Also some additions to our CI for SPIR-V.
SPIR-V: warp reduce/scan, ab9dc0a.
SPIR-V: block scan/reduce/RLD, 9f0dcf1.
Temporarily stop running device_partition test for SPIR-V due to hanging, 75820ee.
SPIR-V: warp exchange/load/store, d3a8911.
SPIR-V: block exchange/load/store (and funcs), cff88f8.
SPIR-V: block radix rank/sort, cc1c028.
SPIR-V: lookback_scan_state, 1ed863a.
Added generic pragmas and created fallback for atomics, 4784b69.

General bug fixes:

Fix the warning: explicit specialization cannot have a storage class, 0249ce1.
Fix compilation failure in hipCUB/rocThrust to rocPRIM, 6c747d3.
Fix some compile issues introduced in rocm 6.4, 87b473f.
Fix compile warning in thread_load for the new compiler, a69a9bb.
Call non-static method properly in data generation utility, 6ad6c2d.
device_merge_sort custom_huge_type failing test, 7caf280.
Fix for device_run_length_encode failing test, b2bb04c.
fix: skip including the init value in block aggregate for warp and block scan, https://projects.streamhpc.com/amd/libraries/rocPRIM/-/commit/5bad310091a66bf84057ce1bb1e5e11ddf79f40c.

Testing coverage:
Prepare to move 'lookback_scan' to public API, 7da3c6a.

…the type trait interface"

…class""

…irs"

…)CUB"

…source files

…n_length_encode"

…ock scan

spolifroni-amd

The changelog needs some changes and a few clarifications.

spolifroni-amd · 2025-05-12T14:25:17Z

    * past default: `class AccType = detail::input_type_t<InitValueType>>`
    * new default: `class AccType = rocprim::invoke_result_binary_op_t<rocprim::detail::input_type_t<InitValueType>, BinaryFunction>`

+* Added function `is_build_in` into `rocprim::traits::get`.


Move to the "added" section.

Moved it to the added section.

spolifroni-amd · 2025-05-12T14:26:10Z

+* Added function `is_build_in` into `rocprim::traits::get`.
+* Changed the parameters `long_radix_bits` and `LongRadixBits` from `segmented_radix_sort` to `radix_bits` and `RadixBits` respectively.
+* Marked the initialisation constructor of `rocprim::reverse_iterator<Iter>` `explicit`, use `rocprim::make_reverse_iterator`.
+* Dropped `c++14` support for rocprim.


In the hipCUB changelog, it clarifies that there's only C++17 support now. Do you want to specify that here too? If not, move this to "removed".

Suggested change

* Dropped `c++14` support for rocprim.

* Dropped `C++14` support for rocprim.

I both moved it to removed and clarified that we only support C++17 now.

spolifroni-amd · 2025-05-12T14:26:27Z

+* Merged `radix_key_codec` into type_traits system.
+* Renamed `type_traits_interface.hpp` to `type_traits.hpp`, rename the original `type_traits.hpp` to `type_traits_functions.hpp`.
+* Removed `radix_sort.hpp`
+* Added virtual shared memory as a fallback option in `rocprim::device_merge` when it exceeds shared memory capacity, similar to `rocprim::device_select`, `rocprim::device_partition`, and `rocprim::device_merge_sort`, which already include this feature.


move to the "added" section.

Moved it to the added section.

spolifroni-amd · 2025-05-12T14:26:42Z

+* Renamed `type_traits_interface.hpp` to `type_traits.hpp`, rename the original `type_traits.hpp` to `type_traits_functions.hpp`.
+* Removed `radix_sort.hpp`
+* Added virtual shared memory as a fallback option in `rocprim::device_merge` when it exceeds shared memory capacity, similar to `rocprim::device_select`, `rocprim::device_partition`, and `rocprim::device_merge_sort`, which already include this feature.
+* Added initial value support to device level inclusive scans.


Move to the "added" section

Moved it to the added section.

spolifroni-amd · 2025-05-12T14:27:52Z

+* Removed `radix_sort.hpp`
+* Added virtual shared memory as a fallback option in `rocprim::device_merge` when it exceeds shared memory capacity, similar to `rocprim::device_select`, `rocprim::device_partition`, and `rocprim::device_merge_sort`, which already include this feature.
+* Added initial value support to device level inclusive scans.
+* Made new optimization for `device_transform` when the input and output are pointers.


I think this should be in "added". Not sure though because the sentence isn't clear. Is this a backend change where optimization has been added, or is this something a user would need to make changes for?

Tried to clarify the statement and moved it to added.

spolifroni-amd · 2025-05-12T14:28:15Z

+* Added virtual shared memory as a fallback option in `rocprim::device_merge` when it exceeds shared memory capacity, similar to `rocprim::device_select`, `rocprim::device_partition`, and `rocprim::device_merge_sort`, which already include this feature.
+* Added initial value support to device level inclusive scans.
+* Made new optimization for `device_transform` when the input and output are pointers.
+* Added `LoadType` to `transform_config`, which is used for the `device_transform` when the input and output are pointers.


Move to "added"

Moved it to the added section.

spolifroni-amd · 2025-05-12T14:29:12Z

+
+### Resolved issues
+
+* Fixed `device_batch_memcpy` its reported benchmarking throughput being 2x lower than it was in reality.


Suggested change

* Fixed `device_batch_memcpy` its reported benchmarking throughput being 2x lower than it was in reality.

* Fixed an issue where `device_batch_memcpy` reported benchmarking throughput being 2x lower than it was in reality.

Made the suggested changes.

spolifroni-amd · 2025-05-12T14:29:40Z

+### Resolved issues
+
+* Fixed `device_batch_memcpy` its reported benchmarking throughput being 2x lower than it was in reality.
+* Fixed `device_segmented_reduce` its reported autotuning throughput being 5x lower than it was in reality.


Suggested change

* Fixed `device_segmented_reduce` its reported autotuning throughput being 5x lower than it was in reality.

* Fixed an issue where `device_segmented_reduce` reported autotuning throughput being 5x lower than it was in reality.

Made the suggested changes.

NB4444 · 2025-05-14T14:02:40Z

I am currently looking into the device_scan failure on gfx90a. I found two issues so far, one of them I already have a fix. The block_scan_reduce_then_scan had some missing wave_barrier. The other one seems to be an issue in the thread_scan_inclusive. It might be a compiler issue, but no luck so far in creating a smaller reproducer.

NB4444 · 2025-05-19T12:23:35Z

I have added a fix and a temp fix for the device_scan failure.

jayhawk-commits · 2025-05-22T00:35:03Z

Closing the pull request in this repo. Please refer to the migrated pull request for updates.

@NB4444

The changes we made to the benchmarks: - Applied new benchmark abstraction to increase maintainability. - Introduced device_ptr to the benchmarks. - Output more data in JSON output of the benchmarks. Deprecation and removal related to major release: - Remove [[deprecated]]float_bit_mask and all uses of it from the type trait interface, ROCm/rocPRIM@a472e80. - Remove short_radix_bits in segmented_radix_sort_config_params, ROCm/rocPRIM@38ac5a5. - Remove already deprecated functions, ROCm/rocPRIM@4d265b7. - Move rocprim::detail::radix_key_codec_base into traits system, ROCm/rocPRIM@7c664b3. CCCL 2.7 parity changes: - Add virtual shared memory fallback to device_merge, ROCm/rocPRIM@ca3b0cd. - Add device-level inclusive_scan with initial value support, ROCm/rocPRIM@cff2e16. - Change default scan accumulator type to be in line with (hip) CUB, ROCm/rocPRIM@5f7accb. - Match CUB's behavior in rocPRIM for device merge, ROCm/rocPRIM@4239ef5. - Create tests for rocPRIM's bit_cast, ROCm/rocPRIM@668f913. Performance optimizations: - Implement tuning for rocprim::search_n, ROCm/rocPRIM@79b4655. - Make use of vectorized load in rocprim::transform, ROCm/rocPRIM@a755431. SPIR-V support changes: - ROCPRIM_WAVEFRONT_SIZE' from architecture defines, ROCm/rocPRIM@68354b7. - Implement mechanism for wavefront size-based dispatching, ROCm/rocPRIM@ff1b0c5. - Fix various compile issues when targeting spir-v, ROCm/rocPRIM@53a1bea. - SPIR-V: warp sort, ROCm/rocPRIM@1ed975a. - Disable dispatching with macro for usage with spir-v, ROCm/rocPRIM@960461d. - Additional fix for warp_sort for SPIR-v, ROCm/rocPRIM@c50cd7a. - fix(intrinsics/atomics.hpp): fix atomics when compiler to spirv, ROCm/rocPRIM@527c24c. - fix: improve compatibility with spir-v target in algorithms using 'lane_mask_type', ROCm/rocPRIM@5815656. - Also some additions to our CI for SPIR-V. - SPIR-V: warp reduce/scan, ROCm/rocPRIM@ab9dc0a. - SPIR-V: block scan/reduce/RLD, ROCm/rocPRIM@9f0dcf1. - Temporarily stop running device_partition test for SPIR-V due to hanging, ROCm/rocPRIM@75820ee. - SPIR-V: warp exchange/load/store, ROCm/rocPRIM@d3a8911. - SPIR-V: block exchange/load/store (and funcs), ROCm/rocPRIM@cff88f8. - SPIR-V: block radix rank/sort, ROCm/rocPRIM@cc1c028. - SPIR-V: lookback_scan_state, ROCm/rocPRIM@1ed863a. - Added generic pragmas and created fallback for atomics, ROCm/rocPRIM@4784b69. General bug fixes: - Fix the warning: explicit specialization cannot have a storage class, ROCm/rocPRIM@0249ce1. - Fix compilation failure in hipCUB/rocThrust to rocPRIM, ROCm/rocPRIM@6c747d3. - Fix some compile issues introduced in rocm 6.4, ROCm/rocPRIM@87b473f. - Fix compile warning in thread_load for the new compiler, ROCm/rocPRIM@a69a9bb. - Call non-static method properly in data generation utility, ROCm/rocPRIM@6ad6c2d. - device_merge_sort custom_huge_type failing test, ROCm/rocPRIM@7caf280. - Fix for device_run_length_encode failing test, ROCm/rocPRIM@b2bb04c. - fix: skip including the init value in block aggregate for warp and block scan, https://projects.streamhpc.com/amd/libraries/rocPRIM/-/commit/5bad310091a66bf84057ce1bb1e5e11ddf79f40c. Testing coverage: Prepare to move 'lookback_scan' to public API, ROCm/rocPRIM@7da3c6a. --- 🔁 Imported from [ROCm/rocPRIM#731](ROCm/rocPRIM#731) 🧑‍💻 Originally authored by @NB4444 --------- Co-authored-by: Sander Bos <sander@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: Apoorva Kalyani <apoorva@streamhpc.com> Co-authored-by: Balint Csala <balint.csala@streamhpc.com> Co-authored-by: Jaap Blok <jaap@streamhpc.com> Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: Borys Petrov <borys@streamhpc.com> Co-authored-by: Balint Siklosi <balint.siklosi@streamhpc.com> Co-authored-by: Saiyang Zhang <saiyang@streamhpc.com> Co-authored-by: Bence Parajdi <bence@streamhpc.com> Co-authored-by: assistant-librarian[bot] <assistant-librarian[bot]@users.noreply.github.com> Co-authored-by: Joseph Macaranas <Joseph.Macaranas@amd.com>

@NB4444

The changes we made to the benchmarks: - Applied new benchmark abstraction to increase maintainability. - Introduced device_ptr to the benchmarks. - Output more data in JSON output of the benchmarks. Deprecation and removal related to major release: - Remove [[deprecated]]float_bit_mask and all uses of it from the type trait interface, a472e80. - Remove short_radix_bits in segmented_radix_sort_config_params, 38ac5a5. - Remove already deprecated functions, 4d265b7. - Move rocprim::detail::radix_key_codec_base into traits system, 7c664b3. CCCL 2.7 parity changes: - Add virtual shared memory fallback to device_merge, ca3b0cd. - Add device-level inclusive_scan with initial value support, cff2e16. - Change default scan accumulator type to be in line with (hip) CUB, 5f7accb. - Match CUB's behavior in rocPRIM for device merge, 4239ef5. - Create tests for rocPRIM's bit_cast, 668f913. Performance optimizations: - Implement tuning for rocprim::search_n, 79b4655. - Make use of vectorized load in rocprim::transform, a755431. SPIR-V support changes: - ROCPRIM_WAVEFRONT_SIZE' from architecture defines, 68354b7. - Implement mechanism for wavefront size-based dispatching, ff1b0c5. - Fix various compile issues when targeting spir-v, 53a1bea. - SPIR-V: warp sort, 1ed975a. - Disable dispatching with macro for usage with spir-v, 960461d. - Additional fix for warp_sort for SPIR-v, c50cd7a. - fix(intrinsics/atomics.hpp): fix atomics when compiler to spirv, 527c24c. - fix: improve compatibility with spir-v target in algorithms using 'lane_mask_type', 5815656. - Also some additions to our CI for SPIR-V. - SPIR-V: warp reduce/scan, ab9dc0a. - SPIR-V: block scan/reduce/RLD, 9f0dcf1. - Temporarily stop running device_partition test for SPIR-V due to hanging, 75820ee. - SPIR-V: warp exchange/load/store, d3a8911. - SPIR-V: block exchange/load/store (and funcs), cff88f8. - SPIR-V: block radix rank/sort, cc1c028. - SPIR-V: lookback_scan_state, 1ed863a. - Added generic pragmas and created fallback for atomics, 4784b69. General bug fixes: - Fix the warning: explicit specialization cannot have a storage class, 0249ce1. - Fix compilation failure in hipCUB/rocThrust to rocPRIM, 6c747d3. - Fix some compile issues introduced in rocm 6.4, 87b473f. - Fix compile warning in thread_load for the new compiler, a69a9bb. - Call non-static method properly in data generation utility, 6ad6c2d. - device_merge_sort custom_huge_type failing test, 7caf280. - Fix for device_run_length_encode failing test, b2bb04c. - fix: skip including the init value in block aggregate for warp and block scan, https://projects.streamhpc.com/amd/libraries/rocPRIM/-/commit/5bad310091a66bf84057ce1bb1e5e11ddf79f40c. Testing coverage: Prepare to move 'lookback_scan' to public API, 7da3c6a. --- 🔁 Imported from [#731](#731) 🧑‍💻 Originally authored by @NB4444 --------- Co-authored-by: Sander Bos <sander@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: Apoorva Kalyani <apoorva@streamhpc.com> Co-authored-by: Balint Csala <balint.csala@streamhpc.com> Co-authored-by: Jaap Blok <jaap@streamhpc.com> Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: Borys Petrov <borys@streamhpc.com> Co-authored-by: Balint Siklosi <balint.siklosi@streamhpc.com> Co-authored-by: Saiyang Zhang <saiyang@streamhpc.com> Co-authored-by: Bence Parajdi <bence@streamhpc.com> Co-authored-by: assistant-librarian[bot] <assistant-librarian[bot]@users.noreply.github.com> Co-authored-by: Joseph Macaranas <Joseph.Macaranas@amd.com>

MyNameIsTrez and others added 30 commits May 6, 2025 13:37

Abstract benchmarking loop

1e13119

Apply new benchmark abstraction - Part 2

2c0800c

Apply new benchmark abstraction - Part 1

7348a1d

Resolve "Introduce device_ptr to Benchmarks Part 3"

43da4c0

Resolve "Introduce device_ptr to Benchmarks Part 5"

f6f2195

Apply new benchmark abstraction - Part 3

29a6d12

Apply new benchmark abstraction - Part 4

8bffaad

Resolve "Remove [[deprecated]]float_bit_mask and all uses of it from …

a472e80

…the type trait interface"

Resolve "Remove short_radix_bits in segmented_radix_sort_config_params"

38ac5a5

Remove deprecations

4d265b7

Resolve "Fix "warning: explicit specialization cannot have a storage …

0249ce1

…class""

Resolve "Move rocprim::detail::radix_key_codec_base into traits system"

7c664b3

Apply new benchmark abstraction - Part 8

e642d8e

Apply new benchmark abstraction - Part 6

d6ba6da

Fix merge issues and clang format

06ce1d9

Resolve "Implement tuning for rocprim::search_n"

79b4655

Resolve "Apply new benchmark abstraction - Part 7"

95ad769

Resolve "Introduce device_ptr to BenchmarksPart 8"

eed8369

Resolve "Introduce device_ptr to Benchmarks Part 2"

0c5e3dd

Resolve "Add virtual shared memory fallback to device_merge"

ca3b0cd

Resolve "Add device-level inclusive_scan with initial value support"

cff2e16

Resolve "Make use of vectorized load in rocprim::transform"

a755431

Fix autotuning in benchmark_device_transform

300de6e

Apply new benchmark abstraction - Part 9

ffd4887

Resolve "Introduce device_ptr to Benchmarks Part 4"

dacfb1e

Resolve "Introduce device_ptr to Benchmarks Part 6"

a88ed02

Resolve "Autotune failure on benchmark_device_segmented_radix_sort_pa…

2148abe

…irs"

Derive 'ROCPRIM_WAVEFRONT_SIZE' from architecture defines

68354b7

Resolve "Change default scan accumulator type to be in line with (hip…

5f7accb

…)CUB"

Apply new benchmark abstraction - Part 10

b135286

Resolve "device_run_length_encode failing test"

b2bb04c

ScottTodd mentioned this pull request May 6, 2025

Warning triage ROCm/TheRock#47

Open

Naraenda and others added 7 commits May 12, 2025 12:47

test: fix clangd language server errors for tests that use generated …

9fc6fff

…source files

Resolve "SPIR-V: block radix rank/sort"

cc1c028

Resolve "REVERT: fallback to host side input generation for device_ru…

5daf2de

…n_length_encode"

Added generic pragmas and created fallback for atomics

4784b69

ci(.gitlab-ci.yml): add timeout to spirv tests

2db9f4e

Resolve "SPIR-V: lookback_scan_state"

1ed863a

Resolve "Prepare to move 'lookback_scan' to public API"

7da3c6a

NB4444 marked this pull request as ready for review May 12, 2025 12:50

NB4444 requested review from a team, RobsonRLemos, lawruble13, stanleytsang-amd and umfranzw as code owners May 12, 2025 12:50

NB4444 and others added 2 commits May 12, 2025 12:58

Clang format

9d6dd68

fix: skip including the init value in block aggregate for warp and bl…

ab1ef9e

…ock scan

spolifroni-amd requested changes May 12, 2025

View reviewed changes

CHANGELOG update

578f5a1

jayhawk-commits mentioned this pull request May 15, 2025

Develop stream 2025 05 06 ROCm/rocm-libraries#30

Merged

Fix failing test device_scan

cff99e3

NB4444 added 2 commits May 20, 2025 07:28

update rocprim version to 4.0.0

d210bc2

Fix build error in benchmark_utils

8ed4a9a

stanleytsang-amd requested a review from spolifroni-amd May 21, 2025 20:09

jayhawk-commits closed this May 22, 2025

	* Dropped `c++14` support for rocprim.
	* Dropped `C++14` support for rocprim.


		### Resolved issues

		* Fixed `device_batch_memcpy` its reported benchmarking throughput being 2x lower than it was in reality.

	* Fixed `device_batch_memcpy` its reported benchmarking throughput being 2x lower than it was in reality.
	* Fixed an issue where `device_batch_memcpy` reported benchmarking throughput being 2x lower than it was in reality.

	* Fixed `device_segmented_reduce` its reported autotuning throughput being 5x lower than it was in reality.
	* Fixed an issue where `device_segmented_reduce` reported autotuning throughput being 5x lower than it was in reality.

Conversation

NB4444 commented May 6, 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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

NB4444 commented May 14, 2025

Uh oh!

NB4444 commented May 19, 2025

Uh oh!

jayhawk-commits commented May 22, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

13 participants

NB4444 commented May 6, 2025 •

edited

Loading