---
id: pr-tilelang-2063
repo: tile-ai/tilelang
pr: 2063
title: '[CUDA] Support int4 `T.gemm`'
author: LeiWang1999
date: '2026-04-20'
url: https://github.com/tile-ai/tilelang/pull/2063
source_category: upstream-code
architectures:
- sm100
tags:
- MLA
- attention_kv
- benchmark_test
- compiler_runtime
- copy lowering
- cuda backend
- fp4
- fp8
- gemm
- gemm_quant
- intrinsic
- matmul
- memory_primitives
techniques:
- fine-grained-quantization
- pipeline-stages
- vectorized-loads
hardware_features:
- fp4
- fp8
kernel_types:
- attention
- compiler
- decode
- gemm
- memory
- mla
languages:
- cuda-cpp
- cute-dsl
- ptx
captured_at: '2026-05-16'
status: merged
merge_sha: 99ee74b3437611ce610f62804a71e0174fed0a44
inclusion_reason: 'Treat compiler/runtime integration as part of the kernel: launch wrapper, cache/JIT behavior, generated code, and build flags need tests.'
changed_paths:
- examples/deepseek_mla/example_mla_decode_ws.py
- examples/deepseek_v32/sparse_mla_fwd_pipelined.py
- examples/deepseek_v32/sparse_mla_fwd_seesaw.py
- examples/gemm_int4/example_tilelang_gemm_int4.py
- src/op/builtin.h
- src/op/copy.cc
- src/target/codegen_cuda.cc
- src/target/codegen_cutedsl.cc
- src/target/codegen_hip.cc
- src/transform/loop_vectorize.cc
- src/transform/lower_ptx_async_copy.cc
- src/transform/merge_shared_memory_allocations.cc
- src/transform/vectorize_loop.cc
- testing/python/kernel/test_tilelang_kernel_int4_gemm.py
- testing/python/kernel/test_tilelang_kernel_int4_gemm_mma.py
- testing/python/language/test_tilelang_language_access_ptr_codegen.py
- tilelang/engine/phase.py
- tilelang/intrinsics/mma_macro_generator.py
- tilelang/language/dtypes.py
- tilelang/language/tir/op.py
- tilelang/tileop/gemm/gemm_mma.py
- tilelang/utils/tensor.py
- testing/python/transform/test_tilelang_transform_lower_ptx_async_copy.py
- testing/python/transform/test_tilelang_transform_lower_tile_op.py
- testing/python/transform/test_tilelang_transform_pipeline_planning.py
artifact_dir: evidence/pull-bundles/tilelang/gh-2063
head_ref_oid: eaca5b749829c2c0bcfa462f3057827664b99fb5
---
# [CUDA] Support int4 `T.gemm`

- Repository: `tile-ai/tilelang`
- PR: [2063](https://github.com/tile-ai/tilelang/pull/2063)
- Decision: `Treat compiler/runtime integration as part of the kernel: launch wrapper, cache/JIT behavior, generated code, and build flags need tests.`
- Reason: Treat compiler/runtime integration as part of the kernel: launch wrapper, cache/JIT behavior, generated code, and build flags need tests.

## Changed Paths

- `examples/deepseek_mla/example_mla_decode_ws.py`
- `examples/deepseek_v32/sparse_mla_fwd_pipelined.py`
- `examples/deepseek_v32/sparse_mla_fwd_seesaw.py`
- `examples/gemm_int4/example_tilelang_gemm_int4.py`
- `src/op/builtin.h`
- `src/op/copy.cc`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_cutedsl.cc`
- `src/target/codegen_hip.cc`
- `src/transform/loop_vectorize.cc`
- `src/transform/lower_ptx_async_copy.cc`
- `src/transform/merge_shared_memory_allocations.cc`
- `src/transform/vectorize_loop.cc`
- `testing/python/kernel/test_tilelang_kernel_int4_gemm.py`
- `testing/python/kernel/test_tilelang_kernel_int4_gemm_mma.py`
- `testing/python/language/test_tilelang_language_access_ptr_codegen.py`
- `tilelang/engine/phase.py`
- `tilelang/intrinsics/mma_macro_generator.py`
- `tilelang/language/dtypes.py`
- `tilelang/language/tir/op.py`
- `tilelang/tileop/gemm/gemm_mma.py`
- `tilelang/utils/tensor.py`
- `testing/python/transform/test_tilelang_transform_lower_ptx_async_copy.py`
- `testing/python/transform/test_tilelang_transform_lower_tile_op.py`
- `testing/python/transform/test_tilelang_transform_pipeline_planning.py`

## Evidence Bundle

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