Merge pull request #3573 from ROCm/jshumway/builder-readme

Author: shumway

experimental/builder/README.md

@@ -2,13 +2,13 @@
 
 This directory contains the experimental builder feature for composable_kernel.
 
-* Status: In development (October - December 2025)
+* Status: In development (October 2025 - March 2026)
 
 ## 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.
+This project is a prototype for a more general builder pattern for all of composable_kernel (CK) and CK Tile, but is currently limited to formalizing the interface between MIOpen and CK.
 
 ## Design descriptions
 

experimental/builder/include/ck_tile/builder/README.md

@@ -100,8 +100,8 @@ concept ConvSignatureDescriptor = requires(T t) {
   - `FORWARD`: Standard forward convolution
   - `BACKWARD_DATA`: Gradient computation w.r.t. input
   - `BACKWARD_WEIGHT`: Gradient computation w.r.t. weights
-- **`data_type`**: Default data type for all tensors (FP32, FP16, BF16, FP8, I8, U8). (Optional, defaults to UNDEFINED_DATA_TYPE, may be overridden by tensors)
-- **`operation`**: Default Operation (Optional, defaults to PASS_THROUGH, may be overridden by tensors)
+- **`data_type`**: Default data type for all tensors (FP32, FP16, BF16, FP8, I8, U8). (Optional, defaults to UNDEFINED_DATA_TYPE which indicates the type should be inferred or specified per-tensor, may be overridden by individual tensors)
+- **`elementwise_operation`**: Default elementwise operation for all tensors (Optional, defaults to PASS_THROUGH, may be overridden by individual tensors via their `operation` field)
 - **`accumulation_data_type`**: Type used for internal accumulation
 
 #### 2. Tensor Level
@@ -133,6 +133,9 @@ concept TensorConfigDescriptor = requires(T t) {
 ```
 
 **Layout Types** (dimension-specific):
+- **Special Values**:
+  - `UNDEFINED_TENSOR_LAYOUT`: Placeholder value indicating layout is not yet specified or should be inferred
+
 - **1D Convolution**:
   - Input: `GNCW`, `GNWC`, `NWGC`, `NGCW`, `G_NW_C_strided`
   - Weight: `GKXC`, `GKCX`, `KXGC`, `G_K_X_C_strided`
@@ -148,6 +151,9 @@ concept TensorConfigDescriptor = requires(T t) {
   - Weight: `GKZYXC`, `GKCZYX`, `KZYXGC`, `G_K_ZYX_C_strided`
   - Output: `GNKDHW`, `GNDHWK`, `NDHWGK`, `NGKDHW`, `G_NDHW_K_strided`
 
+- **Bias Tensors**:
+  - `GC`, `G_C_strided`, `G_K_strided`
+
 Where:
 - `G` = Groups
 - `N` = Batch size

experimental/builder/include/ck_tile/builder/testing/README.md

@@ -53,7 +53,7 @@ struct ConvSignature {
     ck_tile::builder::DataType data_type =
         ck_tile::builder::DataType::FP16;
     ck_tile::builder::ElementwiseOperation elementwise_operation =
-        ck_tile::builder::ElementwiseOperation::NONE;
+        ck_tile::builder::ElementwiseOperation::PASS_THROUGH;
 };
 
 // Double-check that out structure is well-defined according to the CK-Builder API.
@@ -66,7 +66,7 @@ constexpr auto SIGNATURE = ConvSignature{
     .direction = ck_tile::builder::ConvDirection::FORWARD,
     .layout = ck_tile::builder::GroupConvLayout2D::NHWGC_GKYXC_NHWGK,
     .data_type = ck_tile::builder::DataType::FP16,
-    .elementwise_operation = ck_tile::builder::ElementwiseOperation::NONE,
+    .elementwise_operation = ck_tile::builder::ElementwiseOperation::PASS_THROUGH,
 };
 ```
 
@@ -243,15 +243,15 @@ struct ConvSignature {
     ck_tile::builder::DataType data_type =
         ck_tile::builder::DataType::FP16;
     ck_tile::builder::ElementwiseOperation elementwise_operation =
-        ck_tile::builder::ElementwiseOperation::NONE;
+        ck_tile::builder::ElementwiseOperation::PASS_THROUGH;
 };
 static_assert(ck_tile::builder::ConvSignatureDescriptor<ConvSignature>);
 constexpr auto SIGNATURE = ConvSignature{
     .spatial_dim = 2,
     .direction = ck_tile::builder::ConvDirection::FORWARD,
     .layout = ck_tile::builder::GroupConvLayout2D::NHWGC_GKYXC_NHWGK,
     .data_type = ck_tile::builder::DataType::FP16,
-    .elementwise_operation = ck_tile::builder::ElementwiseOperation::NONE,
+    .elementwise_operation = ck_tile::builder::ElementwiseOperation::PASS_THROUGH,
 };
 
 // Define the convolution algorithm