Add Ampere (sm_80/sm_86) extension

AMPERE-EXTENSION.md

@@ -0,0 +1,36 @@
+# Ampere Extension (sm_80 / sm_86)
+
+The knowledge base is Blackwell/Hopper-first. This extension makes the same pages actionable on Ampere hardware: A100 (sm_80) and GA10x consumer/workstation cards (sm_86: RTX 3090/3080, RTX A6000, A40). It was written against a 4x RTX 3090 rig, and numbers marked "measured" come from that machine.
+
+## Design
+
+All changes are **additive** (new files + vocabulary entries) so upstream merges stay trivial:
+
+- `data/tags.yaml` — added architectures `sm80`, `sm86`; hardware features `cp-async`, `mma-sync`, `l2-persistence`.
+- `data/aliases.yaml` — added alias groups for `sm80`, `sm86`, `cp-async`, `mma-sync`, `l2-persistence`.
+- `sources/docs/` — 4 new source pages: `doc-ampere-tuning-guide`, `doc-ga102-whitepaper`, `doc-ptx-isa-ampere`, `doc-cutlass-ampere`.
+- `wiki/hardware/` — 3 new pages: `hw-cp-async`, `hw-mma-sync-ampere`, `hw-ampere-memory-model`.
+- `wiki/migration/` — 1 new page: `migration-hopper-to-ampere` (the entry point: instruction replacement table, capacity re-planning, scheduling paradigm, ncu checklist).
+- `references/primer.md` — "Ampere Extension" section + alias cheat-sheet rows.
+- `SKILL.md` — trigger description now includes Ampere/SM86/RTX 3090 and backport questions.
+
+The upstream validator rules are untouched and still pass: Ampere pages don't trip the Blackwell-first rule (it only constrains Hopper-only pages), and every new page carries the full required frontmatter for its type.
+
+## Query examples
+
+```bash
+python3 scripts/query.py --architecture sm86 --compact          # aliases: "RTX 3090", GA102 …
+python3 scripts/query.py --tag cp-async --type hardware
+python3 scripts/get_page.py migration-hopper-to-ampere
+python3 scripts/grep_wiki.py "wait_group" --only wiki
+```
+
+## Verification status
+
+Snippets on the new pages are compiled against `nvcc -arch=sm_86` on a real 4x RTX 3090 machine (see per-page `reproducibility`). Facts sourced from: NVIDIA Ampere tuning guide (CC 8.0/8.6 table), GA102 whitepaper v2.1 (tensor throughput incl. the FP32-accumulate half-rate), PTX ISA (instruction availability), CUTLASS (SM80 mainloop idiom).
+
+## Candidate follow-ups (not done)
+
+- `source-pr` pages for canonical Ampere kernels in tracked repos (e.g. vLLM Marlin W4A16, exllama kernels) — would upgrade wiki-page confidence to `verified` (needs an `upstream-code` source).
+- `kernel-` case-study pages with benchmarked claims from the 4x3090 rig (FlashAttention-2, Marlin, Triton GEMM autotune points).
+- sm_89 (Ada) column where it differs (FP8 tensor cores present, SMEM 100 KB like sm_86).

README.md

@@ -9,6 +9,8 @@
 
 > **Last repository update: 2026-06-09.** Information after this date is not included in KernelWiki yet.
 
+> **Downstream branch note:** this clone carries an additive **Ampere (sm_80/sm_86) extension** — see [AMPERE-EXTENSION.md](AMPERE-EXTENSION.md). Entry point: `migration-hopper-to-ampere`.
+
 A structured knowledge base of NVIDIA Blackwell (SM100, B200) and Hopper (SM90, H100) GPU kernel optimization, packaged as a Claude Code skill. The repository root **is** the skill directory — clone it directly into `~/.claude/skills/` and it works out of the box.
 
 ## Install as a Claude Code Skill

SKILL.md

@@ -1,6 +1,6 @@
 ---
 name: KernelWiki
-description: Use when the user asks about optimizing NVIDIA Blackwell (SM100, B200) or Hopper (SM90, H100) GPU kernels — tcgen05/TMEM/CLC/NVFP4/2-SM cooperative, warp specialization, FlashAttention-4, DeepGEMM, FlashMLA, MoE, grouped GEMM, CuTe-DSL/PTX/Triton on Blackwell, or wants concrete PR references from CUTLASS/SGLang/vLLM/FlashInfer/PyTorch. Do NOT use for generic CUDA Q&A that is not Blackwell/Hopper-specific, host-side framework integration, or distributed systems (DeepEP/EPLB/DualPipe).
+description: Use when the user asks about optimizing NVIDIA Blackwell (SM100, B200), Hopper (SM90, H100), or Ampere (SM80/SM86, A100, RTX 3090) GPU kernels — tcgen05/TMEM/CLC/NVFP4/2-SM cooperative, warp specialization, FlashAttention-4, DeepGEMM, FlashMLA, MoE, grouped GEMM, CuTe-DSL/PTX/Triton on Blackwell, cp.async/mma.sync/ldmatrix on Ampere, backporting Hopper/Blackwell kernel designs to sm_86 rigs, or wants concrete PR references from CUTLASS/SGLang/vLLM/FlashInfer/PyTorch. Do NOT use for generic CUDA Q&A that is not architecture-specific, host-side framework integration, or distributed systems (DeepEP/EPLB/DualPipe).
 argument-hint: "[natural-language-question] | [--tag foo --type kernel] | [page-id]"
 allowed-tools: "Bash Read Grep Glob"
 ---
@@ -18,6 +18,7 @@ Trigger this skill when the user asks about:
 - **Performance patterns** — low SM utilization, memory-bound, register pressure, compute-bound, tail effects, pipeline stalls
 - **DSLs for Blackwell** — CuTe DSL, CUDA C++ with PTX inline, Triton on Blackwell
 - **Hopper → Blackwell migration** — wgmma → tcgen05, register → TMEM accumulators
+- **Ampere (SM80/SM86) kernel programming and backports** — cp.async pipelines, mma.sync/ldmatrix, GA10x quirks (99 KB SMEM, FP32-accumulate half-rate), porting Hopper/Blackwell designs to A100/RTX 3090 rigs — start at `migration-hopper-to-ampere`
 - **PR references** — "how did vLLM/SGLang/FlashInfer/CUTLASS/PyTorch implement X for SM100?"
 - **Competition solutions** — GPU Mode NVFP4 hackathon, FlashInfer MLSys 2026 submissions
 

data/aliases.yaml

@@ -52,6 +52,30 @@ wgmma:
   - "wgmma.mma_async"
   - WGMMA
 
+cp-async:
+  - "cp.async"
+  - LDGSTS
+  - "async copy"
+  - "cp.async.ca"
+  - "cp.async.cg"
+  - "cuda::memcpy_async"
+  - memcpy_async
+
+mma-sync:
+  - "mma.sync"
+  - HMMA
+  - m16n8k16
+  - m16n8k8
+  - "mma.sync.aligned"
+  - "warp MMA"
+
+l2-persistence:
+  - "L2 persistence"
+  - cudaAccessPolicyWindow
+  - "persisting L2 cache"
+  - "L2 set-aside"
+  - "L2 residency control"
+
 block-scale:
   - "block scaling"
   - "block-scaled"
@@ -77,6 +101,25 @@ sm90:
   - H800
   - "SM90"
 
+sm80:
+  - Ampere
+  - ampere
+  - A100
+  - A800
+  - GA100
+  - "SM80"
+
+sm86:
+  - GA102
+  - "RTX 3090"
+  - "RTX 3090 Ti"
+  - "3090"
+  - "RTX 3080"
+  - "RTX A6000"
+  - A40
+  - A10
+  - "SM86"
+
 # Kernel types
 moe:
   - MoE

data/tags.yaml

@@ -8,6 +8,9 @@ architectures:
   - sm90
   - sm90a
   - sm120
+  # Ampere extension (backport target: A100 / GA10x consumer cards)
+  - sm80
+  - sm86
 
 hardware_features:
   - tcgen05
@@ -27,6 +30,10 @@ hardware_features:
   - mbarrier
   - ldmatrix
   - stmatrix
+  # Ampere extension
+  - cp-async
+  - mma-sync
+  - l2-persistence
 
 techniques:
   - warp-specialization

references/primer.md

@@ -101,6 +101,23 @@ All page IDs below resolve via `get_page.py <id>`. All paths are relative to the
 
 ---
 
+## Ampere Extension (SM80 / SM86 backport target)
+
+Added for running this wiki's knowledge on Ampere hardware (A100, RTX 3090-class rigs). Start at `migration-hopper-to-ampere` when the question is "how do I do X from this wiki on a 3090/A100".
+
+| Topic | Page ID | Path | Notes |
+|---|---|---|---|
+| Hopper/Blackwell → Ampere backport map | `migration-hopper-to-ampere` | `wiki/migration/hopper-to-ampere-backport.md` | Instruction replacement table, capacity re-planning, scheduling paradigm shift, ncu checklist. THE entry point. |
+| cp.async (async global→shared) | `hw-cp-async` | `wiki/hardware/cp-async.md` | What TMA becomes on sm_8x: per-thread 16 B copies, commit-groups, zfill OOB, multistage pipeline snippet. |
+| mma.sync + ldmatrix | `hw-mma-sync-ampere` | `wiki/hardware/mma-sync-ampere.md` | Warp-scope synchronous tensor core; register economics; GA10x FP32-accumulate half-rate quirk; no stmatrix. |
+| Ampere memory model | `hw-ampere-memory-model` | `wiki/hardware/ampere-memory-model.md` | sm_80 vs sm_86 capacity card (99 KB SMEM/block, 1536 thr/SM, 6 MB L2 on 3090), L2 persistence window, roofline position. |
+
+Ampere source docs: `doc-ampere-tuning-guide`, `doc-ga102-whitepaper`, `doc-ptx-isa-ampere`, `doc-cutlass-ampere`.
+
+Techniques that transfer to Ampere as-is: swizzling, double-buffering, pipeline-stages (cp.async flavor), persistent-kernel (atomic queue instead of CLC), tile-scheduling (more important — small L2), vectorized-loads, cache-policy, register-budgeting, kernel-fusion. Techniques that do NOT transfer: anything tagged tcgen05/tmem/clc/2sm-cooperative/tma-multicast — consult the migration table for each.
+
+---
+
 ## Source Repositories (PR coverage)
 
 | Repo | PR pages | Ledger |
@@ -150,6 +167,11 @@ When the user types one of these, match to the canonical term shown:
 | MLA, multi-head latent attention | `mla` |
 | GDN, GatedDeltaNet, gated delta rule | `gated-delta-net` |
 | NSA, native sparse attention | `sparse-attention` |
+| Ampere, A100, A800, GA100, SM80 | `sm80` |
+| GA102, RTX 3090, RTX 3090 Ti, RTX 3080, RTX A6000, A40, A10, SM86 | `sm86` |
+| cp.async, LDGSTS, async copy, cuda::memcpy_async | `cp-async` |
+| mma.sync, HMMA, m16n8k16, warp MMA | `mma-sync` |
+| L2 persistence, cudaAccessPolicyWindow, L2 set-aside | `l2-persistence` |
 
 The `query.py` tool applies these automatically when scoring and when using `--tag`.
 

scripts/query.py

@@ -119,10 +119,13 @@ def score_keyword_match(fm, body, keywords):
     """
     score = 0
     title_text = str(fm.get("title", "")).lower()
+    def _as_list(v):
+        return v if isinstance(v, list) else ([v] if v else [])
     tag_text = " ".join(
         str(v) for k in ("tags", "techniques", "hardware_features", "kernel_types",
-                          "languages", "aliases", "symptoms")
-        for v in (fm.get(k) or [])
+                          "languages", "aliases", "symptoms", "architectures",
+                          "from_arch", "to_arch")
+        for v in _as_list(fm.get(k))
     ).lower()
     body_lower = body.lower()
     for kw in keywords:

sources/docs/cutlass-ampere-support.md

@@ -0,0 +1,51 @@
+---
+id: doc-cutlass-ampere
+title: "CUTLASS Ampere (SM80/SM86) Support: MmaMultistage, CuTe SM80 Atoms, 3xTF32"
+url: https://docs.nvidia.com/cutlass/
+source_category: official-doc
+architectures: [sm80, sm86]
+tags: [cp-async, mma-sync, ldmatrix, swizzling, double-buffering, pipeline-stages]
+retrieved_at: 2026-08-03
+---
+
+# CUTLASS Ampere (SM80/SM86) Support
+
+## Overview
+
+CUTLASS is the reference implementation for high-performance Ampere GEMM. Its SM80 code path is the canonical public example of the full Ampere idiom: multi-stage `cp.async` software pipeline feeding warp-scope `mma.sync` tensor core ops via `ldmatrix`, with swizzled shared-memory layouts. Anyone backporting Hopper/Blackwell kernel structure to sm_86 should read the CUTLASS SM80 mainloop first — it IS the target shape of the port.
+
+## The SM80 mainloop: MmaMultistage
+
+CUTLASS 2.x introduced `cutlass::gemm::threadblock::MmaMultistage` — the Ampere replacement for the Volta/Turing 2-stage `MmaPipelined`:
+
+- **N-stage circular SMEM buffer** (typically `Stages = 3..5`, bounded by SMEM: 99 KB/block on sm_86, 163 KB on sm_80).
+- Global→shared operand movement issued with `cp.async` (`CacheOperation::Global` → `.cg` 16B for main operands, `.ca` for small fragments).
+- Completion via `cp.async.wait_group`/commit-group counting; no mbarriers in the classic mainloop.
+- Warp-level tile compute in `cutlass::gemm::warp::MmaTensorOp`, which lowers to `mma.sync.aligned.m16n8k16` (arch `cutlass::arch::Mma<GemmShape<16,8,16>, 32, ...>`) with `ldmatrix`-based operand fetch (`cutlass::arch::LdMatrix`).
+- Shared-memory layouts use XOR-swizzles (`TensorOpMultiplicandCrosswise` etc.) sized to keep `ldmatrix` bank-conflict-free.
+
+## CuTe (C++) SM80 atoms
+
+CUTLASS 3.x re-expresses the same hardware in CuTe:
+
+- MMA atoms: `SM80_16x8x16_F32F16F16F32_TN`, `SM80_16x8x16_F32BF16BF16F32_TN`, `SM80_16x8x8_F32TF32TF32F32_TN`, INT8 variants — all wrap `mma.sync`.
+- Copy atoms: `SM80_CP_ASYNC_CACHEALWAYS<T>` / `SM80_CP_ASYNC_CACHEGLOBAL<T>` (cp.async .ca/.cg), `SM75_U32x4_LDSM_N` / `SM75_U16x8_LDSM_T` (ldmatrix, non-transposed/transposed).
+- Tutorial `examples/cute/tutorial/sgemm_sm80.cu` walks the full pattern; `cute::cp_async_fence()` / `cute::cp_async_wait<N>()` expose commit-group semantics.
+- CollectiveBuilder path: SM80 collectives use the multistage schedule (no warp specialization — contrast with SM90 `KernelTmaWarpSpecialized*` schedules).
+
+## Ampere-specific examples in the CUTLASS tree
+
+| Example | What it demonstrates |
+|---|---|
+| `14_ampere_tf32_tensorop_gemm` | TF32 tensor-core GEMM on FP32 data |
+| `15_ampere_sparse_tensorop_gemm` | 2:4 structured sparsity (`mma.sp`) |
+| `27_ampere_3xtf32_fast_accurate_tensorop_gemm` | **3xTF32**: emulate FP32 GEMM with 3 TF32 MMAs (error-compensated), ~2x+ over FP32 CUDA cores at near-FP32 accuracy |
+| `sgemm_sm80.cu` (CuTe tutorial) | Minimal readable multistage cp.async + mma.sync mainloop |
+
+Note: the newer Python CuTe DSL examples tree is Hopper/Blackwell-first; check the current `examples/python/CuTeDSL/` tree before assuming SM80 DSL coverage. The mature, performance-proven Ampere path is the C++ one above.
+
+## Practical notes for sm_86 specifically
+
+- CUTLASS kernels compiled for sm_80 run on sm_86, but tile shapes tuned for A100's 164 KB SMEM will fail to launch or spill on GA10x's 99 KB — pick smaller `Stages` or tile sizes (the profiler's sm_86 presets do this).
+- Threadblock shapes that assume 2048 threads/SM occupancy targets should be re-tuned for 1536.
+- FP16 accumulation (`ElementAccumulator = half_t`) doubles tensor throughput on GA10x (see doc-ga102-whitepaper) at the usual numerical risk; CUTLASS supports it natively.

sources/docs/nvidia-ampere-tuning-guide.md

@@ -0,0 +1,76 @@
+---
+id: doc-ampere-tuning-guide
+title: "NVIDIA Ampere GPU Architecture Tuning Guide (CC 8.0 / 8.6)"
+url: https://docs.nvidia.com/cuda/ampere-tuning-guide/
+source_category: official-doc
+architectures: [sm80, sm86]
+tags: [cuda-cpp, cp-async, l2-persistence, mbarrier]
+retrieved_at: 2026-08-03
+---
+
+# NVIDIA Ampere GPU Architecture Tuning Guide (CC 8.0 / 8.6)
+
+## Overview
+
+The official NVIDIA Ampere tuning guide covers both Ampere compute capabilities: **8.0 (GA100: A100/A800)** and **8.6 (GA10x: RTX 3090/3080, RTX A6000, A40, A10)**. The two are NOT the same architecture from a tuning standpoint — occupancy ceilings, shared memory capacity, and FP32 issue rate all differ. This page records the facts that matter when writing or backporting kernels for sm_86.
+
+## CC 8.0 vs CC 8.6 — the table that matters
+
+| Resource | CC 8.0 (A100) | CC 8.6 (GA10x) |
+|---|---|---|
+| Max concurrent warps / SM | 64 (2048 threads) | **48 (1536 threads)** |
+| Max thread blocks / SM | 32 | **16** |
+| Unified L1/SMEM capacity / SM | 192 KB | **128 KB** |
+| Max SMEM carveout / SM | 164 KB | **100 KB** |
+| Max SMEM per thread block | 163 KB | **99 KB** |
+| FP32 ops per cycle per SM | 64 | **128 (2x)** |
+
+Implications:
+
+- **Occupancy math changes on sm_86.** A block size that yields 100% occupancy on A100 (e.g., 2048 threads via 2x1024) caps at 75% of the sm_80 figure on GA10x. Blocks of 256/512 threads with ≤3 blocks resident hit the 1536-thread ceiling exactly.
+- **SMEM-heavy Hopper-style kernels do not fit.** 99 KB per block is the hard ceiling on sm_86 (vs 163 KB on A100, 227 KB on H100). Multi-stage pipelines must shrink stage count or tile size accordingly.
+- **The 2x FP32 pipe** (both datapaths can issue FP32; on GA100 one is FP32-only, the other INT32) means non-tensor-core elementwise/epilogue code is comparatively cheap on GA10x.
+
+## Asynchronous Data Copies (cp.async)
+
+Ampere introduces asynchronous copy from global to shared memory, bypassing the register file and (optionally) L1:
+
+- `cp.async.ca.shared.global` — 4, 8, or 16 bytes, caches in L1.
+- `cp.async.cg.shared.global` — 16 bytes only, caches only in L2 (bypasses L1). Recommended for tensor-core operand staging.
+- Completion managed either via **commit-group semantics** (`cp.async.commit_group` / `cp.async.wait_group N`) or via **asynchronous barriers** (`cuda::barrier`, PTX `mbarrier`, available since sm_80).
+- The C++ surface is `cuda::memcpy_async` + `cuda::pipeline` (libcudax / `<cuda/pipeline>`).
+
+This is the Ampere ancestor of Hopper TMA: per-thread addressing, no bulk tensor descriptors, no multicast, no swizzle-on-the-fly — but the same "load directly to SMEM without burning registers" principle.
+
+## L2 Persistence (Access Policy Window)
+
+CC 8.0+ can set aside a portion of L2 for **persisting** accesses via `cudaAccessPolicyWindow` (per-stream or per-graph-node attribute):
+
+```cpp
+cudaStreamAttrValue attr;
+attr.accessPolicyWindow.base_ptr  = kv_cache_ptr;
+attr.accessPolicyWindow.num_bytes = window_bytes;          // <= cudaLimitMaxL2FetchGranularity window
+attr.accessPolicyWindow.hitRatio  = 0.6f;                   // fraction treated as persisting
+attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
+attr.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
+cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);
+```
+
+`cudaDeviceProp::persistingL2CacheMaxSize` reports the maximum set-aside. Useful for small hot operands re-read across many CTAs (e.g., decode-time KV heads, router weights). Streaming accesses can be marked `evict_first` via PTX cache policies to avoid polluting the persisting region.
+
+## Other sm_80+ features recorded by the guide
+
+- **Asynchronous barriers** (`cuda::barrier<cuda::thread_scope_block>`): arrive/wait split, enables producer/consumer overlap without `__syncthreads()` full-block convergence.
+- **Warp reduce instructions** (`redux.sync`) for int operands.
+- **Improved L2 residency management + larger L2** vs Volta/Turing.
+- **Third-generation tensor cores**: BF16 and TF32 inputs; `mma.sync` shapes up to m16n8k16 (FP16/BF16) and m16n8k8 (TF32).
+
+## What Ampere does NOT have (forward-looking)
+
+Recorded here because this KB is used for Hopper/Blackwell backports; all of the following first appear in sm_90 or sm_100:
+
+- No TMA (`cp.async.bulk.tensor`), no multicast loads.
+- No thread block clusters, no distributed shared memory (DSMEM).
+- No `wgmma` / `tcgen05` asynchronous tensor core instructions — only synchronous warp-scope `mma.sync`.
+- No TMEM, no `setmaxnreg` register reallocation, no `stmatrix` (sm_90+), no PDL/GDC kernel-launch overlap.
+- No FP8/FP6/FP4 tensor core datatypes (INT8/INT4 via `mma.sync` are available).