[CK_TILE] Add the GEMM Memory pipeline to Stream-K tests (#5242)

## Motivation

We want to extend our Stream-K coverage to include other GEMM pipeline
since our current tests only test the CompV3 pipeline.

## Technical Details

All Stream-K unit tests currently only tests one pipeline: CompV3. These
changes extend the test support to also test the Memory pipeline. Future
work will add support for additional GEMM pipelines.

The major changes are as follows:
- **Remove of fp8 and bf8 extended tests for gfx90a**: gfx90a does not
have native support for fp8 and bf8 and emulate the behavior with fp32
mfma instruction sizes. We've observed extremely long compile times for
fp8 and bf8 on gfx90a (exceeding 15 minutes), hence we've opted to
disable these tests.
- **Add the memory pipeline to the Stream-K tile engine tests**: Now our
smoke tests covers compv3 and memory pipelines.
- **Add the memory pipeline to the Stream-K extended tests**: These
changes modify the test kernel types to include the appropriate
pipeline. Each pipeline is contained within a separate kernel type to
help avoid large increases in build time.

## Test Plan

- Ran existing and added tests on all architectures.

## Test Result

- All local tests pass. 

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Author: ecamartins

projects/composablekernel/test/ck_tile/gemm_streamk/CMakeLists.txt

@@ -23,16 +23,31 @@ if(GPU_TARGETS MATCHES "gfx90a|gfx942|gfx950")
     #TODO: support all arches
     #TODO: current c-shuffle only supports C layout as R
     add_gtest_executable(test_ck_tile_streamk_tile_partitioner test_streamk_tile_partitioner.cpp)
-    add_gtest_executable(test_ck_tile_streamk_extended
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp16_persistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf16_persistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp8_persistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_persistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp16_nonpersistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf16_nonpersistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp8_nonpersistent.cpp
-                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_nonpersistent.cpp
-                         test_gemm_streamk_util.cpp)
+    set(STREAMK_EXTENDED_SOURCES
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp16_persistent_compv3.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp16_persistent_mem.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf16_persistent_compv3.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf16_persistent_mem.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp16_nonpersistent_compv3.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp16_nonpersistent_mem.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf16_nonpersistent_compv3.cpp
+        ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf16_nonpersistent_mem.cpp
+        test_gemm_streamk_util.cpp)
+
+    # We only test fp8 and bf8 on gfx942 and gfx950 since these types are not natively supported on gfx90a
+    if(GPU_TARGETS MATCHES "gfx942|gfx950")
+        list(APPEND STREAMK_EXTENDED_SOURCES
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp8_persistent_compv3.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp8_persistent_mem.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_persistent_compv3.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_persistent_mem.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp8_nonpersistent_compv3.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_fp8_nonpersistent_mem.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_nonpersistent_compv3.cpp
+            ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_nonpersistent_mem.cpp)
+    endif()
+
+    add_gtest_executable(test_ck_tile_streamk_extended ${STREAMK_EXTENDED_SOURCES})
     target_compile_options(test_ck_tile_streamk_extended PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
 
     # Collect all test targets for umbrella label

projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_nonpersistent_compv3.cpp

@@ -0,0 +1,18 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "test_gemm_streamk_common_includes.hpp"
+
+template <typename Tuple>
+class TestCkTileStreamKBf16NonPersistentCompV3 : public TestCkTileStreamK<Tuple>
+{
+};
+
+#define TEST_SUITE_NAME TestCkTileStreamKBf16NonPersistentCompV3
+
+TYPED_TEST_SUITE(TestCkTileStreamKBf16NonPersistentCompV3,
+                 KernelTypesStreamKBf16NonPersistentCompV3);
+
+#include "test_gemm_streamk_extended_cases.inc"
+
+#undef TEST_SUITE_NAME

projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_nonpersistent_mem.cpp

@@ -0,0 +1,17 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "test_gemm_streamk_common_includes.hpp"
+
+template <typename Tuple>
+class TestCkTileStreamKBf16NonPersistentMem : public TestCkTileStreamK<Tuple>
+{
+};
+
+#define TEST_SUITE_NAME TestCkTileStreamKBf16NonPersistentMem
+
+TYPED_TEST_SUITE(TestCkTileStreamKBf16NonPersistentMem, KernelTypesStreamKBf16NonPersistentMem);
+
+#include "test_gemm_streamk_extended_cases.inc"
+
+#undef TEST_SUITE_NAME

projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_persistent_compv3.cpp

@@ -0,0 +1,17 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "test_gemm_streamk_common_includes.hpp"
+
+template <typename Tuple>
+class TestCkTileStreamKBf16PersistentCompV3 : public TestCkTileStreamK<Tuple>
+{
+};
+
+#define TEST_SUITE_NAME TestCkTileStreamKBf16PersistentCompV3
+
+TYPED_TEST_SUITE(TestCkTileStreamKBf16PersistentCompV3, KernelTypesStreamKBf16PersistentCompV3);
+
+#include "test_gemm_streamk_extended_cases.inc"
+
+#undef TEST_SUITE_NAME

…test_gemm_streamk_bf16_nonpersistent.cpp …est_gemm_streamk_bf16_persistent_mem.cppprojects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_nonpersistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_persistent_mem.cpp projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_nonpersistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_persistent_mem.cpp

@@ -4,13 +4,13 @@
 #include "test_gemm_streamk_common_includes.hpp"
 
 template <typename Tuple>
-class TestCkTileStreamKBf16NonPersistent : public TestCkTileStreamK<Tuple>
+class TestCkTileStreamKBf16PersistentMem : public TestCkTileStreamK<Tuple>
 {
 };
 
-#define TEST_SUITE_NAME TestCkTileStreamKBf16NonPersistent
+#define TEST_SUITE_NAME TestCkTileStreamKBf16PersistentMem
 
-TYPED_TEST_SUITE(TestCkTileStreamKBf16NonPersistent, KernelTypesStreamKBf16NonPersistent);
+TYPED_TEST_SUITE(TestCkTileStreamKBf16PersistentMem, KernelTypesStreamKBf16PersistentMem);
 
 #include "test_gemm_streamk_extended_cases.inc"
 

projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_nonpersistent_compv3.cpp

@@ -0,0 +1,17 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "test_gemm_streamk_common_includes.hpp"
+
+template <typename Tuple>
+class TestCkTileStreamKBf8NonPersistentCompV3 : public TestCkTileStreamK<Tuple>
+{
+};
+
+#define TEST_SUITE_NAME TestCkTileStreamKBf8NonPersistentCompV3
+
+TYPED_TEST_SUITE(TestCkTileStreamKBf8NonPersistentCompV3, KernelTypesStreamKBf8NonPersistentCompV3);
+
+#include "test_gemm_streamk_extended_cases.inc"
+
+#undef TEST_SUITE_NAME

…ts/test_gemm_streamk_bf16_persistent.cpp …t_gemm_streamk_bf8_nonpersistent_mem.cppprojects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_persistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_nonpersistent_mem.cpp projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf16_persistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_nonpersistent_mem.cpp

@@ -4,13 +4,13 @@
 #include "test_gemm_streamk_common_includes.hpp"
 
 template <typename Tuple>
-class TestCkTileStreamKBf16Persistent : public TestCkTileStreamK<Tuple>
+class TestCkTileStreamKBf8NonPersistentMem : public TestCkTileStreamK<Tuple>
 {
 };
 
-#define TEST_SUITE_NAME TestCkTileStreamKBf16Persistent
+#define TEST_SUITE_NAME TestCkTileStreamKBf8NonPersistentMem
 
-TYPED_TEST_SUITE(TestCkTileStreamKBf16Persistent, KernelTypesStreamKBf16Persistent);
+TYPED_TEST_SUITE(TestCkTileStreamKBf8NonPersistentMem, KernelTypesStreamKBf8NonPersistentMem);
 
 #include "test_gemm_streamk_extended_cases.inc"
 

…sts/test_gemm_streamk_bf8_persistent.cpp …t_gemm_streamk_bf8_persistent_compv3.cppprojects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_persistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_persistent_compv3.cpp projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_persistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_persistent_compv3.cpp

@@ -4,13 +4,13 @@
 #include "test_gemm_streamk_common_includes.hpp"
 
 template <typename Tuple>
-class TestCkTileStreamKBf8Persistent : public TestCkTileStreamK<Tuple>
+class TestCkTileStreamKBf8PersistentCompV3 : public TestCkTileStreamK<Tuple>
 {
 };
 
-#define TEST_SUITE_NAME TestCkTileStreamKBf8Persistent
+#define TEST_SUITE_NAME TestCkTileStreamKBf8PersistentCompV3
 
-TYPED_TEST_SUITE(TestCkTileStreamKBf8Persistent, KernelTypesStreamKBf8Persistent);
+TYPED_TEST_SUITE(TestCkTileStreamKBf8PersistentCompV3, KernelTypesStreamKBf8PersistentCompV3);
 
 #include "test_gemm_streamk_extended_cases.inc"
 

…/test_gemm_streamk_fp8_nonpersistent.cpp …test_gemm_streamk_bf8_persistent_mem.cppprojects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_fp8_nonpersistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_persistent_mem.cpp projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_fp8_nonpersistent.cpp renamed to projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_bf8_persistent_mem.cpp

@@ -4,13 +4,13 @@
 #include "test_gemm_streamk_common_includes.hpp"
 
 template <typename Tuple>
-class TestCkTileStreamKFp8NonPersistent : public TestCkTileStreamK<Tuple>
+class TestCkTileStreamKBf8PersistentMem : public TestCkTileStreamK<Tuple>
 {
 };
 
-#define TEST_SUITE_NAME TestCkTileStreamKFp8NonPersistent
+#define TEST_SUITE_NAME TestCkTileStreamKBf8PersistentMem
 
-TYPED_TEST_SUITE(TestCkTileStreamKFp8NonPersistent, KernelTypesStreamKFp8NonPersistent);
+TYPED_TEST_SUITE(TestCkTileStreamKBf8PersistentMem, KernelTypesStreamKBf8PersistentMem);
 
 #include "test_gemm_streamk_extended_cases.inc"
 

projects/composablekernel/test/ck_tile/gemm_streamk/extended_tests/test_gemm_streamk_fp16_nonpersistent_compv3.cpp

@@ -0,0 +1,18 @@
+// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include "test_gemm_streamk_common_includes.hpp"
+
+template <typename Tuple>
+class TestCkTileStreamKFp16NonPersistentCompV3 : public TestCkTileStreamK<Tuple>
+{
+};
+
+#define TEST_SUITE_NAME TestCkTileStreamKFp16NonPersistentCompV3
+
+TYPED_TEST_SUITE(TestCkTileStreamKFp16NonPersistentCompV3,
+                 KernelTypesStreamKFp16NonPersistentCompV3);
+
+#include "test_gemm_streamk_extended_cases.inc"
+
+#undef TEST_SUITE_NAME