composable_kernel
composable_kernel copied to clipboard
Dlejeune/ck tile 2d multiple reductions
Proposed changes
Implementation of multi reduce ops, both in a threadwise or multiblock (aka blockwise) fashion. It migrate most of the feature already present in old CK with a few exception:
- The alpha and beta terms have not been added. The examples in the old CK are using with values 1 and 0, respectively. As this feature is very easy to implement I suggest we wait the need to have them brought in.
Also some notable limitations are to be noted:
- The multi block version makes use of atomic operation when performing the inter-block reduction/update. While it works for the atomic add for a reasonable collection of types and thread tile size, the other atomic operations (such as MAX) is pretty limited (e.g. no fp16, tile size of 1). While these limitations could be improved on, I suggest it to be part of another PR, if necessary
- Unit testing. While some tests and examples are present it is worth noting that the GPU and CPU reference outputs experience a discrepancy making it difficult to match the two. The nature of the reduction (stochastic order, atomic operation behavior, down casting) make small error to accumulate on the GPU side compared to the deterministic, sequential execution of the CPU reference. Testing large reduction size makes it very apparent. Unfortunately I haven't been able to mitigate much these issues: a generous error tolerance (absolute tolerance of 0.1, relative tolerance of 0.01) has been used, along with small to moderate input sizes.
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.
- [X] I have added tests relevant to the introduced functionality, and the unit tests are passing locally
- [ ] I have added the test to REGRESSION_TESTS list defined at the top of CMakeLists.txt in tests/CMakeLists.txt, IF the test takes more than 30 seconds to run.
- [X] I have added inline documentation which enables the maintainers with understanding the motivation
- [ ] 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-formaton all changed files - [X] Any dependent changes have been merged