Hadacore

PyTorch Labs' fast Hadamard transform CUDA kernel, packaged so it installs with uv / pip instead of requiring a manual git clone && python setup.py. The kernel itself is unmodified from upstream; this repo is about build reproducibility.

Original work

Documented on the PyTorch blog: https://pytorch.org/blog/hadacore/ Source: https://github.com/meta-pytorch/applied-ai/tree/main/kernels/cuda/inference/hadamard_transform

Commit history was preserved using git-filter-repo when initializing this repo.

Supported configurations

• GPUs: SM 8.0 (Ampere A100), 8.6 (Ampere GeForce, e.g. RTX 30xx), 8.9 (Ada), 9.0 (Hopper). The kernel requires at least SM 8.0 for cp.async and the bf16/fp16 m16n8k16 tensor-core MMA. Override via TORCH_CUDA_ARCH_LIST if you need a narrower/wider set.
• dtypes: torch.float16 and torch.bfloat16. The kernel is hard-wired to the m16n8k16 MMA family and will refuse other dtypes. Extending to fp32/tf32 or fp8 would be a separate kernel.
• Hadamard sizes: powers of two from 2 up to $2^{15}$ (32768) along the last dimension.

Install

Requires Python 3.12 and a CUDA toolkit (13.0 or newer recommended; see "Known build-environment issue" below if on CUDA 13.1 with glibc ≥ 2.42).

uv sync
uv pip install -e . --no-build-isolation

--no-build-isolation is required because the CUDA extension's setup.py imports torch to generate the right compile flags; the isolated build env would install a fresh (CPU-only) torch and produce a broken .so.

Smoke test

uv run python -c "
import torch, math, scipy.linalg
import faster_hadamard_transform as fht
H = torch.tensor(scipy.linalg.hadamard(1024), device='cuda', dtype=torch.float32) * math.sqrt(1/1024)
x = torch.randn(8, 1024, device='cuda', dtype=torch.float16)
assert torch.allclose((x.float() @ H).half(), fht.hadamard_transform(x.clone(), inplace=True), atol=1e-2)
print('ok')
"

test.py runs a larger correctness grid. It allocates up to 4 GiB of fp32 reference matrices on the GPU for the 32768-size case, so run it with a largely idle GPU.

Known build-environment issue: CUDA 13.1 + glibc ≥ 2.42

glibc 2.42 added noexcept(true) to the declarations of rsqrt/rsqrtf in <bits/mathcalls.h>. CUDA 13.1's crt/math_functions.h declares the same functions without an exception specifier, and nvcc's EDG frontend rejects the mismatch as a hard error. NVIDIA already provides __NV_IEC_60559_FUNCS_EXCEPTION_SPECIFIER and uses it for sinpi/cospi but missed rsqrt/rsqrtf — expected to be fixed upstream in CUDA 13.2.

setup.py detects this case at build time and, instead of patching the system CUDA install, materializes a symlink-farm "shadow" of the toolkit under build/cuda_shadow/ with the single offending header replaced by a patched copy (see scripts/build_cuda_shadow.py). Nothing outside the project tree is modified.

If you don't want this behavior, set CUDA_HOME (and PATH) to point at a CUDA install that has a fixed crt/math_functions.h, or upgrade to a glibc / CUDA combination that doesn't trigger the conflict.

Hardware tested

GPU	Arch	Notes
RTX 3080	sm_86	fp16 + bf16 correctness up to size 8192 verified; size 32768 requires ~4 GiB of free GPU memory for the reference matrix

Additional reports welcome.

Future work

See TODO.md.