perf(rocprim): reduce false sharing in lookback scan by Naraenda · Pull Request #11702 · ROCm/rocm-libraries

Naraenda · 2026-09-04T09:31:02Z

Motivation

Look-back scan is an integral part of many primitive algorithms. It relies heavily on atomic load/store operations. These operate on cachelines, and may cause issues akin to false sharing. This PR introduces a tunable mechanism to pad atomic flags to reduce this false sharing.

Partially addresses the following tickets:

JIRA ID : EXSWSTRHPC-330
JIRA ID : EXSWSTRHPC-331

Technical Details

Padding is a compile-time mechanism.
This is added as a template parameter to lookback_scan_state.
On the host-side we must resolve the required lookback scan state which we only can know on runtime.
This is solved by leveraging the config launching mechanism, were use a visitor pattern to get compile time access to the config.
The padding is introduced by wrapping the (packed) flag type.
The padding parameter is added to the device scan config.
No additional tuning is done.
- Manually modifying some of the configs did result in a performance increase.

Test Plan

Performance should not change, as configs should be generated in a later PR.

Test Result

Verified no change in kernels before and after these changes:

# Compile benchmark_device_scan with --save-temps
grep -v '^[[:space:]]*[^[:alpha:]]' *.s.old | md5sum
# 4ae345b846f0b8074204d72151e3d1e7 -- gfx908, hash may very if rebased.
grep -v '^[[:space:]]*[^[:alpha:]]' *.s.new | md5sum
# 4ae345b846f0b8074204d72151e3d1e7 -- gfx908, hash may very if rebased.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

therock-pr-bot · 2026-09-04T09:37:20Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/rocprim/rocprim/include/rocprim/device/config_types.hpp`, `projects/rocprim/rocprim/include/rocprim/device/detail/device_config_helper.hpp`, `projects/rocprim/rocprim/include/rocprim/device/detail/lookback_scan_state.hpp`, `projects/rocprim/rocprim/include/rocprim/device/device_scan.hpp`; no test file found
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

therock-pr-bot · 2026-09-04T09:37:24Z

🎉 All checks passed! This PR is ready for review.

Naraenda · 2026-09-17T09:50:14Z

Rebased.

amd-jmahovsky

Some comments. We might also want to add a test that uses non-zero alignment to ensure correct output.

…riable

…variable to params

…arams

amd-jmahovsky

LGTM

amd-jmahovsky · 2026-09-24T21:19:24Z

@Naraenda This is legitimately failing CI with a compile error:

[hipCUB] /__w/rocm-libraries/rocm-libraries/external-rocm-libraries/projects/hipcub/hipcub/include/hipcub/agent/../backend/rocprim/agent/single_pass_scan_operators.hpp:226:38: error: too many template arguments for class template 'lookback_scan_state'
[hipCUB]   226 |     using NativeT = rocprim::detail::lookback_scan_state<T, true, 0, SINGLE_WORD>;
[hipCUB]       |                                      ^                               ~~~~~~~~~~~~

plus other errors that follow.

Naraenda · 2026-09-25T07:05:22Z

@Naraenda This is legitimately failing CI with a compile error:

[hipCUB] /__w/rocm-libraries/rocm-libraries/external-rocm-libraries/projects/hipcub/hipcub/include/hipcub/agent/../backend/rocprim/agent/single_pass_scan_operators.hpp:226:38: error: too many template arguments for class template 'lookback_scan_state'
[hipCUB]   226 |     using NativeT = rocprim::detail::lookback_scan_state<T, true, 0, SINGLE_WORD>;
[hipCUB]       |                                      ^                               ~~~~~~~~~~~~

plus other errors that follow.

AFAIK, this happens because TheRockCI does not handle depedent builds well. E.g. for hipCUB it will use the rocPRIM from develop, and not from source.

I do think I can find a workaround for this.

…ring

amd-jmahovsky · 2026-09-28T22:47:11Z

@Naraenda You mentioned finding a workaround for the failing CI. Any update on this? I suppose I can have this gardened and just have it merged in, but will it then break TheRock?

Naraenda · 2026-09-29T07:38:46Z

@Naraenda You mentioned finding a workaround for the failing CI. Any update on this? I suppose I can have this gardened and just have it merged in, but will it then break TheRock?

It's already included: f25f562

…ring

amd-jmahovsky · 2026-10-05T16:45:52Z

Closing because there is a newer PR: #12972

@Naraenda

…to hipccl3) (#12972) Port of #11702 (by @Naraenda) to the hipccl3 layout, so the change lands in `projects/hipccl/hipccl3/` instead of the retired `projects/` location. The original PR will be closed once this one is accepted. ## Original description ## Motivation Look-back scan is an integral part of many primitive algorithms. It relies heavily on atomic load/store operations. These operate on cachelines, and may cause issues akin to false sharing. This PR introduces a tunable mechanism to pad atomic flags to reduce this false sharing. Partially addresses the following tickets: JIRA ID : EXSWSTRHPC-330 JIRA ID : EXSWSTRHPC-331 ## Technical Details * Padding is a compile-time mechanism. This is added as a template parameter to `lookback_scan_state`. * On the host-side we must resolve the required lookback scan state which we only can know on runtime. This is solved by leveraging the config launching mechanism, were use a visitor pattern to get compile time access to the config. * The padding is introduced by wrapping the (packed) flag type. * The padding parameter is added to the device scan config. * **No additional tuning is done.** * Manually modifying some of the configs did result in a performance increase. ## Test Plan Performance should not change, as configs should be generated in a later PR. ## Test Result Verified no change in kernels before and after these changes: ```sh # Compile benchmark_device_scan with --save-temps grep -v '^[[:space:]]*[^[:alpha:]]' *.s.old | md5sum # 4ae345b846f0b8074204d72151e3d1e7 -- gfx908, hash may very if rebased. grep -v '^[[:space:]]*[^[:alpha:]]' *.s.new | md5sum # 4ae345b846f0b8074204d72151e3d1e7 -- gfx908, hash may very if rebased. ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## About this port - Commits are the original 13 commit(s), with their authors and messages unchanged. Only the paths were remapped (`projects/<name>/` -> `projects/hipccl/hipccl3/<name>/`). The commit IDs differ from the originals, as they must. - Result: 4 files changed, 414 insertions(+), 261 deletions(-). - `git am --3way` applied it cleanly on top of `develop`. - **Not built or tested here.** No CI covers `projects/hipccl` yet, so the checks on this PR do not exercise the change. Please build the affected hipccl3 component and run its relevant tests before merging. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Nara Prasetya <nara@streamhpc.com> Co-authored-by: amd-jmahovsky <jeffrey.mahovsky@amd.com>

## Motivation #11702 did add optimizations, but did not include them to the tuning mechanisms. This PR addresses that and tunes said algorithms for MI300X. JIRA ID : EXSWSTRHPC-330 JIRA ID : EXSWSTRHPC-331 ## Technical Details * 345fe33: #11702 (should be dropped) * 96b6d8f: adds the new parameters to the search space * 36cee5f: allows tuning with mismatching default parameters * 3fd8430: the resulting configs ## Test Plan * Tests pass. * Performance improved. ## Test Result * CI passes. * Performance verified on MI300X node. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: amd-jmahovsky <jeffrey.mahovsky@amd.com>

…ring

## Motivation #11702 did add optimizations, but did not include them to the tuning mechanisms. This PR addresses that and tunes said algorithms for MI300X. This PR should not be merged before #11702. After the aforementioned PR is merged, commit 345fe33 should be dropped from this PR. JIRA ID : EXSWSTRHPC-330 JIRA ID : EXSWSTRHPC-331 ## Technical Details 345fe33: #11702 (should be dropped) 96b6d8f: adds the new parameters to the search space 36cee5f: allows tuning with mismatching default parameters 3fd8430: the resulting configs ## Test Plan Tests pass. Performance improved. ## Test Result CI passes. Performance verified on MI300X node. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Ed D'Azevedo <98369794+EdDAzevedo@users.noreply.github.com> Co-authored-by: Mark Gates <mark.gates@amd.com> Co-authored-by: Mahovsky <jmahovsk@amd.com> Co-authored-by: Daming Feng <dmfeng8898@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>

NguyenNhuDi · 2026-10-15T18:12:04Z

closing in favour of #13756

## Motivation #11702 did add optimizations, but did not include them to the tuning mechanisms. This PR addresses that and tunes said algorithms for MI300X. This PR should not be merged before #11702. After the aforementioned PR is merged, commit 345fe33 should be dropped from this PR. JIRA ID : EXSWSTRHPC-330 JIRA ID : EXSWSTRHPC-331 ## Technical Details 345fe33: #11702 (should be dropped) 96b6d8f: adds the new parameters to the search space 36cee5f: allows tuning with mismatching default parameters 3fd8430: the resulting configs ## Test Plan Tests pass. Performance improved. ## Test Result CI passes. Performance verified on MI300X node. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Ed D'Azevedo <98369794+EdDAzevedo@users.noreply.github.com> Co-authored-by: Mark Gates <mark.gates@amd.com> Co-authored-by: Mahovsky <jmahovsk@amd.com> Co-authored-by: Daming Feng <dmfeng8898@gmail.com> Co-authored-by: Claude <noreply@anthropic.com>

Naraenda requested a review from a team as a code owner September 4, 2026 09:31

github-actions Bot added project: hipcub project: rocprim labels Sep 4, 2026

assistant-librarian Bot added the organization: streamhpc contributors from streamhpc label Sep 4, 2026

Naraenda changed the title ~~Users/naraenda/rocprim lookback false sharing~~ perf(rocprim): reduce false sharing in lookback scan Sep 4, 2026

Naraenda self-assigned this Sep 9, 2026

Naraenda force-pushed the users/naraenda/rocprim-lookback-false-sharing branch from 3c82a3c to e0b503a Compare September 17, 2026 09:50

amd-jmahovsky requested changes Sep 18, 2026

View reviewed changes

amd-jmahovsky reviewed Sep 18, 2026

View reviewed changes

Comment thread projects/rocprim/rocprim/include/rocprim/device/config_types.hpp Resolved

Naraenda added 12 commits September 24, 2026 09:02

perf(rocprim): reduce false sharing in lookback scan

a1f0759

perf(rocprim): align false sharing padding to 128 bytes

ccdbd92

perf(rocprim): reduce false sharing for large types in look back scan

b84b0ff

refactor(rocprim): specify lookback atomic alignment through macro va…

8b534b0

…riable

fix(rocprim): always pad lookback scan state atomics

4c2c307

refactor: allow passing of custom alignment size to lookback scan state

1dd1b5b

feat(rocprim): add state alignment to scan config

abd79d0

feat(hipcub): update type of lookbackscan state

1d3fd29

fix(rocprim): fix config and target not being passed as compile time …

60a19a2

…variable to params

docs(rocprim): improve comments and docs

8a052f3

fix(rocprim): add extra static assertions

269800a

fix(rocprim): use unsigned int for lookback scan state aligbment in p…

728de81

…arams

Naraenda force-pushed the users/naraenda/rocprim-lookback-false-sharing branch from e0b503a to 728de81 Compare September 24, 2026 09:03

amd-jmahovsky approved these changes Sep 24, 2026

View reviewed changes

fix(rocprim): use using alias to hide changes to hipcub

f25f562

Naraenda force-pushed the users/naraenda/rocprim-lookback-false-sharing branch from 67dab51 to f25f562 Compare September 25, 2026 07:31

Merge branch 'develop' into users/naraenda/rocprim-lookback-false-sha…

2a32e46

…ring

amd-jmahovsky added 3 commits September 29, 2026 09:32

Merge branch 'develop' into users/naraenda/rocprim-lookback-false-sha…

6edb908

…ring

Merge branch 'develop' into users/naraenda/rocprim-lookback-false-sha…

e914d87

…ring

Merge branch 'develop' into users/naraenda/rocprim-lookback-false-sha…

92cdf13

…ring

stanleytsang-amd mentioned this pull request Oct 2, 2026

perf(rocprim): reduce false sharing in lookback scan (port of #11702 to hipccl3) #12972

Merged

1 task done

amd-jmahovsky closed this Oct 5, 2026

amd-jmahovsky reopened this Oct 15, 2026

github-actions Bot added the project: rocprim-deprecated label Oct 15, 2026

Merge branch 'develop' into users/naraenda/rocprim-lookback-false-sha…

d3dc6a3

…ring

amd-jmahovsky mentioned this pull request Oct 15, 2026

feat(rocprim): tune new device scan params on gfx942 #13738

Closed

1 task

NguyenNhuDi added this pull request to stack #13742 October 15, 2026 16:43

NguyenNhuDi mentioned this pull request Oct 15, 2026

feat(rocprim): tune new device scan params on gfx942 #13743

Closed

1 task

NguyenNhuDi removed this pull request from stack #13742 October 15, 2026 16:56

NguyenNhuDi added this pull request to stack #13748 October 15, 2026 17:00

NguyenNhuDi removed this pull request from stack #13748 October 15, 2026 17:01

NguyenNhuDi mentioned this pull request Oct 15, 2026

feat(rocprim): tune new device scan params on gfx942 #13750

Merged

1 task

Merge branch 'develop' into users/naraenda/rocprim-lookback-false-sha…

4063480

…ring

amd-jmahovsky mentioned this pull request Oct 15, 2026

feat(rocprim): backport of kernel tuner related changes from hipccl #13756

Merged

1 task done

NguyenNhuDi closed this Oct 15, 2026

NguyenNhuDi deleted the users/naraenda/rocprim-lookback-false-sharing branch October 16, 2026 19:24

NguyenNhuDi restored the users/naraenda/rocprim-lookback-false-sharing branch October 16, 2026 19:24

Conversation

Naraenda commented Sep 4, 2026

Motivation

Technical Details

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Sep 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Sep 4, 2026

Uh oh!

Naraenda commented Sep 17, 2026

Uh oh!

amd-jmahovsky left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-jmahovsky left a comment

Choose a reason for hiding this comment

Uh oh!

amd-jmahovsky commented Sep 24, 2026

Uh oh!

Naraenda commented Sep 25, 2026

Uh oh!

amd-jmahovsky commented Sep 28, 2026

Uh oh!

Naraenda commented Sep 29, 2026

Uh oh!

amd-jmahovsky commented Oct 5, 2026

Uh oh!

NguyenNhuDi commented Oct 15, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

therock-pr-bot Bot commented Sep 4, 2026 •

edited

Loading