2b54a86c04
* ck-builder: remove SPATIAL_DIM parameter from ConvTensorLayouts
This information is already in the SIGNATURE, so its pointless to pass it
separately. This streamlines the interface of those functions a bit. Also
touches up the style of those files in general.
* ck-builder: implement reference conv using old ck
The old ck implementation is more featureful and better tested.
* ck-builder: replace test_reference_execution reference with old ck
This strips out the ck-tile gpu reference implementation completely.
* ck-builder: clean up test_reference_execution
- Remove unneccesary messages
- Replace EXPECT_TRUE(true) with EXPECT_NO_THROW()
[ROCm/composable_kernel commit: 1040d9b1f5]
Convolution Builder Factory Directory
This directory implements compile-time dispatch from high-level signature algorithm descriptors to our exisitng specialized convolution kernel implementations.
See the main builder documentation for an overview.
Design Overview
The factory system operates in two phases:
-
Algorithm Classification: The function
make_conv_instanceinconv_dispatcher.hppinspects the signature and algorithm descriptors to determine which kernel variant they satisfy (XDL V3, XDL, WMMA, DL, or Large Tensor) -
Factory Instantiation: Each factory (
conv_fwd_*_factory.hpp) transforms builder descriptors into CK device operation template parameters and instantiates the corresponding kernel device operation.
Key Files
-
conv_dispatcher.hpp: Entry point withmake_conv_instance()function. Contains dispatch logic and algorithm classification predicates. Start here to understand the overall flow. -
conv_fwd_*_factory.hpp: Individual factories for each kernel variant. Each extracts configuration from descriptors, validates parameters, and instantiates the underlying CK device operation. -
helpers/: Transformation utilities that map builder types to CK device operation parameters (layouts, data types, elementwise ops, block configurations, etc.)
Usage
#include "ck_tile/builder/factory/conv_dispatcher.hpp"
using Factory = decltype(make_conv_instance<signature, algorithm, "v1">());
The dispatcher automatically selects the appropriate factory following explicit logic.