Optimize Mixture-of-Experts (MoE) CUDA kernels on NVIDIA B200 (Blackwell, SM100). Use when writing, migrating, debugging, or optimizing FP8/FP16 MoE operators that involve grouped GEMM, per-expert routing, SwiGLU activation, gather/scatter, or Blackwell-specific features (TMA, tcgen05, TMEM). Covers PyTorch→CUDA migration, an optimization ladder ordered by measured ROI, FP8 correctness failure-mode catalog, cuBLAS/CUTLASS/tcgen05 backend selection, and agent-team patterns for multi-round optimization.
Multi-agent parallel CUDA kernel optimization. Use this skill when the user wants to run two (or more) optimization agents in parallel — typically on separate GPUs — to explore different optimization directions concurrently, or when the single-agent loop has plateaued and the user asks for parallel exploration. Also use when the user says things like "spawn agents on GPUs 6 and 7", "have two workers try different directions", or "Agent 1 focus on X, Agent 2 focus on Y". Extends `cuda-kernel-autodev` (which defines the single-agent loop) with orchestration protocol: isolated working directories, per-agent GPU pinning, sync/merge cadence, direction assignment, and main-agent role. Typically activated after you've already been running `cuda-kernel-autodev` for a while and the curve has flattened.
CUDA kernel development and optimization for NVIDIA B200 (Blackwell, compute capability 10.0) in Claude Code. Use when writing, reviewing, debugging, profiling, or optimizing CUDA kernels for B200/Blackwell. Covers kernel structure, memory hierarchy, shared memory and cluster constraints, PTX/TensorCore inspection, nsys/ncu profiling, compute-sanitizer debugging, and use of local PTX/Runtime/Driver reference docs. Prefer this skill for CUDA, PTX, Tensor Core, WGMMA/TMA/tcgen05, shared memory, occupancy, coalescing, register pressure, launch configuration, and Blackwell-specific tuning.
Autonomous end-to-end CUDA kernel development and optimization workflow. Use this skill whenever the user wants to write, optimize, or speed up a CUDA kernel/operator against a reference — whether they say "optimize this kernel", "improve GPU throughput", "beat PyTorch baseline", "speed up my CUDA op", "auto-tune this kernel", are iterating on kernel.cu / kernel.cpp / kernel.py in a benchmark framework (FlashInfer, KernelBench, MLSys contests), or are asking how to structure kernel optimization experiments. Covers the full 3-phase workflow: setup/baseline → hypothesize-implement-eval-commit loop with keep/revert discipline → submission. Explicitly pairs with `cuda-roofline-strategy` for choosing what to try next, `cuda-kernel-techniques` for the catalog of proven techniques, `cuda-agent-team` for multi-agent parallel mode, and `ncu-cuda-profiling` for collecting NCU data.
Catalog of proven CUDA kernel optimization techniques, broken by sub-topic (memory access, data placement, parallelism, compute, control flow, occupancy, numerical, anti-patterns). Use this skill whenever you need to look up *how* a specific technique works or *when* it applies — e.g., "how does Split-K FlashDecoding work", "when should I use cp.async vs TMA", "how do I implement online softmax", "what is register-centric data path", "what bank-conflict padding should I use", "when does warp specialization help", or before committing to a technique suggested by `cuda-roofline-strategy`. The main SKILL.md is an index; each sub-topic reference has the implementation details, when-it-helps, when-it-hurts, and code sketches. Each technique includes hardware-general guidance with B200 (SM100) and older-arch examples where they differ.
Roofline-driven strategy selection for CUDA kernel optimization. Use this skill during the "what do I try next" moment in a kernel optimization loop — when the user has a profile (NCU report or similar) and needs to decide which class of technique to try, or when experiments are plateauing and the user asks "what now?", or when diagnosing whether a kernel is compute-bound / bandwidth-bound / occupancy-limited / latency-limited. Maps (roofline position, iteration phase) → technique tier so the next experiment has a principled justification rather than being a guess. Complements `cuda-kernel-autodev` (which drives the loop) and `cuda-kernel-techniques` (which details each technique).
Automated NCU (Nsight Compute) profiling workflow with full metrics collection and persistent storage