Write-only benchmark exceeds 100% bandwidth
Nobody has claimed this yet.
- Dominant language
- Cuda
- Stars
- 927
- Forks
- 123
- Avg merge
- 2d 10h
- Merged PRs (30d)
- 2
Description
@gevtushenko and I have run into this curious case: When benchmarking cub::DeviceTransform::Fill to just fill a buffer with values (see source), we sometimes get results like this:
## generate
### [0] NVIDIA GeForce RTX 5090
| T{ct} | Elements{io} | Samples | CPU Time | Noise | GPU Time | Noise | Elem/s | GlobalMem BW | BWUtil | Samples | Batch GPU |
|-------|------------------|---------|------------|--------|------------|--------|----------|--------------|---------|---------|------------|
| I8 | 2^16 = 65536 | 314x | 9.475 us | 13.30% | 2.836 us | 34.84% | 23.108G | 23.108 GB/s | 1.29% | 230528x | 2.169 us |
| I8 | 2^20 = 1048576 | 414x | 9.200 us | 10.08% | 2.683 us | 34.71% | 390.860G | 390.860 GB/s | 21.81% | 235432x | 2.124 us |
| I8 | 2^24 = 16777216 | 262x | 15.828 us | 4.94% | 8.182 us | 1.94% | 2.050T | 2.050 TB/s | 114.41% | 100812x | 4.961 us |
| I8 | 2^28 = 268435456 | 370x | 166.190 us | 0.79% | 158.758 us | 0.73% | 1.691T | 1.691 TB/s | 94.35% | 3338x | 158.649 us |
Notice the 114.41% BWUtil for the third run. Trying to figure out what happens, I run the benchmark under ncu and the memory workload analysis shows:
It looks like the kernel did not write back the data from L2 to GMEM. This may explain why the kernel could finish faster than it would take to do the full write back to GMEM (which is what I guess nvbench assumes).
ncu shows me that smaller runs write back zero bytes from L2 to GMEM, and the larger 2^28 run writes back everything except about 62MB of data, which remains in L2. The experiment was done on an RTX 5090 with 92MB L2.
It seems nvbench should incorporate a mechanism that forces the benchmarked kernel to do the full write back to GMEM, so bandwidth measurements on writes are accurate.
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
The issue points to cub/benchmarks/bench/transform/fill.cu and nvbench bandwidth calculations; start by reproducing the write benchmark under ncu and comparing L2-to-GMEM writeback. Done means the benchmark reports full global-memory traffic rather than over 100% BWUtil, with the writeback behavior validated on the reported case.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 35/100