Skip to content

Conversation

fzyzcjy
Copy link
Contributor

@fzyzcjy fzyzcjy commented Jun 20, 2025

Quickly tested:

before

[layout] Kernel performance: 0.170 ms
[layout] Kernel performance: 0.178 ms

after

[layout] Kernel performance: 0.045 ms
[layout] Kernel performance: 0.046 ms

core change is simple, just change num SM.

The code diff is huge since I inherited code change from #218 but that's mostly unnecessary.

I will try to find out some time later to cleanup the code (eg still use few SM when need two stream overlap, and only use this one when do not enable two stream overlap etc), and here just quickly PR to know whether you think this looks acceptable and if anyone happens to need this then can copy-paste.

@LyricZhao
Copy link
Collaborator

I did some similar change according to your PR in d4f3497 and 77ddb01. Thanks! 👍🏻

Different from your change:

  1. __launch_bounds__ should be removed, as for compute-comm overlapping, if computation launches first and only left, e.g. 20 SMs for comm, we better make all the blocks into 20 SMs together but not by several waves.
  2. The lower bound of kNumRanksPerSM should be 8, it is related to the number of nodes.

@LyricZhao LyricZhao closed this Jul 2, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants