Suggestion on additional memory benchmarks and fix to UMOV 0x400

[bdh0404]

As seen in NVIDIA tuning guides

https://docs.nvidia.com/cuda/ampere-tuning-guide/index.html#unified-shared-memory-l1-texture-cache
https://docs.nvidia.com/cuda/ada-tuning-guide/index.html#unified-shared-memory-l1-texture-cache
https://docs.nvidia.com/cuda/hopper-tuning-guide/index.html#unified-shared-memory-l1-texture-cache
https://docs.nvidia.com/cuda/blackwell-tuning-guide/index.html#unified-shared-memory-l1-texture-cache

each thread block reserves 1KB of shared memory, which is the same size of 0x400.

This is an open hypothesis that these reserved 1KB of shared memory is related to the triplet @!PT LDS RZ, [RZ].

What I remember in sm90 SASS is that these triplets are used for synchronizing shared memory accesses between ld.acquire and st.release, or synchronizing general proxy and async proxy by fence.proxy.async.shared. (I could not provide access to hopper GPU and SASS code of my project as it is not my property, but my workplace).

@!PT LDS RZ, [RZ] could be used for safely flushing data into shared memory as it is fixed latency instructions. My hypothesis is that 1KB padding is required for instructions like @!PT LDS RZ, [RZ], with maintaining the memory consistency without RAW/WAR/WAW hazard AND avoiding shared memory bank conflicts.

I suggest making a microbenchmark for ld.acquire, st.release, ld.relaxed, st.relaxed into shared memory and mbarrier based instructions for future Hopper and Blackwell server GPU benchmarks with Async Tensor Core/Tensor Memory Accelerators.

You could use mbarrier based async memory access benchmarks on sm120 through LDGSTS & cp.async.mbarrier.arrive pairs.

[florianmattana]

Thanks, this maps directly to an existing open SASS King gap.

We already observe 3 × @!PT LDS RZ, [RZ] before LDGSTS in ch.18 and ch.24, currently tracked as GAP-18-1. The role of that triplet is not decoded yet.

The 1 KB reserved shared-memory connection is a useful hypothesis. NVIDIA documents that CUDA reserves 1 KB of shared memory per thread block on these architectures, and 0x400 matching that
reservation is worth testing. But SASS King cannot promote the link to [OBS] without a controlled variant showing that the triplet depends on that reserved region or on a specific shared-memory
ordering/proxy mechanism.

I would track this as an extension of GAP-18-1:

• [OBS] ch.18/ch.24 emit 3 × @!PT LDS RZ, [RZ] before LDGSTS.
• [OBS] NVIDIA tuning guides document 1 KB shared-memory reservation per thread block.
• [HYP] the false-predicated LDS RZ, [RZ] triplet may interact with that reserved region or with shared-memory ordering/proxy behavior.
• [GAP] no controlled test yet distinguishes padding/alignment, proxy synchronization, fixed-latency scheduling, or bank-conflict avoidance.

On maintainer hardware I can test the SM120 side first through LDGSTS, LDGDEPBAR, DEPBAR, and cp.async-style paths. Hopper/Blackwell server features such as TMA or full proxy-fence behavior would
need later hardware or external dumps.

Lett me know if thats sounds good to you

[bdh0404]

Oh, what I suggest for new benchmarks is about the mbarrier-based asynchronous operations with LDGSTS.

The mbarrier instructions are usually paired with TMA (SM90) and UMMA (SM100) ops, but it can be paired with LDGSTS too.

The PTX documentation shows an example with LDGSTS(cp.async)-mbarrier paired examples.

The following is noted in https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-cp-async-mbarrier-arrive

// Example 1: no .noinc
mbarrier.init.shared.b64 [shMem], threadCount;
....
cp.async.ca.shared.global [shard1], [gbl1], 4;
cp.async.cg.shared.global [shard2], [gbl2], 16;
....
// Absence of .noinc accounts for arrive-on from completion of prior cp.async operations.
// So mbarrier.init must only account for arrive-on from mbarrier.arrive.
cp.async.mbarrier.arrive.shared.b64 [shMem];
....
mbarrier.arrive.shared.b64 state, [shMem];

waitLoop:
mbarrier.test_wait.shared.b64 p, [shMem], state;
@!p bra waitLoop;



// Example 2: with .noinc

// Tracks arrive-on from mbarrier.arrive and cp.async.mbarrier.arrive.

// All threads participating in the mbarrier perform cp.async
mov.b32 copyOperationCnt, threadCount;

// 3 arrive-on operations will be triggered per-thread
mul.lo.u32 copyArrivalCnt, copyOperationCnt, 3;

add.u32 totalCount, threadCount, copyArrivalCnt;

mbarrier.init.shared.b64 [shMem], totalCount;
....
cp.async.ca.shared.global [shard1], [gbl1], 4;
cp.async.cg.shared.global [shard2], [gbl2], 16;
...
// Presence of .noinc requires mbarrier initalization to have accounted for arrive-on from cp.async
cp.async.mbarrier.arrive.noinc.shared.b64 [shMem]; // 1st instance
....
cp.async.ca.shared.global [shard3], [gbl3], 4;
cp.async.ca.shared.global [shard4], [gbl4], 16;
cp.async.mbarrier.arrive.noinc.shared::cta.b64 [shMem]; // 2nd instance
....
cp.async.ca.shared.global [shard5], [gbl5], 4;
cp.async.cg.shared.global [shard6], [gbl6], 16;
cp.async.mbarrier.arrive.noinc.shared.b64 [shMem]; // 3rd and last instance
....
mbarrier.arrive.shared.b64 state, [shMem];

waitLoop:
mbarrier.test_wait.shared.b64 p, [shMem], state;
@!p bra waitLoop;

From these examples, you can test the mbarrier ptx instructions first and analyze the SASS-equivalent ones, with LDGSTS asynchronous execution.