Skip to content

Commit 13fa245

Browse files
committed
Add describe() method to device ops for runtime introspection
Introduces a polymorphic describe() method to BaseOperator that enables runtime introspection of kernel configurations through a unified interface. Key changes: - Add virtual describe() method to BaseOperator returning Description objects - Implement describe() in 6 device operation classes (conv fwd/bwd variants) - Create conv_describe.hpp with factory function for ConvDescription - Extract type definitions to conv_types.hpp to resolve circular dependencies - Add InstanceStringDescription for kernels without full ConvDescription support Improvements: - Add ODD_C to ConvFwdSpecialization enum and fix OddC mapping - Add missing types and missing elementwise ops - Add missing layouts to conv_trits.hpp - Replace silent fallbacks in conv_traits.hpp to compile-time errors - Remove circular dependency include from conv_traits.hpp - Update tests to use describe() instead of GetInstanceString() TODO: The enums for types and elementwise ops should be simplified with a better design, but that's out of scope for this PR. This provides a foundation for runtime kernel introspection and better tooling support for analyzing and debugging kernel configurations.
1 parent ce99cab commit 13fa245

22 files changed

+567
-229
lines changed

experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -153,6 +153,7 @@ consteval ck::tensor_operation::device::ConvolutionForwardSpecialization SetFwdC
153153
case ConvFwdSpecialization::FILTER_1X1_PAD0: return ck_conv_spec::Filter1x1Pad0;
154154
case ConvFwdSpecialization::FILTER_1X1_STRIDE1_PAD0: return ck_conv_spec::Filter1x1Stride1Pad0;
155155
case ConvFwdSpecialization::FILTER_3x3: return ck_conv_spec::Filter3x3;
156+
case ConvFwdSpecialization::ODD_C: return ck_conv_spec::OddC;
156157
default: throw "Unknown ConvFwdSpecialization";
157158
}
158159
}
Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
1+
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
2+
// SPDX-License-Identifier: MIT
3+
4+
/// @file
5+
/// @brief Implementation of the describe() function template for convolution kernels
6+
7+
#pragma once
8+
9+
#include "ck_tile/builder/reflect/conv_description.hpp"
10+
#include "ck_tile/builder/reflect/conv_traits.hpp"
11+
12+
namespace ck_tile::reflect {
13+
14+
/// @brief Factory function to create ConvDescription from a convolution instance type
15+
/// @tparam Instance The convolution instance type (must have ConvTraits)
16+
/// @return A ConvDescription object populated with the instance's configuration details
17+
template <conv::HasConvTraits Instance>
18+
conv::ConvDescription describe()
19+
{
20+
using Traits = conv::ConvTraits<Instance>;
21+
22+
return conv::ConvDescription(
23+
conv::ConvSignatureInfo{
24+
.spatial_dim = Traits::spatial_dim,
25+
.direction = Traits::direction,
26+
.input_layout = Traits::layout[0],
27+
.weight_layout = Traits::layout[1],
28+
.output_layout = Traits::layout[2],
29+
.data_type = Traits::data_type,
30+
.input_element_op = Traits::input_element_op,
31+
.weight_element_op = Traits::weight_element_op,
32+
.output_element_op = Traits::output_element_op,
33+
},
34+
conv::GemmAlgorithmInfo{
35+
.thread_block_size = Traits::thread_block_size,
36+
.tile_dims = Traits::tile_dims,
37+
.warp_gemm = Traits::warp_gemm,
38+
.a_tile_transfer = Traits::a_tile_transfer,
39+
.b_tile_transfer = Traits::b_tile_transfer,
40+
.c_tile_transfer = Traits::c_tile_transfer,
41+
.pipeline_version = Traits::pipeline_version,
42+
.pipeline_scheduler = Traits::pipeline_scheduler,
43+
.conv_specialization = Traits::conv_specialization,
44+
.padding = Traits::gemm_padding,
45+
},
46+
[]() { return reflect::instance_string<Instance>(); });
47+
}
48+
49+
} // namespace ck_tile::reflect

experimental/builder/include/ck_tile/builder/reflect/conv_description.hpp

Lines changed: 2 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@
2525
#include <functional>
2626

2727
#include <ck_tile/builder/conv_signature_concepts.hpp>
28-
#include <ck_tile/builder/reflect/conv_traits.hpp>
28+
#include <ck_tile/builder/reflect/conv_types.hpp>
2929
#include <ck_tile/builder/reflect/description.hpp>
3030
#include <ck_tile/builder/reflect/instance_traits.hpp>
3131
#include <ck_tile/builder/reflect/tree_formatter.hpp>
@@ -249,41 +249,7 @@ class ConvDescription : public Description
249249
GemmAlgorithmInfo algorithm_;
250250
std::function<std::string()> instance_string_getter_;
251251
};
252-
} // namespace conv
253-
254-
/// @brief Factory function to create ConvDescription from a convolution instance type
255-
/// @tparam Instance The convolution instance type (must have ConvTraits specialization)
256-
/// @return A ConvDescription object populated with the instance's configuration details
257-
template <conv::HasConvTraits Instance>
258-
conv::ConvDescription describe()
259-
{
260-
using Traits = conv::ConvTraits<Instance>;
261252

262-
return conv::ConvDescription(
263-
conv::ConvSignatureInfo{
264-
.spatial_dim = Traits::spatial_dim,
265-
.direction = Traits::direction,
266-
.input_layout = Traits::layout[0],
267-
.weight_layout = Traits::layout[1],
268-
.output_layout = Traits::layout[2],
269-
.data_type = Traits::data_type,
270-
.input_element_op = Traits::input_element_op,
271-
.weight_element_op = Traits::weight_element_op,
272-
.output_element_op = Traits::output_element_op,
273-
},
274-
conv::GemmAlgorithmInfo{
275-
.thread_block_size = Traits::thread_block_size,
276-
.tile_dims = Traits::tile_dims,
277-
.warp_gemm = Traits::warp_gemm,
278-
.a_tile_transfer = Traits::a_tile_transfer,
279-
.b_tile_transfer = Traits::b_tile_transfer,
280-
.c_tile_transfer = Traits::c_tile_transfer,
281-
.pipeline_version = Traits::pipeline_version,
282-
.pipeline_scheduler = Traits::pipeline_scheduler,
283-
.conv_specialization = Traits::conv_specialization,
284-
.padding = Traits::gemm_padding,
285-
},
286-
[]() { return reflect::instance_string<Instance>(); });
287-
}
253+
} // namespace conv
288254

289255
} // namespace ck_tile::reflect

0 commit comments

Comments
 (0)