# 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