65df23d170
## Summary
Add HBM bandwidth microbenchmark for GPU hardware characterization (M6
module).
- `benchmarks/hardware/memory/hbm_bandwidth.py` — measures
read/write/copy bandwidth across sizes and dtypes, computes calibration
factor against theoretical peak
- Reads theoretical specs directly from `tileops/perf/profiles/` via
`load_profile()` (single source of truth)
- `benchmarks/hardware/utils/bench.py` — thin wrapper around
`tilelang.profiler.do_bench` with bandwidth calculator
- `docs/roofline.md` — relax "zero external dependencies" to "no
third-party dependencies beyond the project itself"
No CSV output, no env.py, no global state — minimal code focused on
producing the `hbm.calibration` value for GPU profiles.
## Test plan
- [ ] `python -c "import ast;
ast.parse(open('benchmarks/hardware/memory/hbm_bandwidth.py').read())"`
— syntax OK
- [ ] Run `python benchmarks/hardware/memory/hbm_bandwidth.py` on H200 —
verify theoretical from profile, measured peak, and calibration factor
🤖 Generated with [Claude Code](https://claude.com/claude-code)
---------
Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
|
||
|---|---|---|
| .. | ||
| memory | ||
| README.md | ||
| __init__.py | ||
README.md
Hardware Microbenchmarks
GPU hardware characterization benchmarks that produce calibration factors for tileops/perf/profiles/.
Prerequisites
- NVIDIA GPU with CUDA toolkit (
nvccin PATH) - TileOPs installed (
pip install -e .from project root) - Root/sudo access for clock locking (recommended)
HBM Bandwidth
Measures peak HBM bandwidth using vectorized CUDA kernels (float4 load/store) with cudaEvent timing. The calibration factor is derived from the STREAM Triad kernel (a[i] = b[i] + s*c[i], 2 reads + 1 write), the industry-standard pattern for roofline bandwidth calibration.
Triad's 2:1 read:write ratio is closer to real compute kernels than pure copy (1:1), which suffers worst-case HBM bus turnaround overhead. Copy, read-only, and write-only results are included as reference measurements.
Lock GPU clocks (recommended)
GPU boost clocks fluctuate during benchmarks. Lock memory and SM clocks to their maximum for stable, reproducible results:
# Lock clocks (requires root/sudo)
sudo nvidia-smi -lgc $(nvidia-smi --query-gpu=clocks.max.sm --format=csv,noheader,nounits)
sudo nvidia-smi -lmc $(nvidia-smi --query-gpu=clocks.max.mem --format=csv,noheader,nounits)
After benchmarking, reset to default:
sudo nvidia-smi -rgc
sudo nvidia-smi -rmc
Run
Run from project root:
python benchmarks/hardware/memory/hbm_bandwidth.py --profile h200 --arch sm_90
Options:
| Flag | Default | Description |
|---|---|---|
--profile |
h200 |
GPU profile name (reads theoretical peak from tileops/perf/profiles/) |
--arch |
sm_90 |
CUDA compute capability for nvcc |
--size-mb |
2048 |
Working set size in MB |
Output
Measured peak (triad vec4): 4070.44 GB/s
Theoretical: 4800.0 GB/s
Calibration: 0.8480
Update tileops/perf/profiles/h200.yaml:
hbm.calibration: 0.8480
Methodology
- Calibration kernel: STREAM Triad
a[i] = b[i] + s*c[i](2 reads + 1 write,float4vectorized) - Reference kernels: Copy (1:1 read:write), Read-only, Write-only
- Timing:
cudaEvent(GPU-side, no host overhead) - Warmup: 100 iterations per config (ensures boost clocks stabilize)
- Measurement: 200 iterations × 5 runs, report best and median
- Working set: 2 GB default (>> L2 cache, ensures HBM is measured)
- Calibration source: best Triad bandwidth across block size sweep (128/256/512)
Adding a new GPU profile
- Create
tileops/perf/profiles/<gpu>.yamlwith theoretical specs from the datasheet - Lock GPU clocks (see above)
- Run
python benchmarks/hardware/memory/hbm_bandwidth.py --profile <gpu> --arch <sm_XX> - Update
<gpu>.yamlwith the measured calibration factor - Reset GPU clocks