[CK-Tile][Async gemm] add missing sync and f8 inputs test cases (#3000)

tenpercent · web-flow · commit 86d542f66320 · 2025-10-24T12:16:01.000-07:00
* add missing sync and f8 test cases

* reformat test cases

* comment failing cases

* bump

* reintroduce compv4 shapes
diff --git a/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp b/include/ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_comp_async.hpp
@@ -472,6 +472,8 @@ struct GemmPipelineAgBgCrCompAsync : public BaseGemmPipelineAgBgCrCompAsync<Prob
                     block_gemm(c_block_tile, a_block_tile0, b_block_tile0);
                 }
                 {
+                    // write to LDS window(0) must complete before the local prefetch
+                    block_sync_lds_direct_load();
                     // read A(num_loop), B(num_loop) from LDS window(0) to pipeline registers(0)
                     Base::LocalPrefetch(a_block_tile0, a_lds_ld_window0, is_a_load_tr_v);
                     Base::LocalPrefetch(b_block_tile0, b_lds_ld_window0, is_b_load_tr_v);
diff --git a/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp b/test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp
@@ -124,12 +124,59 @@ using KernelTypesCompV3Wmma = ::testing::Types<
     std::tuple<    Col,     Col,     Row,       BF8,       BF8,         F32,       F16,        I64,         I64,          I32,        I16,        I16,        I16, Intrawave,        CompV3>
 >;
 
-using KernelTypesCompV4 = ::testing::Types<
-    std::tuple<    Row,     Row,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompV4>,
-    std::tuple<    Row,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompV4>,
-    std::tuple<    Col,     Row,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompV4>,
-    std::tuple<    Col,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompV4>
->;
+// clang-format on
+template <typename ALayout, typename BLayout, typename CLayout, typename InputType>
+using CompV4Config = std::tuple<ALayout,
+                                BLayout,
+                                CLayout,
+                                InputType, // AType
+                                InputType, // BType
+                                F32,       // AccType
+                                F16,       // OutputType
+                                I256,      // MBlockTileSize
+                                I256,      // NBlockTileSize
+                                I32,       // KBlockTileSize
+                                I32,       // MWarpTileSize
+                                I32,       // NWarpTileSize
+                                I16,       // KWarpTileSize
+                                Intrawave,
+                                CompV4>;
+
+using KernelTypesCompV4 = ::testing::Types<CompV4Config<Row, Row, Row, F16>,
+                                           CompV4Config<Row, Col, Row, F16>,
+                                           CompV4Config<Col, Row, Row, F16>,
+                                           CompV4Config<Col, Col, Row, F16>,
+                                           CompV4Config<Row, Row, Row, F8>,
+                                           CompV4Config<Row, Col, Row, F8>,
+                                           CompV4Config<Col, Row, Row, F8>,
+                                           CompV4Config<Col, Col, Row, F8>>;
+
+template <typename ALayout, typename BLayout, typename CLayout, typename InputType>
+using CompAsyncConfig = std::tuple<ALayout,
+                                   BLayout,
+                                   CLayout,
+                                   InputType, // AType
+                                   InputType, // BType
+                                   F32,       // AccType
+                                   F16,       // OutputType
+                                   I256,      // MBlockTileSize
+                                   I256,      // NBlockTileSize
+                                   I32,       // KBlockTileSize
+                                   I32,       // MWarpTileSize
+                                   I32,       // NWarpTileSize
+                                   I16,       // KWarpTileSize
+                                   Intrawave,
+                                   CompAsync>;
+
+using KernelTypesCompAsync = ::testing::Types<CompAsyncConfig<Row, Row, Row, F16>,
+                                              CompAsyncConfig<Row, Col, Row, F16>,
+                                              CompAsyncConfig<Col, Row, Row, F16>,
+                                              CompAsyncConfig<Col, Col, Row, F16>,
+                                              CompAsyncConfig<Row, Row, Row, F8>,
+                                              CompAsyncConfig<Row, Col, Row, F8>,
+                                              CompAsyncConfig<Col, Row, Row, F8>,
+                                              CompAsyncConfig<Col, Col, Row, F8>>;
+// clang-format off
 
 using KernelTypesCompV6 = ::testing::Types<
     std::tuple<    Row,     Row,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I64,        I32,        I32,        I16, Intrawave,        CompV6>,
@@ -153,12 +200,6 @@ using KernelTypesCompV6 = ::testing::Types<
     std::tuple<    Col,     Row,     Row,       INT8,      INT8,        INT32,     INT32,      I256,        I256,         I64,        I32,        I32,        I16, Intrawave,        CompV6>,
     std::tuple<    Col,     Col,     Row,       INT8,      INT8,        INT32,     INT32,      I256,        I256,         I64,        I32,        I32,        I16, Intrawave,        CompV6>
 >;
-using KernelTypesCompAsync = ::testing::Types<
-    std::tuple<    Row,     Row,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompAsync>,
-    std::tuple<    Row,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompAsync>,
-    std::tuple<    Col,     Row,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompAsync>,
-    std::tuple<    Col,     Col,     Row,       F16,       F16,         F32,       F16,        I256,        I256,         I32,        I32,        I32,        I16, Intrawave,        CompAsync>
->;
 
 using KernelTypesCompV4Wmma = ::testing::Types<
     std::tuple<    Row,     Row,     Row,       F16,       F16,         F32,       F16,        I64,         I64,          I32,        I16,        I16,        I16, Intrawave,        CompV4>,

Original file line number	Diff line number	Diff line change
`@@ -472,6 +472,8 @@ struct GemmPipelineAgBgCrCompAsync : public BaseGemmPipelineAgBgCrCompAsync<Prob`
`472`	`472`	`block_gemm(c_block_tile, a_block_tile0, b_block_tile0);`
`473`	`473`	`}`
`474`	`474`	`{`
	`475`	`+ // write to LDS window(0) must complete before the local prefetch`
	`476`	`+ block_sync_lds_direct_load();`
`475`	`477`	`// read A(num_loop), B(num_loop) from LDS window(0) to pipeline registers(0)`
`476`	`478`	`Base::LocalPrefetch(a_block_tile0, a_lds_ld_window0, is_a_load_tr_v);`
`477`	`479`	`Base::LocalPrefetch(b_block_tile0, b_lds_ld_window0, is_b_load_tr_v);`