diff --git a/README.md b/README.md index f1ab58d..d15b4c0 100644 --- a/README.md +++ b/README.md @@ -94,6 +94,24 @@ Model weights consume 56.9 GB — only 23 GB headroom remaining for training com Real Model MoE Benchmark

+### 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. diff --git a/benchmarks/benchmark_qwen3_moe_real_model.py b/benchmarks/benchmark_qwen3_moe_real_model.py new file mode 100644 index 0000000..46f85b3 --- /dev/null +++ b/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()) diff --git a/docs/benchmarking/hardware-dashboard.md b/docs/benchmarking/hardware-dashboard.md index 044d5ac..813c35a 100644 --- a/docs/benchmarking/hardware-dashboard.md +++ b/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: diff --git a/reports/benchmark_linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txt b/reports/benchmark_linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txt new file mode 100644 index 0000000..5109ed0 --- /dev/null +++ b/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 | diff --git a/reports/benchmark_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txt b/reports/benchmark_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txt new file mode 100644 index 0000000..2081051 --- /dev/null +++ b/reports/benchmark_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txt @@ -0,0 +1,31 @@ +# RL-Kernel real-model logprob benchmark — raw output (Qwen3-30B-A3B, real weights + real forward pass) +# GPU: NVIDIA H100 80GB HBM3 (SXM5) | Driver 535.309.01 | CUDA 12.4 | PyTorch 2.6.0+cu124 | transformers 4.51.1 +# Extension built with KERNEL_ALIGN_FORCE_SM90=1 +# Command: python benchmarks/benchmark_qwen3_moe_real_model.py +# Date: 2026-07-19T08:47:54Z + +INFO 07-19 04:47:57 [RL-Kernel]: RL-Engine initialized with NVIDIA CUDA backend (Version: 12.4) +INFO 07-19 04:47:57 [RL-Kernel]: KernelRegistry initialized for cuda +INFO 07-19 04:47:58 [RL-Kernel]: Loading Qwen/Qwen3-30B-A3B in torch.bfloat16 ... + Loading checkpoint shards: 0%| | 0/16 [00:00 headroom 22.24 GB +INFO 07-19 04:48:26 [RL-Kernel]: hidden_size=2048 vocab_size=151936 lm_head shape=(151936, 2048) +INFO 07-19 04:48:26 [RL-Kernel]: Running a real forward pass to obtain real hidden states... +INFO 07-19 04:48:27 [RL-Kernel]: Real hidden states shape: (29, 2048) +INFO 07-19 04:48:27 [RL-Kernel]: Successfully linked to precompiled _C.fused_linear_logp_sm90 kernel. +INFO 07-19 04:48:27 [RL-Kernel]: Dispatched linear_logp backend: FusedLinearLogpSM90Op + +==================================================================================================== +Qwen/Qwen3-30B-A3B REAL MODEL LOGPROB BENCHMARK on NVIDIA H100 80GB HBM3 +Weight VRAM: 56.87 GB | Total: 79.11 GB | Headroom: 22.24 GB +linear_logp backend: FusedLinearLogpSM90Op +==================================================================================================== +| N tokens | Native extra VRAM | RL-Kernel extra VRAM | Native ms | RL-Kernel ms | note | +|------------|---------------------|------------------------|-------------|----------------|------------------------| +| 2048 | 2.90 GB | 0.00 GB | 4.21 ms | 8.06 ms | current alloc: 56.9 GB | +| 4096 | 5.80 GB | 0.00 GB | 8.36 ms | 16.02 ms | current alloc: 56.9 GB | +| 8192 | 11.59 GB | 0.00 GB | 16.71 ms | 33.02 ms | current alloc: 56.9 GB | +| 12288 | 17.39 GB | 0.00 GB | 25.08 ms | 48.03 ms | current alloc: 56.9 GB | +| 16384 | OOM | 0.00 GB | N/A | 65.90 ms | current alloc: 57.0 GB | +| 20480 | OOM | 0.00 GB | N/A | 87.83 ms | current alloc: 57.0 GB | +| 24576 | OOM | 0.00 GB | N/A | 105.50 ms | current alloc: 57.0 GB | diff --git a/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt b/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt new file mode 100644 index 0000000..a8beb42 --- /dev/null +++ b/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt @@ -0,0 +1,40 @@ +# RL-Kernel sampling benchmark — raw output +# GPU: NVIDIA H100 80GB HBM3 (SXM5) | Driver 535.309.01 | CUDA 12.4 | PyTorch 2.6.0+cu124 | FlashInfer 0.2.5 +# Command: python benchmarks/benchmark_sampling.py --g-sizes 32,64,128,256 --vocab-size 128256 --top-k 50 --top-p 0.9 +# Date: 2026-07-19T08:44:50Z + +INFO 07-19 04:44:53 [RL-Kernel]: RL-Engine initialized with NVIDIA CUDA backend (Version: 12.4) +INFO 07-19 04:44:53 [RL-Kernel]: Detected NVIDIA GPU (CUDA) - Using FlashInfer backend +2026-07-19 04:44:54,082 - INFO - flashinfer.jit: Prebuilt kernels not found, using JIT backend +INFO 07-19 04:44:54 [RL-Kernel]: FlashInfer kernels loaded successfully. +INFO 07-19 04:44:54 [RL-Kernel]: Starting Sampling Benchmark on cuda +INFO 07-19 04:44:54 [RL-Kernel]: Config: VocabSize=128256, TopK=50, TopP=0.9 +INFO 07-19 04:44:54 [RL-Kernel]: Warming up kernels... +2026-07-19 04:44:54,376 - INFO - flashinfer.jit: Loading JIT ops: sampling +/gpu02home/nkl5280/miniconda3/envs/mlaaware/lib/python3.10/site-packages/torch/utils/cpp_extension.py:2059: UserWarning: TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation. +If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST']. + warnings.warn( +/gpu02home/nkl5280/miniconda3/envs/mlaaware/lib/python3.10/site-packages/torch/utils/cpp_extension.py:2059: UserWarning: TORCH_CUDA_ARCH_LIST is not set, all archs for visible cards are included for compilation. +If this is not desired, please set os.environ['TORCH_CUDA_ARCH_LIST']. + warnings.warn( +2026-07-19 04:44:54,415 - INFO - flashinfer.jit: Finished loading JIT ops: sampling +INFO 07-19 04:44:54 [RL-Kernel]: Testing Batch Size G=32... +INFO 07-19 04:44:55 [RL-Kernel]: Testing Batch Size G=64... +INFO 07-19 04:44:55 [RL-Kernel]: Testing Batch Size G=128... +INFO 07-19 04:44:55 [RL-Kernel]: Testing Batch Size G=256... + +================================================================================ +RL-KERNEL SAMPLING BENCHMARK REPORT (TopK=50, TopP=0.9) +================================================================================ +╒══════════════════╤══════════════════╤═════════════╤═══════════╕ +│ Batch Size (G) │ Native Latency │ RL-Kernel │ Speedup │ +╞══════════════════╪══════════════════╪═════════════╪═══════════╡ +│ 32 │ 711.22 ms │ 0.97 ms │ 732.86x │ +├──────────────────┼──────────────────┼─────────────┼───────────┤ +│ 64 │ 8.73 ms │ 0.84 ms │ 10.33x │ +├──────────────────┼──────────────────┼─────────────┼───────────┤ +│ 128 │ 15.56 ms │ 1.00 ms │ 15.62x │ +├──────────────────┼──────────────────┼─────────────┼───────────┤ +│ 256 │ 29.30 ms │ 1.65 ms │ 17.71x │ +╘══════════════════╧══════════════════╧═════════════╧═══════════╛ +================================================================================ diff --git a/reports/perf_report_NVIDIA_H100_80GB_HBM3.csv b/reports/perf_report_NVIDIA_H100_80GB_HBM3.csv new file mode 100644 index 0000000..96d992a --- /dev/null +++ b/reports/perf_report_NVIDIA_H100_80GB_HBM3.csv @@ -0,0 +1,31 @@ +batch_size,benchmark_name,gpu_architecture,gpu_backend,gpu_compute_capability,gpu_device_index,gpu_driver_version,gpu_name,gpu_total_memory_gb,latency_ms,latency_std_ms,notes,peak_vram_gb,peak_vram_reduction_gb,repeat_iterations,seq_len,status,tflops,timestamp,tokens_per_sec,total_tokens,vocab_size,warmup_iterations +8,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.07048000022768974,0.14564340182830976,,0.0390777587890625,,10,128,pass,0.2975527799694989,2026-07-19T08:13:49.355864+00:00,14528944.33444819,1024,4096,3 +8,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.03223999962210655,0.007813027023014181,,0.015651702880859375,0.03125,10,128,pass,0.6504813972026258,2026-07-19T08:13:49.367397+00:00,31761786.972784463,1024,4096,3 +8,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,1.4682559967041016,0.42242479020318746,,1.2251129150390625,,10,128,pass,0.44724538600494423,2026-07-19T08:13:49.390668+00:00,697426.063505714,1024,128256,3 +8,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.44065600633621216,0.024757200304702666,,0.4892845153808594,0.98046875,10,128,pass,1.4902116629699873,2026-07-19T08:13:49.409014+00:00,2323808.1071762526,1024,128256,3 +8,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.19468799978494644,0.004186119690128024,,0.15631103515625,,10,512,pass,0.4308744251965251,2026-07-19T08:13:49.413953+00:00,21038790.292799078,4096,4096,3 +8,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.048736000433564186,0.0012531813785628062,,0.0626068115234375,0.125,10,512,pass,1.7212343904656602,2026-07-19T08:13:49.416526+00:00,84044647.97195606,4096,4096,3 +8,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,5.7114880084991455,0.543704311402483,,4.89263916015625,,10,512,pass,0.4598946677453035,2026-07-19T08:13:49.498278+00:00,717151.1161197972,4096,128256,3 +8,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,1.6853919625282288,0.08949274100935575,,1.9571380615234375,3.9140625,10,512,pass,1.5584997071302962,2026-07-19T08:13:49.560716+00:00,2430295.201987114,4096,128256,3 +16,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.11575999855995178,0.007184852476117778,,0.078155517578125,,10,128,pass,0.3623275787989736,2026-07-19T08:13:49.567464+00:00,17691776.308543634,2048,4096,3 +16,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.034383999183773994,0.0027051407968610897,,0.03130340576171875,0.0625,10,128,pass,1.2198418158348827,2026-07-19T08:13:49.569606+00:00,59562588.66381264,2048,4096,3 +16,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,2.8949280977249146,0.6106983816844578,,2.446319580078125,,10,128,pass,0.4536697961625152,2026-07-19T08:13:49.614780+00:00,707444.1681675948,2048,128256,3 +16,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.8684479892253876,0.05456398517326904,,0.9785690307617188,1.95703125,10,128,pass,1.5122856593535732,2026-07-19T08:13:49.650578+00:00,2358229.88297401,2048,128256,3 +16,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.356672003865242,0.00521876305426766,,0.3126220703125,,10,512,pass,0.47038219479482263,2026-07-19T08:13:49.658667+00:00,22967880.60521595,8192,4096,3 +16,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.09430400282144547,0.0010523893741398102,,0.125213623046875,0.25,10,512,pass,1.7790566145708429,2026-07-19T08:13:49.662785+00:00,86867998.75834194,8192,4096,3 +16,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,11.36900806427002,0.8880843584656619,,9.7852783203125,,10,512,pass,0.4620777582619568,2026-07-19T08:13:49.823775+00:00,720555.386511285,8192,128256,3 +16,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,3.3190720081329346,0.1793436943975878,,3.914276123046875,7.828125,10,512,pass,1.5827814964927973,2026-07-19T08:13:49.944974+00:00,2468159.7687325305,8192,128256,3 +32,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.19995199888944626,0.006524416779928997,,0.15631103515625,,10,128,pass,0.419531089791109,2026-07-19T08:13:49.959713+00:00,20484916.493706495,4096,4096,3 +32,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.056432001292705536,0.0015660179285981511,,0.0626068115234375,0.125,10,128,pass,1.4864984065493918,2026-07-19T08:13:49.962811+00:00,72582930.00729452,4096,4096,3 +32,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,5.711199998855591,0.5593094263567243,,4.89263916015625,,10,128,pass,0.4599178597363662,2026-07-19T08:13:50.048115+00:00,717187.281275521,4096,128256,3 +32,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,1.7408799529075623,0.10505182946176396,,1.9571380615234375,3.9140625,10,128,pass,1.5088248190881846,2026-07-19T08:13:50.112505+00:00,2352833.1135980925,4096,128256,3 +32,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.6809599995613098,0.005424716325805069,,0.625244140625,,10,512,pass,0.4927518800166903,2026-07-19T08:13:50.125636+00:00,24060150.391439956,16384,4096,3 +32,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,0.17295999825000763,0.008745538835254063,,0.25042724609375,0.5,10,512,pass,1.9400111204614052,2026-07-19T08:13:50.133597+00:00,94727105.49127956,16384,4096,3 +32,logp_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,23.288352012634277,0.8077733488005617,,19.570556640625,,10,512,pass,0.4511582234028385,2026-07-19T08:13:50.450632+00:00,703527.6687294763,16384,128256,3 +32,logp_fused,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,6.934223890304565,0.012915920597149159,,7.82855224609375,15.65625,10,512,pass,1.5151993483640638,2026-07-19T08:13:50.692229+00:00,2362773.4349489515,16384,128256,3 +64,sampling_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,7.564879894256592,0.6454671091540795,,0.40225791931152344,,10,128,pass,0.005425323412095392,2026-07-19T08:14:04.383107+00:00,1082898.884649779,8192,128256,3 +64,sampling_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,7.520159959793091,0.5859296206931399,,0.40225791931152344,,10,512,pass,0.005457586037987578,2026-07-19T08:14:04.490511+00:00,4357354.122145771,32768,128256,3 +128,sampling_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,15.566224098205566,5.571870377029163,,0.6145744323730469,,10,128,pass,0.0052732017400072256,2026-07-19T08:14:04.715619+00:00,1052535.277446552,16384,128256,3 +128,sampling_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,14.841343879699707,1.7508382671291745,,0.6145744323730469,,10,512,pass,0.00553075521093989,2026-07-19T08:14:04.924452+00:00,4415772.623504902,65536,128256,3 +256,sampling_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,30.545663833618164,1.4212846855153107,,1.2249984741210938,,10,128,pass,0.005374500318415708,2026-07-19T08:14:05.315329+00:00,1072754.5545739937,32768,128256,3 +256,sampling_native,Hopper,cuda,9.0,0,12.4,NVIDIA H100 80GB HBM3,79.11,29.640719413757324,2.119837827020084,,1.2249984741210938,,10,512,pass,0.005538586216763817,2026-07-19T08:14:05.709231+00:00,4422024.923563926,131072,128256,3 diff --git a/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json new file mode 100644 index 0000000..1eb2240 --- /dev/null +++ b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json @@ -0,0 +1,687 @@ +{ + "report_version": "1.0", + "generated_at": "2026-07-19T08:13:50.692667+00:00", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "metrics": [ + { + "timestamp": "2026-07-19T08:13:49.355864+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 128, + "vocab_size": 4096, + "total_tokens": 1024, + "latency_ms": 0.07048000022768974, + "latency_std_ms": 0.14564340182830976, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 14528944.33444819, + "tflops": 0.2975527799694989, + "peak_vram_gb": 0.0390777587890625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.367397+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 128, + "vocab_size": 4096, + "total_tokens": 1024, + "latency_ms": 0.03223999962210655, + "latency_std_ms": 0.007813027023014181, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 31761786.972784463, + "tflops": 0.6504813972026258, + "peak_vram_gb": 0.015651702880859375, + "peak_vram_reduction_gb": 0.03125, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.390668+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 1024, + "latency_ms": 1.4682559967041016, + "latency_std_ms": 0.42242479020318746, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 697426.063505714, + "tflops": 0.44724538600494423, + "peak_vram_gb": 1.2251129150390625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.409014+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 1024, + "latency_ms": 0.44065600633621216, + "latency_std_ms": 0.024757200304702666, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 2323808.1071762526, + "tflops": 1.4902116629699873, + "peak_vram_gb": 0.4892845153808594, + "peak_vram_reduction_gb": 0.98046875, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.413953+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 512, + "vocab_size": 4096, + "total_tokens": 4096, + "latency_ms": 0.19468799978494644, + "latency_std_ms": 0.004186119690128024, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 21038790.292799078, + "tflops": 0.4308744251965251, + "peak_vram_gb": 0.15631103515625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.416526+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 512, + "vocab_size": 4096, + "total_tokens": 4096, + "latency_ms": 0.048736000433564186, + "latency_std_ms": 0.0012531813785628062, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 84044647.97195606, + "tflops": 1.7212343904656602, + "peak_vram_gb": 0.0626068115234375, + "peak_vram_reduction_gb": 0.125, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.498278+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 4096, + "latency_ms": 5.7114880084991455, + "latency_std_ms": 0.543704311402483, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 717151.1161197972, + "tflops": 0.4598946677453035, + "peak_vram_gb": 4.89263916015625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.560716+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 8, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 4096, + "latency_ms": 1.6853919625282288, + "latency_std_ms": 0.08949274100935575, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 2430295.201987114, + "tflops": 1.5584997071302962, + "peak_vram_gb": 1.9571380615234375, + "peak_vram_reduction_gb": 3.9140625, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.567464+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 128, + "vocab_size": 4096, + "total_tokens": 2048, + "latency_ms": 0.11575999855995178, + "latency_std_ms": 0.007184852476117778, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 17691776.308543634, + "tflops": 0.3623275787989736, + "peak_vram_gb": 0.078155517578125, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.569606+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 128, + "vocab_size": 4096, + "total_tokens": 2048, + "latency_ms": 0.034383999183773994, + "latency_std_ms": 0.0027051407968610897, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 59562588.66381264, + "tflops": 1.2198418158348827, + "peak_vram_gb": 0.03130340576171875, + "peak_vram_reduction_gb": 0.0625, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.614780+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 2048, + "latency_ms": 2.8949280977249146, + "latency_std_ms": 0.6106983816844578, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 707444.1681675948, + "tflops": 0.4536697961625152, + "peak_vram_gb": 2.446319580078125, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.650578+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 2048, + "latency_ms": 0.8684479892253876, + "latency_std_ms": 0.05456398517326904, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 2358229.88297401, + "tflops": 1.5122856593535732, + "peak_vram_gb": 0.9785690307617188, + "peak_vram_reduction_gb": 1.95703125, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.658667+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 512, + "vocab_size": 4096, + "total_tokens": 8192, + "latency_ms": 0.356672003865242, + "latency_std_ms": 0.00521876305426766, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 22967880.60521595, + "tflops": 0.47038219479482263, + "peak_vram_gb": 0.3126220703125, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.662785+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 512, + "vocab_size": 4096, + "total_tokens": 8192, + "latency_ms": 0.09430400282144547, + "latency_std_ms": 0.0010523893741398102, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 86867998.75834194, + "tflops": 1.7790566145708429, + "peak_vram_gb": 0.125213623046875, + "peak_vram_reduction_gb": 0.25, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.823775+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 8192, + "latency_ms": 11.36900806427002, + "latency_std_ms": 0.8880843584656619, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 720555.386511285, + "tflops": 0.4620777582619568, + "peak_vram_gb": 9.7852783203125, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.944974+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 16, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 8192, + "latency_ms": 3.3190720081329346, + "latency_std_ms": 0.1793436943975878, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 2468159.7687325305, + "tflops": 1.5827814964927973, + "peak_vram_gb": 3.914276123046875, + "peak_vram_reduction_gb": 7.828125, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.959713+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 128, + "vocab_size": 4096, + "total_tokens": 4096, + "latency_ms": 0.19995199888944626, + "latency_std_ms": 0.006524416779928997, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 20484916.493706495, + "tflops": 0.419531089791109, + "peak_vram_gb": 0.15631103515625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:49.962811+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 128, + "vocab_size": 4096, + "total_tokens": 4096, + "latency_ms": 0.056432001292705536, + "latency_std_ms": 0.0015660179285981511, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 72582930.00729452, + "tflops": 1.4864984065493918, + "peak_vram_gb": 0.0626068115234375, + "peak_vram_reduction_gb": 0.125, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:50.048115+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 4096, + "latency_ms": 5.711199998855591, + "latency_std_ms": 0.5593094263567243, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 717187.281275521, + "tflops": 0.4599178597363662, + "peak_vram_gb": 4.89263916015625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:50.112505+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 4096, + "latency_ms": 1.7408799529075623, + "latency_std_ms": 0.10505182946176396, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 2352833.1135980925, + "tflops": 1.5088248190881846, + "peak_vram_gb": 1.9571380615234375, + "peak_vram_reduction_gb": 3.9140625, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:50.125636+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 512, + "vocab_size": 4096, + "total_tokens": 16384, + "latency_ms": 0.6809599995613098, + "latency_std_ms": 0.005424716325805069, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 24060150.391439956, + "tflops": 0.4927518800166903, + "peak_vram_gb": 0.625244140625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:50.133597+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 512, + "vocab_size": 4096, + "total_tokens": 16384, + "latency_ms": 0.17295999825000763, + "latency_std_ms": 0.008745538835254063, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 94727105.49127956, + "tflops": 1.9400111204614052, + "peak_vram_gb": 0.25042724609375, + "peak_vram_reduction_gb": 0.5, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:50.450632+00:00", + "benchmark_name": "logp_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 16384, + "latency_ms": 23.288352012634277, + "latency_std_ms": 0.8077733488005617, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 703527.6687294763, + "tflops": 0.4511582234028385, + "peak_vram_gb": 19.570556640625, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:13:50.692229+00:00", + "benchmark_name": "logp_fused", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 32, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 16384, + "latency_ms": 6.934223890304565, + "latency_std_ms": 0.012915920597149159, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 2362773.4349489515, + "tflops": 1.5151993483640638, + "peak_vram_gb": 7.82855224609375, + "peak_vram_reduction_gb": 15.65625, + "status": "pass", + "notes": "", + "extra": {} + } + ] +} diff --git a/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json new file mode 100644 index 0000000..96aa149 --- /dev/null +++ b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json @@ -0,0 +1,183 @@ +{ + "report_version": "1.0", + "generated_at": "2026-07-19T08:14:05.709615+00:00", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "metrics": [ + { + "timestamp": "2026-07-19T08:14:04.383107+00:00", + "benchmark_name": "sampling_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 64, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 8192, + "latency_ms": 7.564879894256592, + "latency_std_ms": 0.6454671091540795, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 1082898.884649779, + "tflops": 0.005425323412095392, + "peak_vram_gb": 0.40225791931152344, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:14:04.490511+00:00", + "benchmark_name": "sampling_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 64, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 32768, + "latency_ms": 7.520159959793091, + "latency_std_ms": 0.5859296206931399, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 4357354.122145771, + "tflops": 0.005457586037987578, + "peak_vram_gb": 0.40225791931152344, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:14:04.715619+00:00", + "benchmark_name": "sampling_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 128, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 16384, + "latency_ms": 15.566224098205566, + "latency_std_ms": 5.571870377029163, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 1052535.277446552, + "tflops": 0.0052732017400072256, + "peak_vram_gb": 0.6145744323730469, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:14:04.924452+00:00", + "benchmark_name": "sampling_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 128, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 65536, + "latency_ms": 14.841343879699707, + "latency_std_ms": 1.7508382671291745, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 4415772.623504902, + "tflops": 0.00553075521093989, + "peak_vram_gb": 0.6145744323730469, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:14:05.315329+00:00", + "benchmark_name": "sampling_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 256, + "seq_len": 128, + "vocab_size": 128256, + "total_tokens": 32768, + "latency_ms": 30.545663833618164, + "latency_std_ms": 1.4212846855153107, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 1072754.5545739937, + "tflops": 0.005374500318415708, + "peak_vram_gb": 1.2249984741210938, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + }, + { + "timestamp": "2026-07-19T08:14:05.709231+00:00", + "benchmark_name": "sampling_native", + "gpu_target": { + "name": "NVIDIA H100 80GB HBM3", + "architecture": "Hopper", + "total_memory_gb": 79.11, + "driver_version": "12.4", + "backend": "cuda", + "compute_capability": "9.0", + "device_index": 0 + }, + "batch_size": 256, + "seq_len": 512, + "vocab_size": 128256, + "total_tokens": 131072, + "latency_ms": 29.640719413757324, + "latency_std_ms": 2.119837827020084, + "warmup_iterations": 3, + "repeat_iterations": 10, + "tokens_per_sec": 4422024.923563926, + "tflops": 0.005538586216763817, + "peak_vram_gb": 1.2249984741210938, + "peak_vram_reduction_gb": null, + "status": "pass", + "notes": "", + "extra": {} + } + ] +}