tensorcore support - #9283

Draft

abadams wants to merge 76 commits into

mainfrom

abadams/wmma_intrinsics

Commits on Jul 30, 2026

Add CUDA compute capabilities 8.9, 9.0, 10.0 and 12.0

abadams
and
claude
committed
17c0288
Expose the new CUDA capabilities to the Python bindings

abadams
and
claude
committed
398506f
Let the driver choose the CUDA register count

abadams
and
claude
committed
fac4209
Simplify a comment in the cuda runtime

abadams
and
claude
committed
7aa69ba
Move subtile analysis out of the AMX pass and into MultiRamp

abadams
and
claude
committed
2c7aa91
Don't let a cast hide the type AMX multiplies at

abadams
and
claude
committed
435bf78
Recognize a transposing shuffle as just another reshaping shuffle

abadams
and
claude
committed
7d3307b
Merge remote-tracking branch 'origin/main' into abadams/multiramp_subtile
abadams
committed
9a754fc
Simplify division of a nested vector by a broadcast

abadams
and
claude
committed
8e26fb5
Use cp.async for global to shared copies in the PTX backend

abadams
and
claude
committed
5d72da5
Sink redundant GPU thread barriers

abadams
and
claude
committed
474331b
Add MemoryType::GPUSharedAsync

abadams
and
claude
committed
2efd511
Expose the remaining memory types to the Python bindings

abadams
and
claude
committed
565a68e
Test the async copy constraints, including the ways to break them

abadams
and
claude
committed
cd865e9

Commits on Jul 31, 2026

Only copy asynchronously when the schedule asks for it

abadams
and
claude
committed
c7eb195
Say which requirement an async copy failed to meet

abadams
and
claude
committed
97cd537
Test that the destination of an async copy must be dense too

abadams
and
claude
committed
baa9e04
Merge branch 'abadams/cuda_capabilities' into abadams/wmma_intrinsics
abadams
committed
918c454
Merge branch 'abadams/no_cuda_register_cap' into abadams/wmma_intrinsics
abadams
committed
2dd6655
Merge branch 'abadams/multiramp_subtile' into abadams/wmma_intrinsics
abadams
committed
cff963c
Merge branch 'abadams/simplify_nested_div' into abadams/wmma_intrinsics
abadams
committed
4629bc4
Merge branch 'abadams/gpu_async_copy' into abadams/wmma_intrinsics
abadams
committed
5d5cd81
Add NVIDIA tensor core support via a WMMAAccumulator memory type

abadams
and
claude
committed
999bd5f
Let tensor core accumulators live outside the loop over warps

abadams
and
claude
committed
0be212c
Give each tensor core accumulator its own shape, and enable the resize app

abadams
and
claude
committed
0468a09
Rename WMMAAccumulator to WMMAFragment and let operands be staged in it

abadams
and
claude
committed
dd2d9b2
Stop tensor core fragments becoming vectors in the PTX backend

abadams
and
claude
committed
f72d389
Tell ptxas how large a CUDA thread block is

abadams
and
claude
committed
bb7456b
Retune apps/tensorcore_matmul for the cheaper fragment handling

abadams
and
claude
committed
adae5de
Pad the shared B panel to a stride that actually avoids bank conflicts

abadams
and
claude
committed
8f6e08d
Pick the block shape and its padding by problem size

abadams
and
claude
committed
6d6c505
Check the AMX error tests reach the error they are named for

abadams
and
claude
committed
dd504b2
Recognize any permutation of the lanes as a reshaping shuffle

abadams
and
claude
committed
29a287e
Merge branch 'abadams/multiramp_subtile' into abadams/wmma_intrinsics
abadams
committed
ff07bbd
Expose WMMAFragment to the Python bindings

abadams
and
claude
committed
582b342
Make asynchronous copies a property of the store, not the memory type

abadams
and
claude
committed
a083b48
Merge remote-tracking branch 'origin/main' into abadams/gpu_async_copy
abadams
committed
4f74572
Add the asynchronous copy tests to the CMake build

abadams
and
claude
committed
b010531
Move GPU barrier sinking out to its own branch

abadams
and
claude
committed
a1914d6
Merge the tensorcore matmul app into cuda_mat_mul

abadams
and
claude
committed
218c95a
Merge branch 'abadams/gpu_async_copy' into abadams/wmma_intrinsics
abadams
committed
f5967ea
Use Func::in for the matmul operand wrappers

abadams
and
claude
committed
bc1a364
clang-format the tensor core sources

abadams
and
claude
committed
3b2c1c5
Benchmark cublas at half precision too

abadams
and
claude
committed
adf78ff
Retune the tensor core block shapes

abadams
and
claude
committed
5143ece

Commits on Aug 3, 2026

Record the measured times for all four configurations

abadams
and
claude
committed
473ad08
Take align_up's argument by const reference

abadams
and
claude
committed
dd75861
Support bfloat and eight-bit integer tensor core multiplies

abadams
and
claude
committed
10fd10d
Retune the block shapes for every operand type

abadams
and
claude
committed
8573b89
Benchmark cublas at the other operand and accumulator types

abadams
and
claude
committed
c2d8a11
Widen the staged copy for byte operands, and retune them

abadams
and
claude
committed
fc1d0f3
Measure the instruction ceilings, and compare against those

abadams
and
claude
committed
a966cc8
Say why the single precision row is the slow one

abadams
and
claude
committed
4207b52
Tighten the matmul comments, and fix the stale ones

abadams
and
claude
committed
ad40d85
Break the staging schedules back across lines

abadams
and
claude
committed
15ba249

Commits on Aug 4, 2026

Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams
committed
6cb381f
Use the benchmark helper for the batched timing

abadams
and
claude
committed
fa13cb2
Note that the benchmark data was checked for compressibility

abadams
and
claude
committed
e39f114
Hand cublas the same device buffers as the Halide filters

abadams
and
claude
committed
8ba8604
Use Halide's float16 types in the runner, and link libHalide for them

abadams
and
claude
committed
77022a5
Record that cublas is asked to write rather than accumulate

abadams
and
claude
committed
81be1dc
Unify tile storage as MemoryType::Tile and add tile scheduling directives

abadams
and
claude
committed
b437f75
Remove the tensorcore_resize app

abadams
and
claude
committed
1f7875b
Fix include order in two files

abadams
and
claude
committed
5b2d9a6
Merge remote-tracking branch 'origin/abadams/gpu_async_copy' into abadams/wmma_intrinsics
abadams
committed
f6de42b
Use the tile directives in the existing tests

abadams
and
claude
committed
a0bd1f4

Commits on Aug 5, 2026

Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams
committed
9648d72
Merge remote-tracking branch 'origin/main' into abadams/gpu_async_copy
abadams
committed
3f101f3
Merge remote-tracking branch 'origin/abadams/gpu_async_copy' into abadams/wmma_intrinsics
abadams
committed
23b8ef1
Pass the lane to the wmma intrinsics

abadams
and
claude
committed
9a51eaa
Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams
committed
b60a780
Apply pre-commit auto-fixes
halide-ci[bot]
committed
917e9de
Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams
committed
d4060b2
Use the short form of Load::make in the multiramp test

abadams
and
claude
committed
b06c693
Use the short forms of Load::make and Store::make

abadams
and
claude
committed
e692508
Merge remote-tracking branch 'origin/abadams/wmma_intrinsics' into abadams/wmma_intrinsics
abadams
committed
2423dcf