pytorch - 💡(How to fix) Fix bf16 matmul throughput drops on RTX 4090 for shapes where N%16==8 [1 participants]

Official PRs (…)
ON THIS PAGE

Recommended Tools

×6

Utilities matched from this issue’s tags and category — try them while you read without losing context.

GitHub issue graph ai analysis

Paste a GitHub issue URL. We fetch that issue, discover linked issues from bodies/comments/timeline, collect linked pull requests, and produce a structured English report.

The report is written in English Markdown for sharing and archival.

Helpful · Quick feedback

Loading…
GitHub stats
pytorch/pytorch#181718Fetched 2026-04-29 06:11:16
View on GitHub
Comments
0
Participants
1
Timeline
165
Reactions
0
Participants
Timeline (top)
mentioned ×78subscribed ×78labeled ×8added_to_project_v2 ×1

Fix Action

Fix / Workaround

On RTX 4090, cuBLAS dispatches an inefficient kernel for bf16 square matmuls in most cases when N % 16 == 8. As shown in the diagram, these matmuls often have the lowest throughput. Some shapes (e.g. N=2040) aren’t affected and get a faster kernel.

import torch, matplotlib.pyplot as plt, matplotlib.patches as mpatches
from torch.profiler import profile, ProfilerActivity

fig, ax = plt.subplots(figsize=(14, 5))
ax.scatter(xs_gray, ys_gray, c="steelblue", s=2, alpha=0.3)
ax.scatter(xs_bad,  ys_bad,  c="crimson",   s=6, alpha=0.9)
ax.scatter(xs_good, ys_good, c="limegreen", s=6, alpha=0.9)
ax.set_xlabel("N  (N×N @ N×N  bfloat16  torch.mm)")
ax.set_ylabel("Throughput (TF/s)")
ax.set_title("bf16 matmul throughput")
ax.legend(handles=[
    mpatches.Patch(color="steelblue", alpha=0.5, label="N%16 != 8"),
    mpatches.Patch(color="crimson",   label="N%16 == 8, 64×64 kernel (slow)"),
    mpatches.Patch(color="limegreen", label="N%16 == 8, 128×128 kernel (fast)"),
], fontsize=8.5, loc="lower right")
plt.tight_layout()
plt.savefig("matmul_kernel_selection.png", dpi=150)

Code Example

import torch, matplotlib.pyplot as plt, matplotlib.patches as mpatches
from torch.profiler import profile, ProfilerActivity

device, dtype = "cuda", torch.bfloat16

def bench(N, warmup=20, iters=200):
    a = torch.randn(N, N, device=device, dtype=dtype)
    b = torch.randn(N, N, device=device, dtype=dtype)
    for _ in range(warmup): torch.mm(a, b)
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(iters): torch.mm(a, b)
    e.record(); torch.cuda.synchronize()
    return 2 * N**3 / (s.elapsed_time(e) / iters * 1e-3) / 1e12

def get_kernel(N):
    a = torch.randn(N, N, device=device, dtype=dtype)
    b = torch.randn(N, N, device=device, dtype=dtype)
    with profile(activities=[ProfilerActivity.CUDA]) as prof:
        torch.mm(a, b)
    torch.cuda.synchronize()
    evts = prof.events()
    return max(evts, key=lambda e: e.device_time_total).name

all_Ns = list(range(1, 4097))
tf = {N: bench(N) for N in all_Ns}
kern = {N: get_kernel(N) for N in all_Ns}

xs_gray, ys_gray, xs_bad, ys_bad, xs_good, ys_good = [], [], [], [], [], []
for N in all_Ns:
    if N % 16 != 8:
        xs_gray.append(N); ys_gray.append(tf[N])
    elif "64x64" in kern[N]:
        xs_bad.append(N);  ys_bad.append(tf[N])
    else:
        xs_good.append(N); ys_good.append(tf[N])

fig, ax = plt.subplots(figsize=(14, 5))
ax.scatter(xs_gray, ys_gray, c="steelblue", s=2, alpha=0.3)
ax.scatter(xs_bad,  ys_bad,  c="crimson",   s=6, alpha=0.9)
ax.scatter(xs_good, ys_good, c="limegreen", s=6, alpha=0.9)
ax.set_xlabel("N  (N×N @ N×N  bfloat16  torch.mm)")
ax.set_ylabel("Throughput (TF/s)")
ax.set_title("bf16 matmul throughput")
ax.legend(handles=[
    mpatches.Patch(color="steelblue", alpha=0.5, label="N%16 != 8"),
    mpatches.Patch(color="crimson",   label="N%16 == 8, 64×64 kernel (slow)"),
    mpatches.Patch(color="limegreen", label="N%16 == 8, 128×128 kernel (fast)"),
], fontsize=8.5, loc="lower right")
plt.tight_layout()
plt.savefig("matmul_kernel_selection.png", dpi=150)

for N in [2040, 2055, 2056, 2057, 2071, 2072, 2073]:
    print(f"  N={N}: {tf[N]:6.1f} TF/s  {kern[N]}")
RAW_BUFFERClick to expand / collapse

🐛 Describe the bug

On RTX 4090, cuBLAS dispatches an inefficient kernel for bf16 square matmuls in most cases when N % 16 == 8. As shown in the diagram, these matmuls often have the lowest throughput. Some shapes (e.g. N=2040) aren’t affected and get a faster kernel.

<img width="1568" height="556" alt="Image" src="https://github.com/user-attachments/assets/0b085a35-134a-4eba-a0ae-345c971e2fcb" />

Examples

NTF/sKernel
2040155.9ampere_bf16_s1688gemm_bf16_128x128_ldg8_f2f_stages_32x1_nn
2055115.6cutlass::Kernel2<cutlass_75_tensorop_bf16_s1688gemm_bf16_128x128_nn_align1>
205689.6cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x64_32x6_nn_align8>
2057106.8cutlass::Kernel2<cutlass_75_tensorop_bf16_s1688gemm_bf16_128x128_nn_align1>
2071110.9cutlass::Kernel2<cutlass_75_tensorop_bf16_s1688gemm_bf16_128x128_nn_align1>
207283.6cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x64_32x6_nn_align8>
2073112.0cutlass::Kernel2<cutlass_75_tensorop_bf16_s1688gemm_bf16_128x128_nn_align1>

Repro

import torch, matplotlib.pyplot as plt, matplotlib.patches as mpatches
from torch.profiler import profile, ProfilerActivity

device, dtype = "cuda", torch.bfloat16

def bench(N, warmup=20, iters=200):
    a = torch.randn(N, N, device=device, dtype=dtype)
    b = torch.randn(N, N, device=device, dtype=dtype)
    for _ in range(warmup): torch.mm(a, b)
    s, e = torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)
    s.record()
    for _ in range(iters): torch.mm(a, b)
    e.record(); torch.cuda.synchronize()
    return 2 * N**3 / (s.elapsed_time(e) / iters * 1e-3) / 1e12

def get_kernel(N):
    a = torch.randn(N, N, device=device, dtype=dtype)
    b = torch.randn(N, N, device=device, dtype=dtype)
    with profile(activities=[ProfilerActivity.CUDA]) as prof:
        torch.mm(a, b)
    torch.cuda.synchronize()
    evts = prof.events()
    return max(evts, key=lambda e: e.device_time_total).name

all_Ns = list(range(1, 4097))
tf = {N: bench(N) for N in all_Ns}
kern = {N: get_kernel(N) for N in all_Ns}

xs_gray, ys_gray, xs_bad, ys_bad, xs_good, ys_good = [], [], [], [], [], []
for N in all_Ns:
    if N % 16 != 8:
        xs_gray.append(N); ys_gray.append(tf[N])
    elif "64x64" in kern[N]:
        xs_bad.append(N);  ys_bad.append(tf[N])
    else:
        xs_good.append(N); ys_good.append(tf[N])

fig, ax = plt.subplots(figsize=(14, 5))
ax.scatter(xs_gray, ys_gray, c="steelblue", s=2, alpha=0.3)
ax.scatter(xs_bad,  ys_bad,  c="crimson",   s=6, alpha=0.9)
ax.scatter(xs_good, ys_good, c="limegreen", s=6, alpha=0.9)
ax.set_xlabel("N  (N×N @ N×N  bfloat16  torch.mm)")
ax.set_ylabel("Throughput (TF/s)")
ax.set_title("bf16 matmul throughput")
ax.legend(handles=[
    mpatches.Patch(color="steelblue", alpha=0.5, label="N%16 != 8"),
    mpatches.Patch(color="crimson",   label="N%16 == 8, 64×64 kernel (slow)"),
    mpatches.Patch(color="limegreen", label="N%16 == 8, 128×128 kernel (fast)"),
], fontsize=8.5, loc="lower right")
plt.tight_layout()
plt.savefig("matmul_kernel_selection.png", dpi=150)

for N in [2040, 2055, 2056, 2057, 2071, 2072, 2073]:
    print(f"  N={N}: {tf[N]:6.1f} TF/s  {kern[N]}")

Versions

Collecting environment information... PyTorch version: 2.13.0.dev20260426+cu130 Is debug build: False CUDA used to build PyTorch: 13.0 ROCM used to build PyTorch: N/A

OS: Ubuntu 24.04.3 LTS (x86_64) GCC version: (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0 Clang version: 18.1.3 (1ubuntu1) CMake version: version 3.28.3 Libc version: glibc-2.39

Python version: 3.12.3 (main, Mar 23 2026, 19:04:32) [GCC 13.3.0] (64-bit runtime) Python platform: Linux-6.14.0-36-generic-x86_64-with-glibc2.39 Is CUDA available: True CUDA runtime version: 12.0.140 CUDA_MODULE_LOADING set to: GPU models and configuration: GPU 0: NVIDIA GeForce RTX 4090 Nvidia driver version: 580.95.05 cuDNN version: Could not collect Is XPU available: False HIP runtime version: N/A MIOpen runtime version: N/A Is XNNPACK available: True Caching allocator config: N/A

CPU: Architecture: x86_64 CPU op-mode(s): 32-bit, 64-bit Address sizes: 46 bits physical, 48 bits virtual Byte Order: Little Endian CPU(s): 32 On-line CPU(s) list: 0-31 Vendor ID: GenuineIntel Model name: 13th Gen Intel(R) Core(TM) i9-13900K CPU family: 6 Model: 183 Thread(s) per core: 2 Core(s) per socket: 24 Socket(s): 1 Stepping: 1 CPU(s) scaling MHz: 92% CPU max MHz: 5800.0000 CPU min MHz: 800.0000 BogoMIPS: 5990.40 Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault epb ssbd ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid rdseed adx smap clflushopt clwb intel_pt sha_ni xsaveopt xsavec xgetbv1 xsaves split_lock_detect user_shstk avx_vnni dtherm ida arat pln pts hwp hwp_notify hwp_act_window hwp_epp hwp_pkg_req hfi vnmi umip pku ospke waitpkg gfni vaes vpclmulqdq rdpid movdiri movdir64b fsrm md_clear serialize pconfig arch_lbr ibt flush_l1d arch_capabilities Virtualization: VT-x L1d cache: 896 KiB (24 instances) L1i cache: 1.3 MiB (24 instances) L2 cache: 32 MiB (12 instances) L3 cache: 36 MiB (1 instance) NUMA node(s): 1 NUMA node0 CPU(s): 0-31 Vulnerability Gather data sampling: Not affected Vulnerability Ghostwrite: Not affected Vulnerability Indirect target selection: Not affected Vulnerability Itlb multihit: Not affected Vulnerability L1tf: Not affected Vulnerability Mds: Not affected Vulnerability Meltdown: Not affected Vulnerability Mmio stale data: Not affected Vulnerability Reg file data sampling: Mitigation; Clear Register File Vulnerability Retbleed: Not affected Vulnerability Spec rstack overflow: Not affected Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl Vulnerability Spectre v1: Mitigation; usercopy/swapgs barriers and __user pointer sanitization Vulnerability Spectre v2: Mitigation; Enhanced / Automatic IBRS; IBPB conditional; PBRSB-eIBRS SW sequence; BHI BHI_DIS_S Vulnerability Srbds: Not affected Vulnerability Tsa: Not affected Vulnerability Tsx async abort: Not affected Vulnerability Vmscape: Mitigation; IBPB before exit to userspace

Versions of relevant libraries: [pip3] numpy==2.4.4 [pip3] nvidia-cublas==13.1.1.3 [pip3] nvidia-cuda-cupti==13.0.85 [pip3] nvidia-cuda-nvrtc==13.0.88 [pip3] nvidia-cuda-runtime==13.0.96 [pip3] nvidia-cudnn-cu13==9.20.0.48 [pip3] nvidia-cufft==12.0.0.61 [pip3] nvidia-curand==10.4.0.35 [pip3] nvidia-cusolver==12.0.4.66 [pip3] nvidia-cusparse==12.6.3.3 [pip3] nvidia-cusparselt-cu13==0.8.1 [pip3] nvidia-nccl-cu13==2.29.7 [pip3] nvidia-nvjitlink==13.0.88 [pip3] nvidia-nvtx==13.0.85 [pip3] torch==2.13.0.dev20260426+cu130 [pip3] torchvision==0.27.0.dev20260426+cu130 [pip3] triton==3.7.0+git88b227e2 [conda] Could not collect

cc @jerryzh168 @ptrblck @msaroufim @eqy @tinglvv @nWEIdia @csarofeen @jianyuh @nikitaved @mruberry @walterddr @xwang233 @Lezcano

extent analysis

TL;DR

The most likely fix for the inefficient kernel dispatch issue on RTX 4090 for bf16 square matmuls is to pad the matrix size to avoid the condition N % 16 == 8.

Guidance

  • Investigate the cuBLAS kernel selection logic to understand why the inefficient kernel is being dispatched for certain matrix sizes.
  • Consider padding the matrix size to a multiple of 16 to avoid the condition N % 16 == 8 and potentially improve throughput.
  • Verify the kernel selection and throughput for different matrix sizes using the provided benchmarking code.
  • Check for any updates or patches to the cuBLAS library that may address this issue.

Example

No code example is provided as the issue is related to the underlying cuBLAS kernel selection and not a specific code snippet.

Notes

The issue appears to be specific to the RTX 4090 GPU and the cuBLAS library, and may not be reproducible on other hardware or software configurations.

Recommendation

Apply a workaround by padding the matrix size to avoid the condition N % 16 == 8, as this may improve throughput for bf16 square matmuls on RTX 4090.

Vote matrix · Quick signals

Works
Did the solution work? Tap to confirm.
Easy Fix
Was it a quick fix?
Time Saver
Did it save you time?
Blocking
Was it severely blocking?
Common Issue
Are others likely hitting this too?
Flaky / Intermittent
Is it intermittent?
Verified / Reproducible
Can you reproduce it reliably?
Loading…

Still need to ship something?

×6

Another batch ranked right after the header list — different links, same matching logic.

Back to top recommendations

TRENDING

pytorch - 💡(How to fix) Fix bf16 matmul throughput drops on RTX 4090 for shapes where N%16==8 [1 participants]