MoE-Kernel-Agent

Optimize fused cross-entropy loss kernels in Triton for NVIDIA and AMD GPUs. Covers fused log-softmax + NLL, online log-sum-exp, and large vocabulary handling. Use when writing or optimizing cross-entropy, label-smoothed CE, or similar loss kernels.

Optimize FlashAttention-style fused attention kernels in Triton for NVIDIA and AMD GPUs. Covers online softmax, tiled QK/AV GEMM, causal masking, and memory-efficient attention. Use when writing or optimizing self-attention, cross-attention, or any QKV attention kernel.

Optimize dense matrix multiplication (GEMM) kernels in Triton for NVIDIA and AMD GPUs. Covers tiled blocking, L2 cache grouping, tensor core utilization, and dual-platform autotune. Use when writing or optimizing matmul, linear layers, or any GEMM-based kernel.

Optimize RMS Normalization kernels in Triton for NVIDIA and AMD GPUs. Covers row-parallel reduction, vectorized loads, warp shuffle patterns, and rsqrt usage. Use when writing or optimizing RMSNorm, LayerNorm, or similar normalization kernels.

Optimize Rotary Position Embedding (RoPE) kernels in Triton for NVIDIA and AMD GPUs. Covers interleaved/sequential layouts, sincos precomputation, and multi-row processing. Use when writing or optimizing RoPE, position embeddings, or rotary transformations.

Optimize fused softmax kernels in Triton for NVIDIA and AMD GPUs. Covers online max/sum reduction, multi-row processing, numerical stability, and warp-level patterns. Use when writing or optimizing softmax, log-softmax, or similar row-wise normalization kernels.

Chooses block/tile dimensions, warp counts, pipeline stages, and L2-aware block grouping for GPU kernels (Triton/CUDA/HIP). Use when autotuning matmul-like kernels, tuning occupancy, or aligning tile geometry with M/N/K and NVIDIA vs AMD thread models.

Optimizes global, shared/LDS, and cache behavior for GPU kernels: coalescing, pipelining, bank conflicts, vector loads, and register vs shared tradeoffs. Use when NCU/rocprof shows low memory throughput, poor coalescing, or L2/LDS pressure.

Tightens inner-loop math, mixed precision, fused epilogues, and fast math intrinsics for GPU kernels. Use when kernels are compute-bound, epilogue-bound, or dominated by elementwise ops around matmul-like cores.

Applies Split-K, persistent kernels, Stream-K/Stream-K++, warp specialization, and multi-GPU fused schedules. Use when simple tiling leaves tails, load imbalance, or excessive kernel launch overhead on large GPUs.

AMD ROCm/CDNA3 tuning: MI300X topology, LDS/L2, wavefront-64 and Triton num_warps, MFMA, HIP/Triton limitations, rocprof counters, and coalescing. Use when optimizing for gfx942/MI300 or comparing against NVIDIA H100 baselines.

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.

Unified kernel benchmarking protocol producing JSON results with latency, TFLOPS, GBps, and comparison against PyTorch baselines. Covers warmup, timing, and cross-platform reporting. Use when benchmarking kernels, measuring performance, or comparing implementations.

Classifies GPU kernel bottlenecks (memory, compute, latency) from profiling metrics and applies GEAK-style workload guidance: what to prefer, consider, or deprioritize. Use after NCU/rocprof or when choosing between fusion, tiling, and scheduling changes without a clear winner.

Orchestrates continuous kernel optimization by chaining profiling, bottleneck diagnosis, tier-based optimization, verification, and benchmarking into an iterative loop. This is the "main program" that drives the other skills. Use when you want to iteratively optimize a kernel until it meets a performance target, similar to AutoKernel.

Profile GPU kernels using NCU (NVIDIA) or rocprof (AMD) to collect performance metrics. Produces structured metrics.json with throughput, bandwidth, occupancy, and bottleneck classification. Use when profiling kernels, collecting NCU/rocprof data, or diagnosing performance issues.

5-stage kernel correctness verification protocol for Triton and CUDA kernels. Covers numerical correctness, dtype sensitivity, edge cases, determinism, and stress testing. Use when verifying kernel correctness, running validation, or setting up test harnesses.