| name | gpu-opt-tier5-nvidia-arch |
| description | NVIDIA-specific caps and tuning: Hopper SM90 (WGMMA, TMA, cluster SMEM), Ampere SM80 (cp.async, TF32, 2:4 sparsity), Ada SM89 consumer (4090-class). Use when targeting a known compute capability or interpreting NCU metrics for a specific chip. |
Tier 5 (NVIDIA): Architecture-Specific Optimization
AutoKernel Tier 5 maps generic techniques to SKU capabilities. This file is the NVIDIA branch.
Compute capability quick reference
| Arch | CC | Examples | Notes |
|---|
| Ampere | 8.0 | A100 | TF32 TC, cp.async, 2:4 structured sparsity (sparse TC) |
| Ada Lovelace | 8.9 | RTX 4090 | 16384 "CUDA cores" (marketing aggregate), ~1008 GB/s ref bandwidth class |
| Hopper | 9.0 | H100 | WGMMA, TMA, cluster SMEM, up to ~228 KB SMEM/SM (config-dependent) |
Exact SMEM/L1 splits and clocks vary by SKU; always verify in NCU "Launch Statistics" and occupancy calculator.
H100 / SM90 (Hopper)
- WGMMA: prefer tensor-core-native tile multiples; pair with TMA for async bulk loads.
- TMA: reduces load instruction overhead; needs correct tensor descriptors and alignment.
- Cluster: distributed shared memory across blocks in a cluster for multi-block cooperation.
- Autotune tips: sweep cluster size, TMA vs cp.async paths, WGMMA K-dimension tiling; watch register explosion on wide tiles.
A100 / SM80 (Ampere)
- cp.async: pipeline global→shared; tune
num_stages with occupancy.
- TF32: training default for many frameworks; know accuracy implications vs FP16/BF16 TC.
- 2:4 sparsity: only wins when weights are structured-sparse and paths are enabled end-to-end.
4090 / Ada (SM89)
- High FP16/BF16 throughput for consumer bandwidth; often memory-bound on large models despite high FLOPS.
- Prefer vectorized loads, fusion, and L2-friendly traversal; autotune similarly to Ampere but respect smaller L2 vs datacenter parts in some metrics.
Autotune workflow (NVIDIA)
- Fix Tier 1–2 (tiling + memory) before Hopper-only features.
- Use NCU to separate DRAM vs L1TEX vs TC bottlenecks.
- For Hopper, add WGMMA/TMA knobs only when baseline kernel already uses compatible abstractions (e.g. CUTLASS 3.x, cuBLASLt, or hand-written PTX).
NCU metrics cheat sheet (orientation)
| Symptom in NCU | Likely lever |
|---|
Low sm__throughput + low dram__throughput | Occupancy / launch / tail (Tier 1/4) |
| High DRAM, low TC active | Memory path / vectorization (Tier 2) |
| High TC, reg spills | Tile size / WGMMA shape / fewer live tensors |
Blackwell / future SKUs
- Treat unseen CC values as unknown: fall back to prior-gen patterns, then re-measure; do not assume Hopper knobs map 1:1 without vendor release notes.
Agent Instructions
- Detect CC (
nvidia-smi, NCU header) and load SKU-specific peak FLOPS/BW for roofline.
- Choose instruction family: WGMMA+TMA (9.0) vs WMMA/cp.async (8.x) vs legacy paths.
- Record SMEM usage per kernel; on Hopper, include cluster constraints if used.
- Cross-link GEMM/flash skills for op-specific templates; keep this file for capability gating only.
- Treat numbers like 16384 cores and 1008 GB/s as marketing/spec sheet anchors—always measure on the user's card.
