Fully enable cache bypass access for Simple/LL/LL128 by wenkaidu · Pull Request #3812 · ROCm/rocm-systems

wenkaidu · 2026-03-05T21:51:56Z

Motivation

This PR completes what was left from #2847 and implements cache bypassing access for Simple/LL/LL128 protocols

Technical Details

Since compiler only supports __builtin_amdgcn_global_load_b128 and __builtin_amdgcn_global_store_b128 at this time, system scope relaxed atomic is used for other data sizes.

JIRA ID

Test Plan

There should be minimal performance impact as RCCL communication buffers are still allocated as uncached. The PR does nothing with existing use cases as there is no cache to bypass. Real use case will be "Enable user buffer and graph registration feature": #2758

Test Result

No observable change in performance with and without the PR.

Submission Checklist

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

thananon · 2026-03-06T14:48:24Z

This breaks topo_expl. Please try to run make on rccl/tools/topo_expl.

hipify_rccl/include/device.h:101:10: fatal error: 'device/rccl_ptr.h' file not found
  101 | #include "device/rccl_ptr.h"
      |          ^~~~~~~~~~~~~~~~~~~
1 warning and 1 error generated when compiling for gfx942.

wenkaidu · 2026-03-06T21:56:39Z

This breaks topo_expl. Please try to run make on rccl/tools/topo_expl.

hipify_rccl/include/device.h:101:10: fatal error: 'device/rccl_ptr.h' file not found
  101 | #include "device/rccl_ptr.h"
      |          ^~~~~~~~~~~~~~~~~~~
1 warning and 1 error generated when compiling for gfx942.

Fixed

math-ci-jobs · 2026-03-12T21:36:11Z

regression-detection smoke run on commit `6bfcea4`\nArtifacts - Results

math-ci-jobs · 2026-03-12T23:20:23Z

regression-detection smoke run on commit `42be1f5`\nArtifacts - Results

Cherry-picked and squashed from PR #3812 (users/wenkaidu/cache_bypass). When RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS is enabled (gfx942/gfx950): - load128/store128 use __builtin_amdgcn_global_{load,store}_b128 with system scope, ensuring cache-bypassing visibility for LL128 P2P buffers - LL protocol loads/stores use DWORDX4 builtins for system-scope access - Sub-16B ld_volatile/st uses __hip_atomic_{load,store} with address_space(1) pointers and __HIP_MEMORY_SCOPE_SYSTEM for correct cross-GPU visibility and global_* instruction emission - DEFINE_ld_st macros split into hip_atomic (DWORDX4) and fallback (nontemporal) variants; fallback uses address_space(1) casts - ncclLLFifoLine gains v4u member for b128 load/store on supported HW - rccl_ptr.h: define RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS 0 for non-device - topo_expl build fix for rccl_ptr.h include path

…f-include - Cherry-pick PR #3812 (cache bypass for Simple/LL/LL128): dwordx4 system-scope loads/stores via __builtin_amdgcn_global_{load,store}_b128, hip_atomic ld/st for sub-16B accesses, v4u member in ncclLLFifoLine - Remove bundler target string probe (broken with clang 23 .llvm.offloading format); hardcode hip-amdgcn-amd-amdhsa--<arch> which is stable across all relevant compiler versions - Fix lds_pointer_transform.py self-include bug: skip add_rccl_ptr_include fallback when processing rccl_ptr.h itself - Restore CMAKE_CXX_COMPILER override in do-cmake Made-with: Cursor

math-ci-jobs · 2026-03-17T15:11:31Z

regression-detection smoke run on commit `61e01fe`

Artifacts - Results

alex-breslow-amd · 2026-03-17T16:16:21Z

Great! I'm glad this landed. Thanks for the assist @wenkaidu!

…f-include - Cherry-pick PR #3812 (cache bypass for Simple/LL/LL128): dwordx4 system-scope loads/stores via __builtin_amdgcn_global_{load,store}_b128, hip_atomic ld/st for sub-16B accesses, v4u member in ncclLLFifoLine - Remove bundler target string probe (broken with clang 23 .llvm.offloading format); hardcode hip-amdgcn-amd-amdhsa--<arch> which is stable across all relevant compiler versions - Fix lds_pointer_transform.py self-include bug: skip add_rccl_ptr_include fallback when processing rccl_ptr.h itself - Restore CMAKE_CXX_COMPILER override in do-cmake Made-with: Cursor

This PR completes what was left from #2847 and implements cache bypassing access for Simple/LL/LL128 protocols Since compiler only supports __builtin_amdgcn_global_load_b128 and __builtin_amdgcn_global_store_b128 at this time, system scope relaxed atomic is used for other data sizes. There should be minimal performance impact as RCCL communication buffers are still allocated as uncached. The PR does nothing with existing use cases as there is no cache to bypass. Real use case will be "Enable user buffer and graph registration feature": #2758 No observable change in performance with and without the PR. - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. --------- Co-authored-by: Yiltan <ytemucin@amd.com> Co-authored-by: Bertan Dogancay <111835151+BertanDogancay@users.noreply.github.com>

storeData has called store<> since #9336. When RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS is 0 (compiler without __builtin_amdgcn_global_load_b128, e.g. ROCm 7.0.1, or DWORDX4_INTRINSICS=OFF), store<> falls back to a non-temporal store followed, on gfx950 with ROCM_VERSION < 70002, by a release fence: buffer_wbl2 sc0 sc1 plus s_waitcnt vmcnt(0) lgkmcnt(0) after every user-buffer store. The fence works around the ROCr hipHostMallocUncached bug for the comm FIFO, which storeLL still covers. store<> only writes user buffers, had no callers before #9336, and its builtin path has not carried the fence since #3812. Only gfx950 builds against ROCm < 7.0.2 without the builtins change; every other build preprocesses to the same code. JIRA ID : ROCM-31938 Co-authored-by: Cursor <cursoragent@cursor.com>

## Motivation Since #9336, the LL `storeData` writes the user output buffer through `store<>`. Without the b128 builtins (`RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS == 0`), `store<>` uses a non-temporal store. On gfx950 with `ROCM_VERSION < 70002` it also issues a release fence (`buffer_wbl2 sc0 sc1` + `s_waitcnt vmcnt(0) lgkmcnt(0)`) after every store. This makes 4 KB–64 KB `ncclBroadcast` on MI355X about 30% slower in RCCL built against ROCm 7.0.1. ## Technical Details - Removes the gfx950 `ROCM_VERSION < 70002` fence from the `store<>` fallback. - The fence is the ROCr `hipHostMallocUncached` workaround for the comm FIFO, and `storeLL` keeps it. `store<>` only writes user buffers, and its builtin path has had no fence since #3812. - Builds against ROCm 7.0.2 or later, and gfx942 builds, compile to the same code as before. ## Issue Tracking JIRA ID: ROCM-31938 ## Test Plan 4 nodes x 8 MI355X, RCCL built against ROCm 7.0.1: - `broadcast_perf` 1K–256K, interleaved with RCCL 2.27.7 and develop. - `all_reduce_perf` and `reduce_perf` 1K–4M. - Correctness with `NCCL_PROTO=LL -d all -u 1`, 8 B–1 MB, for broadcast, all_reduce, all_gather, reduce_scatter and reduce. ## Test Result `broadcast_perf` out-of-place time, ratio vs RCCL 2.27.7: | Size | develop | this PR | | --- | --- | --- | | 8K | 1.23 | 1.08 | | 16K | 1.26 | 1.08 | | 32K | 1.28 | 1.06 | | 64K | 1.31 | 1.02 | | 128K | 0.99 | 1.00 | - `all_reduce` and `reduce`: no change. - Correctness: `#wrong` = 0 for every collective, data type and size. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Cursor <cursoragent@cursor.com>

Fully enable cache bypass access for Simple/LL/LL128

423aac7

wenkaidu requested a review from a team as a code owner March 5, 2026 21:51

github-actions Bot added the project: rccl label Mar 5, 2026

wenkaidu mentioned this pull request Mar 5, 2026

[AICOMRCCL-98] [ROCM-1974] Enable user buffer and graph registration feature #2758

Merged

1 task

systems-assistant Bot added the organization: ROCm label Mar 5, 2026

mustafabar reviewed Mar 5, 2026

View reviewed changes

Comment thread projects/rccl/src/device/op128.h Resolved

Comment thread projects/rccl/src/device/op128.h Resolved

wenkaidu added 2 commits March 5, 2026 17:05

Restore original fallback path

6658a9d

Fix another messed up path

e08c750

Fix topo_expl build issue due to path of rccl_ptr.h

ec0b906

wenkaidu added 3 commits March 6, 2026 15:54

Merge branch 'develop' into users/wenkaidu/cache_bypass

d17b7a2

Merge branch 'develop' into users/wenkaidu/cache_bypass

8f392da

Merge branch 'develop' into users/wenkaidu/cache_bypass

20f8dd5

thananon approved these changes Mar 9, 2026

View reviewed changes

wenkaidu and others added 4 commits March 10, 2026 07:33

Merge branch 'develop' into users/wenkaidu/cache_bypass

fb214bf

Merge branch 'develop' into users/wenkaidu/cache_bypass

afe982f

Merge branch 'develop' into users/wenkaidu/cache_bypass

b92fd4b

Merge branch 'develop' into users/wenkaidu/cache_bypass

6bfcea4

Add missing address space attribute and remove warnings

42be1f5

mustafabar approved these changes Mar 16, 2026

View reviewed changes

atulkulk approved these changes Mar 16, 2026

View reviewed changes

Merge branch 'develop' into users/wenkaidu/cache_bypass

bfdece7

lfmeadow mentioned this pull request Mar 17, 2026

Split device compile with LDS pointer transform and cache bypass #4140

Closed

3 of 4 tasks

Merge branch 'develop' into users/wenkaidu/cache_bypass

61e01fe

wenkaidu merged commit 5152a3f into develop Mar 17, 2026
16 of 18 checks passed

wenkaidu deleted the users/wenkaidu/cache_bypass branch March 17, 2026 16:06

lfmeadow mentioned this pull request Mar 30, 2026

Split device compile with LDS pointer transform, cache bypass, and selective sharding #4541

Closed

6 tasks done

i-kosarev mentioned this pull request Oct 14, 2026

fix(rccl): drop the hipHostMallocUncached fence from LL store<> #13451

Merged

1 task done

Navigation Menu

Repository navigation

Fully enable cache bypass access for Simple/LL/LL128 - #3812

Fully enable cache bypass access for Simple/LL/LL128#3812
wenkaidu merged 14 commits into
developfrom
users/wenkaidu/cache_bypass

wenkaidu commented Mar 5, 2026

Uh oh!

Uh oh!

Uh oh!

thananon commented Mar 6, 2026 •

edited

Loading

Uh oh!

wenkaidu commented Mar 6, 2026

Uh oh!

math-ci-jobs Bot commented Mar 12, 2026

Uh oh!

math-ci-jobs Bot commented Mar 12, 2026

Uh oh!

math-ci-jobs Bot commented Mar 17, 2026

Uh oh!

Uh oh!

alex-breslow-amd commented Mar 17, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

Conversation

wenkaidu commented Mar 5, 2026

Motivation

Technical Details

JIRA ID

Test Plan

Test Result

Submission Checklist

Uh oh!

Uh oh!

Uh oh!

thananon commented Mar 6, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

wenkaidu commented Mar 6, 2026

Uh oh!

math-ci-jobs Bot commented Mar 12, 2026

regression-detection smoke run on commit 6bfcea4\nArtifacts - Results

Uh oh!

math-ci-jobs Bot commented Mar 12, 2026

regression-detection smoke run on commit 42be1f5\nArtifacts - Results

Uh oh!

math-ci-jobs Bot commented Mar 17, 2026

regression-detection smoke run on commit 61e01fe

Uh oh!

Uh oh!

alex-breslow-amd commented Mar 17, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

6 participants

thananon commented Mar 6, 2026 •

edited

Loading

regression-detection smoke run on commit `6bfcea4`\nArtifacts - Results

regression-detection smoke run on commit `42be1f5`\nArtifacts - Results

regression-detection smoke run on commit `61e01fe`