---
id: pr-tilelang-2165
repo: tile-ai/tilelang
pr: 2165
title: '[Refactor] Move backend-specific GEMM implementations and transforms into backend directories'
author: LeiWang1999
date: '2026-05-08'
url: https://github.com/tile-ai/tilelang/pull/2165
source_category: upstream-code
architectures:
- sm90
tags:
- benchmark kernel
- cuda kernel
- fp8
- gemm
- hopper
- matmul
techniques: []
hardware_features:
- fp8
- hopper
kernel_types:
- gemm
languages:
- cuda-cpp
captured_at: '2026-05-16'
status: merged
merge_sha: a797e5120f347e4460ccb188eb84fb8d017c90b4
inclusion_reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.
changed_paths:
- benchmark/matmul/benchmark_matmul_intrinsic.py
- docs/deeplearning_operators/matmul.md
- examples/bitnet-1.58b/kernel_benchmark/tilelang_bitnet_158_int8xint2_prefill.py
- examples/dequantize_gemm/example_dequant_gemm_fine_grained.py
- examples/gemm/README.md
- examples/gemm/example_gemm_intrinsics.py
- examples/gemm_fp8/example_tilelang_gemm_amd_fp8_preshuffle.py
- examples/gemm_fp8/example_tilelang_gemm_fp8_intrinsic.py
- examples/gemm_fp8/regression_example_gemm_fp8.py
- examples/gemm_fp8/test_example_gemm_fp8.py
- examples/hadamard_transform/example_hadamard.py
- examples/plot_layout/README.md
- examples/plot_layout/fragment_mfma_load_a.py
- examples/plot_layout/fragment_mma_load_a.py
- src/backend/cuda/CMakeLists.txt
- src/backend/cuda/runtime.cc
- src/backend/cuda/runtime.h
- src/transform/lower_hopper_intrin.cc
- src/transform/lower_l2_persistent_annotation.cc
- src/transform/persist_threadblock.cc
- testing/python/amd/test_tilelang_gemm_mfma_intrinsic.py
- testing/python/amd/test_tilelang_gemm_mfma_preshuffle.py
- testing/python/kernel/test_tilelang_kernel_bf16_gemm_mma.py
- testing/python/kernel/test_tilelang_kernel_element_wise_add.py
- testing/python/kernel/test_tilelang_kernel_fp8_gemm_mma.py
- testing/python/kernel/test_tilelang_kernel_gemm_mma_intrinsic.py
- testing/python/kernel/test_tilelang_kernel_gemm_simt.py
- testing/python/language/test_tilelang_language_reshape.py
- testing/python/language/test_tilelang_language_vectorize.py
- testing/python/tilelibrary/test_tilelang_tilelibrary_gemm_sp.py
- testing/python/tilelibrary/test_tilelang_tilelibrary_gemm_sp_v2.py
- testing/python/transform/test_tilelang_transform_lower_hopper_intrin.py
- tilelang/__init__.py
- tilelang/backend/__init__.py
- tilelang/cpu/__init__.py
- tilelang/cpu/op/__init__.py
- tilelang/cpu/op/gemm/__init__.py
- tilelang/cpu/op/gemm/gemm_scalar.py
- tilelang/cuda/__init__.py
- tilelang/cuda/intrinsics/__init__.py
- tilelang/cuda/intrinsics/layout/__init__.py
- tilelang/cuda/intrinsics/layout/mma_layout.py
- tilelang/cuda/intrinsics/layout/mma_sm70_layout.py
- tilelang/cuda/intrinsics/layout/mma_sp_layout.py
- tilelang/cuda/intrinsics/layout/utils.py
- tilelang/cuda/intrinsics/macro/__init__.py
- tilelang/cuda/intrinsics/macro/mma_macro_generator.py
- tilelang/cuda/intrinsics/macro/mma_sm70_macro_generator.py
- tilelang/cuda/intrinsics/macro/mma_sp_macro_generator.py
- tilelang/cuda/intrinsics/macro/tcgen05_macro_generator.py
- tilelang/cuda/intrinsics/macro/wgmma_macro_generator.py
- tilelang/cuda/op/__init__.py
- tilelang/cuda/op/gemm/__init__.py
- tilelang/cuda/op/gemm/gemm_mma.py
- tilelang/cuda/op/gemm/gemm_mma_sm70.py
- tilelang/cuda/op/gemm/gemm_tcgen05.py
- tilelang/cuda/op/gemm/gemm_wgmma.py
- tilelang/cuda/op/gemm_sp/__init__.py
- tilelang/cuda/op/gemm_sp/gemm_sp_mma.py
- tilelang/cuda/transform/__init__.py
- tilelang/engine/phase.py
- tilelang/intrinsics/__init__.py
- tilelang/language/gemm_op.py
- tilelang/rocm/__init__.py
- tilelang/rocm/intrinsics/__init__.py
- tilelang/rocm/intrinsics/mfma_layout.py
- tilelang/rocm/intrinsics/mfma_macro_generator.py
- tilelang/rocm/intrinsics/utils.py
- tilelang/rocm/intrinsics/wmma_layout.py
- tilelang/rocm/intrinsics/wmma_macro_generator.py
- tilelang/rocm/op/__init__.py
- tilelang/rocm/op/gemm/__init__.py
- tilelang/rocm/op/gemm/gemm_mfma.py
- tilelang/rocm/op/gemm/gemm_wmma.py
- tilelang/tileop/gemm/__init__.py
- tilelang/tileop/gemm/registry.py
- tilelang/tileop/gemm_sp/__init__.py
- tilelang/tileop/gemm_sp/registry.py
- tilelang/transform/__init__.py
artifact_dir: evidence/pull-bundles/tilelang/gh-2165
head_ref_oid: 9752f55b143c22b050242266ebf307f9c5523b14
---
# [Refactor] Move backend-specific GEMM implementations and transforms into backend directories

- Repository: `tile-ai/tilelang`
- PR: [2165](https://github.com/tile-ai/tilelang/pull/2165)
- Decision: `Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.`
- Reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.

## Changed Paths

- `benchmark/matmul/benchmark_matmul_intrinsic.py`
- `docs/deeplearning_operators/matmul.md`
- `examples/bitnet-1.58b/kernel_benchmark/tilelang_bitnet_158_int8xint2_prefill.py`
- `examples/dequantize_gemm/example_dequant_gemm_fine_grained.py`
- `examples/gemm/README.md`
- `examples/gemm/example_gemm_intrinsics.py`
- `examples/gemm_fp8/example_tilelang_gemm_amd_fp8_preshuffle.py`
- `examples/gemm_fp8/example_tilelang_gemm_fp8_intrinsic.py`
- `examples/gemm_fp8/regression_example_gemm_fp8.py`
- `examples/gemm_fp8/test_example_gemm_fp8.py`
- `examples/hadamard_transform/example_hadamard.py`
- `examples/plot_layout/README.md`
- `examples/plot_layout/fragment_mfma_load_a.py`
- `examples/plot_layout/fragment_mma_load_a.py`
- `src/backend/cuda/CMakeLists.txt`
- `src/backend/cuda/runtime.cc`
- `src/backend/cuda/runtime.h`
- `src/transform/lower_hopper_intrin.cc`
- `src/transform/lower_l2_persistent_annotation.cc`
- `src/transform/persist_threadblock.cc`
- `testing/python/amd/test_tilelang_gemm_mfma_intrinsic.py`
- `testing/python/amd/test_tilelang_gemm_mfma_preshuffle.py`
- `testing/python/kernel/test_tilelang_kernel_bf16_gemm_mma.py`
- `testing/python/kernel/test_tilelang_kernel_element_wise_add.py`
- `testing/python/kernel/test_tilelang_kernel_fp8_gemm_mma.py`
- `testing/python/kernel/test_tilelang_kernel_gemm_mma_intrinsic.py`
- `testing/python/kernel/test_tilelang_kernel_gemm_simt.py`
- `testing/python/language/test_tilelang_language_reshape.py`
- `testing/python/language/test_tilelang_language_vectorize.py`
- `testing/python/tilelibrary/test_tilelang_tilelibrary_gemm_sp.py`
- `testing/python/tilelibrary/test_tilelang_tilelibrary_gemm_sp_v2.py`
- `testing/python/transform/test_tilelang_transform_lower_hopper_intrin.py`
- `tilelang/__init__.py`
- `tilelang/backend/__init__.py`
- `tilelang/cpu/__init__.py`
- `tilelang/cpu/op/__init__.py`
- `tilelang/cpu/op/gemm/__init__.py`
- `tilelang/cpu/op/gemm/gemm_scalar.py`
- `tilelang/cuda/__init__.py`
- `tilelang/cuda/intrinsics/__init__.py`
- `tilelang/cuda/intrinsics/layout/__init__.py`
- `tilelang/cuda/intrinsics/layout/mma_layout.py`
- `tilelang/cuda/intrinsics/layout/mma_sm70_layout.py`
- `tilelang/cuda/intrinsics/layout/mma_sp_layout.py`
- `tilelang/cuda/intrinsics/layout/utils.py`
- `tilelang/cuda/intrinsics/macro/__init__.py`
- `tilelang/cuda/intrinsics/macro/mma_macro_generator.py`
- `tilelang/cuda/intrinsics/macro/mma_sm70_macro_generator.py`
- `tilelang/cuda/intrinsics/macro/mma_sp_macro_generator.py`
- `tilelang/cuda/intrinsics/macro/tcgen05_macro_generator.py`
- `tilelang/cuda/intrinsics/macro/wgmma_macro_generator.py`
- `tilelang/cuda/op/__init__.py`
- `tilelang/cuda/op/gemm/__init__.py`
- `tilelang/cuda/op/gemm/gemm_mma.py`
- `tilelang/cuda/op/gemm/gemm_mma_sm70.py`
- `tilelang/cuda/op/gemm/gemm_tcgen05.py`
- `tilelang/cuda/op/gemm/gemm_wgmma.py`
- `tilelang/cuda/op/gemm_sp/__init__.py`
- `tilelang/cuda/op/gemm_sp/gemm_sp_mma.py`
- `tilelang/cuda/transform/__init__.py`
- `tilelang/engine/phase.py`
- `tilelang/intrinsics/__init__.py`
- `tilelang/language/gemm_op.py`
- `tilelang/rocm/__init__.py`
- `tilelang/rocm/intrinsics/__init__.py`
- `tilelang/rocm/intrinsics/mfma_layout.py`
- `tilelang/rocm/intrinsics/mfma_macro_generator.py`
- `tilelang/rocm/intrinsics/utils.py`
- `tilelang/rocm/intrinsics/wmma_layout.py`
- `tilelang/rocm/intrinsics/wmma_macro_generator.py`
- `tilelang/rocm/op/__init__.py`
- `tilelang/rocm/op/gemm/__init__.py`
- `tilelang/rocm/op/gemm/gemm_mfma.py`
- `tilelang/rocm/op/gemm/gemm_wmma.py`
- `tilelang/tileop/gemm/__init__.py`
- `tilelang/tileop/gemm/registry.py`
- `tilelang/tileop/gemm_sp/__init__.py`
- `tilelang/tileop/gemm_sp/registry.py`
- `tilelang/transform/__init__.py`

## Evidence Bundle

`evidence/pull-bundles/tilelang/gh-2165`
