# Lds Optimization

> Optimize LDS (Local Data Share / shared memory) access patterns in FlyDSL GPU kernels. Diagnose bank conflicts and high lgkmcnt stalls from ATT trace data, then apply swizzle or padding layouts to eliminate conflicts. Also increase the distance between LDS write and subsequent LDS read to hide LDS latency. LDS read preceded by write always requires a sync (s_waitcnt lgkmcnt or s_barrier). Use when trace analysis shows ds_read/ds_write/lgkmcnt as a bottleneck. Usage: /lds-optimization

- Skill: `rocm/lds-optimization` (Agent Skill)
- Install (CLI): `npx skillmds@latest add rocm/lds-optimization`
- Raw SKILL.md: https://api.skillmd.com/api/skills/rocm/lds-optimization/raw
- Safety review: pending
- Works with: Claude Code, Claude.ai, OpenAI Codex
- Category: Coding & Dev Tools
- Author: rocm (https://skillmd.com/u/rocm)
- Updated: 2026-09-17
- Page: https://skillmd.com/skills/rocm/lds-optimization

---


# LDS Optimization

Diagnose and fix LDS (shared memory) performance issues in FlyDSL kernels
on AMD CDNA GPUs (MI300X/MI308/MI350).

## When To Use

Run `/kernel-trace-analysis` first. Apply this skill when the trace shows:

| Signal | Threshold | Example |
|--------|-----------|---------|
| `s_waitcnt lgkmcnt(0)` with high stall | > 3000 cycles per instance | `L605: stall=4080 s_waitcnt lgkmcnt(0)` |
| `ds_write` / `ds_read` with high latency | > 500 cycles per instance | `L761: stall=960 ds_write2_b32` |
| Multiple `s_barrier` between `ds_write` and `ds_read` | Barrier stall > 5000 | `L606: stall=17024 s_barrier` |
| Total LDS-related stall > 15% of kernel stall | Sum all lgkmcnt + ds stalls | Softmax reduce phase in PA decode |

## LDS Architecture on CDNA3 (gfx942)

### Hardware Facts

- LDS size: **64 KB per CU** (workgroup-shared)
- LDS is organized into **32 banks**, each **4 bytes wide**
- Bank index = `(byte_address / 4) % 32`
- **Bank conflict**: when 2+ threads in the same wavefront access **different addresses** in the **same bank** in the same cycle, accesses are serialized
- **Broadcast**: when 2+ threads access the **same address** in the same bank, hardware broadcasts (no conflict)
- LDS throughput: **128 bytes/cycle** (peak, no conflicts)
- LDS latency: **~20-40 cycles** (async, hidden if enough work between write and read)
- **VGPR context**: LDS ops use **arch_vgpr** (for addresses/data), not accum_vgpr. But occupancy on CDNA3 is governed by the **combined** 512-entry budget `arch_vgpr + accum_vgpr` (the two physical files share one occupancy budget — see `/kernel-trace-analysis`). So LDS-addressing logic that grows arch_vgpr *can* still cost occupancy even though it never touches the MFMA accumulators. Keep LDS-address VGPR pressure low when the kernel is near a 2-wave boundary.

## LDS Architecture on CDNA4 (gfx950)

### Hardware Facts

- LDS size: **160 KB per CU** (2.5x larger than gfx942's 64 KB)
- LDS is organized into **64 banks**, each **4 bytes wide** (640 DWords per bank)
- Bank index = `(byte_address / 4) % 64`
- **Bank conflict**: same rule as gfx942 — when 2+ threads in the same wavefront access **different addresses** in the **same bank** in the same cycle, accesses are serialized
- **Broadcast**: same as gfx942 — when 2+ threads access the **same address** in the same bank, hardware broadcasts (no conflict)
- LDS throughput: **256 bytes/cycle** (peak, no conflicts; 2x gfx942 due to 64 banks)
- LDS latency: **~2-64 cycles** per operation depending on bank conflicts (2 cycles best case, 64 cycles worst case with all threads conflicting on one bank)
- LDS allocation granularity: **1280 bytes** on **1280-byte alignment**; LDS allocations do not wrap around the LDS storage
- **Wavefront dispatch**: reads across a 64-thread wavefront are dispatched over **4 cycles** in waterfall fashion
- **32 concurrent LDS operations**: hardware can concurrently execute 32 read or write instructions (each 32-bit); extended instructions (read2/write2) can be 64-bit each
- **32 integer atomic units** for unordered atomic operations
- **VGPR context**: same as gfx942 — LDS ops use **arch_vgpr** (for addresses/data), but occupancy is governed by the combined 512-entry budget `arch_vgpr + accum_vgpr`. LDS-addressing arch_vgpr competes with MFMA accumulators for that shared budget.

### Key Differences from gfx942

| Aspect | gfx942 (CDNA3) | gfx950 (CDNA4) |
|--------|----------------|-----------------|
| LDS size per CU | 64 KB | 160 KB |
| Number of banks | 32 | 64 |
| Bank index formula | `(addr/4) % 32` | `(addr/4) % 64` |
| Peak throughput | 128 bytes/cycle | 256 bytes/cycle |
| LDS allocation granularity | 256 bytes | 1280 bytes |
| Max LDS per workgroup | 64 KB | 160 KB |
| MFMA Transpose Load | Not available | `DS_READ_B64_TR_B16/B8/B4`, `DS_READ_B96_TR_B6` |

### Impact on Bank Conflict Analysis

Because gfx950 has **64 banks** instead of 32, the bank conflict patterns change:

- **Stride that causes conflicts**: multiples of 64 banks (256 bytes) instead of 32 banks (128 bytes)
- A stride of 128 bytes that caused **full conflict on gfx942** (all threads hit same bank) will only cause **partial conflict on gfx950** (threads alternate between 2 banks)
- To cause full 64-way conflict on gfx950, the stride must be a multiple of `64 * 4 = 256` bytes
- **XOR swizzle masks may need adjustment** — masks designed for 32-bank gfx942 may be suboptimal on 64-bank gfx950

### MFMA Transpose Load from LDS (gfx950 only)

CDNA4 introduces dedicated instructions for transposing data while loading from LDS to VGPRs, eliminating the need for explicit transpose via `ds_write` + `ds_read` with permuted addresses:

| Instruction | Element Size | VGPRs Written | Description |
|-------------|-------------|---------------|-------------|
| `DS_READ_B64_TR_B16` | 16-bit (fp16/bf16) | 2 VGPRs | Load column-major A or row-major B matrix; two instructions load a complete matrix. Each lane holds 4 consecutive M or N values. |
| `DS_READ_B64_TR_B8` | 8-bit (fp8/bf8) | 2 VGPRs | Same as B16 but for 8-bit elements. First loads K=0..7,16..23,32..39,48..55; second loads remaining K values. |
| `DS_READ_B64_TR_B4` | 4-bit (int4) | 2 VGPRs | Same pattern for 4-bit elements. First loads K=0..15,32..47; second loads remaining K values. |
| `DS_READ_B96_TR_B6` | 6-bit | 3 VGPRs | 6-bit element transpose load into 3 VGPRs. Does NOT require even-VGPR alignment. |

Requirements:
- EXEC mask must be set to all 1's before executing
- LDS address must be aligned to the data size
- DS ops reading/writing 64-bit or larger data must use even-aligned VGPRs (except `DS_READ_B96_TR_B6`)

These instructions are useful for MFMA operand preparation — loading A/B matrices from LDS in the transposed layout needed by MFMA instructions without explicit LDS-based transpose.

### LDS Instruction Model

LDS operations (`ds_read_*`, `ds_write_*`, `ds_bpermute_*`, `ds_swizzle_*`) are **asynchronous**:

```
ds_write_b32 v_addr, v_data    ; issues async write, returns immediately
; ... other instructions ...    ; LDS write completes in background
s_waitcnt lgkmcnt(0)            ; stall until all LDS/SMEM ops complete
ds_read_b32 v_result, v_addr   ; now safe to read
```

Key rules:
1. **Write-before-read requires sync**: any `ds_read` that depends on a prior `ds_write` must have `s_waitcnt lgkmcnt(0)` or `s_barrier` in between
2. **`s_barrier` implies cross-wave sync**: if wave A writes and wave B reads, `s_barrier` is required (not just `lgkmcnt`)
3. **Longer write-read distance = better latency hiding**: more instructions between `ds_write` and the subsequent `s_waitcnt lgkmcnt(0)` allow the write to complete in the background

## Diagnosing LDS Bottlenecks from Trace

### Step 1: Identify LDS-heavy regions

```python
import json

with open('ui_output_agent_XXX_dispatch_YYY/code.json') as f:
    data = json.load(f)
instructions = data['code']
# Columns: [ISA, _, LineNum, Source, Codeobj, Vaddr, Hit, Latency, Stall, Idle]

# Find all LDS-related instructions
lds_insts = [i for i in instructions if i[0].startswith('ds_') or
             ('lgkmcnt' in i[0] and i[8] > 0)]

total_lds_stall = sum(i[8] for i in lds_insts)
total_stall = sum(i[8] for i in instructions)
print(f"LDS stall: {total_lds_stall} / {total_stall} = {100*total_lds_stall/total_stall:.1f}%")

# Show hottest LDS instructions
for i in sorted(lds_insts, key=lambda x: x[8], reverse=True)[:15]:
    print(f"  L{i[2]:>4d}  stall={i[8]:>6d}  idle={i[9]:>6d}  {i[0][:55]}  | :{i[3].split(':')[-1]}")
```

### Step 2: Classify the bottleneck type

**Type A: Bank Conflicts** (high stall on `ds_read`/`ds_write` themselves)

```
L 766  stall=  160  ds_read2_b64 v[44:47], v28 offset1:8        ; <-- bank conflict
L 767  stall=  320  ds_read2_b64 v[36:39], v28 offset0:16 offset1:24  ; <-- bank conflict
```

Signs:
- `ds_read_*` / `ds_write_*` instructions with stall > 100 cycles per hit
- Multiple reads/writes with similar base address but different offsets that map to same banks
- `ds_read2_b64` / `ds_write2_b32` with offsets that are multiples of the bank count:
  - **gfx942**: multiples of 32 (= same bank, 32-bank LDS)
  - **gfx950**: multiples of 64 (= same bank, 64-bank LDS)

**Type B: Write-Read Latency Exposed** (high stall on `s_waitcnt lgkmcnt(0)` after `ds_write`)

```
L 761  stall=  960  ds_write2_b32 v28, v41, v43 offset0:32 offset1:48
L 764  stall= 4560  s_waitcnt lgkmcnt(0)    ; <-- write latency fully exposed
L 765  stall= 1468  s_barrier
L 766  stall=  160  ds_read2_b64 v[44:47], v28 offset1:8
```

Signs:
- `s_waitcnt lgkmcnt(0)` with > 2000 stall cycles immediately after `ds_write`
- Very few instructions between `ds_write` and `s_waitcnt`
- This means the write hasn't completed by the time we need to wait

**Type C: Cross-Wave Reduce Serialization** (high stall on `s_barrier` in reduce chains)

```
L 605  stall= 4080  s_waitcnt lgkmcnt(0)     ; wait for ds_bpermute
L 606  stall=17024  s_barrier                 ; cross-wave sync
L 607  stall=27220  s_waitcnt vmcnt(0)        ; also waiting for global loads
```

Signs:
- `ds_bpermute` -> `lgkmcnt(0)` -> `s_barrier` -> `ds_write LDS` -> `lgkmcnt(0)` -> `s_barrier` -> `ds_read LDS` pattern
- Multiple barriers (> 4) in a reduce region

## Optimization Method 1: Swizzle Layout

### The Problem

When multiple threads access LDS with a stride that is a multiple of the bank count, every access hits the same bank:

- **gfx942 (32 banks)**: stride multiple of 128 bytes causes full conflict
- **gfx950 (64 banks)**: stride multiple of 256 bytes causes full conflict; stride of 128 bytes causes 2-way conflict (threads alternate between 2 banks)

```
# gfx942 (32 banks): stride=128 -> full conflict
Thread 0: addr = base + 0*128  -> bank (0*128/4)%32 = 0
Thread 1: addr = base + 1*128  -> bank (1*128/4)%32 = 0  <- CONFLICT
Thread 2: addr = base + 2*128  -> bank (2*128/4)%32 = 0  <- CONFLICT

# gfx950 (64 banks): stride=128 -> only 2-way conflict (NOT full conflict)
Thread 0: addr = base + 0*128  -> bank (0*128/4)%64 = 0
Thread 1: addr = base + 1*128  -> bank (1*128/4)%64 = 32  <- different bank!
Thread 2: addr = base + 2*128  -> bank (2*128/4)%64 = 0   <- conflict with thread 0

# gfx950 (64 banks): stride=256 -> full conflict
Thread 0: addr = base + 0*256  -> bank (0*256/4)%64 = 0
Thread 1: addr = base + 1*256  -> bank (1*256/4)%64 = 0  <- CONFLICT
...
```

### The Solution: XOR-Based Swizzle

Swizzle XORs bits of the row index into the column index of the LDS address, distributing accesses across different banks:

```
swizzled_col = original_col XOR (row >> shift)
```

This ensures threads accessing the same column in different rows hit different banks.

### FlyDSL XOR Swizzle with SharedAllocator

In FlyDSL, declare the LDS storage as an `@fx.struct` of `fx.Array` fields and
allocate it inside the kernel with `fx.SharedAllocator` (the current LDS API).
To apply swizzle, XOR the row index into the LDS address when computing
store/load offsets:

```python
import flydsl.expr as fx

@fx.struct
class SharedStorage:
    key: fx.Array[fx.Float16, KV_BLOCK_SIZE * HEAD_SIZE]

@flyc.kernel
def my_kernel(...):
    lds = fx.SharedAllocator().allocate(SharedStorage).peek()
    lds_key = lds.key.view(fx.make_layout((KV_BLOCK_SIZE, PADDED_STRIDE), (PADDED_STRIDE, 1)))

    # XOR-swizzle address: distribute bank accesses
    # row_idx and col_idx are the logical 2D coordinates
    # XOR_BITS controls swizzle width (typically 4 for fp16 vec=8)
    swizzled_col = col_idx ^ (row_idx & XOR_MASK)
    fx.memref_store(data, lds_key, [row_idx, swizzled_col])
```

The legacy `flydsl.utils.smem_allocator.SmemAllocator` path remains for
un-migrated kernels but is not recommended for new code.

### Choosing Swizzle Parameters

The goal is to make vectorized access span enough banks:

| Data Type | Element Size | Recommended Vec Width | Banks Covered per Vec |
|-----------|-------------|----------------------|----------------------|
| fp32      | 4 bytes     | 4                    | 4 banks (16 bytes)   |
| fp16/bf16 | 2 bytes     | 8                    | 4 banks (16 bytes)   |
| fp8       | 1 byte      | 16                   | 4 banks (16 bytes)   |

For XOR mask:
- **gfx942 (32 banks)**: use `32 / (vec * element_size / 4) - 1` to ensure full bank coverage
- **gfx950 (64 banks)**: use `64 / (vec * element_size / 4) - 1` — the doubled bank count means wider XOR masks may be needed to fully distribute accesses

### Example: Fix Bank Conflicts in KV Cache Load to LDS

Before (conflict-prone, linear layout):

```python
# Linear shared memory layout — threads in same warp hit same banks
@fx.struct
class SharedStorage:
    key: fx.Array[fx.Float16, KV_BLOCK_SIZE * HEAD_SIZE]

lds = fx.SharedAllocator().allocate(SharedStorage).peek()
lds_key = lds.key.view(fx.make_layout((KV_BLOCK_SIZE, HEAD_SIZE), (HEAD_SIZE, 1)))
# Store key tile: all threads write to column 0,1,2... -> bank conflicts
fx.memref_store(data, lds_key, [row, col])
```

After (swizzled, conflict-free):

```python
# XOR-swizzle distributes accesses across banks
XOR_BITS = 4  # for fp16 vec=8: covers 4 banks per vec
lds = fx.SharedAllocator().allocate(SharedStorage).peek()
lds_key = lds.key.view(fx.make_layout((KV_BLOCK_SIZE, HEAD_SIZE), (HEAD_SIZE, 1)))
swizzled_col = col ^ ((row & 0x7) << XOR_BITS)
fx.memref_store(data, lds_key, [row, swizzled_col])  # now conflict-free
```

## Optimization Method 2: Padding

### The Problem

Same as swizzle — stride-aligned accesses cause bank conflicts. Padding adds extra unused elements to change the effective stride.

### The Solution

Add padding per row to break the alignment. The rule is about the **row stride
in dwords**, not in elements: lane `t` reading row `t` lands on bank
`(t * row_stride_bytes / 4) % NBANKS`, so the wavefront touches
`NBANKS / gcd(row_stride_dwords, NBANKS)` distinct banks — and a *fractional*
dword stride (an odd number of 2-byte elements) walks all of them.

```text
fp16, HEAD_SIZE = 128, gfx950 (64 banks):
  pad 0  -> 128 elems -> 256 B -> 64.0 dw -> gcd(64,64)=64 ->  1 bank  -> 64-way
  pad 1  -> 129 elems -> 258 B -> 64.5 dw -> fractional    -> 64 banks ->  1-way
  pad 2  -> 130 elems -> 260 B -> 65.0 dw -> gcd(65,64)=1  -> 64 banks ->  1-way
  pad 4  -> 132 elems -> 264 B -> 66.0 dw -> gcd(66,64)=2  -> 32 banks ->  2-way
  pad 32 -> 160 elems -> 320 B -> 80.0 dw -> gcd(80,64)=16 ->  4 banks -> 16-way
```

**More padding is not safer.** Only a fractional-dword stride, or one coprime
with the bank count, is conflict-free; a large even padding can be nearly as bad
as none.

### FlyDSL Padding Implementation

```python
PADDING = 1  # or a small number
PADDED_HEAD_SIZE = HEAD_SIZE + PADDING

# Allocate with extra column for padding
@fx.struct
class SharedStorage:
    key: fx.Array[fx.Float16, KV_BLOCK_SIZE * PADDED_HEAD_SIZE]

@flyc.kernel
def my_kernel(...):
    lds = fx.SharedAllocator().allocate(SharedStorage).peek()
    # View with padded stride: physical row stride is PADDED_HEAD_SIZE
    lds_key = lds.key.view(fx.make_layout((KV_BLOCK_SIZE, HEAD_SIZE), (PADDED_HEAD_SIZE, 1)))

    # Write key data using padded stride (padding column stays unused)
    fx.memref_store(data, lds_key, [row, col])

    # Read back using same padded-stride view
    data = fx.memref_load(lds_key, [row, col])
```

### Padding Amount

Pick the smallest padding whose row stride is fractional-dword or coprime with
the bank count — for 2-byte elements that is **+1 or +2 elements**. Do **not**
use `padding = bank_count / element_size_bytes`: for fp16 on gfx950 that gives
+32 elements, an 80-dword stride, and a 16-way conflict.

Measured on gfx950 (MI355X), one wavefront strided-column-reading a
`[64][64+PAD]` fp16 tile, best of 9 x 50 launches, `hipcc -O3`. Note this is a
64-element row (128 B unpadded), so its strides differ from the 128-element
worked example above; the predicted column is `64 / gcd(row_stride_dwords, 64)`
banks:

| PAD | row stride | predicted | measured | vs best |
|-----|-----------|-----------|----------|---------|
| 0 | 128 B (32.0 dw) | 32-way | 7.11 ms | +59% |
| 1 | 130 B (32.5 dw) | 1-way | 4.61 ms | +3% |
| 2 | 132 B (33.0 dw) | 1-way | **4.48 ms** | best |
| 4 | 136 B (34.0 dw) | 2-way | 4.62 ms | +3% |
| 32 | 192 B (48.0 dw) | 16-way | 5.81 ms | +30% |

Unpadded is clearly worst, and the `bank_count / element_size` formula is 26%
slower than a single element of padding. (The microbenchmark is latency-bound on
one wavefront, so absolute ratios are compressed relative to the conflict factor
and it cannot separate 1-way from 2-way; use it for the ordering, not for an
expected speedup.)

The cost of padding is extra LDS usage:
- 1 element padding per row: `KV_BLOCK_SIZE * element_size` extra bytes
- Must ensure total LDS usage stays within **64 KB** per CU (gfx942) or **160 KB** per CU (gfx950)

### Swizzle vs Padding Trade-offs

| Aspect | Swizzle | Padding |
|--------|---------|---------|
| LDS overhead | None (zero extra bytes) | Extra bytes per row |
| Complexity | Need correct XOR mask params (arch-dependent) | Simple: just add 1 to stride |
| Address computation | XOR adds ~1 SALU instruction | Simple offset, no extra compute |
| Risk | Wrong params = silent bank conflicts | Exceeding LDS limit = kernel fail |
| LDS limit | N/A | gfx942: 64 KB, gfx950: 160 KB |
| Preferred when | LDS near capacity, need zero overhead | Simple cases, LDS has headroom |

**Recommendation**: Prefer swizzle (zero overhead). Use padding only when swizzle layout is hard to integrate with the kernel's access pattern. On gfx950, the 160 KB LDS gives much more headroom for padding.

## Optimization Method 3: Increase Write-Read Distance

### The Problem

When `ds_write` is immediately followed by `s_waitcnt lgkmcnt(0)` and then `ds_read`, the ~20-40 cycle LDS write latency is fully exposed as stall:

```
ds_write_b32 ...          ; async write issued
s_waitcnt lgkmcnt(0)      ; STALL: write hasn't completed yet (3000+ cycles)
ds_read_b32 ...           ; read must wait for write
```

### The Solution

Insert useful compute work between the write and the wait:

```
ds_write_b32 ...          ; async write issued
; --- insert independent compute here ---
v_mfma_f32_16x16x32 ...  ; MFMA takes ~64 cycles, overlaps with LDS write
v_add_f32 ...             ; more independent ALU work
v_mul_f32 ...
; --- write has completed by now ---
s_waitcnt lgkmcnt(0)      ; no stall (or minimal stall)
ds_read_b32 ...           ; data ready immediately
```

### FlyDSL-Level Implementation

At the Python/FlyDSL level, you control write-read distance by reordering operations.
The snippet below is schematic — it illustrates *ordering*, and spells the global
load in the raw-intrinsic form. On the current surface that load is `fx.copy` from
a `make_buffer_tensor` view (see **kernel-code-cleanup**); the ordering argument is
the same either way.

```python
# BEFORE: write and read are close together
lds_ptr.store(data, [offset])               # ds_write
gpu.barrier()                                # s_barrier (includes lgkmcnt wait)
result = lds_ptr.load([offset])             # ds_read

# AFTER: insert independent work between write and barrier
lds_ptr.store(data, [offset])               # ds_write (async)

# Do independent compute that doesn't need the LDS data
next_offsets = compute_next_offsets()        # SALU/VALU work
next_data = buffer_ops.buffer_load(rsrc, next_offsets, vec_width=4)  # global load (also async)
scale_factor = buffer_ops.buffer_load(rsrc_scale, scale_off, vec_width=1)

gpu.barrier()                                # by now, LDS write has completed
result = lds_ptr.load([offset])             # ds_read (no stall)
```

### What to Insert Between Write and Read

Prioritize by latency-hiding value:

1. **Global loads for next phase** — also async, ~300+ cycle latency
2. **Address computation** (`compute_offsets`) — SALU/VALU, ~4-8 cycles each
3. **Independent MFMA chains** — if available, ~64 cycles per MFMA
4. **Scalar loads** (`s_load_dword*`) — kernel arguments, ~20 cycles

Avoid inserting:
- Operations that depend on the LDS write result (data dependency)
- More LDS operations (would compete for LDS bandwidth)
- Operations that increase register pressure beyond budget

## Verification Checklist

After applying LDS optimizations:

1. **Correctness**: Run tests. Swizzle changes must be applied consistently to both write and read paths — if the write uses swizzled addresses, the read must use the same swizzle.

2. **Re-profile**: Run `/kernel-trace-analysis` and check:
   - `ds_read_*` / `ds_write_*` stall should decrease
   - `s_waitcnt lgkmcnt(0)` stall after `ds_write` should decrease
   - No new bank conflicts introduced

3. **LDS usage**: Check total LDS consumption:
   ```python
   # Estimate: sum of all @fx.struct fx.Array field byte sizes
   # gfx942: Must be <= 65536 bytes (64 KB) per workgroup
   # gfx950: Must be <= 163840 bytes (160 KB) per workgroup
   # Note: gfx950 allocates LDS in 1280-byte granularity (1280-byte aligned blocks)
   ```

4. **Register pressure**: Swizzle adds ~1-2 SALU instructions for address XOR. Padding doesn't add register pressure but uses more LDS. Neither should significantly impact VGPR count.

## Quick Reference: Common LDS Patterns in Paged Attention

| Pattern | Location | Typical Issue | Fix |
|---------|----------|---------------|-----|
| K/V cache tile in LDS | QK/PV MFMA loop | Bank conflicts from stride=HEAD_SIZE | Swizzle with XOR on row index |
| Softmax reduce via LDS | `ds_write -> barrier -> ds_read` | Write-read latency exposed + too many barriers | Increase write-read distance; replace with `ds_bpermute` chain |
| Cross-wave max/sum broadcast | `ds_write -> barrier -> ds_read` from different wave | Cross-wave sync overhead | Merge max+sum into single reduce pass |
| MFMA accumulator shuffle | `ds_write accum -> barrier -> ds_read permuted` | Bank conflicts if accumulator layout misaligns | Swizzle or use `ds_bpermute` for permutation |

## Output

After optimization, report:
- Which LDS bottleneck type was identified (bank conflict / write-read latency / reduce serialization)
- Which optimization was applied (swizzle / padding / distance increase)
- Before/after `lgkmcnt` stall cycles and `ds_*` instruction stalls
- LDS usage before/after (bytes)
- Any impact on VGPR count or occupancy