diff --git a/experimental/builder/README.md b/experimental/builder/README.md index 1156de0e9c..850bcf136e 100644 --- a/experimental/builder/README.md +++ b/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 diff --git a/experimental/builder/include/ck_tile/builder/README.md b/experimental/builder/include/ck_tile/builder/README.md index af8c4ec01b..0af0cede60 100644 --- a/experimental/builder/include/ck_tile/builder/README.md +++ b/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 diff --git a/experimental/builder/include/ck_tile/builder/testing/README.md b/experimental/builder/include/ck_tile/builder/testing/README.md index 85adc59d80..c6662c2b04 100644 --- a/experimental/builder/include/ck_tile/builder/testing/README.md +++ b/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,7 +243,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; }; static_assert(ck_tile::builder::ConvSignatureDescriptor); constexpr auto SIGNATURE = ConvSignature{ @@ -251,7 +251,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, }; // Define the convolution algorithm