perf(kernel): optimize Blackwell grouped GEMM primitives by spectrometerHBH · Pull Request #19988 · apache/tvm

Commits on Jun 27, 2026

refactor(infra): remove tir-bench agent skill (#2)
spectrometerHBH
committed
feat(lower-tirx): vector PTX ld/st and remove copy_xxb (#6)
spectrometerHBH
committed
fix(arith): add overflow-free uint32/uint64 floormod/floordiv rewrites (#7)
spectrometerHBH
committed
feat(infra): tir-bench helper, rounds, and microsecond timings (#8)
spectrometerHBH
committed
feat(infra): TFLOAT32 descriptor support for copy_async(tma) (#9)
spectrometerHBH
committed
feat(infra): plumb TVM_CUDA_PTXAS_EXTRA_OPTS through nvcc and nvrtc (#10)
spectrometerHBH
committed
feat(op-dispatch): add dense fp8 and tf32 tcgen05 gemm_async paths (#11)
spectrometerHBH
committed
fix(infra): reg-copy R2S/S2R for tcgen05 split-laneid atom layouts (#12)
spectrometerHBH
committed
feat(op): add CUDA elementwise dispatch for maximum tile op (#13)
spectrometerHBH
committed
fix(lower-tirx): elementwise reg slice under sctx.target (B00011) (#14)
spectrometerHBH
committed
refactor(op-dispatch): drop redundant pre-slice canonicalize in ld_stmatrix (#18)
spectrometerHBH
committed
fix(op-dispatch): cache one tcgen05.cp descriptor, patch SMEM addr per cp (#19)
spectrometerHBH
committed
feat(op-dispatch): support uint32 shape extents in TMA copy (#17)
spectrometerHBH
committed
perf(op-dispatch): direct ld/st.shared for shared->shared warp permute (#20)
spectrometerHBH
committed
perf(op-dispatch): per-MMA SMEM descriptor with hoist/recompute in gemm_async (#15)
spectrometerHBH
committed
fix(lower-tirx): support PrimType dtype handling (#21)
spectrometerHBH
authored

Commits on Jul 2, 2026

feat(infra): Triton-standard bench harness with proton default timer (#24)
spectrometerHBH
authored

Commits on Jul 4, 2026

Commits on Jul 5, 2026

feat(layout): buffer dim-surgery views (unflatten/flatten/select/narrow/sub/rearrange) (#28)
spectrometerHBH
authored

Commits on Jul 12, 2026

perf(kernel): optimize Blackwell grouped GEMM primitives
spectrometerHBH
committed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

perf(kernel): optimize Blackwell grouped GEMM primitives#19988

perf(kernel): optimize Blackwell grouped GEMM primitives#19988
spectrometerHBH wants to merge 21 commits into
apache:mainfrom
spectrometerHBH:group_fp8

Commits on Jun 27, 2026

refactor(infra): remove tir-bench agent skill (#2)

feat(lower-tirx): vector PTX ld/st and remove copy_xxb (#6)

fix(arith): add overflow-free uint32/uint64 floormod/floordiv rewrites (#7)

feat(infra): tir-bench helper, rounds, and microsecond timings (#8)

feat(infra): TFLOAT32 descriptor support for copy_async(tma) (#9)

feat(infra): plumb TVM_CUDA_PTXAS_EXTRA_OPTS through nvcc and nvrtc (#10)

feat(op-dispatch): add dense fp8 and tf32 tcgen05 gemm_async paths (#11)

fix(infra): reg-copy R2S/S2R for tcgen05 split-laneid atom layouts (#12)

feat(op): add CUDA elementwise dispatch for maximum tile op (#13)

fix(lower-tirx): elementwise reg slice under sctx.target (B00011) (#14)

refactor(op-dispatch): drop redundant pre-slice canonicalize in ld_stmatrix (#18)

fix(op-dispatch): cache one tcgen05.cp descriptor, patch SMEM addr per cp (#19)

feat(op-dispatch): support uint32 shape extents in TMA copy (#17)

perf(op-dispatch): direct ld/st.shared for shared->shared warp permute (#20)

perf(op-dispatch): per-MMA SMEM descriptor with hoist/recompute in gemm_async (#15)

fix(lower-tirx): support PrimType dtype handling (#21)

Commits on Jul 2, 2026

feat(infra): Triton-standard bench harness with proton default timer (#24)

Commits on Jul 4, 2026

refactor(infra): always build and run all reference impls (#27)

Remove dead async_structs, merge tile-primitive headers, rename Predicate to LambdaExpr (#26)

Commits on Jul 5, 2026

feat(layout): buffer dim-surgery views (unflatten/flatten/select/narrow/sub/rearrange) (#28)

Commits on Jul 12, 2026

perf(kernel): optimize Blackwell grouped GEMM primitives