[Naraenda]

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: perf(rocprim): reduce false sharing in lookback scan #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.

[therock-pr-bot]

✅ 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/hipccl/hipccl3/rocprim/benchmark/benchmark_device_scan.cpp, projects/hipccl/hipccl3/rocprim/benchmark/benchmark_device_scan.hpp, projects/hipccl/hipccl3/rocprim/rocprim/include/rocprim/device/detail/config/device_scan.hpp, projects/hipccl/hipccl3/rocprim/tuning/tuner/base_tuner.py, projects/hipccl/hipccl3/rocprim/tuning/tuner/tuning_device_scan.py; 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]

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

[amd-jmahovsky]

Two potential bugs that I am not 100% sure of.

[amd-jmahovsky]

LGTM, but there are merge conflicts and this needs to be ported to the new rocprim location.

[NguyenNhuDi]

@Naraenda tried auto tuning your device scan changes on a 950 and I am getting these errors:

(env) root@bg-1w300-e1-2a:/dockerx/zenguyen/rocm-libraries/projects/rocprim/build# python3 /dockerx/zenguyen/rocm-libraries/projects/rocprim/tuning/run_tuning.py --algo-regex device_scan --arch-name gfx950 --output-dir /dockerx/zenguyen/pr-12642/
Running tuning for algorithms matching 'device_scan':
  - device_scan
  - device_scan_by_key

Running tuning for device_scan...
Working directory: /dockerx/zenguyen/rocm-libraries/projects/rocprim/tuning/tuner

Tuning for rocprim::int128_t
Using size: 134217728 bytes
Using strategy: dual_annealing, max fevals: 100
Running warmup:0:/__w/rockrel/rockrel/rocm-systems/projects/clr/hipamd/src/hip_global.cpp:66  : 19273293000045 us:  Cannot find Symbol with name: _ZN7rocprim26ROCPRIM_ROCPRIM_VERSION_NS6detail17trampoline_kernelINS0_11scan_configILj256ELj8ELNS0_17block_load_methodE3ELNS0_18block_store_methodE3ELNS0_20block_scan_algorithmE0ELj4294967295ELj64EEENS1_20scan_config_selectorInEEZZZNS1_9scan_implILNS1_25lookback_scan_determinismE0ELb0ELb0ES7_PnSC_nNS0_4plusInEEnEEDaPvRmT3_T4_T5_mT6_P12ihipStream_tbENKUlT_T0_E_clISt17integral_constantIbLb0EESR_IbLb1EEEEDaSN_SO_ENKUlSN_SO_T1_T2_E_clINS1_14launch_managerIS7_NS1_11comp_targetILNS1_3genE0ELNS1_11target_archE4294967295ELNS1_3gpuE0ELNS1_3repE0EEEEES7_S14_SR_INS0_4arch9wavefront6targetELS18_1EEEEDaSN_SO_SU_SV_EUlSN_E_S14_NS1_30default_config_static_selectorELS18_2EEEvSU_
Aborted (core dumped)

[amd-jmahovsky]

@Naraenda #11702 has been merged which unblocks this PR.

[amd-jmahovsky]

@Naraenda tried auto tuning your device scan changes on a 950 and I am getting these errors:

(env) root@bg-1w300-e1-2a:/dockerx/zenguyen/rocm-libraries/projects/rocprim/build# python3 /dockerx/zenguyen/rocm-libraries/projects/rocprim/tuning/run_tuning.py --algo-regex device_scan --arch-name gfx950 --output-dir /dockerx/zenguyen/pr-12642/
Running tuning for algorithms matching 'device_scan':
  - device_scan
  - device_scan_by_key

Running tuning for device_scan...
Working directory: /dockerx/zenguyen/rocm-libraries/projects/rocprim/tuning/tuner

Tuning for rocprim::int128_t
Using size: 134217728 bytes
Using strategy: dual_annealing, max fevals: 100
Running warmup:0:/__w/rockrel/rockrel/rocm-systems/projects/clr/hipamd/src/hip_global.cpp:66  : 19273293000045 us:  Cannot find Symbol with name: _ZN7rocprim26ROCPRIM_ROCPRIM_VERSION_NS6detail17trampoline_kernelINS0_11scan_configILj256ELj8ELNS0_17block_load_methodE3ELNS0_18block_store_methodE3ELNS0_20block_scan_algorithmE0ELj4294967295ELj64EEENS1_20scan_config_selectorInEEZZZNS1_9scan_implILNS1_25lookback_scan_determinismE0ELb0ELb0ES7_PnSC_nNS0_4plusInEEnEEDaPvRmT3_T4_T5_mT6_P12ihipStream_tbENKUlT_T0_E_clISt17integral_constantIbLb0EESR_IbLb1EEEEDaSN_SO_ENKUlSN_SO_T1_T2_E_clINS1_14launch_managerIS7_NS1_11comp_targetILNS1_3genE0ELNS1_11target_archE4294967295ELNS1_3gpuE0ELNS1_3repE0EEEEES7_S14_SR_INS0_4arch9wavefront6targetELS18_1EEEEDaSN_SO_SU_SV_EUlSN_E_S14_NS1_30default_config_static_selectorELS18_2EEEvSU_
Aborted (core dumped)

@Naraenda It looks like this symbol lookup error is also happening in #11702, but it slipped through CI due to some tests not being run.

[Naraenda]

@Naraenda tried auto tuning your device scan changes on a 950 and I am getting these errors:
...
@Naraenda It looks like this symbol lookup error is also happening in #11702, but it slipped through CI due to some tests not being run.

Hmm weird. I don't think I touched anything that should impact the gfx950 path. I'll rebase this branch and investigate it.

[amd-jmahovsky]

@Naraenda We found and fixed the problem with gfx950. #13251

[assistant-librarian]

TheRock Submodule Bump Activity

Newest first

• 2026-10-12 — Included in TheRock via ROCm/TheRock#8933.