diff --git a/test/ck_tile/CMakeLists.txt b/test/ck_tile/CMakeLists.txt index d9324119919..82f4ee0ee7a 100644 --- 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_ - Build specific test executable +# - ./build/bin/test_ - 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) diff --git a/test/ck_tile/README.md b/test/ck_tile/README.md new file mode 100644 index 00000000000..027e6b60ddc --- /dev/null +++ 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 +``` + + + + + + diff --git a/test/ck_tile/fmha/CMakeLists.txt b/test/ck_tile/fmha/CMakeLists.txt index e591d5066f0..a9f658652f9 100644 --- a/test/ck_tile/fmha/CMakeLists.txt +++ b/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..." +) diff --git a/test/ck_tile/gemm/CMakeLists.txt b/test/ck_tile/gemm/CMakeLists.txt index ee23ad2f632..e7fb72ab41c 100644 --- 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() diff --git a/test/ck_tile/gemm_block_scale/CMakeLists.txt b/test/ck_tile/gemm_block_scale/CMakeLists.txt index 8e005d588e7..2080fc185f7 100644 --- a/test/ck_tile/gemm_block_scale/CMakeLists.txt +++ b/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() diff --git a/test/ck_tile/gemm_streamk/CMakeLists.txt b/test/ck_tile/gemm_streamk/CMakeLists.txt index 1390e5ee07f..3aa5f327647 100644 --- a/test/ck_tile/gemm_streamk/CMakeLists.txt +++ b/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() diff --git a/test/ck_tile/grouped_gemm_quant/CMakeLists.txt b/test/ck_tile/grouped_gemm_quant/CMakeLists.txt index 5959e44f481..04db7b0c4b8 100644 --- 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() diff --git a/test/ck_tile/reduce/CMakeLists.txt b/test/ck_tile/reduce/CMakeLists.txt index 0e6b4fcd49a..73046499eb1 100644 --- 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()