017d96faaa
* Grouped Conv Bwd Weight Direct Load
* Update gridwise_gemm_xdl_cshuffle_conv_v3.hpp
* Implement group merging for bwd_weight and add instances
* Link direct load instances
* builder fixes
* fix
* fixes
* fix
---------
Co-authored-by: Graner, Johannes <johannes.graner@amd.com>
[ROCm/composable_kernel commit: 83b58bb0c3]
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.