Skip to content

bench(softmax): use profile_with_l2flush for more realistic bandwidth measurements#126

Merged
hannahli-nv merged 1 commit into
NVIDIA:mainfrom
liqiangxl:softmax-profile-fix
May 14, 2026
Merged

bench(softmax): use profile_with_l2flush for more realistic bandwidth measurements#126
hannahli-nv merged 1 commit into
NVIDIA:mainfrom
liqiangxl:softmax-profile-fix

Conversation

@liqiangxl
Copy link
Copy Markdown
Collaborator

Switch from triton.testing.do_bench_cudagraph to profile_with_l2flush which flushes L2 cache between repetitions, giving production-representative numbers instead of inflated warm-cache throughput at small problem sizes.

Perf changes (M=4096, GB/s):

Baseline (no TMA, no chunked):
  N=1024:  CuTile 4277→2745 (-35.8%), PyTorch 5055→3827 (-24.3%)
  N=2048:  CuTile 5383→3427 (-36.3%), PyTorch 4423→3158 (-28.6%)
  N=4096:  CuTile 4010→3866 (-3.6%),  PyTorch 2869→2745 (-4.3%)
  N=8192:  CuTile 4377→4273 (-2.4%),  PyTorch 2339→2242 (-4.1%)
  N=16384: CuTile 3221→3209 (-0.4%),  PyTorch 4199→4116 (-2.0%)

TMA softmax:
  N=1024:  CuTile 2577→2131 (-17.3%), PyTorch 5068→3827 (-24.5%)
  N=2048:  CuTile 3428→2713 (-20.9%), PyTorch 4387→3149 (-28.2%)
  N=4096:  CuTile 3652→3570 (-2.3%),  PyTorch 2878→2740 (-4.8%)
  N=8192:  CuTile 3345→3305 (-1.2%),  PyTorch 2339→2238 (-4.3%)
  N=16384: CuTile 3575→3567 (-0.2%),  PyTorch 4200→4117 (-2.0%)

Chunked softmax:
  N=1024:  CuTile 3927→2325 (-40.8%), PyTorch 5012→3826 (-23.7%)
  N=2048:  CuTile 5242→3410 (-34.9%), PyTorch 4384→3158 (-28.0%)
  N=4096:  CuTile 3998→3852 (-3.6%),  PyTorch 2879→2741 (-4.8%)
  N=8192:  CuTile 4370→4286 (-1.9%),  PyTorch 2339→2240 (-4.2%)
  N=16384: CuTile 3301→3307 (+0.2%),  PyTorch 4199→4113 (-2.1%)

Small N (1024-2048) drops 20-40% as L2 cache warmth is removed; large N (8192+) barely changes since working set already exceeds L2.

Description

CI Configuration

config:
  build: true
  # valid options are "ops", "benchmark", and "sanity"
  test: [benchmark]

Checklist

  • Code formatted and imports sorted via repo specifications (./format.sh)
  • Documentation updated (if needed)
  • CI configuration reviewed

… measurements

Switch from triton.testing.do_bench_cudagraph to profile_with_l2flush which
flushes L2 cache between repetitions, giving production-representative numbers
instead of inflated warm-cache throughput at small problem sizes.

Perf changes (M=4096, GB/s):

Baseline (no TMA, no chunked):
  N=1024:  CuTile 4277→2745 (-35.8%), PyTorch 5055→3827 (-24.3%)
  N=2048:  CuTile 5383→3427 (-36.3%), PyTorch 4423→3158 (-28.6%)
  N=4096:  CuTile 4010→3866 (-3.6%),  PyTorch 2869→2745 (-4.3%)
  N=8192:  CuTile 4377→4273 (-2.4%),  PyTorch 2339→2242 (-4.1%)
  N=16384: CuTile 3221→3209 (-0.4%),  PyTorch 4199→4116 (-2.0%)

TMA softmax:
  N=1024:  CuTile 2577→2131 (-17.3%), PyTorch 5068→3827 (-24.5%)
  N=2048:  CuTile 3428→2713 (-20.9%), PyTorch 4387→3149 (-28.2%)
  N=4096:  CuTile 3652→3570 (-2.3%),  PyTorch 2878→2740 (-4.8%)
  N=8192:  CuTile 3345→3305 (-1.2%),  PyTorch 2339→2238 (-4.3%)
  N=16384: CuTile 3575→3567 (-0.2%),  PyTorch 4200→4117 (-2.0%)

Chunked softmax:
  N=1024:  CuTile 3927→2325 (-40.8%), PyTorch 5012→3826 (-23.7%)
  N=2048:  CuTile 5242→3410 (-34.9%), PyTorch 4384→3158 (-28.0%)
  N=4096:  CuTile 3998→3852 (-3.6%),  PyTorch 2879→2741 (-4.8%)
  N=8192:  CuTile 4370→4286 (-1.9%),  PyTorch 2339→2240 (-4.2%)
  N=16384: CuTile 3301→3307 (+0.2%),  PyTorch 4199→4113 (-2.1%)

Small N (1024-2048) drops 20-40% as L2 cache warmth is removed;
large N (8192+) barely changes since working set already exceeds L2.
@copy-pr-bot
Copy link
Copy Markdown

copy-pr-bot Bot commented May 13, 2026

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@hannahli-nv
Copy link
Copy Markdown
Collaborator

/ok to test ca9c7a7

@hannahli-nv hannahli-nv merged commit f729995 into NVIDIA:main May 14, 2026
17 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants