Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
76 commits
Select commit Hold shift + click to select a range
17c0288
Add CUDA compute capabilities 8.9, 9.0, 10.0 and 12.0
abadams Jul 29, 2026
398506f
Expose the new CUDA capabilities to the Python bindings
abadams Jul 30, 2026
fac4209
Let the driver choose the CUDA register count
abadams Jul 30, 2026
7aa69ba
Simplify a comment in the cuda runtime
abadams Jul 30, 2026
2c7aa91
Move subtile analysis out of the AMX pass and into MultiRamp
abadams Jul 30, 2026
435bf78
Don't let a cast hide the type AMX multiplies at
abadams Jul 30, 2026
7d3307b
Recognize a transposing shuffle as just another reshaping shuffle
abadams Jul 30, 2026
9a754fc
Merge remote-tracking branch 'origin/main' into abadams/multiramp_sub…
abadams Jul 30, 2026
8e26fb5
Simplify division of a nested vector by a broadcast
abadams Jul 30, 2026
5d72da5
Use cp.async for global to shared copies in the PTX backend
abadams Jul 29, 2026
474331b
Sink redundant GPU thread barriers
abadams Jul 29, 2026
2efd511
Add MemoryType::GPUSharedAsync
abadams Jul 30, 2026
565a68e
Expose the remaining memory types to the Python bindings
abadams Jul 30, 2026
cd865e9
Test the async copy constraints, including the ways to break them
abadams Jul 30, 2026
c7eb195
Only copy asynchronously when the schedule asks for it
abadams Jul 31, 2026
97cd537
Say which requirement an async copy failed to meet
abadams Jul 31, 2026
baa9e04
Test that the destination of an async copy must be dense too
abadams Jul 31, 2026
918c454
Merge branch 'abadams/cuda_capabilities' into abadams/wmma_intrinsics
abadams Jul 31, 2026
2dd6655
Merge branch 'abadams/no_cuda_register_cap' into abadams/wmma_intrinsics
abadams Jul 31, 2026
cff963c
Merge branch 'abadams/multiramp_subtile' into abadams/wmma_intrinsics
abadams Jul 31, 2026
4629bc4
Merge branch 'abadams/simplify_nested_div' into abadams/wmma_intrinsics
abadams Jul 31, 2026
5d5cd81
Merge branch 'abadams/gpu_async_copy' into abadams/wmma_intrinsics
abadams Jul 31, 2026
999bd5f
Add NVIDIA tensor core support via a WMMAAccumulator memory type
abadams Jul 28, 2026
0be212c
Let tensor core accumulators live outside the loop over warps
abadams Jul 28, 2026
0468a09
Give each tensor core accumulator its own shape, and enable the resiz…
abadams Jul 29, 2026
dd2d9b2
Rename WMMAAccumulator to WMMAFragment and let operands be staged in it
abadams Jul 29, 2026
f72d389
Stop tensor core fragments becoming vectors in the PTX backend
abadams Jul 29, 2026
bb7456b
Tell ptxas how large a CUDA thread block is
abadams Jul 29, 2026
adae5de
Retune apps/tensorcore_matmul for the cheaper fragment handling
abadams Jul 29, 2026
8f6e08d
Pad the shared B panel to a stride that actually avoids bank conflicts
abadams Jul 29, 2026
6d6c505
Pick the block shape and its padding by problem size
abadams Jul 30, 2026
dd504b2
Check the AMX error tests reach the error they are named for
abadams Jul 31, 2026
29a287e
Recognize any permutation of the lanes as a reshaping shuffle
abadams Jul 31, 2026
ff07bbd
Merge branch 'abadams/multiramp_subtile' into abadams/wmma_intrinsics
abadams Jul 31, 2026
582b342
Expose WMMAFragment to the Python bindings
abadams Jul 31, 2026
a083b48
Make asynchronous copies a property of the store, not the memory type
abadams Jul 31, 2026
4f74572
Merge remote-tracking branch 'origin/main' into abadams/gpu_async_copy
abadams Jul 31, 2026
b010531
Add the asynchronous copy tests to the CMake build
abadams Jul 31, 2026
a1914d6
Move GPU barrier sinking out to its own branch
abadams Jul 31, 2026
218c95a
Merge the tensorcore matmul app into cuda_mat_mul
abadams Jul 31, 2026
f5967ea
Merge branch 'abadams/gpu_async_copy' into abadams/wmma_intrinsics
abadams Jul 31, 2026
bc1a364
Use Func::in for the matmul operand wrappers
abadams Jul 31, 2026
3b2c1c5
clang-format the tensor core sources
abadams Jul 31, 2026
adf78ff
Benchmark cublas at half precision too
abadams Jul 31, 2026
5143ece
Retune the tensor core block shapes
abadams Jul 31, 2026
473ad08
Record the measured times for all four configurations
abadams Aug 3, 2026
dd75861
Take align_up's argument by const reference
abadams Aug 3, 2026
10fd10d
Support bfloat and eight-bit integer tensor core multiplies
abadams Aug 3, 2026
8573b89
Retune the block shapes for every operand type
abadams Aug 3, 2026
c2d8a11
Benchmark cublas at the other operand and accumulator types
abadams Aug 3, 2026
fc1d0f3
Widen the staged copy for byte operands, and retune them
abadams Aug 3, 2026
a966cc8
Measure the instruction ceilings, and compare against those
abadams Aug 3, 2026
4207b52
Say why the single precision row is the slow one
abadams Aug 3, 2026
ad40d85
Tighten the matmul comments, and fix the stale ones
abadams Aug 3, 2026
15ba249
Break the staging schedules back across lines
abadams Aug 3, 2026
6cb381f
Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams Aug 4, 2026
fa13cb2
Use the benchmark helper for the batched timing
abadams Aug 4, 2026
e39f114
Note that the benchmark data was checked for compressibility
abadams Aug 4, 2026
8ba8604
Hand cublas the same device buffers as the Halide filters
abadams Aug 4, 2026
77022a5
Use Halide's float16 types in the runner, and link libHalide for them
abadams Aug 4, 2026
81be1dc
Record that cublas is asked to write rather than accumulate
abadams Aug 4, 2026
b437f75
Unify tile storage as MemoryType::Tile and add tile scheduling direct…
abadams Aug 4, 2026
1f7875b
Remove the tensorcore_resize app
abadams Aug 4, 2026
5b2d9a6
Fix include order in two files
abadams Aug 4, 2026
f6de42b
Merge remote-tracking branch 'origin/abadams/gpu_async_copy' into aba…
abadams Aug 4, 2026
a0bd1f4
Use the tile directives in the existing tests
abadams Aug 4, 2026
9648d72
Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams Aug 5, 2026
3f101f3
Merge remote-tracking branch 'origin/main' into abadams/gpu_async_copy
abadams Aug 5, 2026
23b8ef1
Merge remote-tracking branch 'origin/abadams/gpu_async_copy' into aba…
abadams Aug 5, 2026
9a51eaa
Pass the lane to the wmma intrinsics
abadams Aug 5, 2026
b60a780
Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams Aug 5, 2026
917e9de
Apply pre-commit auto-fixes
halide-ci[bot] Aug 5, 2026
d4060b2
Merge remote-tracking branch 'origin/main' into abadams/wmma_intrinsics
abadams Aug 5, 2026
b06c693
Use the short form of Load::make in the multiramp test
abadams Aug 5, 2026
e692508
Use the short forms of Load::make and Store::make
abadams Aug 5, 2026
2423dcf
Merge remote-tracking branch 'origin/abadams/wmma_intrinsics' into ab…
abadams Aug 5, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -507,6 +507,7 @@ SOURCE_FILES = \
Error.cpp \
Expr.cpp \
ExtractTileOperations.cpp \
ExtractWMMAOperations.cpp \
FastIntegerDivide.cpp \
FindCalls.cpp \
FindIntrinsics.cpp \
Expand Down Expand Up @@ -709,6 +710,7 @@ HEADER_FILES = \
Extern.h \
ExternFuncArgument.h \
ExtractTileOperations.h \
ExtractWMMAOperations.h \
FastIntegerDivide.h \
FindCalls.h \
FindIntrinsics.h \
Expand Down
47 changes: 44 additions & 3 deletions apps/cuda_mat_mul/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -27,12 +27,53 @@ find_package(Halide REQUIRED)
# Generator
add_halide_generator(mat_mul.generator SOURCES mat_mul_generator.cpp)

# Filters
add_halide_library(mat_mul FROM mat_mul.generator FEATURES cuda cuda_capability_50 PARAMS size=1024)
# Filters. The operand type picks the schedule: half precision gets the tensor
# cores, which need compute capability 7.0 or above, and float gets a schedule
# that accumulates in ordinary registers.
add_halide_library(
mat_mul
FROM mat_mul.generator
FEATURES cuda cuda_capability_50
PARAMS size=1024 A.type=float32 B.type=float32 out.type=float32
)
add_halide_library(
mat_mul_f16
FROM mat_mul.generator
GENERATOR mat_mul
FEATURES cuda cuda_capability_80
PARAMS size=1024 A.type=float16 B.type=float16 out.type=float32
)

add_halide_library(
mat_mul_f16_acc16
FROM mat_mul.generator
GENERATOR mat_mul
FEATURES cuda cuda_capability_80
PARAMS size=1024 A.type=float16 B.type=float16 out.type=float16
)
add_halide_library(
mat_mul_bf16
FROM mat_mul.generator
GENERATOR mat_mul
FEATURES cuda cuda_capability_80
PARAMS size=1024 A.type=bfloat16 B.type=bfloat16 out.type=float32
)
add_halide_library(
mat_mul_u8
FROM mat_mul.generator
GENERATOR mat_mul
FEATURES cuda cuda_capability_80
PARAMS size=1024 A.type=uint8 B.type=uint8 out.type=int32
)

# Main executable
add_executable(runner runner.cpp)
target_link_libraries(runner PRIVATE mat_mul Halide::Tools CUDA::cudart CUDA::cublas)
target_link_libraries(
runner
PRIVATE
mat_mul mat_mul_f16 mat_mul_f16_acc16 mat_mul_bf16 mat_mul_u8 Halide::Halide Halide::Tools
CUDA::cudart CUDA::cublas
)

# Test that the app actually works!
add_test(NAME mat_mul COMMAND runner)
Expand Down
41 changes: 37 additions & 4 deletions apps/cuda_mat_mul/Makefile
Original file line number Diff line number Diff line change
Expand Up @@ -2,11 +2,16 @@ include ../support/Makefile.inc

MATRIX_SIZE ?= 1024

CUDA_SDK ?= /usr/local/cuda-10.0
CUDA_SDK ?= /usr/local/cuda

CXXFLAGS += -I $(CUDA_SDK)/include
LDFLAGS += -L $(CUDA_SDK)/lib64 -Wl,-rpath,$(CUDA_SDK)/lib64

# The float variant runs anywhere, but the half variant is scheduled onto the
# tensor cores, which need compute capability 7.0 or above.
FLOAT_TARGET ?= host-cuda-cuda_capability_50
HALF_TARGET ?= host-cuda-cuda_capability_80

all: $(BIN)/$(HL_TARGET)/runner

$(GENERATOR_BIN)/mat_mul.generator: mat_mul_generator.cpp $(GENERATOR_DEPS)
Expand All @@ -15,11 +20,39 @@ $(GENERATOR_BIN)/mat_mul.generator: mat_mul_generator.cpp $(GENERATOR_DEPS)

$(BIN)/%/mat_mul.a: $(GENERATOR_BIN)/mat_mul.generator
@mkdir -p $(@D)
$^ -g mat_mul -e $(GENERATOR_OUTPUTS) -o $(@D) target=host-cuda-cuda_capability_50 size=$(MATRIX_SIZE)
$^ -g mat_mul -f mat_mul -e $(GENERATOR_OUTPUTS) -o $(@D) \
target=$(FLOAT_TARGET) size=$(MATRIX_SIZE) \
A.type=float32 B.type=float32 out.type=float32

$(BIN)/%/mat_mul_f16.a: $(GENERATOR_BIN)/mat_mul.generator
@mkdir -p $(@D)
$^ -g mat_mul -f mat_mul_f16 -e $(GENERATOR_OUTPUTS) -o $(@D) \
target=$(HALF_TARGET) size=$(MATRIX_SIZE) \
A.type=float16 B.type=float16 out.type=float32

$(BIN)/%/mat_mul_f16_acc16.a: $(GENERATOR_BIN)/mat_mul.generator
@mkdir -p $(@D)
$^ -g mat_mul -f mat_mul_f16_acc16 -e $(GENERATOR_OUTPUTS) -o $(@D) \
target=$(HALF_TARGET) size=$(MATRIX_SIZE) \
A.type=float16 B.type=float16 out.type=float16

$(BIN)/%/mat_mul_bf16.a: $(GENERATOR_BIN)/mat_mul.generator
@mkdir -p $(@D)
$^ -g mat_mul -f mat_mul_bf16 -e $(GENERATOR_OUTPUTS) -o $(@D) \
target=$(HALF_TARGET) size=$(MATRIX_SIZE) \
A.type=bfloat16 B.type=bfloat16 out.type=float32

$(BIN)/%/mat_mul_u8.a: $(GENERATOR_BIN)/mat_mul.generator
@mkdir -p $(@D)
$^ -g mat_mul -f mat_mul_u8 -e $(GENERATOR_OUTPUTS) -o $(@D) \
target=$(HALF_TARGET) size=$(MATRIX_SIZE) \
A.type=uint8 B.type=uint8 out.type=int32

$(BIN)/%/runner: runner.cpp $(BIN)/%/mat_mul.a
$(BIN)/%/runner: runner.cpp $(BIN)/%/mat_mul.a $(BIN)/%/mat_mul_f16.a \
$(BIN)/%/mat_mul_bf16.a $(BIN)/%/mat_mul_u8.a \
$(BIN)/%/mat_mul_f16_acc16.a
@mkdir -p $(@D)
$(CXX) $(CXXFLAGS) -I$(BIN)/$* -Wall $^ -o $@ $(LDFLAGS) -lcudart -lcublas
$(CXX) $(CXXFLAGS) -I$(BIN)/$* -Wall $^ -o $@ $(LDFLAGS) $(LIBHALIDE_LDFLAGS) -lcudart -lcublas

test: $(BIN)/$(HL_TARGET)/runner
$^ $(MATRIX_SIZE)
Expand Down
Loading
Loading