Introduces the new partitioner to implement the reduction StreamK kernel.

[CongMa13]

Proposed changes

Implemented classic StreamK reduction.

Streamk https://arxiv.org/abs/2301.03598

• The important changes are in include/ck_tile/ops/gemm/kernel/streamk_gemm_kernel.hpp
• They implemented the logic in the red box in the following Algorithm5.

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
• 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.
• 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
• I have run clang-format on all changed files
• 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

[Copilot]

Pull Request Overview

This PR implements the classic StreamK reduction strategy for GEMM operations, extending the existing atomic reduction approach. The implementation follows the StreamK algorithm from the referenced paper (https://arxiv.org/abs/2301.03598).

Key changes:

• Added support for the Reduction strategy alongside the existing Atomic strategy
• Implemented partial result accumulation across workgroups with synchronization mechanisms
• Modified kernel APIs to return intermediate results for flexible epilogue execution

Reviewed Changes

Copilot reviewed 7 out of 7 changed files in this pull request and generated 2 comments.

Show a summary per file

File	Description
streamk_gemm_tile_partitioner.hpp	Changed buffer size calculation methods from host-only to host-device accessible
streamk_gemm_tile_partitioner_impl.hpp	Updated function qualifiers to enable device-side workspace size calculations
streamk_gemm_kernel.hpp	Core reduction logic implementation with partial storage, accumulation, and synchronization
kernel_launch.hpp	Added preprocessing support in cold iteration loops
streamk_gemm_basic.cpp	Updated to use new partitioner and handle workspace initialization for reduction strategy
run_gemm_example.inc	Removed unused num_sk_blocks parameter and updated args structure
gemm_utils.hpp	Increased tile sizes for improved performance characteristics

---

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

[ecamartins]

LGTM! Great job!

[cgmillette]

LGTM