test: Add umbrella test targets for CK Tile operations

test/ck_tile/CMakeLists.txt

@@ -1,6 +1,32 @@
 # Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
 # SPDX-License-Identifier: MIT
 
+################################################################################
+# CK Tile Test Organization
+################################################################################
+# CK Tile tests can be run using several methods:
+#
+# 1. Global test labels (run tests across all operations):
+#    - ninja smoke       - Fast tests (~30s on gfx90a)
+#    - ninja regression  - Slower comprehensive tests
+#    - ninja check       - All available tests
+#
+# 2. Operation-specific umbrella targets (run all tests for a specific operation):
+#    - ninja ck_tile_gemm_tests                  - All basic GEMM tests
+#    - ninja ck_tile_gemm_block_scale_tests      - All GEMM with block-scale quantization tests
+#    - ninja ck_tile_gemm_streamk_tests          - All GEMM StreamK tests
+#    - ninja ck_tile_grouped_gemm_quant_tests    - All grouped GEMM quantization tests
+#    - ninja ck_tile_reduce_tests                - All reduce operation tests
+#    - ninja ck_tile_fmha_tests                  - All FMHA (Flash Attention) tests
+#
+# 3. Individual test executables:
+#    - ninja test_<test_name>                    - Build specific test executable
+#    - ./build/bin/test_<test_name>              - Run specific test directly
+#
+# These umbrella targets are useful when working on specific operations to quickly
+# validate all related tests without running the entire test suite.
+################################################################################
+
 add_subdirectory(image_to_column)
 add_subdirectory(gemm)
 add_subdirectory(gemm_persistent_async_input)

test/ck_tile/README.md

@@ -0,0 +1,133 @@
+# CK Tile Testing Guide
+
+This document describes the test organization and available test targets for CK Tile operations.
+
+## Overview
+
+CK Tile tests are organized with multiple levels of granularity to support different development workflows:
+
+1. **Global test labels** - Run tests across all operations
+2. **Operation-specific umbrella targets** - Run all tests for a specific operation
+3. **Individual test executables** - Run specific tests
+
+## Global Test Labels
+
+These targets run tests across all CK operations (not just CK Tile):
+
+### `ninja smoke`
+Run fast smoke tests (tests that complete within ~30 seconds on gfx90a).
+```bash
+ninja smoke
+```
+
+### `ninja regression`
+Run slower, more comprehensive regression tests.
+```bash
+ninja regression
+```
+
+### `ninja check`
+Run ALL available tests in the entire codebase.
+```bash
+ninja check
+```
+
+## Operation-Specific Umbrella Targets
+
+These targets allow you to run all tests for a specific CK Tile operation. This is useful when making changes to a particular operation and wanting to validate all related tests without running the entire test suite.
+
+### GEMM Operations
+
+#### `ck_tile_gemm_tests`
+Run all basic GEMM pipeline tests (memory, compute variants, persistent, etc.)
+```bash
+ninja ck_tile_gemm_tests
+```
+**Test executables included:**
+- `test_ck_tile_gemm_pipeline_mem`
+- `test_ck_tile_gemm_pipeline_compv3`
+- `test_ck_tile_gemm_pipeline_compv4`
+- `test_ck_tile_gemm_pipeline_persistent`
+- `test_ck_tile_gemm_pipeline_compv6`
+- `test_ck_tile_gemm_pipeline_comp_async` (gfx95 only)
+- `test_ck_tile_gemm_pipeline_*_wmma` variants (gfx11/gfx12 only)
+
+#### `ck_tile_gemm_block_scale_tests`
+Run all GEMM tests with block-scale quantization (AQuant, BQuant, ABQuant, etc.)
+```bash
+ninja ck_tile_gemm_block_scale_tests
+```
+**Test executables included:** 29 test executables covering:
+- AQuant tests (memory pipelines, base layouts, prefill, preshuffle, transpose)
+- ABQuant tests (base, padding, preshuffle)
+- BQuant tests (1D/2D variants, transpose)
+- BQuant with PreshuffleB (decode/prefill, 1D/2D)
+- BQuant with PreshuffleQuant (decode/prefill, 1D/2D)
+- RowColQuant and TensorQuant tests
+
+#### `ck_tile_gemm_streamk_tests`
+Run all GEMM StreamK tests (tile partitioner, reduction, smoke, extended)
+```bash
+ninja ck_tile_gemm_streamk_tests
+```
+**Test executables included:**
+- `test_ck_tile_streamk_tile_partitioner`
+- `test_ck_tile_streamk_reduction`
+- `test_ck_tile_streamk_smoke`
+- `test_ck_tile_streamk_extended`
+
+#### `ck_tile_grouped_gemm_quant_tests`
+Run all grouped GEMM quantization tests
+```bash
+ninja ck_tile_grouped_gemm_quant_tests
+```
+**Test executables included:**
+- `test_ck_tile_grouped_gemm_quant_rowcol`
+- `test_ck_tile_grouped_gemm_quant_tensor`
+- `test_ck_tile_grouped_gemm_quant_aquant`
+- `test_ck_tile_grouped_gemm_quant_bquant`
+- `test_ck_tile_grouped_gemm_quant_bquant_preshuffleb`
+
+### Other Operations
+
+#### `ck_tile_fmha_tests`
+Run all FMHA (Flash Multi-Head Attention) tests
+```bash
+ninja ck_tile_fmha_tests
+```
+**Test executables included:** Forward and backward tests for fp16, bf16, fp8bf16, fp32
+
+#### `ck_tile_reduce_tests`
+Run all reduce operation tests
+```bash
+ninja ck_tile_reduce_tests
+```
+**Test executables included:**
+- `test_ck_tile_reduce2d`
+- `test_ck_tile_multi_reduce2d_threadwise`
+- `test_ck_tile_multi_reduce2d_multiblock`
+
+## Individual Test Executables
+
+You can also build and run individual test executables:
+
+### Build a specific test
+```bash
+ninja test_ck_tile_gemm_pipeline_mem
+```
+
+### Run a specific test directly
+```bash
+./build/bin/test_ck_tile_gemm_pipeline_mem
+```
+
+### Run a specific test through ctest
+```bash
+ctest -R test_ck_tile_gemm_pipeline_mem --output-on-failure
+```
+
+
+
+
+
+

test/ck_tile/fmha/CMakeLists.txt

@@ -23,7 +23,7 @@ function(add_gtest_fwd test_group)
         set(name "${test_group}_${type}")
         add_gtest_executable(${name} test_fmha_fwd.cpp)
         get_test_property(${name} LABELS COMMON_LABELS)
-        set_tests_properties(${name} PROPERTIES LABELS "${COMMON_LABELS};${TEST_NAME};${test_group}")
+        set_tests_properties(${name} PROPERTIES LABELS "${COMMON_LABELS};${TEST_NAME};${test_group};CK_TILE_FMHA_TESTS")
         target_compile_definitions(${name} PRIVATE DataTypeConfig=${CPP_TYPE_${type}})
         target_link_libraries(${name} PRIVATE ${FMHA_FWD_INSTANCES})
         list(APPEND all_tests ${name})
@@ -43,7 +43,7 @@ function(add_gtest_bwd test_group)
         set(name "${test_group}_${type}")
         add_gtest_executable(${name} test_fmha_bwd.cpp)
         get_test_property(${name} LABELS COMMON_LABELS)
-        set_tests_properties(${name} PROPERTIES LABELS "${COMMON_LABELS};${TEST_NAME};${test_group}")
+        set_tests_properties(${name} PROPERTIES LABELS "${COMMON_LABELS};${TEST_NAME};${test_group};CK_TILE_FMHA_TESTS")
         target_compile_definitions(${name} PRIVATE DataTypeConfig=${CPP_TYPE_${type}})
         target_link_libraries(${name} PRIVATE ${FMHA_BWD_INSTANCES})
         list(APPEND all_tests ${name})
@@ -56,3 +56,12 @@ endfunction()
 add_gtest_fwd(${TEST_NAME}_fwd)
 add_gtest_bwd(${TEST_NAME}_bwd)
 add_custom_target(${TEST_NAME} DEPENDS ${TEST_NAME}_fwd ${TEST_NAME}_bwd)
+
+# Umbrella target to build and run all ck_tile fmha tests
+# Usage: ninja ck_tile_fmha_tests
+add_custom_target(ck_tile_fmha_tests
+    COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "CK_TILE_FMHA_TESTS"
+    DEPENDS ${TEST_NAME}
+    USES_TERMINAL
+    COMMENT "Running all ck_tile fmha tests..."
+)

test/ck_tile/gemm/CMakeLists.txt

@@ -17,6 +17,9 @@ list(APPEND EXAMPLE_GEMM_COMPILE_COMPUTE_V4_OPTIONS
 set(EXAMPLE_GEMM_COMPILE_COMPUTE_ASYNC_OPTIONS ${EXAMPLE_GEMM_COMPILE_COMPUTE_V4_OPTIONS})
 
 if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx90a|gfx11|gfx12")
+    # Collect all test targets for umbrella label
+    set(CK_TILE_GEMM_TEST_TARGETS)
+
     if(GPU_TARGETS MATCHES "gfx94|gfx95")
         add_gtest_executable(test_ck_tile_gemm_pipeline_mem test_gemm_pipeline_mem.cpp)
         add_gtest_executable(test_ck_tile_gemm_pipeline_compv3 test_gemm_pipeline_compv3.cpp)
@@ -29,11 +32,23 @@ if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx90a|gfx11|gfx12")
         target_compile_options(test_ck_tile_gemm_pipeline_compv4 PRIVATE ${EXAMPLE_GEMM_COMPILE_COMPUTE_V4_OPTIONS})
         target_compile_options(test_ck_tile_gemm_pipeline_persistent PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
         target_compile_options(test_ck_tile_gemm_pipeline_compv6 PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
+
+        list(APPEND CK_TILE_GEMM_TEST_TARGETS
+            test_ck_tile_gemm_pipeline_mem
+            test_ck_tile_gemm_pipeline_compv3
+            test_ck_tile_gemm_pipeline_compv4
+            test_ck_tile_gemm_pipeline_persistent
+            test_ck_tile_gemm_pipeline_compv6
+        )
     endif()
 
     if(GPU_TARGETS MATCHES "gfx95")
         add_gtest_executable(test_ck_tile_gemm_pipeline_comp_async test_gemm_pipeline_comp_async.cpp)
         target_compile_options(test_ck_tile_gemm_pipeline_comp_async PRIVATE ${EXAMPLE_GEMM_COMPILE_COMPUTE_ASYNC_OPTIONS})
+
+        list(APPEND CK_TILE_GEMM_TEST_TARGETS
+            test_ck_tile_gemm_pipeline_comp_async
+        )
     endif()
 
     if(GPU_TARGETS MATCHES "gfx11|gfx12")
@@ -55,7 +70,28 @@ if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx90a|gfx11|gfx12")
         target_compile_options(test_ck_tile_gemm_pipeline_compv3_wmma PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
         target_compile_options(test_ck_tile_gemm_pipeline_compv4_wmma PRIVATE ${EXAMPLE_GEMM_COMPILE_COMPUTE_V4_OPTIONS})
         target_compile_options(test_ck_tile_gemm_pipeline_persistent_wmma PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
+
+        list(APPEND CK_TILE_GEMM_TEST_TARGETS
+            test_ck_tile_gemm_pipeline_mem_wmma
+            test_ck_tile_gemm_pipeline_compv3_wmma
+            test_ck_tile_gemm_pipeline_compv4_wmma
+            test_ck_tile_gemm_pipeline_persistent_wmma
+        )
     endif()
+
+    # Label all ck_tile gemm tests with CK_TILE_GEMM_TESTS for selective execution
+    foreach(test_target ${CK_TILE_GEMM_TEST_TARGETS})
+        set_tests_properties(${test_target} PROPERTIES LABELS "CK_TILE_GEMM_TESTS")
+    endforeach()
+
+    # Umbrella target to build and run all ck_tile gemm tests
+    # Usage: ninja ck_tile_gemm_tests
+    add_custom_target(ck_tile_gemm_tests
+        COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "CK_TILE_GEMM_TESTS"
+        DEPENDS ${CK_TILE_GEMM_TEST_TARGETS}
+        USES_TERMINAL
+        COMMENT "Running all ck_tile gemm tests..."
+    )
 else()
     message(DEBUG "Skipping ck_tile_gemm tests for current target test_ck_tile_gemm_pipeline") 
 endif()

test/ck_tile/gemm_block_scale/CMakeLists.txt

@@ -187,17 +187,8 @@ if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx12")
     )
     target_compile_options(test_tile_gemm_quant_tensor PRIVATE ${TEST_GEMM_COMPILE_OPTIONS})
 
-    # Target to build only AQuant memory pipeline tests
-    add_custom_target(test_tile_gemm_aquant_mem_all)
-    add_dependencies(test_tile_gemm_aquant_mem_all
-        test_tile_gemm_quant_aquant_mem_prefill_interwave
-        test_tile_gemm_quant_aquant_mem_decode_intrawave
-        test_tile_gemm_quant_aquant_mem_decode_interwave
-    )
-
-    # Umbrella target to build all gemm quant tests
-    add_custom_target(test_tile_gemm_quant_all)
-    add_dependencies(test_tile_gemm_quant_all
+    # Collect all test targets for umbrella label
+    set(CK_TILE_GEMM_BLOCK_SCALE_TEST_TARGETS
         # AQuant tests
         test_tile_gemm_quant_aquant_mem_prefill_interwave
         test_tile_gemm_quant_aquant_mem_decode_intrawave
@@ -235,6 +226,32 @@ if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx12")
         test_tile_gemm_quant_rowcol
         test_tile_gemm_quant_tensor
     )
+
+    # Label all ck_tile gemm_block_scale tests with CK_TILE_GEMM_BLOCK_SCALE_TESTS for selective execution
+    foreach(test_target ${CK_TILE_GEMM_BLOCK_SCALE_TEST_TARGETS})
+        set_tests_properties(${test_target} PROPERTIES LABELS "CK_TILE_GEMM_BLOCK_SCALE_TESTS")
+    endforeach()
+
+    # Target to build only AQuant memory pipeline tests
+    add_custom_target(test_tile_gemm_aquant_mem_all)
+    add_dependencies(test_tile_gemm_aquant_mem_all
+        test_tile_gemm_quant_aquant_mem_prefill_interwave
+        test_tile_gemm_quant_aquant_mem_decode_intrawave
+        test_tile_gemm_quant_aquant_mem_decode_interwave
+    )
+
+    # Umbrella target to build all gemm quant tests
+    add_custom_target(test_tile_gemm_quant_all)
+    add_dependencies(test_tile_gemm_quant_all ${CK_TILE_GEMM_BLOCK_SCALE_TEST_TARGETS})
+
+    # Umbrella target to build and run all ck_tile gemm_block_scale tests
+    # Usage: ninja ck_tile_gemm_block_scale_tests
+    add_custom_target(ck_tile_gemm_block_scale_tests
+        COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "CK_TILE_GEMM_BLOCK_SCALE_TESTS"
+        DEPENDS ${CK_TILE_GEMM_BLOCK_SCALE_TEST_TARGETS}
+        USES_TERMINAL
+        COMMENT "Running all ck_tile gemm_block_scale tests..."
+    )
 else()
     message(DEBUG "Skipping ck_tile quant gemm tests for current target")
 endif()

test/ck_tile/gemm_streamk/CMakeLists.txt

@@ -48,6 +48,28 @@ if(GPU_TARGETS MATCHES "gfx90a|gfx942|gfx950")
                          test_gemm_streamk_util.cpp)
     target_compile_options(test_ck_tile_streamk_smoke PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
     target_compile_options(test_ck_tile_streamk_extended PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
+
+    # Collect all test targets for umbrella label
+    set(CK_TILE_GEMM_STREAMK_TEST_TARGETS
+        test_ck_tile_streamk_tile_partitioner
+        test_ck_tile_streamk_reduction
+        test_ck_tile_streamk_smoke
+        test_ck_tile_streamk_extended
+    )
+
+    # Label all ck_tile gemm_streamk tests with CK_TILE_GEMM_STREAMK_TESTS for selective execution
+    foreach(test_target ${CK_TILE_GEMM_STREAMK_TEST_TARGETS})
+        set_tests_properties(${test_target} PROPERTIES LABELS "CK_TILE_GEMM_STREAMK_TESTS")
+    endforeach()
+
+    # Umbrella target to build and run all ck_tile gemm_streamk tests
+    # Usage: ninja ck_tile_gemm_streamk_tests
+    add_custom_target(ck_tile_gemm_streamk_tests
+        COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "CK_TILE_GEMM_STREAMK_TESTS"
+        DEPENDS ${CK_TILE_GEMM_STREAMK_TEST_TARGETS}
+        USES_TERMINAL
+        COMMENT "Running all ck_tile gemm_streamk tests..."
+    )
 else()
     message(DEBUG "Skipping test_ck_tile_streamk unit tests for current target")
 endif()

test/ck_tile/grouped_gemm_quant/CMakeLists.txt

@@ -22,5 +22,28 @@ if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx12")
 
     add_gtest_executable(test_ck_tile_grouped_gemm_quant_bquant_preshuffleb test_grouped_gemm_quant_bquant_preshuffleb.cpp)
     target_compile_options(test_ck_tile_grouped_gemm_quant_bquant_preshuffleb PRIVATE ${EXAMPLE_GEMM_COMPILE_OPTIONS})
+
+    # Collect all test targets for umbrella label
+    set(CK_TILE_GROUPED_GEMM_QUANT_TEST_TARGETS
+        test_ck_tile_grouped_gemm_quant_rowcol
+        test_ck_tile_grouped_gemm_quant_tensor
+        test_ck_tile_grouped_gemm_quant_aquant
+        test_ck_tile_grouped_gemm_quant_bquant
+        test_ck_tile_grouped_gemm_quant_bquant_preshuffleb
+    )
+
+    # Label all ck_tile grouped_gemm_quant tests with CK_TILE_GROUPED_GEMM_QUANT_TESTS for selective execution
+    foreach(test_target ${CK_TILE_GROUPED_GEMM_QUANT_TEST_TARGETS})
+        set_tests_properties(${test_target} PROPERTIES LABELS "CK_TILE_GROUPED_GEMM_QUANT_TESTS")
+    endforeach()
+
+    # Umbrella target to build and run all ck_tile grouped_gemm_quant tests
+    # Usage: ninja ck_tile_grouped_gemm_quant_tests
+    add_custom_target(ck_tile_grouped_gemm_quant_tests
+        COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "CK_TILE_GROUPED_GEMM_QUANT_TESTS"
+        DEPENDS ${CK_TILE_GROUPED_GEMM_QUANT_TEST_TARGETS}
+        USES_TERMINAL
+        COMMENT "Running all ck_tile grouped_gemm_quant tests..."
+    )
 endif()
 

test/ck_tile/reduce/CMakeLists.txt

@@ -10,5 +10,26 @@ if(GPU_TARGETS MATCHES "gfx9|gfx11|gfx12")
         target_link_libraries(test_ck_tile_multi_reduce2d_threadwise PRIVATE utility)
         target_link_libraries(test_ck_tile_multi_reduce2d_multiblock PRIVATE utility)
     endif()
+
+    # Collect all test targets for umbrella label
+    set(CK_TILE_REDUCE_TEST_TARGETS
+        test_ck_tile_reduce2d
+        test_ck_tile_multi_reduce2d_threadwise
+        test_ck_tile_multi_reduce2d_multiblock
+    )
+
+    # Label all ck_tile reduce tests with CK_TILE_REDUCE_TESTS for selective execution
+    foreach(test_target ${CK_TILE_REDUCE_TEST_TARGETS})
+        set_tests_properties(${test_target} PROPERTIES LABELS "CK_TILE_REDUCE_TESTS")
+    endforeach()
+
+    # Umbrella target to build and run all ck_tile reduce tests
+    # Usage: ninja ck_tile_reduce_tests
+    add_custom_target(ck_tile_reduce_tests
+        COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR} -L "CK_TILE_REDUCE_TESTS"
+        DEPENDS ${CK_TILE_REDUCE_TEST_TARGETS}
+        USES_TERMINAL
+        COMMENT "Running all ck_tile reduce tests..."
+    )
 endif()