[FEA] make_tma_copy doesnot support shared-to-shared copy
Open
Nobody has claimed this yet.
feature request
inactive-30d
inactive-90d
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
Describe the bug
In PTX, I noticed that
cp.async.bulk.dst.src.completion_mechanism [dstMem], [srcMem], size, [mbar]
.dst = { .shared::cluster }
.src = { .shared::cta }
.completion_mechanism = { .mbarrier::complete_tx::bytes }
supports shared-to-shared copy. But in cute, make_tma_copy has to be global to shared. Why? Could you modify it?
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Start by locating make_tma_copy and reviewing how its source and destination memory spaces are represented. Compare the supported forms with the PTX shared-to-shared cp.async.bulk instruction shown in the issue; done means make_tma_copy supports that shared-to-shared case without regressing existing global-to-shared behavior.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- hpc
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Mostly clear
- Newbie friendliness
- 35/100