mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-04-19 22:39:03 +00:00
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.
This commit is contained in:
@@ -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<Instance>()` 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<MyConvFwdInstance>();
|
||||
|
||||
// Print the detailed description
|
||||
std::cout << description.detailed() << std::endl;
|
||||
```
|
||||
@@ -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 <concepts>
|
||||
@@ -13,11 +29,13 @@
|
||||
#include <ck_tile/builder/reflect/conv_traits.hpp>
|
||||
#include <ck_tile/builder/reflect/tree_formatter.hpp>
|
||||
|
||||
/// @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 <typename T>
|
||||
concept HasInstanceTraits = requires { typename InstanceTraits<T>; };
|
||||
|
||||
// Helper concept to detect ConvBuilder types
|
||||
template <typename T>
|
||||
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 <typename Instance>
|
||||
requires HasInstanceTraits<Instance>
|
||||
ConvDescription Describe()
|
||||
@@ -255,14 +278,4 @@ ConvDescription Describe()
|
||||
.padding = Traits::gemm_padding}};
|
||||
}
|
||||
|
||||
// Backward compatibility: Create ConvDescription from Builder type
|
||||
template <typename Builder>
|
||||
requires IsConvBuilder<Builder> && (!HasInstanceTraits<Builder>)
|
||||
ConvDescription Describe()
|
||||
{
|
||||
// Delegate to Instance-based version
|
||||
using Instance = typename Builder::Instance;
|
||||
return Describe<Instance>();
|
||||
}
|
||||
|
||||
} // namespace ck_tile::reflect::conv
|
||||
|
||||
@@ -109,16 +109,16 @@ TEST(ConvDescriptionTest, DefaultInstanceHasBriefDescription)
|
||||
{
|
||||
static constexpr const ConvSignature SIGNATURE;
|
||||
static constexpr const DefaultAlgorithm ALGORITHM;
|
||||
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
|
||||
EXPECT_THAT(ckr::Describe<Builder>().brief(), ckt::StringEqWithDiff("2D Forward convolution"));
|
||||
using Instance = ckb::ConvBuilder<SIGNATURE, ALGORITHM>::Instance;
|
||||
EXPECT_THAT(ckr::Describe<Instance>().brief(), ckt::StringEqWithDiff("2D Forward convolution"));
|
||||
}
|
||||
|
||||
TEST(ConvDescriptionTest, DefaultInstanceHasDetailedDescription)
|
||||
{
|
||||
static constexpr const ConvSignature SIGNATURE;
|
||||
static constexpr const DefaultAlgorithm ALGORITHM;
|
||||
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
|
||||
EXPECT_THAT(ckr::Describe<Builder>().detailed(),
|
||||
using Instance = ckb::ConvBuilder<SIGNATURE, ALGORITHM>::Instance;
|
||||
EXPECT_THAT(ckr::Describe<Instance>().detailed(),
|
||||
ckt::StringEqWithDiff( //
|
||||
"2D Forward Convolution Kernel\n"
|
||||
"├─ Signature\n"
|
||||
|
||||
Reference in New Issue
Block a user