From f645d827c800f78edb4a29764470dbd192ff091a Mon Sep 17 00:00:00 2001 From: John Shumway Date: Mon, 1 Dec 2025 01:03:58 -0800 Subject: [PATCH] Cleanup convolution description (#3329) Remove obsolete feature for extracting a description from a builder, since this should apply directly to the instance type. Also add some documentation, including a README.md for reflection. [ROCm/composable_kernel commit: abd6a4b3fc535772ecff047b02f2af666987f859] --- .../include/ck_tile/builder/reflect/README.md | 38 +++++++++++ .../builder/reflect/conv_description.hpp | 67 +++++++++++-------- .../builder/test/test_conv_description.cpp | 8 +-- 3 files changed, 82 insertions(+), 31 deletions(-) create mode 100644 experimental/builder/include/ck_tile/builder/reflect/README.md diff --git a/experimental/builder/include/ck_tile/builder/reflect/README.md b/experimental/builder/include/ck_tile/builder/reflect/README.md new file mode 100644 index 0000000000..6ef9ebe87a --- /dev/null +++ b/experimental/builder/include/ck_tile/builder/reflect/README.md @@ -0,0 +1,38 @@ +# Convolution Reflection Directory + +This directory contains tools for "reflecting" on convolution kernel instances. It allows developers to inspect the compile-time configuration of a kernel and generate detailed, human-readable descriptions. + +See the [main builder documentation](../README.md) for an overview. + +## Design Overview + +The reflection system works by extracting properties from a convolution kernel *type* and formatting them into a string. This is useful for debugging, performance tuning, and generating documentation. + +1. **Trait Extraction**: The `ConvTraits` template (in `conv_traits.hpp`) is specialized for each kernel instance. It extracts low-level details like tile sizes, data layouts, and pipeline versions from the kernel's type definition. + +2. **Description Generation**: The `Describe()` function (in `conv_description.hpp`) uses `ConvTraits` to populate a `ConvDescription` struct. + +3. **Formatting**: The `ConvDescription` struct contains methods like `brief()` and `detailed()` that format the extracted properties into well-structured strings for display. + +## Key Files + +- **`conv_description.hpp`**: The main entry point. Contains the `ConvDescription` struct and the `Describe()` factory function. +- **`conv_traits.hpp`**: Home of the `ConvTraits` template, which is the core of the property extraction mechanism. +- **`tree_formatter.hpp`**: A simple utility for generating the indented, tree-like format used in the `detailed()` description. + +## Usage + +To get a description of a convolution kernel instance, use the `Describe` function and call one of its formatting methods: + +```cpp +#include "ck_tile/builder/reflect/conv_description.hpp" + +// Assume MyConvFwdInstance is a type alias for a specific kernel instance +using MyConvFwdInstance = /* ... some kernel type ... */; + +// Describe the instance +const auto description = ck_tile::reflect::conv::Describe(); + +// Print the detailed description +std::cout << description.detailed() << std::endl; +``` diff --git a/experimental/builder/include/ck_tile/builder/reflect/conv_description.hpp b/experimental/builder/include/ck_tile/builder/reflect/conv_description.hpp index 08e506b614..4682f636eb 100644 --- a/experimental/builder/include/ck_tile/builder/reflect/conv_description.hpp +++ b/experimental/builder/include/ck_tile/builder/reflect/conv_description.hpp @@ -1,6 +1,22 @@ // Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT +/** + * @file + * @brief Provides utilities to reflect on convolution kernel instances and generate + * human-readable descriptions of their configuration. + * + * This file contains the necessary components to transform a convolution kernel's + * compile-time properties into a structured, descriptive format. This is primarily + * used for debugging, logging, and generating documentation. + * + * Key components: + * - ck_tile::reflect::conv::ConvDescription: A struct that holds the extracted + * properties and provides methods to format them into strings. + * - ck_tile::reflect::conv::Describe(): A factory function that creates a + * ConvDescription from a given kernel instance type. + */ + #pragma once #include @@ -13,11 +29,13 @@ #include #include -/// @file conv_description.hpp -/// @brief Provides human-readable descriptions of ConvBuilder configurations +/// @brief Provides human-readable descriptions of convolution kernel instances namespace ck_tile::reflect::conv { +/// @brief Signature information for a convolution operation +/// Contains high-level properties that define the convolution's interface, +/// including dimensionality, data layout, data types, and elementwise operations. struct ConvSignatureInfo { int spatial_dim; @@ -30,7 +48,9 @@ struct ConvSignatureInfo builder::ElementwiseOperation output_element_op; }; -// Algorithm information - groups all algorithm-related configuration +/// @brief Algorithm configuration for a convolution kernel +/// Contains low-level implementation details including thread block configuration, +/// tile dimensions, memory access patterns, and pipeline settings. struct GemmAlgorithmInfo { int thread_block_size; @@ -48,13 +68,16 @@ struct GemmAlgorithmInfo builder::GemmPadding padding; }; -// Provides human-readable descriptions of ConvBuilder configurations. +/// @brief Provides human-readable descriptions of convolution kernel instances +/// Generates formatted text descriptions at various levels of detail for +/// understanding and documenting convolution kernel configurations. struct ConvDescription { ConvSignatureInfo signature; GemmAlgorithmInfo algorithm; - // Brief one-line summary + /// @brief Generate a brief one-line summary of the convolution + /// @return A concise description (e.g., "2D Forward convolution") std::string brief() const { std::ostringstream oss; @@ -62,7 +85,8 @@ struct ConvDescription return oss.str(); } - // Detailed hierarchical description + /// @brief Generate a detailed hierarchical description of the convolution + /// @return A multi-line tree-formatted description covering signature and algorithm details std::string detailed() const { TreeFormatter f; @@ -200,7 +224,9 @@ struct ConvDescription return f.getString(); } - // Educational explanation of optimization choices + /// @brief Generate an educational explanation of optimization choices + /// @return Educational content explaining why certain algorithm choices were made + /// @note Currently unimplemented - reserved for future enhancement std::string explain() const { std::ostringstream oss; @@ -208,7 +234,9 @@ struct ConvDescription return oss.str(); } - // Performance characteristics and use case guidance + /// @brief Generate performance characteristics and use case guidance + /// @return Guidance on when this configuration is optimal and expected performance + /// @note Currently unimplemented - reserved for future enhancement std::string suggest() const { std::ostringstream oss; @@ -217,18 +245,13 @@ struct ConvDescription } }; -// Helper concept to detect if a type has InstanceTraits specialization +/// @brief Helper concept to detect if a type has InstanceTraits specialization template concept HasInstanceTraits = requires { typename InstanceTraits; }; -// Helper concept to detect ConvBuilder types -template -concept IsConvBuilder = requires { - typename T::Factory; - typename T::Instance; -}; - -// Primary factory function: Create ConvDescription from Instance type directly +/// @brief Factory function to create ConvDescription from a convolution instance type +/// @tparam Instance The convolution instance type (must have InstanceTraits specialization) +/// @return A ConvDescription object populated with the instance's configuration details template requires HasInstanceTraits ConvDescription Describe() @@ -255,14 +278,4 @@ ConvDescription Describe() .padding = Traits::gemm_padding}}; } -// Backward compatibility: Create ConvDescription from Builder type -template - requires IsConvBuilder && (!HasInstanceTraits) -ConvDescription Describe() -{ - // Delegate to Instance-based version - using Instance = typename Builder::Instance; - return Describe(); -} - } // namespace ck_tile::reflect::conv diff --git a/experimental/builder/test/test_conv_description.cpp b/experimental/builder/test/test_conv_description.cpp index 0d48a91738..b83abe9f43 100644 --- a/experimental/builder/test/test_conv_description.cpp +++ b/experimental/builder/test/test_conv_description.cpp @@ -109,16 +109,16 @@ TEST(ConvDescriptionTest, DefaultInstanceHasBriefDescription) { static constexpr const ConvSignature SIGNATURE; static constexpr const DefaultAlgorithm ALGORITHM; - using Builder = ckb::ConvBuilder; - EXPECT_THAT(ckr::Describe().brief(), ckt::StringEqWithDiff("2D Forward convolution")); + using Instance = ckb::ConvBuilder::Instance; + EXPECT_THAT(ckr::Describe().brief(), ckt::StringEqWithDiff("2D Forward convolution")); } TEST(ConvDescriptionTest, DefaultInstanceHasDetailedDescription) { static constexpr const ConvSignature SIGNATURE; static constexpr const DefaultAlgorithm ALGORITHM; - using Builder = ckb::ConvBuilder; - EXPECT_THAT(ckr::Describe().detailed(), + using Instance = ckb::ConvBuilder::Instance; + EXPECT_THAT(ckr::Describe().detailed(), ckt::StringEqWithDiff( // "2D Forward Convolution Kernel\n" "├─ Signature\n"