Repository navigation
Conversation
refer to presentation for motivation details
…ulticast, tma bulk variants, etc) and improve organization.
…sor loads, prefetching, and reductions. Add new files for UTMALDG, UTMAPF, and UTMAREDG tests, enhancing coverage for TMA operations on rank-2 and rank-3 tensors. Update Makefile to support new benchmarks and streamline build process.
…D" and "FAILED" so CI respects output
New ubench under tma/bar.sync that exercises the CTA-wide barrier
(PTX `bar.sync 0;`, i.e. __syncthreads()) via explicit inline PTX and
verifies the result on the host.
The kernel runs 4 blocks x 256 threads (8 warps/block) and builds two
cross-warp data dependencies that are only correct if the barrier
synchronized every warp in the CTA:
- Phase 1: each thread writes global_tid+1 into its shared-memory slot,
bar.sync, then reads the mirrored slot smem[blockDim-1-tid] (owned by
a different warp) out to global memory.
- Phase 2: block-wide tree reduction with a bar.sync between every step.
A per-warp clock64 skew forces the warps out of lockstep so the barrier
is genuinely required; a USE_BARRIER compile toggle strips the barriers
for a control run. Results are copied back and checked on the host, and
printed as PASSED/FAILED so CI can parse the outcome.
Adds a dedicated Makefile (built by `make gpu_microbenchmark`, auto-
discovered via the top-level ubench/*/*/ wildcard) that compiles
bar_sync.cu into the bin/ directory. ARCH defaults to the tma/ suite set
(sm_80 sm_90a sm_100a sm_101 sm_120) since bar.sync is universal.
Co-authored-by: Cursor <cursoragent@cursor.com>
New ubench under tma/bar_warp_sync that exercises the warp-level barrier
(PTX `bar.warp.sync`, i.e. __syncwarp()) via explicit inline PTX and
verifies an intra-warp exchange on the host.
What the kernel does:
- Each lane writes produce(global_tid) = global_tid+1 into its slot of
shared memory, issues `bar.warp.sync 0xffffffff;` (full membermask),
then reads the MIRRORED lane WITHIN THE SAME WARP (lane l reads lane
31-l) and writes it to global memory. Runs 2 blocks x 256 threads
(16 warps). The host copies the buffer back and checks every lane got
its within-warp mirror value, printing PASSED/FAILED.
How it is NOP'd (and why that is correct):
- bar.warp.sync only synchronizes the 32 lanes within one warp. GPGPU-Sim
already executes each warp in lockstep and reconverges divergent lanes
with the SIMT/PDOM stack, so that intra-warp synchronization is already
guaranteed by the model. The instruction therefore has no work to do
and is NOP'd in the PTX front-end (src/cuda-sim/ptx.l maps the
`bar.warp.*` family to NOP_OP). Its real purpose on silicon is
Independent Thread Scheduling (Volta+), which the SIMT-stack model does
not implement.
- A passing run validates both that the simulator parses and drops
bar.warp.sync cleanly (a missing lexer rule would abort PTX parsing)
and that intra-warp lockstep ordering holds without the barrier doing
anything. Note there is deliberately NO cross-warp dependency and NO
warp skew (unlike the bar.sync test), because intra-warp ordering is
implicit in the model.
Adds a dedicated Makefile (built by `make gpu_microbenchmark`, auto-
discovered via the top-level ubench/*/*/ wildcard) that compiles
bar_warp_sync.cu into the bin/ directory. ARCH defaults to the tma/ suite
set (sm_80 sm_90a sm_100a sm_101 sm_120) since bar.warp.sync is universal.
Co-authored-by: Cursor <cursoragent@cursor.com>
Relocate the bar.sync and bar.warp.sync microbenchmarks to the system/ category. Folder move only; no source or Makefile changes. Co-authored-by: Cursor <cursoragent@cursor.com>
…A pure functional mode
New ubench under l2_cache/ld_global_l2_128b that exercises the CUTLASS-
style global load `ld.global.L2::128B.u32` via explicit inline PTX and
verifies from the host that the load still returns correct data — proving
the .L2::128B prefetch hint is functionally inert while the underlying
global load performs a real, CPU-observable memory transaction.
Instruction under test:
- Base opcode: ordinary ld.global.u32 (reads 4 bytes from global memory
into a register). Already implemented in GPGPU-Sim.
- Qualifier: .L2::128B is a PTX .level::prefetch_size hint (PTX ISA
7.4+, SM_75+). It asks the hardware to also prefetch a 128-byte
region into L2 alongside the load. It does NOT change the value
returned by the load — only (maybe) timing on silicon.
How the hint is NOP'd in GPGPU-Sim (no simulator patch required):
- ptx.l already lexes `.L2::64B` / `.L2::128B` / `.L2::256B` /
`.L2::cache_hint` / `.L2` as L2_OPTION, so PTX carrying the qualifier
parses cleanly instead of aborting.
- ptx_ir.cc already accepts L2_OPTION in the instruction-option switch
and takes the "recognized but no special action needed" path — i.e.
the hint is ignored for functional correctness while ld.global still
executes through the normal load path.
- This matches the team's reading of the CUTLASS / Accel-Sim coverage
table: ld.global is in; L2::128B prefetching can be ignored for now;
confirming that unblocks CUTLASS-48 functional runs.
What the kernel does:
- Issues EXPLICIT inline PTX for both paths so the exercised instruction
is unambiguous:
path A (under test): ld.global.L2::128B.u32 %0, [%1];
path B (control): ld.global.u32 %0, [%1];
- Host fills a global input buffer with unique nonzero values
(tid*3+7). Each of 4 blocks x 256 threads loads in[tid] via both
paths, then stores (loaded+1) into two separate output buffers.
Depending the store on the loaded value prevents DCE and guarantees
a real GPU->CPU-visible memory transaction seeded by the load.
- Host copies both buffers back and checks:
(1) every hinted load returned input+1 (functional load worked);
(2) hinted == plain for every tid (L2::128B did not change results).
Prints PASSED/FAILED so CI respects the result.
Why a passing run validates the NOP'd hint:
1. Parsing: a missing L2_OPTION lexer/option rule would abort PTX
parsing on `.L2::128B`. A completed run proves the qualifier is
recognized and ignored cleanly.
2. Functional load: if the whole instruction were wrongly NOP'd (not
just the hint), outputs would keep their 0xFF sentinel and Check 1
would fail. Passing Check 1 proves a real global load occurred.
3. Hint inertness: Check 2 shows the hinted path is bit-identical to
plain ld.global.u32 — the functional-mode contract for a prefetch-
size hint.
Adds a dedicated Makefile (built by `make GPU_Microbenchmark`, auto-
discovered via the top-level ubench/*/*/ wildcard) that compiles
ld_global_l2_128b.cu into bin/ with shared cudart for silicon runs.
ARCH defaults to sm_80 sm_90a sm_100a sm_101 sm_120 (Ampere+; L2::128B
is SM_75+). Override with e.g. `make ARCH=sm_90a` for a single GPU.
Verified under GPGPU-Sim (1024 threads): both checks PASS; cuobjdump
confirms the binary embeds `ld.global.L2::128B.u32` alongside the plain
control load.
Co-authored-by: Cursor <cursoragent@cursor.com>
…chmark New ubench under mem/mapa_shared_cluster that exercises the Hopper+ PTX instruction `mapa.shared::cluster.u32` (PTX ISA 7.8+, sm_90+) via explicit inline PTX and verifies end-to-end cluster shared-memory remapping from the host. Built through the existing GPU_Microbenchmark recursion (`make GPU_Microbenchmark`); installs mapa_shared_cluster into bin/ with shared cudart for silicon and simulator runs. Instruction under test: mapa.shared::cluster.u32 d, a, b; Remaps shared-window address `a` onto cluster CTA rank `b`. Does NOT move data; a later ld/st.shared::cluster through `d` touches CTA b's copy of that offset. Cluster launch uses __cluster_dims__(2,1,1) with barrier.cluster arrive/wait around remapped stores. Why the gates are strict (a stub/identity mapa must fail): Gate A — Rank sensitivity (address-only): from the same local address, mapa(., 0) and mapa(., 1) MUST differ. A stub that ignores operand b and copies a→a fails this check hard. Gate B — Remote store / local load: each CTA writes its magic into the *peer* slot via st.shared::cluster [mapa(local, peer)], then reads its *own* slot with an ordinary shared load (no mapa). Expected local value is the peer's magic; identity mapa writes self and fails. Gate C — Self-map: mapa(local, own_rank) must still name this CTA's slot (load through that address returns the peer magic after Gate B). Host D2H checks all gates; device buffers are poison-filled so a silent no-op is obvious. Makefile targets Hopper+ arches (sm_90a / sm_100a / sm_101 / sm_120) with SASS+PTX gencode, matching the elect_sync-style per-instruction ubench pattern. Co-authored-by: Cursor <cursoragent@cursor.com>
Cluster ctarank ubenchmark
The wgmma microbenchmarks write their A/B operand tiles into shared memory with a device-side smem_off() helper and then run wgmma.mma_async, which the simulator interprets. The simulator's wgmma_smem_offset() was corrected to CuTe's canonical GMMA Major-K layout, so these helpers -- which encoded the old (non-hardware) formula -- must be updated to stay consistent with the reader. smem_off() now computes, in uint128_t units, u128 = (leading/T)*(LBO/16) + (stride%8)*Wsw + (stride/8)*(SBO/16) off = u128*16 + (leading%T)*E matching the simulator. For the SW=0 tiles used by most benches the leading byte offset becomes LBO = MN*16 (dense); wgmma_verify additionally exercises SW=1/2/3 with the sparse swizzled layout (SBO=128*Wsw, LBO=16) and larger shared-memory buffers. The Swizzle<B,4,3> XOR is omitted to match the reader. All built wgmma benches (verify, dtypes, multik, shapes, scale, sync, fence, timing_edge) pass under functional simulation. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Two companion ubenches that verify the Hopper+ mapa.shared::cluster.u32 address-remapping instruction and end-to-end distributed shared memory (DSMEM) communication between CTAs of a thread-block cluster. mapa_shared_cluster (2-CTA cluster): - Gate A rank sensitivity: mapa(local,0) != mapa(local,1). - Gate B remote store / local load: each CTA writes its magic into its peer's slot via st.shared::cluster[mapa(local,peer)], then reads its own slot with an ordinary shared load and must see the peer's magic. - Gate C self-map identity: mapa(local,own_rank) still names this CTA's slot. mapa_shared_cluster_extended (8-CTA cluster, mirror pairing 0<->7,1<->6,...): - Keeps Gates A/B/C across all 8 CTAs (every slot has exactly one remote writer), and adds - Gate A+ full rank round-trip: getctarank(mapa(local,k)) == k for every rank k, proving mapa lands on the correct CTA rather than just "some other CTA" - a distinction the 2-CTA test cannot make. Both poison-fill result buffers so a silent no-op mapa is caught, use cluster barriers to order the zero-fill before remote stores, and build to standalone executables that pass under GPGPU-Sim and on Hopper silicon. Co-authored-by: Cursor <cursoragent@cursor.com>
This reverts commit 350a767.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
8 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.