composable_kernel/test/ck_tile/pooling_tile_engine

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

# 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

# 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