tile-ai / tile-ai/tilelang

[Feature Request] check async copy is used for auto T.Pipelined

Open
#1,201 0 comments 0 reactions 0 assignees View on GitHub
enhancement
Dominant language
Python
Stars
7.4k
Forks
745
Avg merge
1d 1h
Merged PRs (30d)
104

Description

### Required prerequisites

- [x] I have searched the [Issue Tracker](https://github.com/tile-ai/tilelang/issues) that this hasn't already been reported. (comment there if it has.)

### Motivation

Tilelang will use ldg.256+`tl::cp_async_commit()`+`tl::cp_async_wait`, when auto vectorized 256 bits vector, T.Pipelined is opened, and tma&warpspecialized are disabled. The expected behavior is to use 128bits `cp.async.copy`.

### Solution

_No response_

### Alternatives

_No response_

### Additional context

_No response_

Contributor guide

Open the contributing guide

Research direction

Start by locating the auto-vectorization and T.Pipelined code paths that emit ldg.256, tl::cp_async_commit(), and tl::cp_async_wait when TMA and warp specialization are disabled. Trace the generated async-copy width and verify that the 128-bit cp.async.copy behavior is selected instead of the current 256-bit path.

Written by the indexing model from the issue text.

Assessment

Tech stack
python
Domain
compilers
Issue type
Feature
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.