docs(bench): add H100 SXM5 benchmark results alongside A100 baseline

README.md

@@ -94,6 +94,24 @@ Model weights consume 56.9 GB — only 23 GB headroom remaining for training com
   <img src="docs/assets/3. moe .png" alt="Real Model MoE Benchmark">
 </p>
 
+### 4. H100 SXM5 Independent Validation
+
+The numbers above were measured on A100. We re-ran a comparable set of checks on **NVIDIA H100
+80GB HBM3 (SM90)** with SM90 kernels built in (`KERNEL_ALIGN_FORCE_SM90=1`), including the
+**Qwen3-30B-A3B** memory claim against real downloaded weights (not synthetic tensors). Full
+methodology, exact commands, and raw CSVs are in the
+[Hardware Benchmark Dashboard](docs/benchmarking/hardware-dashboard.md); headline results:
+
+- **Qwen3-30B-A3B weight footprint, confirmed on real weights**: 56.87 GB / 79.11 GB total → 22.24 GB headroom, matching the A100 claim above. RL-Kernel's fused logprob path keeps ~0 GB extra VRAM through 24,576 tokens where the naive path OOMs past 12,288.
+- **Fused `logp` kernel (generic CUDA, vocab=128256, seq=512, batch=32)**: 23.29 ms / 19.57 GB (native) → 6.93 ms / 7.83 GB (fused), a 3.4x speedup *and* lower memory — no tradeoff here.
+- **FlashInfer sampling (vocab=128256, batch=256)**: 29.30 ms (native) → 1.65 ms (RL-Kernel), ~18x. Raw output: `reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt`.
+- **`linear_logp` SM90 kernel — reported honestly, not cherry-picked**: it beats the Triton path by 1.7–1.9x at ~600–2500x less memory than the naive materializing path, but does **not** uniformly beat naive on raw latency (naive wins forward at 2 of 3 vocab sizes tested, and wins backward across the board — the SM90 kernel trades FLOPs for memory via tile recomputation, same as Triton's approach). Its win is fitting at all in constrained memory, not raw speed everywhere.
+
+See the dashboard doc for the full sweep, exact reproduction commands, and two known limitations
+found while benchmarking (a profiler gap: no `sampling-fused` workload registered; and a bug: the
+experimental, off-by-default `fused_logp_sm90` standalone kernel currently crashes the process on
+a failed `cuTensorMapEncodeTiled` call instead of raising).
+
 # Key Features
 
 - **Zero-Growth Memory Pool**: Uses pre-allocated buffers and micro-chunking to prevent VRAM spikes during advantage calculation.

benchmarks/benchmark_qwen3_moe_real_model.py

@@ -0,0 +1,196 @@
+# SPDX-License-Identifier: Apache-2.0
+# Copyright (c) 2026 RL-Kernel Contributors
+
+"""Real-model logprob memory/latency benchmark against Qwen3-30B-A3B.
+
+Downloads real weights from the HF Hub (default: Qwen/Qwen3-30B-A3B) and runs a real
+forward pass to obtain real hidden states and the model's real lm_head weight -- unlike
+benchmarks/benchmark_linear_logp.py, which uses synthetic random tensors. Compares native
+(materializing log_softmax + gather) against the dispatched rl_engine linear_logp kernel
+as token count scales into the model's remaining VRAM headroom.
+
+Usage:
+    python benchmarks/benchmark_qwen3_moe_real_model.py
+    python benchmarks/benchmark_qwen3_moe_real_model.py --model Qwen/Qwen3-30B-A3B \
+        --n-configs 2048,4096,8192,12288,16384
+"""
+
+import argparse
+import time
+
+import torch
+from tabulate import tabulate
+from transformers import AutoModelForCausalLM, AutoTokenizer
+
+from rl_engine.kernels.registry import kernel_registry
+from rl_engine.utils.logger import logger
+
+DEFAULT_N_CONFIGS = [2048, 4096, 8192, 12288, 16384, 20480, 24576]
+DEFAULT_PROMPT = (
+    "Explain the tradeoffs between memory usage and compute latency when "
+    "computing log probabilities over a large vocabulary in reinforcement "
+    "learning post-training for large language models."
+)
+
+
+def gb(x):
+    return x / (1024**3)
+
+
+def native_logprob(hidden, weight, target):
+    """Standard full log_softmax + gather -- O(N*V) extra memory."""
+    logits = torch.nn.functional.linear(hidden, weight)
+    log_probs = torch.log_softmax(logits.float(), dim=-1)
+    return torch.gather(log_probs, dim=-1, index=target.unsqueeze(-1)).squeeze(-1)
+
+
+def measure(fn, warmup=2, iters=5):
+    torch.cuda.synchronize()
+    torch.cuda.empty_cache()
+    base = torch.cuda.memory_allocated()
+    for _ in range(warmup):
+        fn()
+    torch.cuda.synchronize()
+    torch.cuda.reset_peak_memory_stats()
+    t0 = time.perf_counter()
+    for _ in range(iters):
+        fn()
+    torch.cuda.synchronize()
+    t1 = time.perf_counter()
+    peak = torch.cuda.max_memory_allocated()
+    return gb(peak - base), (t1 - t0) / iters * 1000
+
+
+def run_benchmark(args):
+    if not torch.cuda.is_available():
+        raise RuntimeError("this benchmark requires a CUDA GPU.")
+    dtype = getattr(torch, args.dtype)
+
+    torch.cuda.reset_peak_memory_stats()
+    baseline_before_load = torch.cuda.memory_allocated()
+
+    t0 = time.time()
+    logger.info(f"Loading {args.model} in {dtype} ...")
+    tokenizer = AutoTokenizer.from_pretrained(args.model)
+    model = AutoModelForCausalLM.from_pretrained(args.model, torch_dtype=dtype, device_map={"": 0})
+    model.eval()
+    torch.cuda.synchronize()
+    load_s = time.time() - t0
+
+    weight_gb = gb(torch.cuda.memory_allocated() - baseline_before_load)
+    total_gb = gb(torch.cuda.get_device_properties(0).total_memory)
+    headroom_gb = total_gb - weight_gb
+    logger.info(
+        f"Loaded in {load_s:.1f}s. Weight VRAM: {weight_gb:.2f} GB / {total_gb:.2f} GB "
+        f"total -> headroom {headroom_gb:.2f} GB"
+    )
+
+    hidden_size = model.config.hidden_size
+    vocab_size = model.config.vocab_size
+    lm_head_weight = model.lm_head.weight.detach()
+    logger.info(
+        f"hidden_size={hidden_size} vocab_size={vocab_size} "
+        f"lm_head shape={tuple(lm_head_weight.shape)}"
+    )
+
+    enc = tokenizer(args.prompt, return_tensors="pt").to("cuda")
+
+    logger.info("Running a real forward pass to obtain real hidden states...")
+    with torch.no_grad():
+        out = model.model(**enc, output_hidden_states=False, use_cache=False)
+    real_hidden = out.last_hidden_state.detach()  # [1, seq, H]
+    real_hidden = real_hidden.reshape(-1, hidden_size).to(dtype).contiguous()
+    logger.info(f"Real hidden states shape: {tuple(real_hidden.shape)}")
+
+    linear_logp_op = kernel_registry.get_op("linear_logp")
+    logger.info(f"Dispatched linear_logp backend: {type(linear_logp_op).__name__}")
+
+    def make_batch(n):
+        reps = (n + real_hidden.shape[0] - 1) // real_hidden.shape[0]
+        hidden = real_hidden.repeat(reps, 1)[:n].clone()
+        target = torch.randint(0, vocab_size, (n,), device="cuda")
+        return hidden, target
+
+    rows = []
+    for n in args.n_configs:
+        try:
+            hidden, target = make_batch(n)
+        except torch.cuda.OutOfMemoryError:
+            rows.append([n, "OOM (input alloc)", "OOM (input alloc)", "N/A", "N/A", "N/A"])
+            torch.cuda.empty_cache()
+            continue
+
+        try:
+            native_extra, native_ms = measure(
+                lambda: native_logprob(hidden, lm_head_weight, target)
+            )
+            native_str, native_ms_str = f"{native_extra:.2f} GB", f"{native_ms:.2f} ms"
+        except torch.cuda.OutOfMemoryError:
+            native_str, native_ms_str = "OOM", "N/A"
+        torch.cuda.empty_cache()
+
+        try:
+            kernel_extra, kernel_ms = measure(
+                lambda: linear_logp_op(hidden, lm_head_weight, target)
+            )
+            kernel_str, kernel_ms_str = f"{kernel_extra:.2f} GB", f"{kernel_ms:.2f} ms"
+        except torch.cuda.OutOfMemoryError:
+            kernel_str, kernel_ms_str = "OOM", "N/A"
+        torch.cuda.empty_cache()
+
+        rows.append(
+            [
+                n,
+                native_str,
+                kernel_str,
+                native_ms_str,
+                kernel_ms_str,
+                f"current alloc: {gb(torch.cuda.memory_allocated()):.1f} GB",
+            ]
+        )
+        torch.cuda.empty_cache()
+
+    print("\n" + "=" * 100)
+    print(f"{args.model} REAL MODEL LOGPROB BENCHMARK on {torch.cuda.get_device_name(0)}")
+    print(
+        f"Weight VRAM: {weight_gb:.2f} GB | Total: {total_gb:.2f} GB | "
+        f"Headroom: {headroom_gb:.2f} GB"
+    )
+    print(f"linear_logp backend: {type(linear_logp_op).__name__}")
+    print("=" * 100)
+    print(
+        tabulate(
+            rows,
+            headers=[
+                "N tokens",
+                "Native extra VRAM",
+                "RL-Kernel extra VRAM",
+                "Native ms",
+                "RL-Kernel ms",
+                "note",
+            ],
+            tablefmt="github",
+        )
+    )
+
+
+def parse_args():
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--model", type=str, default="Qwen/Qwen3-30B-A3B")
+    parser.add_argument("--dtype", type=str, default="bfloat16", choices=["bfloat16", "float16"])
+    parser.add_argument("--prompt", type=str, default=DEFAULT_PROMPT)
+    parser.add_argument(
+        "--n-configs",
+        type=str,
+        default=None,
+        help="Comma-separated token counts, e.g. '2048,4096,8192'.",
+    )
+    args = parser.parse_args()
+    args.n_configs = (
+        [int(x) for x in args.n_configs.split(",")] if args.n_configs else DEFAULT_N_CONFIGS
+    )
+    return args
+
+
+if __name__ == "__main__":
+    run_benchmark(parse_args())

docs/benchmarking/hardware-dashboard.md

@@ -30,23 +30,99 @@ A fallback backend result must not be presented as a fused-kernel result.
 
 | Environment ID | GPU | Architecture | Driver | Runtime | PyTorch | RL-Kernel Commit | Date |
 | --- | --- | --- | --- | --- | --- | --- | --- |
-| `h100-template` | H100 SXM5 | Hopper | pending | CUDA pending | pending | pending | pending |
+| `h100-sxm5` | H100 SXM5 80GB HBM3 | Hopper (SM90) | 535.309.01 | CUDA 12.4 | 2.6.0+cu124 | `6df029a` | 2026-07-19 |
 | `mi300-template` | MI300X | CDNA 3 | pending | ROCm pending | pending | pending | pending |
 
 ## Selected LogP Results
 
 | Environment | Backend | Batch | Sequence Length | Vocabulary | Dtype | Latency (ms) | Tokens/s | Peak VRAM (GB) | Status | Command |
 | --- | --- | ---: | ---: | ---: | --- | ---: | ---: | ---: | --- | --- |
-| `h100-template` | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending |
+| `h100-sxm5` | `logp_native` (log_softmax + gather) | 32 | 512 | 128256 | float16 | 23.29 | 703,528 | 19.57 | pass | `python scripts/run_profile_suite.py --device cuda --dtype float16 --batch-sizes 32 --seq-lens 512 --vocab-sizes 128256 --workloads logp-native,logp-fused` |
+| `h100-sxm5` | `logp_fused` (generic CUDA `fused_logp`) | 32 | 512 | 128256 | float16 | 6.93 | 2,362,773 | 7.83 | pass | same as above |
+| `h100-sxm5` | `logp_native` | 16 | 512 | 128256 | float16 | 11.37 | 720,555 | 9.79 | pass | `--batch-sizes 16 --seq-lens 512 --vocab-sizes 128256` |
+| `h100-sxm5` | `logp_fused` | 16 | 512 | 128256 | float16 | 3.32 | 2,468,160 | 3.91 | pass | same as above |
 | `mi300-template` | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending |
 
+Full 24-row sweep (batch ∈ {8,16,32} × seq_len ∈ {128,512} × vocab ∈ {4096,128256}):
+`reports/perf_report_NVIDIA_H100_80GB_HBM3.csv` (this PR). `logp_fused` selected the generic
+CUDA kernel (`FusedLogpGenericOp`), not the experimental SM90 TMA kernel — see the linear-logp
+note below for why SM90-specific dispatch is more nuanced.
+
 ## Sampling Results
 
 | Environment | Backend | Batch | Vocabulary | Top-k | Top-p | Temperature | Latency (ms) | Tokens/s | Peak VRAM (GB) | Status | Command |
 | --- | --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | --- | --- |
-| `h100-template` | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending |
+| `h100-sxm5` | native (topk→topp→softmax→multinomial) | 64 | 128256 | 50 | 0.9 | 1.0 | 8.73 | n/a | n/a | pass | `python benchmarks/benchmark_sampling.py --g-sizes 32,64,128,256 --vocab-size 128256 --top-k 50 --top-p 0.9` |
+| `h100-sxm5` | FlashInfer (`RL_Sampler`) | 64 | 128256 | 50 | 0.9 | 1.0 | 0.84 | n/a | n/a | pass | same as above |
+| `h100-sxm5` | native | 256 | 128256 | 50 | 0.9 | 1.0 | 29.30 | n/a | n/a | pass | same as above |
+| `h100-sxm5` | FlashInfer (`RL_Sampler`) | 256 | 128256 | 50 | 0.9 | 1.0 | 1.65 | n/a | n/a | pass | same as above |
 | `mi300-template` | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending |
 
+Raw artifact backing the rows above: `reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt`
+(this PR) — full console output of the exact command listed, including all four batch sizes
+(32/64/128/256), not just the two rows excerpted here.
+
+Limitation: `benchmarks/profiler.py`'s `WORKLOAD_REGISTRY` only registers `sampling-native`
+(no `sampling-fused` workload), so the FlashInfer rows above come from
+`benchmarks/benchmark_sampling.py` directly rather than `run_profile_suite.py`, and it does not
+report tokens/s or peak VRAM. Wiring a `sampling-fused` workload into the profiler would close
+this gap (tracked as a follow-up, not done in this PR).
+
+## Linear-LogP: SM90 warp specialization vs generic paths (H100, bf16)
+
+`rl_engine/kernels/ops/cuda/loss/linear_logp.py` has **no generic (SM86) CUDA kernel** — on
+Ampere, `linear_logp` falls back to Triton or the native materializing path. The SM90 TMA/WGMMA
+kernel is therefore the only path that is both memory-efficient *and* GPU-accelerated for this
+op; the question "does SM90 warp specialization help vs SM86" reduces to "does the SM90 kernel
+help vs SM86's only options (Triton / native)".
+
+Command: `python benchmarks/benchmark_linear_logp.py` (built with `KERNEL_ALIGN_FORCE_SM90=1`).
+Raw artifact: `reports/benchmark_linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txt` (this PR).
+
+| shape (N×H×V) | native fwd ms | triton fwd ms | sm90 fwd ms | native fwd MB | triton fwd MB | sm90 fwd MB | sm90 vs triton | sm90 vs native |
+| --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: |
+| 4096×2048×32768 | 1.80 | 6.29 | 3.65 | 1280 | 0 | 2 | 1.72x faster | 0.49x (slower) |
+| 4096×2048×50257 | 10.27 | 9.88 | 5.40 | 1965 | 0 | 2 | 1.83x faster | 1.90x faster |
+| 4096×2048×131072 | 7.20 | 24.76 | 14.11 | 5120 | 0 | 2 | 1.75x faster | 0.51x (slower) |
+
+Status: pass. Limitation (report honestly, not cherry-picked): SM90 **consistently beats Triton**
+(1.7–1.9x) at ~600–2500x less memory than the native materializing path, but does **not**
+uniformly beat native on raw latency — native is faster in forward at two of three vocab sizes
+tested, and backward is slower on SM90 across the board (tile-recompute trades FLOPs for
+memory). This also reproduces on the real Qwen3-30B-A3B `lm_head` (vocab 151,936) below.
+
+Separately: the standalone (non-`linear_logp`) `fused_logp_sm90` kernel, gated behind
+`RL_KERNEL_ENABLE_EXPERIMENTAL_SM90_LOGP=1`, currently aborts the process
+(`cuTensorMapEncodeTiled` fails, and `csrc/utils/tma_utils.cuh:49` calls `exit(EXIT_FAILURE)`
+instead of raising) at vocab=32768/128256 — filed as a follow-up, out of scope for this PR since
+the flag defaults off and isn't part of any existing benchmark claim.
+
+## Real-model validation: Qwen3-30B-A3B (H100, bf16)
+
+Real weights (`Qwen/Qwen3-30B-A3B`, 56.9 GB, downloaded from the HF Hub) and a real forward pass
+(not synthetic tensors) were used for the `lm_head` weight and hidden-state distribution.
+
+Command: `python benchmarks/benchmark_qwen3_moe_real_model.py` (script added in this PR).
+Raw artifact: `reports/benchmark_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txt` (this PR).
+
+| Metric | Value |
+| --- | --- |
+| Weight VRAM | 56.87 GB |
+| Total H100 VRAM | 79.11 GB |
+| Headroom | 22.24 GB |
+
+| N tokens | native extra VRAM | RL-Kernel (SM90) extra VRAM | native ms | RL-Kernel ms | status |
+| ---: | ---: | ---: | ---: | ---: | --- |
+| 12,288 | 17.39 GB | 0.00 GB | 25.08 | 48.03 | pass |
+| 16,384 | OOM | 0.00 GB | n/a | 65.90 | native: oom |
+| 24,576 | OOM | 0.00 GB | n/a | 105.50 | native: oom |
+
+Status: pass (memory claim), with the same latency caveat as above — RL-Kernel is ~1.9–2x
+slower per call than native at shapes where native still fits; its advantage here is fitting at
+all within the 22.24 GB headroom, not raw speed. Hidden states came from one real 29-token
+prompt, replicated to build larger N (memory/latency at this stage don't depend on token
+content, only shape/dtype) — flagging so this isn't read as N independently-sampled completions.
+
 ## Reproduction
 
 Run the profiler from the repository root:

reports/benchmark_linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txt

@@ -0,0 +1,15 @@
+# RL-Kernel linear_logp benchmark — raw output (native vs Triton vs SM90 TMA/WGMMA)
+# GPU: NVIDIA H100 80GB HBM3 (SXM5) | Driver 535.309.01 | CUDA 12.4 | PyTorch 2.6.0+cu124
+# Extension built with KERNEL_ALIGN_FORCE_SM90=1
+# Command: python benchmarks/benchmark_linear_logp.py
+# Date: 2026-07-19T08:45:44Z
+
+INFO 07-19 04:45:46 [RL-Kernel]: RL-Engine initialized with NVIDIA CUDA backend (Version: 12.4)
+INFO 07-19 04:45:46 [RL-Kernel]: Successfully linked to precompiled _C.fused_linear_logp_sm90 kernel.
+INFO 07-19 04:45:46 [RL-Kernel]: linear_logp benchmark on cuda (dtype=torch.bfloat16); SM90 TMA+MMA backend enabled
+INFO 07-19 04:45:50 [RL-Kernel]: Using fused_linear_logp_sm90_backward fast path.
+| shape (N x H x V)   |   native fwd ms |   triton fwd ms | fwd speedup   |   native f+b ms |   triton f+b ms | f+b speedup   |   native fwd MB |   triton fwd MB |   sm90 fwd ms | sm90 vs native   | sm90 vs triton   |   sm90 f+b ms |   sm90 fwd MB |
+|---------------------|-----------------|-----------------|---------------|-----------------|-----------------|---------------|-----------------|-----------------|---------------|------------------|------------------|---------------|---------------|
+| 4096x2048x32768     |           1.798 |           6.288 | 0.29x         |           4.347 |          47.181 | 0.09x         |            1280 |               0 |         3.653 | 0.49x            | 1.72x            |         7.818 |             2 |
+| 4096x2048x50257     |          10.268 |           9.878 | 1.04x         |          24.005 |          85.23  | 0.28x         |            1965 |               0 |         5.398 | 1.90x            | 1.83x            |        26.307 |             2 |
+| 4096x2048x131072    |           7.199 |          24.755 | 0.29x         |          17.097 |         264.012 | 0.06x         |            5120 |               0 |        14.107 | 0.51x            | 1.75x            |        30.434 |             2 |