B300 (SM103) numerical correctness: nvcc 13.x sm_103 cubin produces wrong output for fwd/update kernels

[vai-minzhou]

Summary

causal_conv1d_fwd and causal_conv1d_update produce numerically incorrect output on B300 (SM103, "Blackwell-Ultra") when built from source with CUDA 13.x and the SM103 gencode that setup.py enables for bare_metal_version >= 13.0. Forward output drifts by max_abs_diff ≈ 0.4–0.6 on BF16 (well above any reasonable tolerance), and causal_conv1d_update loses its in-place conv_state update entirely (state error ≈ 0.98). Same source compiled for sm_100 only is bit-exact correct on the same hardware (max_abs_diff = 0).

Repro

import torch, torch.nn.functional as F, causal_conv1d_cuda

torch.manual_seed(42)
x = torch.randn(1, 64, 32, dtype=torch.bfloat16, device="cuda")
w = torch.randn(64, 4,  dtype=torch.bfloat16, device="cuda")
o = torch.zeros_like(x)

causal_conv1d_cuda.causal_conv1d_fwd(x, w, None, None, None, o, None, False)

ref = F.conv1d(F.pad(x, (3, 0)), w.unsqueeze(1), bias=None, groups=64)
print(f"max_abs_diff = {(o.float() - ref.float()).abs().max().item():.4e}")

Build configuration	max_abs_diff	Verdict
Wheel causal-conv1d==1.6.1 (no sm_103 cubin, falls back to sm_100 on B300)	0.0e+00	✅ correct
Source build, CUDA 13.2, default arch list (incl. compute_103,sm_103)	~6.0e-01	❌ wrong
Source build, CUDA 13.2, arch list with only compute_100,sm_100	0.0e+00	✅ correct
Source build, CUDA 13.2, arch list with only compute_103,sm_103	~6.0e-01	❌ wrong

So when both cubins are present in the fatbin, the CUDA loader picks sm_103 on B300 (exact major.minor match), and that cubin produces wrong output. The sm_100 cubin is forward-compatible to sm_103 within the Blackwell family and is correct.

Environment

• GPU: NVIDIA B300 SXM6 AC (compute capability 10.3, "Blackwell-Ultra")
• Driver: 580.x
• CUDA toolkit: 13.2 (V13.2.51)
• nvcc: cuda_13.2.r13.2/compiler.37434383_0
• causal-conv1d: 0d2252d (current main, includes [NVIDIA] Add support Thor, Spark and GB300 #71's SM103 gencode)

cuobjdump --list-elf on the locally-built .so:

causal_conv1d_fwd.sm_80.cubin
causal_conv1d_fwd.sm_89.cubin
causal_conv1d_fwd.sm_90.cubin
causal_conv1d_fwd.sm_100.cubin
causal_conv1d_fwd.sm_103.cubin   ← loader picks this on B300, produces wrong output
causal_conv1d_fwd.sm_120.cubin

Failing test cases

In our test suite (which mirrors the upstream tests/test_causal_conv1d.py at smaller shapes):

fwd_bf16_width4_no_init                                  max_abs_err = 6.05e-01
fwd_bf16_silu_no_init                                    max_abs_err = 4.09e-01
update_bf16_width4                       out_err = 4.47e-01, state_err = 9.79e-01
fwd_bf16_initial_states_silently_ignored_in_channel_first  drift_vs_no_init = 6.22e-01

update_bf16_width4's state_err ≈ 0.98 is particularly damning — conv_state gets shifted left by one and the new token appended, but with sm_103 it looks like the in-place store is being dropped or scrambled.

Workarounds

The simplest workaround is to drop the arch=compute_103,code=sm_103 gencode from setup.py and let the sm_100 cubin handle B300 via Blackwell forward-compatibility. Patch:

--- a/setup.py
+++ b/setup.py
@@ -190,12 +190,9 @@ def get_extensions(...):
         if bare_metal_version >= Version("12.8"):
             cc_flag.append("-gencode")
             cc_flag.append("arch=compute_100,code=sm_100")
             cc_flag.append("-gencode")
             cc_flag.append("arch=compute_120,code=sm_120")
         if bare_metal_version >= Version("13.0"):
-            cc_flag.append("-gencode")
-            cc_flag.append("arch=compute_103,code=sm_103")
             cc_flag.append("-gencode")
             cc_flag.append("arch=compute_110,code=sm_110")
             cc_flag.append("-gencode")
             cc_flag.append("arch=compute_121,code=sm_121")

Verified that this restores correctness on B300 in our full test matrix (fwd_bf16_*, update_bf16_*, causal_conv1d_varlen_*).

Hypothesis

Either nvcc 13.2's SM103 codegen has a bug for this specific kernel (the warp-shuffle / cub::BlockLoad<..., BLOCK_LOAD_WARP_TRANSPOSE> exchange pattern is the most subtle thing in the kernel), or the CUB version shipped with CUDA 13.2 has an SM103 path that miscompiles the WARP_TRANSPOSE algorithm. We didn't dig into PTX/SASS to pin it down — happy to do so if useful.

Next steps

Will open a PR that applies the diff above and adds a B300 test job to CI if you have B300 capacity.