Skip to content

examples.md

Latest commit

 

History

History
96 lines (66 loc) · 3.58 KB

examples.md

File metadata and controls

96 lines (66 loc) · 3.58 KB
Raw

Worked Query Patterns — RocmKernelWiki

Each pattern: user question → command sequence → synthesis. Use these as templates when answering from the KB.

1. "How do I avoid LDS bank conflicts on MI300?"

python3 scripts/query.py "LDS bank conflict" --architecture gfx942
python3 scripts/get_page.py pattern-lds-bank-conflict --follow-sources
python3 scripts/get_page.py technique-lds-swizzling

Synthesis: cite wiki/patterns/lds-bank-conflict.md for symptoms, then wiki/techniques/lds-swizzling.md and wiki/techniques/lds-double-buffer.md for fixes. Reference hw-lds-cdna for hardware background.

2. "Which MFMA variant should I pick for a 64×64 tile in BF16 on MI300X?"

python3 scripts/query.py --tag mfma --architecture MI300X --type hardware
python3 scripts/get_page.py hw-mfma-cdna3

Synthesis: read the MFMA-variant table in wiki/hardware/mfma-cdna3.md; recommend mfma_32x32x8_bf16 or mfma_16x16x16_bf16 based on the surrounding loop structure. Cite confidence level.

3. "How do I tune Triton GEMM on MI300?"

python3 scripts/query.py "triton autotune" --architecture MI300X
python3 scripts/get_page.py technique-triton-autotune-amd
python3 scripts/get_page.py kernel-gemm-triton-amd --follow-sources

Synthesis: name the AMD-specific knobs (matrix_instr_nonkdim, kpack, waves_per_eu, num_stages, num_warps); link to the GEMM case study for representative configs.

4. "What's different on MI350 vs MI300 for FP4 GEMM?"

python3 scripts/get_page.py migration-mi300-to-mi350
python3 scripts/get_page.py hw-mx-formats
python3 scripts/get_page.py hw-mfma-cdna4

Synthesis: highlight CDNA4 MX format support and any new MFMA opcodes; call out cdna_relevance if it matters.

5. "Show me how the AMD Triton backend implements FlashAttention."

python3 scripts/query.py "flash attention" --language triton-amd
python3 scripts/get_page.py kernel-flash-attention-triton-amd --follow-sources
python3 scripts/grep_wiki.py "matrix_instr_nonkdim" --only wiki

Synthesis: walk through the tile shape choice, K-packing, and pipeline depth; cite source PRs from the triton-amd candidate ledger.

6. "What does s_waitcnt vmcnt(0) cost me?"

python3 scripts/grep_wiki.py "s_waitcnt vmcnt"
python3 scripts/get_page.py technique-s-waitcnt-tuning

Synthesis: explain the counter and the fence semantics; show the pipelining technique that minimizes how often you wait on vmcnt(0).

7. "How do I profile a HIP kernel?"

python3 scripts/query.py --tag rocprof-compute
python3 scripts/get_page.py technique-rocprof-compute

Synthesis: outline rocprof-compute (omniperf) command flow, mention rocprofv3 for instruction-level traces.

8. "I'm spilling registers — what now?"

python3 scripts/get_page.py pattern-register-pressure
python3 scripts/get_page.py technique-register-budgeting
python3 scripts/get_page.py technique-occupancy-tuning

Synthesis: diagnose with --save-temps ISA dump or rocprof-compute VGPR-pressure section; mitigations from the technique pages.

9. "List all merged PRs in the AMD Triton backend touching MFMA scheduling."

python3 scripts/query.py --repo triton-amd --tag mfma-scheduling

10. "I want to port a CUDA tcgen05 kernel to MI350."

python3 scripts/get_page.py migration-wgmma-to-mfma
python3 scripts/get_page.py migration-tma-to-buffer-loads
python3 scripts/get_page.py migration-cuda-to-hip

Synthesis: TMA has no direct equivalent — use buffer_load_dwordx4 + async direct-to-LDS; WGMMA → MFMA tile mapping table from the migration page.