composable_kernel icon indicating copy to clipboard operation
composable_kernel copied to clipboard

Cross GPU Reduce Operator Initial Development

Open ThomasNing opened this issue 11 months ago • 1 comments

Proposed changes

It is a new Composable Kernel Tile Operator that enables the cross GPU connection reduce on block thread level. There is one (multiple in future) master GPU and one slave GPU that will transfer a block of data (in block thread) from the master to slave GPU. It uses mscclpp to awake the Infinity Fabric of AMD GPU to complete this feature.

Checklist

Please put an x into the boxes that apply. You can also fill these out after creating the PR. If you're not sure, please don't hesitate to ask.

  • [ ] I have added tests relevant to the introduced functionality, and the unit tests are passing locally
  • [x] I have added inline documentation which enables the maintainers with understanding the motivation
  • [x] I have removed the stale documentation which is no longer relevant after this pull request
  • [ ] (If this change is user-facing) I have added release notes which provide the end users with a brief summary of the improvement from this pull request
  • [x] I have run clang-format on all changed files
  • [x] Any dependent changes have been merged

Discussion

If this is a relatively large or complex change, feel free to start a discussion by explaining why you chose the solution you did and what alternatives you considered

ThomasNing avatar Jan 06 '25 16:01 ThomasNing

Great job! It is good to see kick-off of Compute+Communication in CK.

zjing14 avatar Feb 06 '25 02:02 zjing14