ckTileEngine pooling

test/ck_tile/CMakeLists.txt

@@ -42,3 +42,4 @@ add_subdirectory(gemm_tile_engine)
 add_subdirectory(pooling)
 add_subdirectory(grouped_conv)
 add_subdirectory(gemm_streamk_tile_engine)
+add_subdirectory(pooling_tile_engine)

test/ck_tile/pooling_tile_engine/CMakeLists.txt

@@ -0,0 +1,298 @@
+# Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
+# SPDX-License-Identifier: MIT
+
+# ============================================================================
+# Pooling Tile Engine Unit Tests
+#
+# This CMake file creates unit tests for tile_engine generated pooling kernels.
+# Each kernel configuration gets its own test executable.
+# ============================================================================
+
+# Locate tile_engine pooling scripts directory
+set(TILE_ENGINE_POOLING_DIR "${PROJECT_SOURCE_DIR}/tile_engine/ops/pooling")
+
+if(NOT EXISTS ${TILE_ENGINE_POOLING_DIR})
+    message(WARNING "Tile engine pooling directory not found: ${TILE_ENGINE_POOLING_DIR}")
+    return()
+endif()
+
+# ============================================================================
+# create_individual_pool_test_target
+#
+# Creates a single test executable for a specific pooling kernel configuration.
+#
+# Parameters:
+#   datatype     - Data type (fp16, fp32, bf16)
+#   config_name  - Configuration file name without .json extension
+#   trait        - Kernel trait combination string
+#   tile_config  - Tile configuration parameters
+#   config_json  - Full path to JSON configuration file
+# ============================================================================
+function(create_individual_pool_test_target datatype config_name trait tile_config config_json)
+    set(target_name "test_pooling_tile_engine_${datatype}_${config_name}_${trait}_${tile_config}")
+    set(working_path "${CMAKE_CURRENT_BINARY_DIR}/${datatype}/${config_name}")
+
+    # Generated header path (already created during cmake configuration)
+    set(test_header "${working_path}/pooling_single_pool_${datatype}_${trait}_${tile_config}.hpp")
+    set(test_params_header "${working_path}/test_params.hpp")
+
+    # Verify header exists
+    if(NOT EXISTS ${test_header})
+        message(WARNING "Generated header not found: ${test_header}")
+        return()
+    endif()
+
+    # Verify test parameters header exists
+    if(NOT EXISTS ${test_params_header})
+        message(WARNING "Test parameters header not found: ${test_params_header}")
+        return()
+    endif()
+
+    # Create GTest executable for this kernel configuration
+    add_gtest_executable(${target_name}
+        ${CMAKE_CURRENT_SOURCE_DIR}/test_pooling_simple.cpp
+    )
+
+    # Configure GPU architectures for HIP compilation
+    set_property(TARGET ${target_name} PROPERTY HIP_ARCHITECTURES ${POOLING_TEST_GPU_TARGETS})
+
+    # Define preprocessor macros for generated header location and test parameters
+    target_compile_definitions(${target_name} PRIVATE
+        POOLING_SINGLE_INSTANCE_HPP="${test_header}"
+        POOLING_TEST_PARAMS_HPP="${test_params_header}"
+    )
+
+    # Include directories for headers and dependencies
+    target_include_directories(${target_name} PRIVATE
+        ${PROJECT_SOURCE_DIR}/include
+        ${PROJECT_BINARY_DIR}/include
+        ${PROJECT_SOURCE_DIR}  # Root directory for tile_engine access
+        ${GTEST_INCLUDE_DIRS}
+    )
+
+    # Compiler options matching tile_engine requirements
+    target_compile_options(${target_name} PRIVATE
+        -Wno-undefined-func-template
+        -Wno-float-equal
+        --offload-compress
+        -include ${test_header}
+    )
+
+    # Add FP8 format definitions for proper data type interpretation
+    if(CK_USE_OCP_FP8)
+        target_compile_options(${target_name} PRIVATE -DCK_TILE_USE_OCP_FP8)
+    endif()
+
+    message(STATUS "  Created test target: ${target_name}")
+endfunction()
+
+# ============================================================================
+# build_pool_test_targets
+#
+# Builds all test targets for a specific datatype/config combination.
+# Uses tile_engine's two-step process: list kernels, then generate tests.
+#
+# Parameters:
+#   datatype     - Data type (fp16, fp32, bf16)
+#   config_name  - Configuration file name without .json extension
+# ============================================================================
+function(build_pool_test_targets datatype config_name)
+    set(working_path "${CMAKE_CURRENT_BINARY_DIR}/${datatype}/${config_name}")
+
+    # Locate and validate configuration file
+    set(config_filename "${config_name}.json")
+    set(json_blob "${CMAKE_CURRENT_SOURCE_DIR}/configs/${config_filename}")
+
+    if(NOT EXISTS ${json_blob})
+        message(WARNING "Test config file not found: ${json_blob}")
+        return()
+    endif()
+
+    # Prepare build directory for this configuration
+    file(MAKE_DIRECTORY ${working_path})
+
+    # STEP 1: Discovery phase - list all valid kernel configurations
+    execute_process(
+        COMMAND ${Python3_EXECUTABLE} -u ${TILE_ENGINE_POOLING_DIR}/pooling_instance_builder.py
+                --working_path ${working_path}
+                --datatype ${datatype}
+                --config_json ${json_blob}
+                --list_kernels
+        WORKING_DIRECTORY ${TILE_ENGINE_POOLING_DIR}
+        RESULT_VARIABLE ret
+        OUTPUT_VARIABLE list_output
+        ERROR_VARIABLE list_error
+    )
+
+    if(NOT ret EQUAL 0)
+        message(WARNING "Failed to list pooling kernels for ${datatype}_${config_name}: ${list_error}")
+        return()
+    endif()
+
+    # Verify kernel list file was generated
+    if(NOT EXISTS ${working_path}/pool_kernel_list.txt)
+        message(STATUS "No pooling kernels found for ${datatype}_${config_name}")
+        return()
+    endif()
+
+    message(STATUS "Building pooling tests for ${datatype}_${config_name}")
+
+    # STEP 2a: Determine pooling dimension from config
+    # Read the trait config to find pooling_dim
+    file(READ ${json_blob} config_content)
+    string(FIND "${config_content}" "\"3d\"" found_3d)
+    if(found_3d GREATER -1)
+        set(pooling_dim "3d")
+    else()
+        set(pooling_dim "2d")
+    endif()
+
+    # STEP 2b: Extract test parameters from config
+    set(test_params_file "${working_path}/test_params.hpp")
+    execute_process(
+        COMMAND ${Python3_EXECUTABLE} -u ${CMAKE_CURRENT_SOURCE_DIR}/extract_test_params.py
+                --config_file ${json_blob}
+                --output_file ${test_params_file}
+                --pooling_dim ${pooling_dim}
+        WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}
+        RESULT_VARIABLE extract_ret
+        OUTPUT_VARIABLE extract_output
+        ERROR_VARIABLE extract_error
+    )
+
+    if(NOT extract_ret EQUAL 0)
+        message(WARNING "Failed to extract test parameters for pooling ${datatype}: ${extract_error}")
+        return()
+    endif()
+
+    # STEP 2c: Header generation phase - generate headers using --gen_single
+    message(STATUS "  Generating pooling headers using --gen_single...")
+
+    file(STRINGS ${working_path}/pool_kernel_list.txt kernel_lines)
+    set(gen_count 0)
+
+    foreach(line IN LISTS kernel_lines)
+        # Parse kernel specification format: kernel_name|tile_config|trait_combo
+        string(REPLACE "|" ";" parts "${line}")
+        list(LENGTH parts parts_len)
+        if(parts_len EQUAL 3)
+            list(GET parts 0 kernel_name)
+            list(GET parts 1 tile_config)
+            list(GET parts 2 trait_combo)
+
+            # Generate header using --gen_single
+            execute_process(
+                COMMAND ${Python3_EXECUTABLE} -u ${TILE_ENGINE_POOLING_DIR}/pooling_instance_builder.py
+                        --working_path ${working_path}
+                        --datatype ${datatype}
+                        --config_json ${json_blob}
+                        --gen_single
+                        --kernel_name "${kernel_name}"
+                        --tile_config "${tile_config}"
+                        --trait_combo "${trait_combo}"
+                WORKING_DIRECTORY ${TILE_ENGINE_POOLING_DIR}
+                RESULT_VARIABLE gen_ret
+                OUTPUT_VARIABLE gen_output
+                ERROR_VARIABLE gen_error
+            )
+
+            if(NOT gen_ret EQUAL 0)
+                message(WARNING "Failed to generate pooling header for ${kernel_name}: ${gen_error}")
+            else()
+                math(EXPR gen_count "${gen_count} + 1")
+            endif()
+        endif()
+    endforeach()
+
+    message(STATUS "  Generated ${gen_count} pooling headers for ${datatype}")
+
+    # STEP 3: Target creation phase - create test targets
+    message(STATUS "  Creating pooling test targets...")
+    file(STRINGS ${working_path}/pool_kernel_list.txt kernel_lines)
+    set(test_count 0)
+    foreach(line IN LISTS kernel_lines)
+        string(REPLACE "|" ";" parts "${line}")
+        list(LENGTH parts parts_len)
+        if(parts_len EQUAL 3)
+            list(GET parts 0 kernel_name)
+            list(GET parts 1 tile_config)
+            list(GET parts 2 trait_combo)
+
+            create_individual_pool_test_target("${datatype}" "${config_name}" "${trait_combo}" "${tile_config}" "${json_blob}")
+            math(EXPR test_count "${test_count} + 1")
+        endif()
+    endforeach()
+    message(STATUS "  Created ${test_count} pooling test targets for ${datatype}")
+endfunction()
+
+# ============================================================================
+# MAIN EXECUTION - Test Target Generation
+# ============================================================================
+
+message(STATUS "=== Starting Pooling Tile Engine Test Configuration ===")
+message(STATUS "SUPPORTED_GPU_TARGETS: ${SUPPORTED_GPU_TARGETS}")
+
+# GPU architecture filtering - only build tests for supported architectures
+set(POOLING_TEST_GPU_TARGETS "")
+set(DESIRED_TARGETS "gfx90a;gfx942")
+
+foreach(target IN LISTS SUPPORTED_GPU_TARGETS)
+    if(target IN_LIST DESIRED_TARGETS)
+        list(APPEND POOLING_TEST_GPU_TARGETS ${target})
+        message(STATUS "  Adding GPU target for pooling tests: ${target}")
+    endif()
+endforeach()
+
+# Early exit if no compatible GPU architectures are available
+if(NOT POOLING_TEST_GPU_TARGETS)
+    message(WARNING "Skipping Pooling Tile Engine tests: No supported GPU targets (gfx90a, gfx942) found in SUPPORTED_GPU_TARGETS: ${SUPPORTED_GPU_TARGETS}")
+    return()
+endif()
+
+message(STATUS "Building Pooling tile engine tests for GPU targets: ${POOLING_TEST_GPU_TARGETS}")
+
+# Enable parallel compilation optimizations
+set_property(GLOBAL PROPERTY JOB_POOLS
+    compile_heavy=4
+    compile_normal=16
+)
+
+# Enable compiler cache if available and explicitly requested
+option(ENABLE_CCACHE_TESTS "Enable ccache for test compilation" OFF)
+if(ENABLE_CCACHE_TESTS)
+    find_program(CCACHE_PROGRAM ccache)
+    if(CCACHE_PROGRAM)
+        set(CMAKE_CXX_COMPILER_LAUNCHER ${CCACHE_PROGRAM})
+        message(STATUS "Using ccache for faster test compilation")
+    else()
+        message(WARNING "ccache requested but not found")
+    endif()
+else()
+    message(STATUS "ccache disabled for tests (use -DENABLE_CCACHE_TESTS=ON to enable)")
+endif()
+
+# ============================================================================
+# Test Configuration Matrix
+# ============================================================================
+
+set(TEST_DATATYPES "fp16;fp32")
+
+# ============================================================================
+# Test Target Generation
+# ============================================================================
+
+# SIMPLE TEST: Basic functionality validation
+set(SIMPLE_TEST_CONFIG "simple_test_config")
+set(SIMPLE_TEST_CONFIG_FILE "${CMAKE_CURRENT_SOURCE_DIR}/configs/${SIMPLE_TEST_CONFIG}.json")
+
+if(EXISTS ${SIMPLE_TEST_CONFIG_FILE})
+    message(STATUS "Processing pooling simple test config: ${SIMPLE_TEST_CONFIG}")
+    foreach(datatype IN LISTS TEST_DATATYPES)
+        build_pool_test_targets("${datatype}" "${SIMPLE_TEST_CONFIG}")
+    endforeach()
+else()
+    message(WARNING "Pooling simple test config file not found: ${SIMPLE_TEST_CONFIG_FILE}")
+endif()
+
+message(STATUS "Pooling tile engine tests configured:")
+message(STATUS "  - Simple test: fp16/fp32")

test/ck_tile/pooling_tile_engine/README.md

@@ -0,0 +1,88 @@
+# Pooling Tile Engine Tests
+
+Unit tests for pooling kernels generated by the tile_engine pooling codegen system.
+
+## Overview
+
+These tests validate pooling kernels that are generated at CMake configuration time
+by `pooling_instance_builder.py`. Each kernel configuration (tile shape + traits)
+gets its own GTest executable that verifies correctness against a CPU reference
+implementation.
+
+## Architecture
+
+
+```
+test/ck_tile/pooling_tile_engine/
+├── CMakeLists.txt                # Build infrastructure
+├── configs/
+│   └── simple_test_config.json   # Test configuration with problem sizes
+├── extract_test_params.py        # Extracts problem sizes to C++ header
+├── test_pooling_simple.cpp       # GTest driver (parameterized)
+└── README.md                     # This file
+```
+
+### Build Flow
+
+1. **CMake configuration**: `CMakeLists.txt` invokes `pooling_instance_builder.py --list_kernels`
+   to discover valid kernel configurations from the JSON config.
+2. **Parameter extraction**: `extract_test_params.py` generates `test_params.hpp` with
+   problem sizes from the JSON config.
+3. **Header generation**: For each kernel, `pooling_instance_builder.py --gen_single`
+   generates a C++ header defining `SelectedKernel` with the specific tile configuration.
+4. **Compilation**: Each kernel gets a separate test executable compiled with the
+   generated header via `-include`.
+5. **Execution**: GTest runs each problem size as a separate test case, comparing
+   device results against the CPU reference.
+
+## Configuration
+
+### `simple_test_config.json`
+
+Defines:
+- **tile_config**: Block/warp/thread tile dimensions for PoolShape
+- **trait_config**: Reduce op (max/avg), output_index, propagate_nan, pooling_dim (2d/3d)
+- **test_params**: Problem sizes (N, H, W, C, window, stride, dilation, padding)
+
+### Supported configurations
+
+- **Data types**: fp16, fp32
+- **Reduce operations**: max (with index output)
+- **Pooling dimensions**: 2D (NHWC), 3D (NDHWC)
+- **GPU targets**: gfx90a, gfx942
+
+## Building
+
+```bash
+# From the build directory:
+cmake --build . --target test_pooling_tile_engine_fp16_simple_test_config_max_true_false_2d_128x1_1x1_128x1_2x1
+
+# Or build all pooling tests:
+cmake --build . --target tests
+```
+
+## Running
+
+```bash
+# Run a specific test:
+./test_pooling_tile_engine_fp16_simple_test_config_max_true_false_2d_128x1_1x1_128x1_2x1
+
+# Run with GTest filters:
+./test_pooling_tile_engine_fp16_simple_test_config_max_true_false_2d_128x1_1x1_128x1_2x1 --gtest_filter="*BasicFunctionality*"
+```
+
+## Relationship to tile_engine
+
+The tile_engine pooling op lives at `tile_engine/ops/pooling/` and provides:
+- `pooling_instance_builder.py` - Codegen for kernel headers
+- `pooling_validation_utils.py` - Configuration validation
+- `pooling_common.hpp` - Shared trait definitions
+- `pooling_benchmark.hpp` - Problem/metric definitions
+- `pooling_profiler.hpp` - Benchmark profiling
+- `pooling_benchmark_single.cpp` - Single-kernel benchmark entry point
+
+The underlying ck_tile pooling kernel lives at `include/ck_tile/ops/pooling/` and provides:
+- `PoolKernel` - GPU kernel implementation
+- `PoolProblem` - Problem parameterization
+- `PoolShape` - Tile shape specification
+- `PoolDefaultPolicy` - Tile distribution and reduction policies