composable_kernel icon indicating copy to clipboard operation
composable_kernel copied to clipboard

[CK_TILE] Add supported grouped and batched GEMM test cases

Open aledudek opened this issue 3 months ago • 4 comments

Proposed changes

Add supported grouped and batched GEMM test cases

Checklist

  • [x] I have added tests relevant to the introduced functionality, and the unit tests are passing locally
  • [x] 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
  • [x] I have removed the stale documentation which is no longer relevant after this pull request
  • [x] (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

aledudek avatar Sep 16 '25 10:09 aledudek

The new tests appear to be failing on gfx12.

illsilin avatar Sep 16 '25 18:09 illsilin

All new tests are also failing on gfx950: [----------] Global test environment tear-down [==========] 20 tests from 20 test suites ran. (11542 ms total) [ PASSED ] 3 tests. [ FAILED ] 17 tests, listed below: [ FAILED ] TestCkTileBatchedGemm/0.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,_Float16,_Float16,float,_Float16> [ FAILED ] TestCkTileBatchedGemm/1.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,_Float16,_Float16,float,_Float16> [ FAILED ] TestCkTileBatchedGemm/3.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::RowMajor,_Float16,_Float16,float,_Float16> [ FAILED ] TestCkTileBatchedGemm/4.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,std::bfloat16_t,std::bfloat16_t,float,_Float16> [ FAILED ] TestCkTileBatchedGemm/5.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,std::bfloat16_t,std::bfloat16_t,float,_Float16> [ FAILED ] TestCkTileBatchedGemm/7.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::RowMajor,std::bfloat16_t,std::bfloat16_t,float,_Float16> [ FAILED ] TestCkTileBatchedGemm/8.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm8RowMajorES3_S3_DU8_S4_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/9.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm11ColumnMajorENS2_8RowMajorES4_DU8_S5_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/10.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm8RowMajorENS2_11ColumnMajorES3_DU8_S5_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/11.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm11ColumnMajorES3_NS2_8RowMajorEDU8_S5_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/12.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm8RowMajorES3_S3_DB8_S4_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/13.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm11ColumnMajorENS2_8RowMajorES4_DB8_S5_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/14.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm8RowMajorENS2_11ColumnMajorES3_DB8_S5_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/15.Basic, where TypeParam = St5tupleIJN7ck_tile13tensor_layout4gemm11ColumnMajorES3_NS2_8RowMajorEDB8_S5_fDF16_EE [ FAILED ] TestCkTileBatchedGemm/16.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,signed char,signed char,int,int> [ FAILED ] TestCkTileBatchedGemm/17.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::RowMajor,ck_tile::tensor_layout::gemm::RowMajor,signed char,signed char,int,int> [ FAILED ] TestCkTileBatchedGemm/19.Basic, where TypeParam = std::tuple<ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::ColumnMajor,ck_tile::tensor_layout::gemm::RowMajor,signed char,signed char,int,int>

17 FAILED TESTS

Please make sure to run CI on gfx950 before merging!

illsilin avatar Sep 16 '25 19:09 illsilin

Please remove duplicates

These are not duplicated cases, one case has Persistent=True, the other Persistent=False

aledudek avatar Sep 30 '25 11:09 aledudek

@aledudek Please address the CI test failures

cgmillette avatar Oct 30 '25 17:10 cgmillette

@aledudek Marking as stale as no response in over two weeks. Risking PR closure if comments are not addressed!

cgmillette avatar Nov 13 '25 18:11 cgmillette