87 lines (74 loc) · 3.94 KB

Primer — RocmKernelWiki Topic Map

Use this as the first stop when a question is broad. Maps hardware features, techniques, symptoms, kernels, and DSLs to canonical page IDs. Page IDs resolve via python3 scripts/get_page.py <id>.

Hardware features

Concept	Page id	File
CDNA3 matrix core (MFMA)	`hw-mfma-cdna3`	`wiki/hardware/mfma-cdna3.md`
CDNA4 matrix core (MFMA + new dtype support)	`hw-mfma-cdna4`	`wiki/hardware/mfma-cdna4.md`
Structured 4:2 sparsity via SMFMAC	`hw-smfmac-sparsity`	`wiki/hardware/smfmac-sparsity.md`
Dual register files (VGPR + AccVGPR)	`hw-accvgpr`	`wiki/hardware/accvgpr.md`
LDS layout, banking, swizzling	`hw-lds-cdna`	`wiki/hardware/lds-cdna.md`
Buffer resource descriptors / `buffer_load*`	`hw-buffer-resource`	`wiki/hardware/buffer-resource.md`
Async direct-to-LDS copy (`global_load_lds_*`)	`hw-async-copy-cdna3`	`wiki/hardware/async-copy-cdna3.md`
Global wave sync (GWS)	`hw-gws`	`wiki/hardware/gws.md`
Wave64 semantics, EXEC mask	`hw-wave64`	`wiki/hardware/wave64.md`
OCP FP8 / BF8 formats	`hw-fp8-bf8`	`wiki/hardware/fp8-bf8.md`
CDNA4 microscaling (MX-FP4/FP6/FP8)	`hw-mx-formats`	`wiki/hardware/mx-formats.md`

Techniques

Concept	Page id
Coalesced linear access	`technique-coalesced-access`
Vectorized load/store (`float4`, `dwordx4`)	`technique-vectorized-loads`
Grid-stride loop	`technique-grid-stride-loop`
Software pipelining (mfma + global + lds)	`technique-software-pipelining`
LDS double-buffering (ping-pong)	`technique-lds-double-buffer`
LDS swizzling / bank-conflict avoidance	`technique-lds-swizzling`
Register budgeting	`technique-register-budgeting`
Occupancy tuning (`waves_per_eu`)	`technique-occupancy-tuning`
Kernel fusion	`technique-kernel-fusion`
Epilogue fusion	`technique-epilogue-fusion`
Persistent kernel	`technique-persistent-kernel`
Stream-K / Split-K scheduling	`technique-stream-k`
Triton autotune knobs (`num_stages`, `num_warps`, `matrix_instr_nonkdim`, `kpack`)	`technique-triton-autotune-amd`
`s_waitcnt` tuning (vmcnt / lgkmcnt)	`technique-s-waitcnt-tuning`
Cache control (GLC/SLC)	`technique-cache-policy`

Patterns (symptom → candidate techniques)

Symptom	Page id
Memory-bound, low HBM bandwidth	`pattern-memory-bound`
Compute-bound, low MFMA issue rate	`pattern-compute-bound`
LDS bank conflict	`pattern-lds-bank-conflict`
Register spilling	`pattern-register-pressure`
Low CU utilization	`pattern-low-cu-utilization`
Poor coalescing / scattered global access	`pattern-poor-coalescing`
Pipeline stall (high `vmcnt` / `lgkmcnt`)	`pattern-pipeline-stall`
Tail effect	`pattern-tail-effect`
MoE load imbalance	`pattern-moe-load-imbalance`

Kernels

Kernel	Page id
GEMM in HIP using MFMA	`kernel-gemm-hip-mfma`
GEMM in Triton AMD	`kernel-gemm-triton-amd`
GEMM in Composable Kernel	`kernel-gemm-ck`
FlashAttention (HIP) on MI300/MI350	`kernel-flash-attention-hip`
FlashAttention (Triton-AMD)	`kernel-flash-attention-triton-amd`
Paged Attention (HIP)	`kernel-paged-attention-hip`
Fused MoE on ROCm	`kernel-fused-moe-rocm`
RMSNorm (HIP)	`kernel-rmsnorm-hip`

Languages / DSLs

Language	Page id
HIP C++	`lang-hip-cpp`
Triton — ROCm backend	`lang-triton-amd`
Composable Kernel	`lang-composable-kernel`
Gluon-AMD (gfx942/gfx950)	`lang-gluon-amd`
FlyDSL	`lang-flydsl`
AMDGCN inline asm	`lang-amdgcn-asm`

Migration

Direction	Page id
CUDA → HIP	`migration-cuda-to-hip`
Triton-NVIDIA → Triton-AMD	`migration-triton-nvidia-to-amd`
WGMMA → MFMA	`migration-wgmma-to-mfma`
TMA → buffer-load + async-copy	`migration-tma-to-buffer-loads`
MI300 → MI350 (gfx942 → gfx950)	`migration-mi300-to-mi350`