# Gpu Opt Tier5 Nvidia Arch

> 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.

- Skill: `zjli2013/gpu-opt-tier5-nvidia-arch` (Agent Skill)
- Install (CLI): `npx skillmds@latest add zjli2013/gpu-opt-tier5-nvidia-arch`
- Raw SKILL.md: https://api.skillmd.com/api/skills/zjli2013/gpu-opt-tier5-nvidia-arch/raw
- Safety review: pending
- Works with: Claude Code, Claude.ai, OpenAI Codex
- Category: Coding & Dev Tools
- Author: ZJLi2013 (https://skillmd.com/u/zjli2013)
- Updated: 2026-09-22
- Page: https://skillmd.com/skills/zjli2013/gpu-opt-tier5-nvidia-arch

---


# 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)

1. Fix **Tier 1–2** (tiling + memory) before Hopper-only features.
2. Use **NCU** to separate **DRAM vs L1TEX vs TC** bottlenecks.
3. 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

1. Detect **CC** (`nvidia-smi`, NCU header) and load **SKU-specific peak** FLOPS/BW for roofline.
2. Choose instruction family: **WGMMA+TMA (9.0)** vs **WMMA/cp.async (8.x)** vs legacy paths.
3. Record **SMEM usage** per kernel; on Hopper, include **cluster** constraints if used.
4. Cross-link **GEMM/flash** skills for op-specific templates; keep this file for **capability gating** only.
5. Treat numbers like **16384 cores** and **1008 GB/s** as **marketing/spec sheet** anchors—always **measure** on the user's card.