Addition of Stream-K tests using Tile Engine #3514
Conversation
| {"m": 256, "n": 256, "k": 128, "split_k": 1}, | ||
| {"m": 512, "n": 256, "k": 256, "split_k": 1}, | ||
| {"m": 256, "n": 512, "k": 256, "split_k": 1} |
There was a problem hiding this comment.
Should these sizes cover the existing cases in test/ck_tile/gemm_streamk?
If so, did we want to remove the old test files and replace them with these to avoid adding redundant tests?
arai713
force-pushed
the
arai/ck_tile/streamk_tile_engine_test
branch
2 times, most recently
from
7db8519 to
0e6e6a7
Compare
arai713
requested review from
a team,
ThomasNing,
afagaj,
andriy-ca,
aosewski,
asleepzzz,
bartekxk,
carlushuang,
cgmillette,
coderfeli,
ddembeckAMD,
geyyer,
illsilin,
poyenc,
qianfengz,
shumway,
tenpercent and
vidyasagar-amd
as code owners
| - **SKIPPED**: Kernel validation returned "Arguments not supported" (expected for certain problem sizes/configurations) ⚠️ | ||
| - **FAILED**: Actual error or incorrect computation results ❌ | ||
|
|
||
| When a kernel's `IsSupportedArgument()` check fails (e.g., due to vector alignment requirements, dimension constraints, or padding limitations), the test is automatically skipped rather than failed. This allows comprehensive testing across various problem sizes while gracefully handling configurations that don't meet specific kernel requirements. |
There was a problem hiding this comment.
My opinion is there should never be a set of arguments passed to the test suite that are not supported.
We should have a set of black and white tests that we know will either absolutely pass or absolutely fail.
Because in a scenario, where IsSupportedAruguments()'s implementation is changed, for instance someone unintentionally reduces the vector alignment requirements, it will result in valid kernels just being skipped rather than failing. So, we might miss this regression.
IsSupportedArgument() is pretty fragile in a sense because it calls a lot of getter functions from all over the place.
arai713
force-pushed
the
arai/ck_tile/streamk_tile_engine_test
branch
from
0e6e6a7 to
89e4aba
Compare
|
|
||
| ## Data Type Support | ||
| - ✅ **fp16, bf16**: Fully supported - all layouts (rcr, rrr, ccr, crr) | ||
| - ❌ **fp64**: Not supported (hardware MFMA limitation) |
There was a problem hiding this comment.
fp64 mfma is supported on gfx9 (minus gfx908). Is this more of a CK limitation?
There was a problem hiding this comment.
Pull request overview
This PR introduces Stream-K GEMM unit tests using the Tile Engine infrastructure, establishing a test generation framework that mirrors tile_engine's kernel generation methodology. The tests cover fp16 and bf16 data types across all matrix layouts (rcr, rrr, ccr, crr) for both atomic and parallel reduction strategies, as well as a newly added TreeReduction strategy.
Changes:
- Added support for TreeReduction strategy across profiler, instance builder, and configuration files
- Created a comprehensive test framework that generates individual test executables for each kernel configuration
- Established test parameter extraction from JSON configurations for flexible test definition
Reviewed changes
Copilot reviewed 9 out of 9 changed files in this pull request and generated 1 comment.
Show a summary per file
| File | Description |
|---|---|
| tile_engine/ops/gemm_streamk/gemm_streamk_profiler.hpp | Added TreeReduction to the reduction strategy mapping for kernel profiling |
| tile_engine/ops/gemm_streamk/gemm_streamk_instance_builder.py | Added TreeReduction mapping and workspace reset logic |
| tile_engine/ops/gemm_streamk/configs/default_config.json | Added "tree" to the list of supported reduction strategies |
| test/ck_tile/gemm_streamk_tile_engine/test_gemm_streamk_simple.cpp | New test implementation using GTest with tile_engine error thresholds |
| test/ck_tile/gemm_streamk_tile_engine/extract_test_params.py | Python utility to extract test parameters from JSON config files |
| test/ck_tile/gemm_streamk_tile_engine/configs/simple_test_config.json | Test configuration defining problem sizes and kernel traits |
| test/ck_tile/gemm_streamk_tile_engine/README.md | Comprehensive documentation of the test framework and its integration |
| test/ck_tile/gemm_streamk_tile_engine/CMakeLists.txt | Build system integration for generating and compiling test targets |
| test/ck_tile/CMakeLists.txt | Added gemm_streamk_tile_engine subdirectory to test build |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
arai713
force-pushed
the
arai/ck_tile/streamk_tile_engine_test
branch
from
89e4aba to
01b28de
Compare
arai713
force-pushed
the
arai/ck_tile/streamk_tile_engine_test
branch
4 times, most recently
from
416ccd6 to
3f48a82
Compare
| ### Config-Specific Test Parameters | ||
|
|
||
| Each test configuration can specify optimized problem sizes in its JSON file: | ||
| - **`test_params.problem_sizes`**: Array of `{m, n, k, split_k}` configurations |
There was a problem hiding this comment.
Maybe we can make a future ticket to remove split-k from the config for stream-k since it may be confusing as we don't support split-k :)
arai713
force-pushed
the
arai/ck_tile/streamk_tile_engine_test
branch
from
3f48a82 to
86a5a56
Compare
This change adds an implementation for generating Stream-K tests using Tile Engine. This will generate various test executables for different combinations based on the config files. This addition has simple tests running for bf16 and fp16, with both atomic and reduction strategies and compv3 pipeline. The tests rely on the implementation of Stream-K in Tile Engine.
…ugs regarding them are being resolved
arai713
force-pushed
the
arai/ck_tile/streamk_tile_engine_test
branch
from
86a5a56 to
1e3844b
Compare
This PR introduces the generation of unit tests for Stream-K using Tile Engine. This will allow us to scale up the unit tests we have for better coverage and maintainability. It establishes a small test targeting fp16 and bf16 data types and covers both atomic and parallel reduction strategies within the compv3 pipeline. It also lays the groundwork for expanding to the full set of Stream-K smoke and extended tests.
These are the supported instances:
Checklist
Please put an
xinto 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.clang-formaton all changed filesDiscussion
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