docs(bench): add H100 SXM5 benchmark results alongside A100 baseline - #231
docs(bench): add H100 SXM5 benchmark results alongside A100 baseline#231Billy1900 wants to merge 5 commits into
Conversation
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 RL-Align#75.
📝 WalkthroughWalkthroughAdds H100 SXM5 benchmark tooling, raw and structured performance reports, and documented LogP, sampling, SM90 kernel, and Qwen3 validation results. ChangesH100 SXM5 benchmark validation
Estimated code review effort: 3 (Moderate) | ~25 minutes Possibly related issues
Possibly related PRs
Suggested labels: Suggested reviewers: Sequence Diagram(s)sequenceDiagram
participant CLI
participant QwenModel
participant CUDA
participant RLKernel
CLI->>QwenModel: Load model and obtain hidden states
QwenModel->>CUDA: Allocate benchmark inputs
CLI->>CUDA: Measure native logprob
CLI->>RLKernel: Dispatch linear_logp
RLKernel->>CUDA: Measure fused logprob
CUDA-->>CLI: Return latency, VRAM, or OOM results
🚥 Pre-merge checks | ✅ 5✅ Passed checks (5 passed)
✨ Finishing Touches🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
There was a problem hiding this comment.
Pull request overview
Adds NVIDIA H100 SXM5 benchmark/validation documentation and accompanying profiler output artifacts to complement the existing A100-focused performance claims, aiming to address Issue #75 by documenting H100 logp, sampling, and linear_logp behavior plus a real-model memory validation.
Changes:
- Adds H100 logp + sampling-native profiler outputs under
reports/(CSV + timestamped JSON). - Updates
docs/benchmarking/hardware-dashboard.mdwith an H100 environment entry and tables summarizing H100 logp/sampling/linear_logp/Qwen3 validation results. - Adds a new “H100 SXM5 Independent Validation” section to
README.mdlinking to the dashboard and reports.
Reviewed changes
Copilot reviewed 5 out of 5 changed files in this pull request and generated 4 comments.
Show a summary per file
| File | Description |
|---|---|
| reports/perf_report_NVIDIA_H100_80GB_HBM3.csv | Adds consolidated CSV profiler output including logp sweep and sampling-native rows on H100. |
| reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json | Adds timestamped JSON report for H100 logp sweep metrics. |
| reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json | Adds timestamped JSON report for H100 sampling-native metrics. |
| README.md | Adds an H100 validation section with headline results and links to dashboard/reports. |
| docs/benchmarking/hardware-dashboard.md | Fills in H100 SXM5 rows and adds summarized tables and methodology notes. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@docs/benchmarking/hardware-dashboard.md`:
- Around line 46-49: Update the artifact reference in the benchmarking
documentation to point to the committed timestamped JSON report files, using the
H100 report naming pattern shown in the review comment instead of the
nonexistent CSV path.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: e2caaeac-eb91-47b1-9023-a8a3fb75fb03
⛔ Files ignored due to path filters (1)
reports/perf_report_NVIDIA_H100_80GB_HBM3.csvis excluded by!**/*.csv
📒 Files selected for processing (4)
README.mddocs/benchmarking/hardware-dashboard.mdreports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.jsonreports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json
…gp/Qwen3 numbers Copilot review on RL-Align#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).
There was a problem hiding this comment.
🧹 Nitpick comments (1)
benchmarks/benchmark_qwen3_moe_real_model.py (1)
116-135: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueResolve static analysis errors for loop variables.
Ruff flags
hiddenandtargetas potentially undefined (F821) and warns about late-binding in the lambdas (B023). While the code executes correctly becausemeasureevaluates the lambda synchronously and thecontinuestatement prevents execution from reaching the lambdas if initialization fails, explicitly binding these variables in the lambdas and pre-initializing them satisfies static analysis.🛠️ Proposed fix to satisfy static analysis
- 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" + hidden = target = None + 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 h=hidden, t=target: native_logprob(h, lm_head_weight, t)) + 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 h=hidden, t=target: linear_logp_op(h, lm_head_weight, t)) + 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"🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@benchmarks/benchmark_qwen3_moe_real_model.py` around lines 116 - 135, Initialize hidden and target before the batch-construction try block so they are defined on every loop path, then explicitly bind both as default arguments in the native_logprob and linear_logp_op lambdas passed to measure. Preserve the existing OOM handling and synchronous measurement behavior.Source: Linters/SAST tools
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@benchmarks/benchmark_qwen3_moe_real_model.py`:
- Around line 116-135: Initialize hidden and target before the
batch-construction try block so they are defined on every loop path, then
explicitly bind both as default arguments in the native_logprob and
linear_logp_op lambdas passed to measure. Preserve the existing OOM handling and
synchronous measurement behavior.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: a1344e09-ef4f-48b1-83f5-5691435ff38c
📒 Files selected for processing (6)
README.mdbenchmarks/benchmark_qwen3_moe_real_model.pydocs/benchmarking/hardware-dashboard.mdreports/benchmark_linear_logp_sm90_NVIDIA_H100_80GB_HBM3.txtreports/benchmark_qwen3_30b_a3b_NVIDIA_H100_80GB_HBM3.txtreports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt
🚧 Files skipped from review as they are similar to previous changes (2)
- README.md
- docs/benchmarking/hardware-dashboard.md
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.
There was a problem hiding this comment.
Actionable comments posted: 1
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (3)
benchmarks/benchmark_qwen3_moe_real_model.py (3)
108-110: 🩺 Stability & Availability | 🟠 Major | ⚡ Quick winAvoid materializing the repeated hidden buffer twice.
repeat(reps, 1)[:n].clone()first allocates the full repeated tensor and then allocates the cloned slice. Near repetition boundaries this temporarily uses almost 2× the requested input memory and can report OOM for a token count that would otherwise fit.Proposed fix
def make_batch(n): - reps = (n + real_hidden.shape[0] - 1) // real_hidden.shape[0] - hidden = real_hidden.repeat(reps, 1)[:n].clone() + indices = torch.arange(n, device=real_hidden.device) % real_hidden.shape[0] + hidden = real_hidden.index_select(0, indices) target = torch.randint(0, vocab_size, (n,), device="cuda") return hidden, target🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@benchmarks/benchmark_qwen3_moe_real_model.py` around lines 108 - 110, Update make_batch so it constructs the requested n-row hidden tensor directly instead of materializing the full real_hidden.repeat(reps, 1) result before cloning the slice. Preserve the existing repeated-row contents and returned shape while ensuring only the needed input buffer is allocated.
69-82: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winReport actual free GPU memory as headroom.
torch.cuda.memory_allocated()only covers tensor allocations belonging to the PyTorch process, sototal_gb - weight_gbcan overstate usable VRAM by ignoring CUDA context, allocator-reserved/cached memory, and any existing usage. Usetorch.cuda.mem_get_info(0)after loading to report headroom, or rename this metric to make clear it is only an allocation estimate.Proposed fix
- total_gb = gb(torch.cuda.get_device_properties(0).total_memory) - headroom_gb = total_gb - weight_gb + free_bytes, total_bytes = torch.cuda.mem_get_info(0) + total_gb = gb(total_bytes) + headroom_gb = gb(free_bytes)🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@benchmarks/benchmark_qwen3_moe_real_model.py` around lines 69 - 82, Update the post-load memory reporting near weight_gb and headroom_gb to use torch.cuda.mem_get_info(0) after model loading, and derive headroom_gb from the returned free-memory value. Preserve weight_gb as the PyTorch allocation estimate, but ensure the reported headroom reflects actual available GPU memory rather than total memory minus allocations.
88-101: 🎯 Functional Correctness | 🟠 Major | ⚡ Quick winDo not expose the real-model benchmark with an unvalidated
--modeloption.
AutoModelForCausalLMis generic; checkpoint-specific attributes vary, butbenchmarks/benchmark_qwen3_moe_real_model.py:90directly accessesmodel.lm_headandbenchmarks/benchmark_qwen3_moe_real_model.py:100directly callsmodel.modelwithout checking whether the loaded model actually provides these names. Restrict--modelto supported Qwen3-compatible checkpoints, or fail fast with validation before accessing architecture-specific modules.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@benchmarks/benchmark_qwen3_moe_real_model.py` around lines 88 - 101, Validate the loaded model in the benchmark before the architecture-specific accesses to model.lm_head and model.model. Restrict --model to supported Qwen3-compatible checkpoints or fail fast with a clear error when those required modules are unavailable, ensuring the benchmark never reaches the lm_head shape logging or real forward pass with an incompatible model.
🧹 Nitpick comments (1)
benchmarks/benchmark_qwen3_moe_real_model.py (1)
132-139: 🩺 Stability & Availability | 🔵 TrivialIsolate the experimental SM90 backend from the benchmark process.
The PR documents that
fused_logp_sm90can terminate the process whencuTensorMapEncodeTiledfails. Catching onlytorch.cuda.OutOfMemoryErrorcannot preserve partial results in that case; run the risky backend in a subprocess if the benchmark must remain resilient and produce a report.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@benchmarks/benchmark_qwen3_moe_real_model.py` around lines 132 - 139, Update the benchmark flow around the fused log-probability measurement and linear_logp_op so the experimental SM90 backend runs in an isolated subprocess when resilience is required. Capture its success, timing, OOM, and abnormal termination as reportable results, allowing the parent benchmark process to continue and preserve partial output instead of relying only on torch.cuda.OutOfMemoryError handling.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@benchmarks/benchmark_qwen3_moe_real_model.py`:
- Around line 124-134: Update both benchmark callbacks passed to measure in the
native and kernel paths to bind the current hidden and target values at lambda
creation time, using default arguments or functools.partial. Preserve the
existing native_logprob and linear_logp_op calls while ensuring each callback is
independent of later loop-variable reassignment and satisfies Ruff B023.
---
Outside diff comments:
In `@benchmarks/benchmark_qwen3_moe_real_model.py`:
- Around line 108-110: Update make_batch so it constructs the requested n-row
hidden tensor directly instead of materializing the full
real_hidden.repeat(reps, 1) result before cloning the slice. Preserve the
existing repeated-row contents and returned shape while ensuring only the needed
input buffer is allocated.
- Around line 69-82: Update the post-load memory reporting near weight_gb and
headroom_gb to use torch.cuda.mem_get_info(0) after model loading, and derive
headroom_gb from the returned free-memory value. Preserve weight_gb as the
PyTorch allocation estimate, but ensure the reported headroom reflects actual
available GPU memory rather than total memory minus allocations.
- Around line 88-101: Validate the loaded model in the benchmark before the
architecture-specific accesses to model.lm_head and model.model. Restrict
--model to supported Qwen3-compatible checkpoints or fail fast with a clear
error when those required modules are unavailable, ensuring the benchmark never
reaches the lm_head shape logging or real forward pass with an incompatible
model.
---
Nitpick comments:
In `@benchmarks/benchmark_qwen3_moe_real_model.py`:
- Around line 132-139: Update the benchmark flow around the fused
log-probability measurement and linear_logp_op so the experimental SM90 backend
runs in an isolated subprocess when resilience is required. Capture its success,
timing, OOM, and abnormal termination as reportable results, allowing the parent
benchmark process to continue and preserve partial output instead of relying
only on torch.cuda.OutOfMemoryError handling.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro Plus
Run ID: 60cdc5bd-c95f-43eb-9c05-1c2c90fe983b
📒 Files selected for processing (4)
benchmarks/benchmark_qwen3_moe_real_model.pyreports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txtreports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.jsonreports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json
🚧 Files skipped from review as they are similar to previous changes (3)
- reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081350.json
- reports/perf_report_NVIDIA_H100_80GB_HBM3_20260719_081405.json
- reports/benchmark_sampling_NVIDIA_H100_80GB_HBM3.txt
| 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) |
There was a problem hiding this comment.
📐 Maintainability & Code Quality | 🟠 Major | ⚡ Quick win
Bind loop values in both benchmark callbacks.
Ruff reports B023 because these lambdas capture loop-scoped hidden and target. Bind them as default arguments (or use functools.partial) so the code passes lint and does not depend on measure() remaining synchronous.
Proposed fix
native_extra, native_ms = measure(
- lambda: native_logprob(hidden, lm_head_weight, target)
+ lambda hidden=hidden, target=target: native_logprob(
+ hidden, lm_head_weight, target
+ )
)
...
kernel_extra, kernel_ms = measure(
- lambda: linear_logp_op(hidden, lm_head_weight, target)
+ lambda hidden=hidden, target=target: linear_logp_op(
+ hidden, lm_head_weight, target
+ )
)📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| 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) | |
| native_extra, native_ms = measure( | |
| lambda hidden=hidden, target=target: 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 hidden=hidden, target=target: linear_logp_op( | |
| hidden, lm_head_weight, target | |
| ) |
🧰 Tools
🪛 Ruff (0.16.0)
[warning] 125-125: Function definition does not bind loop variable hidden
(B023)
[warning] 125-125: Function definition does not bind loop variable target
(B023)
[warning] 134-134: Function definition does not bind loop variable hidden
(B023)
[warning] 134-134: Function definition does not bind loop variable target
(B023)
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@benchmarks/benchmark_qwen3_moe_real_model.py` around lines 124 - 134, Update
both benchmark callbacks passed to measure in the native and kernel paths to
bind the current hidden and target values at lambda creation time, using default
arguments or functools.partial. Preserve the existing native_logprob and
linear_logp_op calls while ensuring each callback is independent of later
loop-variable reassignment and satisfies Ruff B023.
Source: Linters/SAST tools
Summary
pendingh100-sxm5rows indocs/benchmarking/hardware-dashboard.md: logp native/fused (viascripts/run_profile_suite.py), sampling (native vs FlashInfer), andlinear_logpSM90-vs-Triton-vs-native, all measured on a real H100 80GB HBM3 (SM90) with the extension built viaKERNEL_ALIGN_FORCE_SM90=1.README.md(new "H100 SXM5 Independent Validation" section) that links to the dashboard doc for full methodology, raw CSV/JSON reports (reports/perf_report_NVIDIA_H100_80GB_HBM3*), and exact reproduction commands.linear_logpkernel's tradeoffs honestly rather than cherry-picking: 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 — tile recomputation trades FLOPs for memory).benchmarks/profiler.py'sWORKLOAD_REGISTRYhas nosampling-fusedworkload, so the FlashInfer sampling numbers come frombenchmarks/benchmark_sampling.pydirectly instead of the profiler suite; (2) the experimental, off-by-defaultfused_logp_sm90standalone kernel (RL_KERNEL_ENABLE_EXPERIMENTAL_SM90_LOGP=1) currently aborts the process viaexit(EXIT_FAILURE)on a failedcuTensorMapEncodeTiledcall incsrc/utils/tma_utils.cuh, instead of raising a catchable error.Addresses #75.
Test plan
scripts/run_profile_suite.py --device cuda --dtype float16 --batch-sizes 8,16,32 --seq-lens 128,512 --vocab-sizes 4096,128256 --workloads logp-native,logp-fused— run on H100, raw output committed underreports/.benchmarks/benchmark_sampling.py --g-sizes 32,64,128,256 --vocab-size 128256 --top-k 50 --top-p 0.9— run on H100.benchmarks/benchmark_linear_logp.py(SM90 build) — run on H100.Qwen/Qwen3-30B-A3Bweights downloaded and loaded in bf16 on H100; weight VRAM + logprob memory/latency sweep measured directly (script not included in this PR — happy to add it underbenchmarks/if useful as a repeatable workload).pending(no ROCm hardware available for this pass).🤖 Generated with Claude Code
Summary by CodeRabbit