Files
composable_kernel/experimental/builder
Ville Pietilä 91f7d4ac75 [CK_BUILDER] Forward convolution builder improvements (#3179)
Proposed changes
Improve the forward convolution builder implementation and addressed leftover feedback left from PR #3138. Main changes

Refactored tests such that they reflect better the builder pattern. The templates and types for the convolution algorithm concepts are created via factory that facilitates programmatic creation of the device op instances.
Moved tests into anonymous namespace.
The convolution factory had lot of if-else constructs when CK Builder types were converted into CK library types. I had initially trouble in using static_assert in the default branch of switch as the static_assert was evaluated at compile time even for valid types. However, if we change the static_assert to throw "<error message>", it will result in a compile-time error only if the default branch is actually hit. This assumes that the function is consteval. Hence, changed all conversions in the convolution factory to use switch, which is more intuitive.
Removed the explicit device op definition from convolution signature and the corresponding predicate file. The device ops are defined by the corresponding concepts. This allowed to remove lot of boilerplate code from the convolution factory.
Adde inheritance and convolution algorithm specialization to handle device ops that are specialization of a more generic ones. The large tensor support is more naturally expressed by this pattern.
Added support for the FP8 data type.

* WIP: Builder for expected test results.

* Improve ckb fwd conv instance tests.

* clang-format

* Change if-else statements into switch in conv factory.

* Fix clang-formatting.

* Removed unnecessary includes.

* Added missing copyright.

* Remove explicit device op flag from from convolution signature.

* Add missing concept.

* Fix build.

* clang-format

* Add test for building conv fwd FP8 instances.

* Add missing header to instance traits.

* Clean-up recently added instances.

* Introduce inheritance and specialization.

* Use builder to build conv algorithm templates and types.

* clang-format

* Fix conv description tests.

---------

Co-authored-by: John Shumway <john.shumwayjr@gmail.com>

[ROCm/composable_kernel commit: 7d57bc169f]
2025-11-13 08:47:25 -08:00
..

Builder

This directory contains the experimental builder feature for composable_kernel.

  • Status: In development (October - November 2025)

Overview

The builder provides a high-level, semantically-clear interface for constructing composable kernel operations, with an initial focus on convolution kernels for MIOpen. It leverages modern C++20 features (such as POD structs as non-type template parameters, concepts, and designated initializers) to simplify kernel instantiation and improve developer experience.

This project is a prototype for a more general builder pattern for all of composable_kernel (CK) and CKTile, but is currently limited to formalizing the interface between MIOpen and CK.

Directory Structure

  • include/ck_tile/builder/
    Core builder headers and public API.
  • test/
    Unit tests and example usage of the builder pattern.
  • CMakeLists.txt
    CMake configuration for building the experimental builder and its tests.

CMake Configuration

To enable the experimental builder, configure your build with:

cmake                                                                                             \
  -D CMAKE_PREFIX_PATH=/opt/rocm                                                                  \
  -D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc                                                       \
  -D CMAKE_BUILD_TYPE=Release                                                                     \
  -D GPU_TARGETS="gfx942;gfx950"                                                                  \
  -D CK_EXPERIMENTAL_BUILDER=ON                                                                   \
  -D CMAKE_CXX_STANDARD=20                                                                        \
  -G Ninja                                                                                        \
  ..

Building and testing

During development, all CK Builder tests can be built with command

ninja test_ckb_all

To execute all tests, run

ls bin/test_ckb_* | xargs -n1 sh -c

Some tests involve building old CK convolution factories, which will take a long time. Hence, one might want to build only single test targets. For example

ninja test_ckb_conv_builder && bin/test_ckb_conv_builder

When adding new tests, please follow the convention where the CMake build target starts with a prefix test_ckb. This allows us to filter out the CK Builder tests from the set full CK repository tests. Also, the test_ckb_all target that builds all CK Builder tests relies on having the test_ckb prefix on the CMake build targets.