[QST]Why cp.async.ca will influence bank conflcit of Shared Store From Global Load ?
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 10.5k
- Forks
- 2.1k
- Avg merge
- 3d 11h
- Merged PRs (30d)
- 7
Description
when I rewrite cutlass conv with cute api, I encountered an interesting phenomena: During global to smem stage，I can't reduce the bank conflict using cute swizzle.After some research, I find the reason is the instruction cp.async.ca when I copy element in Activation to smem. In cutlass, Activation g2s copy will use cp.async.ca, Filter g2s use cp.async.cg, I can understand this. And I do some experiment below to find out why bank conflict can't decrease:
when I always use cp.async.cg in Activation g2s copy and Filter g2s copy with cute swizzle, bank conflict will decrease as i thought. but why cp.async.ca will block this bank conflict reduction？What happens to the hardware?
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 reproducing the reported CUTLASS convolution experiment comparing cp.async.ca and cp.async.cg for Activation and Filter global-to-shared-memory copies with Cute swizzling. Investigate the observed bank-conflict difference and document what hardware behavior explains it; the issue does not mention a specific source file or test.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- performance
- Issue type
- Bug
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100