NCU Profile: GqaFwdWsPersistentKernel bottleneck analysis (llama8b-4k, H200)

[superAngGao]

Overview

NCU full-set profiling of GqaFwdWsPersistentKernel from tile-ai/TileOPs#871, commit beab3ca.

Shape: B=1 S=4096 H=32 Hkv=8 D=128 causal=True fp16 (llama8b-4k)
Hardware: NVIDIA H200 (132 SMs, 228 KB smem/SM, 65536 regs/SM), CUDA 12.8
Kernel duration: ~310 µs → 66.7% FA3

The kernel is not bandwidth-bound on this shape. The bottleneck is tensor core pipeline utilization — only 57.7% of peak, with 39% idle cycles caused by non-wgmma work (softmax, barriers, smem reads) being too slow to fully overlap with wgmma execution.

---

1. Launch Configuration

Parameter	Value
Grid	(132, 1, 1) — 1 persistent CTA per SM
Block	(384, 1, 1) — 3 warpgroups × 128 threads
Registers/thread	168
Total regs/block	168 × 384 = 64,512 (of 65,536/SM)
Shared memory/block (allocated)	165.9 KB (of 228 KB configurable)
Cluster size	0 (no cluster)
Waves/SM	1.0
Stack size	1024 B

2. Occupancy

Metric	Value
Theoretical occupancy	18.75% (12 warps / 64 max)
Achieved occupancy	18.68%
Achieved active warps/SM	11.96
Block limit — registers	1 (bottleneck)
Block limit — shared mem	1 (bottleneck)
Block limit — warps	5
Block limit — barriers	12
Block limit — SM	32

Both registers and shared memory independently limit the SM to 1 block. The register file is 98.4% full (64,512 / 65,536).

3. Throughput (Speed of Light)

Metric	Value
SM frequency	1.46 GHz
DRAM frequency	3.20 GHz
Compute (SM) throughput	57.7% of peak
Tensor pipe active (of elapsed)	57.7%
Tensor pipe active (of SM active)	60.9%
Memory throughput	38.8%
L1/TEX throughput	40.9%
L2 throughput	29.5%
DRAM throughput	5.17%
SM busy	60.9%
SM active cycles (avg)	430,587
SM elapsed cycles (avg)	455,541

Key observation

DRAM throughput is only 5.17% — the kernel is entirely L2-resident on this shape. Total DRAM traffic is 57.7 MB read + 21.0 MB write = 78.7 MB. The L2 hit rate is 90.7%, confirming that K/V data (B=1 × S=4096 × Hkv=8 × D=128 × 2B = 8 MB per K or V) fits comfortably in H200's 60 MB L2 cache.

The bottleneck is tensor core utilization (57.7%), not memory bandwidth.

4. Warp Stall Breakdown

Measured via smsp__average_warps_issue_stalled_*_per_issue_active.ratio:

Stall Reason	Ratio	Interpretation
long_scoreboard	3.30	Waiting for long-latency memory ops (smem→reg via L1/TEX, TMA completion)
wait (DEPBAR)	1.59	Waiting for wgmma to drain (wait_wgmma(0) / wait_wgmma(1))
barrier (mbarrier)	1.13	Waiting on mbarrier arrive/wait (producer↔consumer sync, ping-pong scheduler)
selected	1.00	Warp was selected and issued (not a stall)
no_instruction	0.42	I-cache miss or warp not yet ready
short_scoreboard	0.34	Waiting for short-latency ops (register deps)
dispatch_stall	0.27	Dispatch unit stall
not_selected	0.26	Warp eligible but not selected
branch_resolving	0.07	Branch resolution
mio_throttle	0.05	Memory I/O throttle
math_pipe_throttle	0.02	Tensor core back-pressure (negligible)
gmma	0.00	GMMA-specific stall
membar	0.00	Memory barrier
sleeping	0.00	—
lg_throttle	0.00	—
tex_throttle	0.00	—
drain	0.00	—
imc_miss	0.00	—
misc	0.00	—

Stall analysis

long_scoreboard (3.30) is the dominant stall. In this kernel, the long-latency operations are:

• Shared memory reads via L1/TEX — consumer warpgroups reading k_smem_* / v_smem_* / q_shared_* to feed wgmma. The swizzled smem layout goes through L1TEX with 72.2% hit rate.
• TMA async copy completion — producer waiting for TMA descriptor-based copies to land in smem.
• This is NOT HBM latency (DRAM is only 5% utilized, L2 hit rate is 90.7%). The stall is on the smem → register file path through L1TEX.

wait / DEPBAR (1.59) is the second largest stall. Consumer warps stall on wait_wgmma(0) or wait_wgmma(1) (WARPGROUP.DEPBAR.LE). With IntraWGOverlap (wait_group<1>), the current softmax/rescale overlaps with the previous wgmma, but 1.59 ratio indicates the overlap is incomplete — the non-wgmma work finishes before the wgmma, so the consumer must idle-wait.

barrier (1.13) is the third stall. The kernel uses 8 mbarriers:

• k_full, k_empty — K pipeline (producer → consumer data ready / consumer → producer buffer free)
• v_full, v_empty — V pipeline
• wg_sched_12, wg_sched_21 — WG1↔WG2 ping-pong scheduler
• q_full_1, q_full_2 — per-sub-tile Q TMA loads

Each K-loop iteration triggers multiple arrive/wait pairs. PR body notes ~2% overhead from using mbarrier-based scheduler vs PTX bar.arrive named barriers (which need an upstream TileLang patch).

5. Instruction Statistics

Metric	Value
Total instructions executed	~80.3M
Executed IPC (active)	1.41
Executed IPC (elapsed)	1.34
Issue slots busy	35.3%
Issued IPC (active)	1.41
Warp cycles per issued instruction	8.46
Avg active threads per warp	31.45
Avg not predicated-off threads per warp	31.06
Branch instructions	~2.28M
Branch efficiency	98.24%
Avg divergent branches	87.5

6. Memory Workload

Metric	Value
Memory throughput	253.9 GB/s
Mem busy	38.8%
Max bandwidth	38.1%
L1/TEX hit rate	72.2%
L2 hit rate	90.7%
L2 compression	0% (disabled)
Mem pipes busy	31.2%
DRAM read	57.7 MB
DRAM write	21.0 MB
DRAM read BW	184.8 GB/s
DRAM write BW	67.2 GB/s

Shared memory (GMMA) instructions

Metric	Value
GMMA instructions executed (smsp)	2,162,688
GMMA instructions (% of peak elapsed)	3.6%

7. Tensor Core Detail

Metric	Value
sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_elapsed	57.68%
sm__pipe_tensor_cycles_active.avg.pct_of_peak_sustained_active	60.92%
sm__pipe_tensor_cycles_active.min.pct_of_peak_sustained_elapsed	44.61%
sm__pipe_tensor_cycles_active.max.pct_of_peak_sustained_elapsed	59.48%
sm__inst_executed_pipe_tensor_op_gmma.avg.pct_of_peak_sustained_active	3.81%
sm__inst_executed_pipe_tensor_op_hmma.avg.pct_of_peak_sustained_active	1.90%
sm__pipe_tensor_cycles_active_realtime.avg.pct_of_peak_sustained_elapsed	27.42%
Achieved FLOPS (hgmma fp16, no sparsity)	57.68% of peak

The min/max spread (44.6%–59.5%) across SMs reflects causal tile pairing imbalance — even with (k, M-1-k) pairing, the two sub-tiles within a pair have different K-loop lengths, so some CTAs finish earlier than others.

8. Workload Distribution

Metric	Value
Avg DRAM active cycles	51,325
Total DRAM elapsed cycles	47.7M
Avg SM active cycles	430,586
Total SM elapsed cycles	59.9M
Avg L2 active cycles	509,492
Total L2 elapsed cycles	48.9M
Avg SMSP active cycles	430,099
Total SMSP elapsed cycles	239.4M

9. Shared Memory Budget Breakdown

Current smem allocation (block_m=128, block_n=128, dim=128, fp16):

q_shared_1:  half_m × dim × 2B  =  64 × 128 × 2  =  16 KB
q_shared_2:  half_m × dim × 2B  =  64 × 128 × 2  =  16 KB
k_smem_0:    block_n × dim × 2B = 128 × 128 × 2  =  32 KB
k_smem_1:    block_n × dim × 2B = 128 × 128 × 2  =  32 KB
v_smem_0:    block_n × dim × 2B = 128 × 128 × 2  =  32 KB
v_smem_1:    block_n × dim × 2B = 128 × 128 × 2  =  32 KB
barriers (8) + driver smem + padding            ≈   6 KB
─────────────────────────────────────────────────────────
Total:                                            ~166 KB
Allocated (NCU reported):                         165.9 KB
Remaining (of 228 KB max):                        ~62 KB

3-stage buffer would need +k_smem_2 + v_smem_2 = +64 KB → ~230 KB > 228 KB limit.

Even if it did fit, adding pipeline depth only helps when the bottleneck is producer-side TMA fill latency. On this shape, DRAM is 5% utilized and L2 hit rate is 91% — the producer is not the bottleneck.

10. Why 3-Stage Buffering Cannot Help

1. Doesn't fit: 166 KB + 64 KB = 230 KB > 228 KB smem limit per SM.
2. Wrong bottleneck: The producer (TMA loads) is not the bottleneck. DRAM throughput is 5.17%, and L2 hit rate is 90.7%. More prefetch depth doesn't help when data is already L2-resident.
3. Consumer-side stall: The top stalls (long_scoreboard 3.30, wait 1.59, barrier 1.13) are all on the consumer side — reading smem, waiting for wgmma, and synchronizing between warpgroups. More pipeline stages don't address these.

11. Bottleneck Model

Tensor core 60.9% active (of SM active) → 39.1% idle
  ├── [3.30] smem→register read latency (long_scoreboard)
  │     └── Consumer warps stall reading k_smem/v_smem through L1TEX
  │         L1 hit rate 72.2%, throughput 40.9% — swizzled layout contention?
  ├── [1.59] wgmma pipeline drain (wait / DEPBAR)
  │     └── Non-wgmma work (softmax, rescale, copy) finishes before wgmma
  │         → consumer idle-waits on DEPBAR
  │         → IntraWGOverlap (wait_group<1>) helps but doesn't fully hide
  └── [1.13] mbarrier synchronization (barrier)
        └── 8 mbarriers × multiple arrive/wait per K-loop iteration
            ~2% overhead from mbarrier scheduler vs named barriers (PR #872)

The two consumers ping-pong on the tensor core port. Ideal utilization would be 100% if each consumer's non-wgmma work (softmax + rescale + barriers + copies) exactly matches the wgmma latency of the other consumer. The 60.9% utilization implies the non-wgmma chain is shorter than wgmma, so consumers finish their non-wgmma work and then idle-wait for wgmma to complete.

12. Optimization Directions

Based on this profile, ranked by expected impact:

Direction	Mechanism	Target stall	Expected impact
TMA multicast	Broadcast K/V via cluster to multiple Q-head CTAs	long_scoreboard (reduces L1/L2 pressure)	High (for high group-ratio shapes)
Thread block cluster	Prerequisite for TMA multicast; also enables distributed smem	long_scoreboard, barrier	High
Named barriers	Replace mbarrier scheduler with PTX bar.arrive	barrier (1.13)	~2-5%
Reduce softmax path length	Fewer instructions in non-wgmma chain → less idle time	wait (1.59)	Medium
block_n=176	Larger wgmma tile → higher arithmetic intensity → wgmma takes longer → better overlap with non-wgmma work	wait (1.59)	Medium (needs TileLang wgmma_macro fix)
Split-K	Parallelize K loop across CTAs for long sequences	N/A for this shape, helps 32k/128k	Medium (long-seq only)
Tile ordering / L2 locality	KV-head-major persistent iteration order	long_scoreboard (L2 miss reduction)	Low-medium
Register pressure reduction	Reduce from 168 → ~160 regs/thread	Unlocks headroom for larger blocks or 3-stage	Low (indirect)

Not useful for this shape:

• 3-stage buffering (doesn't fit; wrong bottleneck)
• Increasing DRAM bandwidth utilization (already L2-resident)

Reproduction

# Profiling script at: ncu_profile_gqa.py in the PR branch
cd /tmp/tileops-pr871
TMPDIR=/home/ga/tmp TILELANG_CLEANUP_TEMP_FILES=1 CUDA_VISIBLE_DEVICES=7 \
  ncu --set full --target-processes all -f \
  -o /home/ga/tmp/gqa_ws_ncu \
  python ncu_profile_gqa.py

# Summary
TMPDIR=/home/ga/tmp ncu -i /home/ga/tmp/gqa_ws_ncu.ncu-rep --print-summary per-kernel

NCU report file: /home/ga/tmp/gqa_ws_ncu.ncu-rep

---

Profiled on 2026-04-10, branch feat/flash-attn/sm90-gqa-ws-persistent @ beab3ca, H200 GPU 7.

[superAngGao]

Update 2: Root cause identified — softmax reduction mechanism is the bottleneck

Clock instrumentation results (measured)

Added clock64() instrumentation to the WS persistent kernel's consumer 1 inner K-loop (clock_profile.py). Shape: B=1 S=4096 H=32 Hkv=8 D=128, causal, fp16, block_m=128, block_n=128.

Per K-loop iteration cycle breakdown (Consumer 1, thread 128):

Stage	Mean cycles/iter	% of total
softmax	1363	65.9%
k_full.wait (mbarrier)	486	23.5%
v_full.wait	83	4.0%
wait_wgmma (DEPBAR)	69	3.3%
wg_sched.wait	65	3.2%
TOTAL instrumented	~2067

Softmax is 66% of each iteration, not barriers or wgmma. The 1363 cycles come from TileLang's AllReduce<MaxOp/SumOp, 4, 1, 128, NamedBarrier<128>> — a cross-warp shared memory reduction using named barriers for 128 threads.

SASS-level PC sampling confirms softmax as bottleneck

The top stall instructions (long_scoreboard) in the NCU SASS source view are SYNCS.PHASECHK.TRANS64.TRYWAIT + @!P0 BRA spin-wait loops. These correspond to mbarrier try_wait inside the AllReduce cross-warp reduction (softmax's max/sum reductions), not the K/V pipeline barriers.

FA3 comparison — warp shuffle vs shared memory reduction

FA3's softmax (flash-attention/hopper/softmax.h) uses a fundamentally different reduction:

Aspect	TileOPs (TileLang)	FA3 (CUTLASS)
Reduction method	AllReduce<..., NamedBarrier<128>> — smem + named barrier, 128 threads	Allreduce<4>::run() — warp shuffle (__shfl_xor_sync), 4 threads
Softmax cycles/iter	~1363 (measured)	~40 (estimated from instruction count)
Cross-warp communication	Shared memory store → barrier → smem load	None — 4-thread quad group is intra-warp
Data path	Register → smem → barrier → smem → register	Register → shuffle → register

FA3's quad_allreduce_ (softmax.h:36-48) reduces within a 4-thread group using __shfl_xor_sync, requiring zero shared memory traffic. This is possible because the wgmma M=64 fragment distributes rows such that each 4-thread group already holds all the data needed for one row's reduction.

Revised bottleneck model

TileOPs per-iteration budget: ~2067 cycles
  ├── softmax: 1363 cycles (66%)
  │     ├── AllReduce<MaxOp>: smem write → barrier → smem read → local max
  │     ├── exp2f: element-wise (fast)
  │     └── AllReduce<SumOp>: smem write → barrier → smem read → local sum
  ├── k_full.wait: 486 cycles (23%) — pipeline turnaround
  ├── v_full.wait: 83 cycles (4%)
  ├── wait_wgmma: 69 cycles (3%)
  └── wg_sched.wait: 65 cycles (3%)

wgmma compute per iter: ~1024 cycles (QK 512 + SV 512)
SV wgmma overlaps with softmax start: ~512 cycles effective overlap
Tensor core idle: 2067 - 1024 - 512(overlap) ≈ 531 cycles/iter/consumer

Two consumers ping-pong → tensor utilization ≈ 1024 / (1024 + 531) ≈ 58%
(matches NCU's 57.7%)

If softmax drops to ~100 cycles (warp shuffle): total ≈ 800 cycles, idle ≈ 0 (softmax fits entirely within SV wgmma overlap window of 512 cycles), utilization → ~80%+.

Optimization priority (revised)

Priority	Direction	Expected impact
1	Replace TileLang's smem-based AllReduce with warp shuffle (quad_allreduce)	+20pp tensor utilization — the single largest optimization
2	Named barrier (bar.arrive) for ping-pong scheduler (issue #872 Gap 1)	+1.9pp (measured)
3	block_n=176 (issue #872 Gap 2)	+6pp (measured on non-causal)
4	4-WG softmax separation	Not recommended — FA3 proves inline softmax is correct if reduction is fast

Conclusion

The 4-WG approach (separating softmax to a dedicated WG) attacks the wrong bottleneck. The real fix is making the softmax reduction faster, not moving it elsewhere. FA3 demonstrates that with warp shuffle reduction (~40 cycles), softmax fits entirely within the SV wgmma overlap window and doesn't create tensor core bubbles.

The key TileLang gap: T.reduce_max / T.reduce_sum should support warp-shuffle-based reduction (like FA3's Allreduce<4>) as an alternative to the current smem + named barrier approach. This would bring TileOPs' tensor core utilization from 58% to ~80% without architectural changes to the kernel.

[superAngGao]

Update 3: FA3 SASS-level profiling — measured comparison

FA3 SASS PC sampling (same shape: B=1 S=4096 H=32 Hkv=8 D=128, causal, fp16, H200)

Top 20 hot SASS instructions from ncu --set full source-level PC sampling:

Samples	Instruction	Top Stalls	Interpretation
1187	HGMMA.64x128x16.F32 R88, R180, gdesc[UR16].tnspB, R88	not_sel=52, dispatch=45	SV wgmma — tensor core is busy (dispatch stall = good)
1008	UMOV UR23, 0x40000040	barrier=907	Named barrier wait (ping-pong scheduler)
441	WARPSYNC.ALL	sleep=391	Producer WG sleeping (nanosleep) — yields SM to consumers
423	WARPGROUP.DEPBAR.LE gsb0, 0x1	barrier=392	wait_wgmma(1) — waiting for QK completion
276	HGMMA.64x128x16.F32 R24, gdesc[UR32], R24	wait=119, dispatch=70, not_sel=70	QK wgmma (ss mode, both operands from smem)
271	HGMMA.64x128x16.F32 R24, gdesc[UR40], R24	wait=119, dispatch=75	QK wgmma continued
266	HGMMA.64x128x16.F32 R24, gdesc[UR48], RZ, !UPT	not_sel=1	QK wgmma first call (zero-init)
242	HGMMA.64x128x16.F32 R24, gdesc[UR20], R24	wait=99, not_sel=64	QK wgmma
237	HGMMA.64x128x16.F32 R24, gdesc[UR28], R24	wait=91, dispatch=63	QK wgmma
230	HGMMA.64x128x16.F32 R24, gdesc[UR24], R24	wait=99, not_sel=61	QK wgmma
219	HGMMA.64x128x16.F32 R24, gdesc[UR44], R24, gsb0	wait=120, dispatch=43	QK wgmma (last call, with scoreboard)
213	HGMMA.64x128x16.F32 R24, gdesc[UR48], RZ, !UPT	not_sel=64, wait=36	QK wgmma (2nd invocation)
207	FMNMX.FTZ R23, R20, R9, !PT	short_sb=187	Softmax max reduction — warp shuffle! short_scoreboard not long_scoreboard
207	FMNMX.FTZ R18, R23, R18, !PT	short_sb=190	Softmax max reduction continued
194	WARPGROUP.DEPBAR.LE gsb0, 0x0	barrier=191	wait_wgmma(0)
193	WARPGROUP.DEPBAR.LE gsb0, 0x0	barrier=192	wait_wgmma(0) (2nd invocation)
183	HGMMA.64x128x16.F32 R88, R176, gdesc[UR20].tnspB, R88	wait=89, dispatch=47	SV wgmma
178	UIADD3 UR4, UR6, 0x1, URZ	no_inst=120	Loop counter increment
165	@P1 BRA 0x...	long_sb=151	Only long_scoreboard in FA3 — 151 samples
163	HGMMA.64x128x16.F32 R88, R164, gdesc[UR16].tnspB, R88	wait=117, mio=32	SV wgmma

Side-by-side: TileOPs vs FA3 stall signatures

Softmax reduction instructions:

	TileOPs (smem AllReduce)	FA3 (warp shuffle)
Instruction	SYNCS.PHASECHK.TRANS64.TRYWAIT + @!P0 BRA (mbarrier spin-wait inside AllReduce)	FMNMX.FTZ (register-to-register max via __shfl_xor_sync)
Stall type	long_scoreboard (smem read latency)	short_scoreboard (register dependency, ~1-2 cycles)
Total samples	5255 (single hottest instruction)	207 (combined max reduction)
Data path	register → smem → barrier → smem → register	register → shuffle → register

Overall stall distribution:

Stall reason	TileOPs samples (top inst)	FA3 samples (top inst)
long_scoreboard	5255	151 (35x less)
barrier (named/mbarrier)	1437	907 + 392 = 1299
wait (DEPBAR)	—	119 (in HGMMA)
short_scoreboard	179 (softmax FMNMX)	187-190 (softmax FMNMX)
dispatch_stall	—	45-75 (tensor core busy)
sleep	0	391 (producer yields SM)

Key observations:

1. FA3's softmax max/sum reduction appears as FMNMX.FTZ with short_scoreboard — confirming warp shuffle implementation. Total ~414 samples vs TileOPs' smem AllReduce at ~5255 samples.

2. FA3's hottest instruction is an HGMMA (1187 samples) — tensor core is the bottleneck, not softmax. The dominant stalls are not_selected and dispatch (warp scheduling / tensor core contention), meaning the tensor core pipeline is saturated.

3. FA3 uses producer nanosleep (WARPSYNC.ALL with sleep=391) — the producer WG sleeps when K/V are loaded, yielding SM resources (warp scheduler slots, register file bandwidth) to consumer WGs. TileOPs' producer spin-waits on mbarrier, consuming warp scheduler slots and contributing to long_scoreboard.

4. FA3's only long_scoreboard is 151 samples on a single BRA — likely a pipeline boundary or epilogue edge case, not a structural bottleneck. TileOPs has 5255+1757+1186+... = 8000+ samples of long_scoreboard concentrated on mbarrier spin-waits inside the softmax AllReduce.

5. FA3's named barrier stall (907 samples on UMOV) is the ping-pong scheduler wait — both TileOPs and FA3 have this, but FA3 uses hardware named barriers (bar.sync/bar.arrive) while TileOPs uses mbarrier (smem-based).

Softmax cycle estimate from SASS

From the FA3 SASS, the softmax-related instructions are:

• FMNMX.FTZ (max reduction): 2 instructions, ~207 samples each
• MUFU.EX2 (exp2f): not in top 20 → very few samples → fast
• FADD/FMUL (sum, scale): not in top 20 → fast

Compared to the total kernel samples (~3200 per invocation), softmax accounts for ~414/3200 ≈ 13% of FA3's execution time. In TileOPs, softmax accounts for ~8000/15000 ≈ 53% of execution time (estimated from PC sampling proportions).

This aligns with the clock instrumentation measurement: TileOPs softmax = 66% of iteration time (1363/2067 cycles).

Producer nanosleep design

FA3's producer WG uses nanosleep after loading data, visible in SASS as WARPSYNC.ALL with sleep stall. This is significant because:

• Sleeping warps don't consume warp scheduler slots
• More scheduler slots available for consumer WGs → better instruction-level parallelism
• TileOPs' producer spin-waits on mbarrier (SYNCS.PHASECHK.TRANS64.TRYWAIT + BRA loop), consuming scheduler cycles that could go to consumers

This is a secondary optimization but worth noting — FA3's producer is designed to be out of the way when not loading data.

[superAngGao]

Update 4: Why two consumers don't fully overlap — the stagger gap

Clarification: acc_s is NOT the bottleneck

The two consumer WGs (WG1, WG2) each have their own acc_s registers — they're split along the Q (M) dimension. WG2's QK wgmma writes acc_s_2, which does NOT conflict with WG1's softmax reading acc_s_1. The ping-pong overlap IS happening.

The real issue: both WGs enter softmax simultaneously

The ping-pong stagger is too small. WG1 signals WG2 (arrive(sched_12)) at step 8 — right after issuing QK+SV, which is very early in the iteration. WG2 then quickly completes its own QK+SV and enters softmax. Timeline:

WG1: [QK+SV issue ~20cy] → arrive(sched) → wait_wgmma(1) → ─── softmax (1363cy) ───
WG2:                        start → [QK+SV issue] → wait_wgmma(1) → ─── softmax (1363cy) ───
TC:  ████ WG1 QK+SV ████████ WG2 QK+SV █████████████  idle (~339cy)  ████ WG1 next ████
                                                        ↑
                                                   Both WGs in softmax
                                                   No one feeding TC

WG2's wgmma (1024 cycles) is shorter than WG1's softmax (1363 cycles). So for ~339 cycles, both WGs are in softmax and TC is idle. Adding the k_full.wait overhead (486 cycles/iter from clock measurement), the total TC idle per consumer iteration is ~825 cycles.

Utilization calculation

Per consumer iteration: wgmma = 1024 cycles, total = 2067 cycles.
Single consumer utilization: 1024 / 2067 = 49.5%.
Two consumers with imperfect stagger: measured 57.7% (NCU), matching this model.

If softmax drops to ~500 cycles (warp shuffle): softmax < WG2's wgmma (1024 cycles), so WG1's softmax would be fully covered by WG2's wgmma. No simultaneous-softmax gap. Both WGs always have wgmma in-flight when the other does softmax → TC utilization ≈ 80%+.

Conclusion

The bottleneck is softmax duration (1363 cycles), not a barrier or register conflict. The two WGs already overlap via ping-pong, but softmax is longer than one WG's wgmma window (1024 cycles), creating an unavoidable gap where both WGs are in softmax simultaneously.

Fix: reduce softmax from 1363 to <1024 cycles (ideally ~500). This makes softmax fit entirely within the other WG's wgmma window, eliminating the simultaneous-softmax gap. FA3 achieves this via warp shuffle (quad_allreduce) instead of smem-based cross-warp AllReduce.