[Epic] Develop CK Builder as an API for CK and CK Tile within MIOpen

[assistant-librarian]

[Migrated from original issue] ROCm/composable_kernel#3574

Original issue author: @shumway

Overview

The CK Builder provides a high-level, semantically-clear interface for constructing Composable Kernel convolution operations. This experimental feature focuses on formalizing the interface between MIOpen and CK, leveraging modern C++20 features (POD structs as non-type template parameters, concepts, and designated initializers) to simplify kernel instantiation and improve developer experience.

This project serves as a prototype for a more general builder pattern that will eventually extend to all CK and CK-Tile operations.

Motivation

The current CK interface requires deep knowledge of complex template parameters, making it difficult for library users like MIOpen to instantiate kernels correctly. The builder pattern addresses this by:

• Providing compile-time validation through C++20 concepts and other best practices
• Offering clear, self-documenting interfaces with designated initializers
• Separating "what to compute" (signature) from "how to compute it" (algorithm)
• Enabling testable, strongly-typed kernel construction with meaningful compiler error messages

Status

Development Period: October 2025 - March 2026

The initial scope of CK Builder is a narrow focus on MIOpen, to deliver all supported CK kernels (convolutions) to MIOpen using CK / CK Tile as a header-only library.

The builder currently supports forward convolution kernels with various implementations (XDL, WMMA, DL, Large Tensor variants). Support for backward data and backward weight convolutions is in progress.

Requirements

• C++20 compiler support
• CMake configuration: -D CK_EXPERIMENTAL_BUILDER=ON

Documentation

• CK Builder Design
• Testing Framework

Success Criteria

MIOpen can instantiate all required convolution kernels using the builder interface, with compile-time validation ensuring configuration correctness.

Future Vision

Extend the builder pattern to all CK and CK-Tile operations beyond convolutions, making it the primary interface for library users.

[assistant-librarian]

Comment by @shumway on 2026-01-21T04:18:23Z

CK Builder Status Update for 2026.01.21

Current Status: Task Breakdown

1. ✅ Builder signature and algorithm concepts with C++20 validation
2. ✅ Factory dispatch system for all kernel variants (XDL, WMMA, DL, Large Tensor)
3. ✅ Forward convolution support (1D, 2D, 3D) with multiple data types
4. ✅ Reflection infrastructure for kernel introspection
5. ✅ Testing framework with Given-When-Then pattern and GPU validation
6. ✅ Comprehensive documentation (design, factory, testing)
7. ✅ CKTile integration foundation (header installation, initial kernel support)
8. 🔄 Backward weight convolution testing and validation
9. 🔄 WMMA kernel reflection and algorithm concept refinements
10. 🔄 CKTile kernel expansion (backward data/weight support)
11. ⏳ MIOpen end-to-end integration (Critical - escalating this week)
12. ⏳ Hardware awareness implementation (using ck_tile/core/config.hpp helpers to distinguish XDL vs WMMA)

Current Work In-Flight (5 Open PRs)

The team is actively working on completing backward convolution support and refining the builder infrastructure for both CK and CKTile kernels. Key efforts include: finalizing backward weight convolution testing (#3618), improving reference implementations for validation (#3604), and extending reflection capabilities for WMMA and backward weight kernels (#3592). Additionally, we're refactoring algorithm concepts to improve type safety and developer experience (#3567), and enhancing testing utilities to support more complex validation scenarios (#3563). These PRs represent the final pieces needed to complete the core builder functionality before focusing on MIOpen integration.

Critical Path Items

🔴 High Priority - This Week:

• MIOpen Integration: Escalating end-to-end integration work to reduce risk and enable multi-developer scaling. The CK builder side is maturing, but we need to close the loop with MIOpen to validate the interface and identify integration issues.

🟡 Medium Priority - Next 2 Weeks:

• Hardware Awareness: Implement hardware detection using helper functions in include/ck_tile/core/config.hpp to automatically distinguish between XDL and WMMA kernels based on GPU architecture.
• CKTile Expansion: Continue extending builder support to cover CKTile kernels alongside traditional CK kernels.

---

Progress: 78 PRs merged, 5 PRs open | Target Completion: March 2026

[assistant-librarian]

Comment by @shumway on 2026-01-28T04:59:57Z

CK Builder Status Update for 2026.01.28

Week's Accomplishments (Jan 20-27)

• ✅ Backward Weight Testing Complete (#3618): Added comprehensive ckt::run integration for backward weight algorithms across reference, CK, and CK Tile implementations. Introduced RunResult struct for better error handling and testing integration with SuccessfulRun() matcher. Reorganized testing structure into testing/conv/ directory with separate files for direction-specific implementations.

• ✅ Reference Implementation Improvements (#3604): Replaced reference convolution implementation with proven CK implementation to improve validation reliability and reduce maintenance burden.

• ✅ CK Tile Profiler (#3518): Added grouped convolution forward profiler for CK Tile, enabling performance analysis and optimization.

• ✅ Transfer Concepts (#3535): Implemented convolution forward transfer concepts to improve type safety and compile-time validation.

• ✅ Documentation Updates (#3573): Updated README files across the builder codebase to reflect recent architectural changes.

• ✅ CI Fix (#3581): Disabled CK Builder for SLES15 in Jenkins CI to address compatibility issues.

Current Work In-Flight (1 Open PR)

• 🔄 Validation Integration (#3649): Integrating builder/testing/validation.hpp with CK's gpu_verification.hpp to reduce code duplication. This PR introduces an iterator adapter for multi-dimensional tensor validation, cleans up the verification infrastructure with abstracted kernel launches and block reductions, and improves NaN handling in test validation. Includes updated documentation for ckt::validate().

Critical Path Items

🔴 High Priority - This Week:

• MIOpen Integration: This is critical for validating the builder interface in production and identifying any remaining integration issues.
• Validation Infrastructure: Complete the integration of validation utilities (#3649) to establish a unified testing framework across CK and CK Builder.

🟡 Medium Priority - Next 2 Weeks:

• Hardware Awareness: Implement hardware detection using helper functions in include/ck_tile/core/config.hpp to automatically distinguish between XDL and WMMA kernels based on GPU architecture.
• Backward Data Support: Extend builder support to cover backward data convolutions, completing the full convolution operation suite.

Next Week's Focus

1. Complete validation infrastructure integration (#3649)
2. Complete end-to-end builder integration in MIOpen.
3. Start hardware awareness implementation for automatic XDL/WMMA selection
4. Begin backward data convolution support

---

Progress: 84 PRs merged, 1 PR open | Target Completion: March 2026

[shumway]

CK Builder Status Update for 2026.03.04

Period Covered: Jan 28 – Mar 4 (5 weeks)

This update covers a longer period due to the composable_kernel → rocm-libraries monorepo migration (completed Feb 3). All CK Builder development now lands in rocm-libraries.

MIOpen Integration

The critical-path MIOpen integration saw major progress this period. DarylHawkinsAMD landed the first end-to-end CK Builder kernels in MIOpen (PR #4123), establishing the integration pattern behind the MIOPEN_CK_BUILDER_EXPERIMENTAL CMake flag with parity validation against legacy instantiations. This was followed by a large refactoring that ported all valid 2D grouped forward convolution instances (FP32, FP16, BFP16, I8) to use CK Builder (PR #4570), mirroring CK's code structure for easier automated conversion. Work is now in progress on BILINEAR and ADD_CLAMP elementwise operation mappings (PR #5026) to expand integration coverage. A CI test naming issue was also resolved (PR #4458).

Backward Data Support

kabrahamAMD added backward data WMMA and XDL kernels to the builder factory (PR #4582) with instance definitions, conv traits, and tests. aosewski fixed convolution forward transfer validation logic (PR #4750).

Multi-Architecture / Hardware Awareness

AnaghaRaoAMD opened PR #4955 implementing CMake-time architecture-based kernel filtering — XDL kernels compile only for CDNA targets (gfx908/gfx90a/gfx942/gfx950), WMMA kernels only for RDNA (gfx1030/gfx1100/gfx1200/gfx1201). This addresses both the hardware awareness task and build time optimization, and directly supports RFC0008 multi-arch packaging.

Infrastructure & Code Quality

• ✅ Fixed staging-compiler errors in CK builder code (PR [CK_BUILDER] Fix two staging-compiler errors in CK builder code #4598, @shumway )
• ✅ Moved GPU-dependent smoke tests to appropriate test targets (PR [CK_BUILDER] Move some smoke tests that require GPU #4593, @shumway )
• ✅ Updated developer notes across CK Builder source directories (PR [CK_BUILDER] Update developer notes in the CK Builder source directories #5038, @shumway )
• 🔄 Reflection traits refactoring to use .inc files across 55 files (PR [CK_BUILDER] Removing forward declarations #5053, @bidlekm)
• 🔄 TreeFormatter simplification (PR [CK_BUILDER] Simplify the TreeFormatter. #5083, @shumway )
• 🔄 ConvDescription output formatting cleanup (PR [CK_BUILDER] Clean up ConvDescription output formatting #5085, @shumway )
• 🔄 Testing-specific descriptor initialization (PR [CK_BUILDER] Testing-specific descriptor initialization  #4343, @JH-Leon-KIM-AMD )
• 🔄 Warp GEMM and conv algorithm concept refactoring (PR [Issue]: TheRock CI not triggering builds for some PRs #3567, @vpietila-amd )

Current Status: Task Breakdown

1. ✅ Builder signature and algorithm concepts with C++20 validation
2. ✅ Factory dispatch system for all kernel variants (XDL, WMMA, DL, Large Tensor)
3. ✅ Forward convolution support (1D, 2D, 3D) with multiple data types
4. ✅ Reflection infrastructure for kernel introspection
5. ✅ Testing framework with Given-When-Then pattern and GPU validation
6. ✅ Comprehensive documentation (design, factory, testing)
7. ✅ CKTile integration foundation (header installation, initial kernel support)
8. ✅ Backward weight convolution testing and validation
9. 🔄 WMMA kernel reflection and algorithm concept refinements (PR [Issue]: TheRock CI not triggering builds for some PRs #3567)
10. 🔄 CKTile kernel expansion — backward data kernels added to factory (PR [CK_Builder] added bwd data kernels to builder factory #4582), further expansion ongoing
11. ✅ MIOpen end-to-end integration — First integration merged (PR [MIOpen] First set of kernels using CK Builder end to end #4123), most 2D fwd convolutions ported (PR [MIOpen][CK Builder] Refactored CK Builder integration and ported most 2D fwd convs to use CK Builder within MIOpen #4570). Elementwise operations in progress (PR [CK BUILDER] Add BILINEAR and ADD_CLAMP elementwise operation mappings to CK builder #5026).
12. 🔄 Hardware awareness / multi-architecture — Architecture-based kernel filtering in progress (PR [MIOpen] CK Builder Multi-Architecture Support #4955)

---

Progress: 93 PRs merged, 7 PRs open | Target Completion: March 2026