---
id: pr-tilelang-1667
repo: tile-ai/tilelang
pr: 1667
title: '[Feature] Support `cp.reduce.async.bulk.tensor`'
author: Rachmanino
date: '2026-01-14'
url: https://github.com/tile-ai/tilelang/pull/1667
source_category: upstream-code
architectures:
- sm90
tags:
- arch_pipeline
- attention
- attention_kv
- benchmark_test
- memory_primitives
- scheduler_autotune
techniques:
- async-copy
- autotuning
- pipeline-stages
- swizzling
- tile-scheduling
hardware_features:
- hopper
- tma
kernel_types:
- attention
- memory
- runtime
- scheduler
languages:
- cuda-cpp
captured_at: '2026-05-16'
status: merged
merge_sha: 732971a8de216099664c4e80ec2879ce26b9d4b5
inclusion_reason: 'Transfer primitive policy carefully: coalescing, shared-memory staging, partial tiles, and determinism are part of the contract.'
changed_paths:
- examples/autodd/tilelang_buggy.py
- examples/autodd/tilelang_minimized_expected.py
- examples/flash_attention/example_gqa_bwd_tma_reduce.py
- examples/flash_attention/example_gqa_bwd_tma_reduce_varlen.py
- src/op/atomic_add.cc
- src/op/atomic_add.h
- src/op/copy.cc
- src/op/utils.cc
- src/op/utils.h
- src/tl_templates/cuda/copy_sm90.h
- testing/python/language/test_tilelang_language_atomic_add.py
artifact_dir: evidence/pull-bundles/tilelang/gh-1667
head_ref_oid: 8f320da14adf5833fb1e253254f4d02db3fb5f44
---
# [Feature] Support `cp.reduce.async.bulk.tensor`

- Repository: `tile-ai/tilelang`
- PR: [1667](https://github.com/tile-ai/tilelang/pull/1667)
- Decision: `Transfer primitive policy carefully: coalescing, shared-memory staging, partial tiles, and determinism are part of the contract.`
- Reason: Transfer primitive policy carefully: coalescing, shared-memory staging, partial tiles, and determinism are part of the contract.

## Changed Paths

- `examples/autodd/tilelang_buggy.py`
- `examples/autodd/tilelang_minimized_expected.py`
- `examples/flash_attention/example_gqa_bwd_tma_reduce.py`
- `examples/flash_attention/example_gqa_bwd_tma_reduce_varlen.py`
- `src/op/atomic_add.cc`
- `src/op/atomic_add.h`
- `src/op/copy.cc`
- `src/op/utils.cc`
- `src/op/utils.h`
- `src/tl_templates/cuda/copy_sm90.h`
- `testing/python/language/test_tilelang_language_atomic_add.py`

## Evidence Bundle

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