[rocm-libraries] ROCm/rocm-libraries#4301 (commit 0821c9f)

test: Add umbrella test targets for CK Tile operations (#4301) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Proposed changes Adds operation-specific umbrella test targets for CK Tile to enable running all tests for a specific operation without running the entire test suite. This improves the development workflow by allowing faster iteration when working on specific operations. ## Motivation Previously, developers working on CK Tile operations could only: - Run individual test executables one at a time - Run global labels (, , ) which test the entire codebase - Build all tests for an operation but had no simple way to run them all This made it cumbersome to validate changes to a specific operation (e.g., GEMM quantization) without either running tests individually or running the entire test suite. ### Documentation - - Comprehensive testing guide with usage examples and implementation details ## Usage Examples # Run all GEMM tests with 256 parallel jobs ninja -j256 ck_tile_gemm_tests # Run all GEMM block scale (quantization) tests ninja -j256 ck_tile_gemm_block_scale_tests # Run all GEMM StreamK tests ninja -j256 ck_tile_gemm_streamk_tests ## Checklist Please put an into 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. - [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 - [ ] (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 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
2026-04-19 22:39:03 +00:00 · 2026-03-03 15:40:50 +00:00
parent 1cd031c21d
commit f00ec5afd9
8 changed files with 298 additions and 13 deletions
--- a/test/ck_tile/CMakeLists.txt
+++ b/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)
--- a/test/ck_tile/README.md
+++ b/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
+```
+
+
+
+
+
+
--- a/test/ck_tile/fmha/CMakeLists.txt
+++ b/test/ck_tile/fmha/CMakeLists.txt
@@ -32,7 +32,7 @@ function(add_gtest_fwd test_group)
        endif()
        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})
@@ -62,7 +62,7 @@ function(add_gtest_bwd test_group)
        endif()
        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})
@@ -75,3 +75,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..."
+)
--- a/test/ck_tile/gemm/CMakeLists.txt
+++ b/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()
--- a/test/ck_tile/gemm_block_scale/CMakeLists.txt
+++ b/test/ck_tile/gemm_block_scale/CMakeLists.txt
@@ -254,17 +254,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
@@ -312,6 +303,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()
--- a/test/ck_tile/gemm_streamk/CMakeLists.txt
+++ b/test/ck_tile/gemm_streamk/CMakeLists.txt
@@ -34,6 +34,26 @@ if(GPU_TARGETS MATCHES "gfx90a|gfx942|gfx950")
                         ${CMAKE_CURRENT_SOURCE_DIR}/extended_tests/test_gemm_streamk_bf8_nonpersistent.cpp
                         test_gemm_streamk_util.cpp)
    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_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()
--- a/test/ck_tile/grouped_gemm_quant/CMakeLists.txt
+++ b/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()

--- a/test/ck_tile/reduce/CMakeLists.txt
+++ b/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()