feat(tf32): add TF32 TensorCore GEMM kernel achieving 27 TFLOPS#45

Merged

m96-chan merged 23 commits intomainfrom

feature/v0.2.3-tf32-tensorcore

Dec 14, 2025

Commits on Dec 13, 2025

wip(tf32): add TF32 TensorCore GEMM kernel (correctness bug)

m96-chan
and
claude
committed
fix(tf32): WMMA store_matrix_sync alignment bug for N % 8 != 0

m96-chan
and
claude
committed
docs(CLAUDE.md): add kernel development workflow

m96-chan
and
claude
committed
wip(tf32): 2-stage cp.async pipeline with transposed B (correctness bug)

m96-chan
and
claude
committed
wip(tf32): BK=32 kernel with manual epilogue fix

m96-chan
and
claude
committed
wip(tf32): BK=16 with Bs[BN][BK] col-major layout

m96-chan
and
claude
committed
wip(tf32): user rewrite v3 - 51KB smem, 0 spills

m96-chan
and
claude
committed
wip(tf32): G3 kernel - 40KB smem, BK=16

m96-chan
and
claude
committed
wip(tf32): cp.async for both A and B, row_major fragments (6 TFLOPS)

m96-chan
and
claude
committed
wip(tf32): revert to simplified kernel structure

m96-chan
and
claude
committed
wip(tf32): add launch function (correctness broken)

m96-chan
and
claude
committed
wip(tf32): restore 32 TFLOPS kernel (correctness bug)

m96-chan
and
claude
committed
wip(tf32): 44 TFLOPS kernel (correctness still broken)

m96-chan
and
claude
committed
wip(tf32): WMMA row_major×row_major verified working

m96-chan
and
claude
committed
docs(tf32): add WMMA 16x16x8 fragment mapping from dump_fragments

m96-chan
and
claude
committed
fix(tf32): correct C fragment mapping for PTX mma.sync m16n8k8

m96-chan
and
claude
committed
feat(tf32): correct cp.async pipeline achieving 27 TFLOPS

m96-chan
and
claude
committed
docs(readme): add v0.2.3 TF32 benchmark comparison table

m96-chan
and
claude
committed
docs(readme): fix benchmark numbers with actual cuBLAS data

m96-chan
and
claude
committed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

feat(tf32): add TF32 TensorCore GEMM kernel achieving 27 TFLOPS#45

feat(tf32): add TF32 TensorCore GEMM kernel achieving 27 TFLOPS#45
m96-chan merged 23 commits intomainfrom
feature/v0.2.3-tf32-tensorcore

Commits on Dec 13, 2025

wip(tf32): add TF32 TensorCore GEMM kernel (correctness bug)

fix(tf32): WMMA store_matrix_sync alignment bug for N % 8 != 0

docs(CLAUDE.md): add kernel development workflow

wip(tf32): 2-stage cp.async pipeline with transposed B (correctness bug)

wip(tf32): BK=32 kernel with manual epilogue fix

wip(tf32): BK=16 with Bs[BN][BK] col-major layout

wip(tf32): user rewrite v3 - 51KB smem, 0 spills

wip(tf32): G3 kernel - 40KB smem, BK=16

wip(tf32): cp.async for both A and B, row_major fragments (6 TFLOPS)

wip(tf32): revert to simplified kernel structure

wip(tf32): add launch function (correctness broken)

wip(tf32): restore 32 TFLOPS kernel (correctness bug)

wip(tf32): 44 TFLOPS kernel (correctness still broken)

wip(tf32): WMMA row_major×row_major verified working

docs(tf32): add WMMA 16x16x8 fragment mapping from dump_fragments

fix(tf32): correct C fragment mapping for PTX mma.sync m16n8k8

feat(tf32): correct cp.async pipeline achieving 27 TFLOPS

docs(readme): add v0.2.3 TF32 benchmark comparison table

docs(readme): fix benchmark numbers with actual cuBLAS data

Commits on Dec 14, 2025

perf(tf32): optimize kernel with A fragment hoisting (+1.35 TFLOPS)

docs: add TF32 kernel optimization summary

docs: remove optimization summary md (moving to Issue #41)

fix(lint): remove extraneous f-string prefixes