---
id: pr-tilelang-1981
repo: tile-ai/tilelang
pr: 1981
title: '[Feature] Support TMA store in T.tma_copy()'
author: LeiWang1999
date: '2026-03-27'
url: https://github.com/tile-ai/tilelang/pull/1981
source_category: upstream-code
architectures:
- sm100
tags:
- TMA
- arch_pipeline
- benchmark_test
- compiler_runtime
- copy lowering
- cuda backend
- matmul
- memory_primitives
techniques:
- async-copy
- pipeline-stages
hardware_features:
- tma
kernel_types:
- compiler
- gemm
- memory
- runtime
languages:
- cuda-cpp
- cute-dsl
captured_at: '2026-05-16'
status: merged
merge_sha: bdf436d4c3ca1d001936fc1a5c883e60fdd4d3cf
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:
- src/op/atomic_add.cc
- src/op/builtin.cc
- src/op/copy.cc
- src/target/codegen_cuda.cc
- src/target/codegen_cutedsl.cc
- testing/python/language/test_tilelang_language_tma_copy.py
- testing/python/language/test_tilelang_language_tma_store.py
- tilelang/language/builtin.py
- tilelang/language/copy_op.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1981
head_ref_oid: ea6af5b141db1235120d9a9a7a75b10d0a8e25be
---
# [Feature] Support TMA store in T.tma_copy()

- Repository: `tile-ai/tilelang`
- PR: [1981](https://github.com/tile-ai/tilelang/pull/1981)
- 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

- `src/op/atomic_add.cc`
- `src/op/builtin.cc`
- `src/op/copy.cc`
- `src/target/codegen_cuda.cc`
- `src/target/codegen_cutedsl.cc`
- `testing/python/language/test_tilelang_language_tma_copy.py`
- `testing/python/language/test_tilelang_language_tma_store.py`
- `tilelang/language/builtin.py`
- `tilelang/language/copy_op.py`

## Evidence Bundle

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