---
id: pr-tilelang-1887
repo: tile-ai/tilelang
pr: 1887
title: '[Refactor] Improve cp.async lowering and add async_copy op'
author: LeiWang1999
date: '2026-03-06'
url: https://github.com/tile-ai/tilelang/pull/1887
source_category: upstream-code
architectures:
- sm100
tags:
- mxfp4
- schedule
techniques: []
hardware_features:
- fp4
- mxfp4
kernel_types: []
languages: []
captured_at: '2026-05-16'
status: merged
merge_sha: 4929ad89ccdb5622c43a42878beeb3438c77188e
inclusion_reason: Search-matched CUDA kernel optimization PR; materialize diff/source before synthesis.
changed_paths:
- docs/programming_guides/instructions.md
- docs/programming_guides/language_basics.md
- examples/deepseek_mhc/example_mhc_post.py
- examples/deepseek_mhc/example_mhc_pre.py
- examples/deepseek_mhc/regression_example_mhc.py
- 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/dequantize_gemm/example_dequant_gemm_bf16_fp4_hopper.py
- examples/gdn/test_example_gdn_compilation.py
- examples/gemm_fp8/example_tilelang_gemm_fp8_intrinsic.py
- examples/minference/example_vertical_slash_sparse_attn.py
- src/op/builtin.cc
- src/op/builtin.h
- src/op/copy.cc
- src/op/copy.h
- src/op/operator.h
- src/op/utils.h
- src/target/codegen_cutedsl.cc
- src/transform/annotate_warp_group_reg_alloc.cc
- src/transform/flatten_buffer.cc
- src/transform/inject_pipeline.cc
- src/transform/inject_ptx_async_copy.cc
- src/transform/inject_tma_barrier.cc
- src/transform/layout_inference.cc
- src/transform/legalize_safe_memory_access.cc
- src/transform/loop_vectorize.cc
- src/transform/lower_ldg_stg.cc
- src/transform/lower_ptx_async_copy.cc
- src/transform/lower_tile_op.cc
- src/transform/merge_shared_memory_allocations.cc
- src/transform/optimize_cp_async_sync.cc
- src/transform/pipeline_planning.cc
- src/transform/ptx_async_copy_injector.h
- src/transform/thread_storage_sync.cc
- src/transform/vectorize_loop.cc
- src/transform/warp_specialized_rewriter.cc
- testing/python/analysis/test_tilelang_nested_loop_checker.py
- testing/python/kernel/test_tilelang_kernel_int4_gemm_mma.py
- testing/python/language/test_tilelang_language_access_ptr_codegen.py
- testing/python/language/test_tilelang_language_annotate_safe_value.py
- testing/python/language/test_tilelang_language_async_copy_gemm_sm80.py
- testing/python/language/test_tilelang_language_vectorize.py
- testing/python/target/test_tilelang_codegen_cutedsl_cp_async.py
- testing/python/transform/test_tilelang_transform_Inject_software_pipeline.py
- testing/python/transform/test_tilelang_transform_flatten_buffer.py
- testing/python/transform/test_tilelang_transform_inject_tma_barrier.py
- testing/python/transform/test_tilelang_transform_legalize_safe_memory_access.py
- testing/python/transform/test_tilelang_transform_lower_ldgstg.py
- testing/python/transform/test_tilelang_transform_lower_ptx_async_copy.py
- testing/python/transform/test_tilelang_transform_optimize_cp_async_sync.py
- testing/python/transform/test_tilelang_transform_pipeline_planning.py
- tilelang/engine/phase.py
- tilelang/language/__init__.py
- tilelang/language/copy_op.py
- tilelang/language/loop.py
- tilelang/transform/__init__.py
- tilelang/transform/pass_config.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1887
head_ref_oid: 734d3faeb8830233e2464df8c42973c25fdb7216
---
# [Refactor] Improve cp.async lowering and add async_copy op

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

- `docs/programming_guides/instructions.md`
- `docs/programming_guides/language_basics.md`
- `examples/deepseek_mhc/example_mhc_post.py`
- `examples/deepseek_mhc/example_mhc_pre.py`
- `examples/deepseek_mhc/regression_example_mhc.py`
- `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/dequantize_gemm/example_dequant_gemm_bf16_fp4_hopper.py`
- `examples/gdn/test_example_gdn_compilation.py`
- `examples/gemm_fp8/example_tilelang_gemm_fp8_intrinsic.py`
- `examples/minference/example_vertical_slash_sparse_attn.py`
- `src/op/builtin.cc`
- `src/op/builtin.h`
- `src/op/copy.cc`
- `src/op/copy.h`
- `src/op/operator.h`
- `src/op/utils.h`
- `src/target/codegen_cutedsl.cc`
- `src/transform/annotate_warp_group_reg_alloc.cc`
- `src/transform/flatten_buffer.cc`
- `src/transform/inject_pipeline.cc`
- `src/transform/inject_ptx_async_copy.cc`
- `src/transform/inject_tma_barrier.cc`
- `src/transform/layout_inference.cc`
- `src/transform/legalize_safe_memory_access.cc`
- `src/transform/loop_vectorize.cc`
- `src/transform/lower_ldg_stg.cc`
- `src/transform/lower_ptx_async_copy.cc`
- `src/transform/lower_tile_op.cc`
- `src/transform/merge_shared_memory_allocations.cc`
- `src/transform/optimize_cp_async_sync.cc`
- `src/transform/pipeline_planning.cc`
- `src/transform/ptx_async_copy_injector.h`
- `src/transform/thread_storage_sync.cc`
- `src/transform/vectorize_loop.cc`
- `src/transform/warp_specialized_rewriter.cc`
- `testing/python/analysis/test_tilelang_nested_loop_checker.py`
- `testing/python/kernel/test_tilelang_kernel_int4_gemm_mma.py`
- `testing/python/language/test_tilelang_language_access_ptr_codegen.py`
- `testing/python/language/test_tilelang_language_annotate_safe_value.py`
- `testing/python/language/test_tilelang_language_async_copy_gemm_sm80.py`
- `testing/python/language/test_tilelang_language_vectorize.py`
- `testing/python/target/test_tilelang_codegen_cutedsl_cp_async.py`
- `testing/python/transform/test_tilelang_transform_Inject_software_pipeline.py`
- `testing/python/transform/test_tilelang_transform_flatten_buffer.py`
- `testing/python/transform/test_tilelang_transform_inject_tma_barrier.py`
- `testing/python/transform/test_tilelang_transform_legalize_safe_memory_access.py`
- `testing/python/transform/test_tilelang_transform_lower_ldgstg.py`
- `testing/python/transform/test_tilelang_transform_lower_ptx_async_copy.py`
- `testing/python/transform/test_tilelang_transform_optimize_cp_async_sync.py`
- `testing/python/transform/test_tilelang_transform_pipeline_planning.py`
- `tilelang/engine/phase.py`
- `tilelang/language/__init__.py`
- `tilelang/language/copy_op.py`
- `tilelang/language/loop.py`
- `tilelang/transform/__init__.py`
- `tilelang/transform/pass_config.py`

## Evidence Bundle

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