From 946edee86f3ee21fc16952f117fc28ebabb9eb20 Mon Sep 17 00:00:00 2001 From: Billy1900 Date: Sun, 19 Jul 2026 04:17:25 -0400 Subject: [PATCH 1/3] docs(bench): add H100 SXM5 benchmark results alongside A100 baseline Fills in the h100-sxm5 rows in the hardware benchmark dashboard (logp native/fused, sampling, linear_logp SM90-vs-Triton-vs-native) and adds a real Qwen3-30B-A3B validation using downloaded weights and a real forward pass instead of synthetic tensors, confirming the 56.9 GB / 23 GB headroom claim on actual hardware. Reports honestly: the SM90 linear_logp kernel beats Triton and uses far less memory than the naive path, but does not uniformly beat naive on raw latency (tile-recompute trades FLOPs for memory). Also documents two gaps found while benchmarking: no sampling-fused workload in the profiler suite, and the experimental (off-by-default) fused_logp_sm90 kernel currently crashes the process on cuTensorMapEncodeTiled failure. Addresses #75. --- README.md | 18 + docs/benchmarking/hardware-dashboard.md | 74 +- reports/perf_report_NVIDIA_H100_80GB_HBM3.csv | 31 + ...NVIDIA_H100_80GB_HBM3_20260719_081350.json | 687 ++++++++++++++++++ ...NVIDIA_H100_80GB_HBM3_20260719_081405.json | 183 +++++ 5 files changed, 990 insertions(+), 3 deletions(-) create mode 100644 reports/perf_report_NVIDIA_H100_80GB_HBM3.csv create mode 100644 reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json create mode 100644 reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json diff --git a/README.md b/README.md index f1ab58d8..66fc8b7c 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)**: 32.43 ms (native) → 1.57 ms (RL-Kernel), ~21x. +- **`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/docs/benchmarking/hardware-dashboard.md b/docs/benchmarking/hardware-dashboard.md index 044d5acf..c4587901 100644 --- a/docs/benchmarking/hardware-dashboard.md +++ b/docs/benchmarking/hardware-dashboard.md @@ -30,23 +30,91 @@ 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 | 10.66 | 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 | 1.09 | n/a | n/a | pass | same as above | +| `h100-sxm5` | native | 256 | 128256 | 50 | 0.9 | 1.0 | 32.43 | n/a | n/a | pass | same as above | +| `h100-sxm5` | FlashInfer (`RL_Sampler`) | 256 | 128256 | 50 | 0.9 | 1.0 | 1.57 | n/a | n/a | pass | same as above | | `mi300-template` | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending | pending | +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`). + +| 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.37 | 3.78 | 1280 | 0 | 2 | 1.69x faster | 0.48x (slower) | +| 4096×2048×50257 | 10.26 | 9.94 | 5.40 | 1965 | 0 | 2 | 1.84x faster | 1.90x faster | +| 4096×2048×131072 | 7.20 | 24.74 | 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. + +| 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.09 | 48.47 | pass | +| 16,384 | OOM | 0.00 GB | n/a | 66.08 | native: oom | +| 24,576 | OOM | 0.00 GB | n/a | 105.98 | 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/perf_report_NVIDIA_H100_80GB_HBM3.csv b/reports/perf_report_NVIDIA_H100_80GB_HBM3.csv new file mode 100644 index 00000000..96d992ae --- /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 00000000..3d08f439 --- /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": {} + } + ] +} \ No newline at end of file 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 00000000..3c5d0777 --- /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": {} + } + ] +} \ No newline at end of file From 55a2b4e3424a8762cd225066e23aedb3ed883ed5 Mon Sep 17 00:00:00 2001 From: Billy1900 Date: Sun, 19 Jul 2026 04:50:37 -0400 Subject: [PATCH 2/3] docs(bench): commit raw artifacts backing the H100 sampling/linear_logp/Qwen3 numbers Copilot review on #231 correctly flagged that the FlashInfer sampling, linear_logp SM90, and Qwen3-30B-A3B numbers in the dashboard/README had no committed raw output to audit against. Fixes that: - Adds benchmarks/benchmark_qwen3_moe_real_model.py (previously only a local scratch script) so the real-model benchmark is itself reviewable and reproducible, not just its output. - Re-runs all three benchmarks and commits their raw console output under reports/, then updates the dashboard/README numbers to match those committed runs exactly (previous figures were from an earlier, uncommitted run and differed by normal run-to-run noise). --- README.md | 2 +- benchmarks/benchmark_qwen3_moe_real_model.py | 190 ++++++++++++++++++ docs/benchmarking/hardware-dashboard.md | 28 ++- ...linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txt | 15 ++ ...rk_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txt | 31 +++ ...nchmark_sampling_NVIDIA_H100_80GB_HBM3.txt | 41 ++++ 6 files changed, 296 insertions(+), 11 deletions(-) create mode 100644 benchmarks/benchmark_qwen3_moe_real_model.py create mode 100644 reports/benchmark_linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txt create mode 100644 reports/benchmark_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txt create mode 100644 reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt diff --git a/README.md b/README.md index 66fc8b7c..d15b4c04 100644 --- a/README.md +++ b/README.md @@ -104,7 +104,7 @@ methodology, exact commands, and raw CSVs are in the - **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)**: 32.43 ms (native) → 1.57 ms (RL-Kernel), ~21x. +- **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 diff --git a/benchmarks/benchmark_qwen3_moe_real_model.py b/benchmarks/benchmark_qwen3_moe_real_model.py new file mode 100644 index 00000000..725c0452 --- /dev/null +++ b/benchmarks/benchmark_qwen3_moe_real_model.py @@ -0,0 +1,190 @@ +# 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", + ] + ) + del hidden, target + 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 | 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 c4587901..813c35a5 100644 --- a/docs/benchmarking/hardware-dashboard.md +++ b/docs/benchmarking/hardware-dashboard.md @@ -52,12 +52,16 @@ note below for why SM90-specific dispatch is more nuanced. | Environment | Backend | Batch | Vocabulary | Top-k | Top-p | Temperature | Latency (ms) | Tokens/s | Peak VRAM (GB) | Status | Command | | --- | --- | ---: | ---: | ---: | ---: | ---: | ---: | ---: | ---: | --- | --- | -| `h100-sxm5` | native (topk→topp→softmax→multinomial) | 64 | 128256 | 50 | 0.9 | 1.0 | 10.66 | 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 | 1.09 | n/a | n/a | pass | same as above | -| `h100-sxm5` | native | 256 | 128256 | 50 | 0.9 | 1.0 | 32.43 | n/a | n/a | pass | same as above | -| `h100-sxm5` | FlashInfer (`RL_Sampler`) | 256 | 128256 | 50 | 0.9 | 1.0 | 1.57 | n/a | n/a | pass | same as above | +| `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 @@ -73,12 +77,13 @@ op; the question "does SM90 warp specialization help vs SM86" reduces to "does t 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.37 | 3.78 | 1280 | 0 | 2 | 1.69x faster | 0.48x (slower) | -| 4096×2048×50257 | 10.26 | 9.94 | 5.40 | 1965 | 0 | 2 | 1.84x faster | 1.90x faster | -| 4096×2048×131072 | 7.20 | 24.74 | 14.11 | 5120 | 0 | 2 | 1.75x faster | 0.51x (slower) | +| 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** @@ -97,6 +102,9 @@ the flag defaults off and isn't part of any existing benchmark claim. 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 | @@ -105,9 +113,9 @@ Real weights (`Qwen/Qwen3-30B-A3B`, 56.9 GB, downloaded from the HF Hub) and a r | 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.09 | 48.47 | pass | -| 16,384 | OOM | 0.00 GB | n/a | 66.08 | native: oom | -| 24,576 | OOM | 0.00 GB | n/a | 105.98 | native: oom | +| 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 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 00000000..5109ed00 --- /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 00000000..20810514 --- /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 00000000..9bd6c55e --- /dev/null +++ b/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt @@ -0,0 +1,41 @@ +# 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 │ +╘══════════════════╧══════════════════╧═════════════╧═══════════╛ +================================================================================ + From 8b69ca422c65f54053633a9bb6fdd72e12153eb3 Mon Sep 17 00:00:00 2001 From: Billy1900 Date: Tue, 28 Jul 2026 23:13:34 -0400 Subject: [PATCH 3/3] fix(lint): resolve pre-commit failures in benchmark script and reports Remove del hidden, target which caused flake8 F821 false positives (pyflakes checks lambda bodies in a deferred pass after the del had already removed the bindings from scope). Apply black formatting and fix trailing whitespace/EOF newlines in report files. --- benchmarks/benchmark_qwen3_moe_real_model.py | 14 ++++++++++---- .../benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt | 5 ++--- ...port_NVIDIA_H100_80GB_HBM3_20260719_081350.json | 2 +- ...port_NVIDIA_H100_80GB_HBM3_20260719_081405.json | 2 +- 4 files changed, 14 insertions(+), 9 deletions(-) diff --git a/benchmarks/benchmark_qwen3_moe_real_model.py b/benchmarks/benchmark_qwen3_moe_real_model.py index 725c0452..46f85b35 100644 --- a/benchmarks/benchmark_qwen3_moe_real_model.py +++ b/benchmarks/benchmark_qwen3_moe_real_model.py @@ -121,14 +121,18 @@ def make_batch(n): continue try: - native_extra, native_ms = measure(lambda: native_logprob(hidden, lm_head_weight, target)) + 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_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" @@ -144,12 +148,14 @@ def make_batch(n): f"current alloc: {gb(torch.cuda.memory_allocated()):.1f} GB", ] ) - del hidden, target 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 | Headroom: {headroom_gb:.2f} GB") + 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( diff --git a/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt b/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt index 9bd6c55e..a8beb425 100644 --- a/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt +++ b/reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt @@ -11,10 +11,10 @@ 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. +/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. +/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 @@ -38,4 +38,3 @@ RL-KERNEL SAMPLING BENCHMARK REPORT (TopK=50, TopP=0.9) │ 256 │ 29.30 ms │ 1.65 ms │ 17.71x │ ╘══════════════════╧══════════════════╧═════════════╧═══════════╛ ================================================================================ - diff --git a/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json index 3d08f439..1eb22404 100644 --- a/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json +++ b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json @@ -684,4 +684,4 @@ "extra": {} } ] -} \ No newline at end of file +} diff --git a/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json index 3c5d0777..96aa1492 100644 --- a/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json +++ b/reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json @@ -180,4 +180,4 @@ "extra": {} } ] -} \ No newline at end of file +}