[CK_BUILDER] Clean up conv_traits.hpp #3354
Conversation
We were getting ugly error messages when we tried to call `describe<Instance>` on an instance that is not a supported (forward XDL) kernel. * Fix concepts and requires clauses to get much better error messages. * Also clean up functions in conv_traits.hpp to make the more concise and readable.
shumway
requested review from
ThomasNing,
afagaj,
andriy-ca,
aosewski,
asleepzzz,
bartekxk,
carlushuang,
cgmillette,
coderfeli,
geyyer,
illsilin,
poyenc,
qianfengz,
tenpercent and
vidyasagar-amd
as code owners
There was a problem hiding this comment.
Pull request overview
This PR enhances error reporting and code readability in conv_traits.hpp by introducing C++20 concepts and refactoring several functions. The changes aim to provide clearer, more descriptive error messages when using unsupported kernel instances with the builder API.
Key changes:
- Added six C++20 concepts (
HasFwdConvLayouts,HasGemmSpec,HasDataTypes,HasElementwiseOps,HasTileParams,IsXdlFwdConv, andHasConvTraits) to validate kernel instance properties and provide better compiler diagnostics - Refactored multiple functions to use switch statements instead of if-constexpr chains for improved readability
- Applied requires clauses to functions that need specific traits, constraining them at the correct level
Reviewed changes
Copilot reviewed 2 out of 2 changed files in this pull request and generated 10 comments.
| File | Description |
|---|---|
| experimental/builder/include/ck_tile/builder/reflect/conv_traits.hpp | Added concepts for trait validation, refactored conversion and extraction functions to use switch statements, added requires clauses to constrain template functions |
| experimental/builder/include/ck_tile/builder/reflect/conv_description.hpp | Updated to use the new HasConvTraits concept from conv namespace instead of local definition |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
| if constexpr(layouts_are<A, B, E, ctl::NGCDHW, ctl::GKCZYX, ctl::NGKDHW>) | ||
| return layouts(NGCDHW, GKCZYX, NGKDHW); | ||
| break; | ||
| } |
There was a problem hiding this comment.
The function may not return a value in all code paths. If none of the layout combinations match or if kSpatialDim is not 1, 2, or 3, the function will have undefined behavior. Consider adding a default return value or static_assert to ensure all cases are handled.
| } | |
| } | |
| static_assert(always_false<Instance>, "Unsupported convolution layout or spatial dimension in conv_layout()"); |
| case MNOPadding: return MNO_PADDING; | ||
| case MKOPadding: return MKO_PADDING; | ||
| case NKOPadding: return NKO_PADDING; | ||
| case MNKOPadding: return MNKO_PADDING; |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. If gemm_spec doesn't match any of the enumerated cases, the function will have undefined behavior. Consider adding a default case to handle unexpected values.
| case MNKOPadding: return MNKO_PADDING; | |
| case MNKOPadding: return MNKO_PADDING; | |
| default: | |
| static_assert([]{return false;}(), "Unhandled GemmSpecialization value in gemm_spec()"); |
| if constexpr(detail::case_insensitive_equal(name, "PassThrough")) | ||
| return PASS_THROUGH; | ||
| if constexpr(detail::case_insensitive_equal(name, "ScaleAddScaleAddRelu")) | ||
| return SCALEADD_SCALEADD_RELU; |
There was a problem hiding this comment.
The function may not return a value in all code paths. If none of the elementwise operation names match, the function will have undefined behavior. Consider adding a default return value or ensuring all cases are exhaustively handled.
| return SCALEADD_SCALEADD_RELU; | |
| return SCALEADD_SCALEADD_RELU; | |
| return PASS_THROUGH; // Default fallback if no match found |
| case v2: return V2; | ||
| case v3: return V3; | ||
| case v4: return V4; | ||
| case v5: return V5; |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. If ck_ver doesn't match any of the enumerated cases (v1-v5), the function will have undefined behavior. Consider adding a default case or ensuring all enum values are covered.
| case v5: return V5; | |
| case v5: return V5; | |
| default: | |
| static_assert(ck_ver == v1 || ck_ver == v2 || ck_ver == v3 || ck_ver == v4 || ck_ver == v5, | |
| "Unsupported ck::BlockGemmPipelineVersion in convert_pipeline_version"); | |
| // The following is just to satisfy the compiler; this line should never be reached. | |
| return V1; |
| case v1: return V1; | ||
| case v2: return V2; | ||
| case v4: return V4; | ||
| case weight_only: return WEIGHT_ONLY; |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. If ck_ver doesn't match any of the enumerated cases (v1, v2, v4, weight_only), the function will have undefined behavior. Consider adding a default case or ensuring all enum values are covered.
| case weight_only: return WEIGHT_ONLY; | |
| case weight_only: return WEIGHT_ONLY; | |
| default: | |
| static_assert(ck_ver == v1 || ck_ver == v2 || ck_ver == v4 || ck_ver == weight_only, | |
| "Unsupported ck::PipelineVersion value in convert_pipeline_version."); | |
| // To satisfy all code paths, return a dummy value (will never be used due to static_assert) | |
| return V1; |
| switch(ck_sched) | ||
| { | ||
| case Intrawave: return INTRAWAVE; | ||
| case Interwave: return INTERWAVE; |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. If ck_sched doesn't match any of the enumerated cases (Intrawave, Interwave), the function will have undefined behavior. Consider adding a default case or ensuring all enum values are covered.
| case Interwave: return INTERWAVE; | |
| case Interwave: return INTERWAVE; | |
| default: | |
| static_assert( | |
| ck_sched == Intrawave || ck_sched == Interwave, | |
| "Unsupported ck::BlockGemmPipelineScheduler value in convert_pipeline_scheduler" | |
| ); | |
| // The following is unreachable, but required to satisfy all code paths. | |
| return INTRAWAVE; |
| switch(InstTraits::kConvForwardSpecialization) | ||
| { | ||
| return builder::ConvFwdSpecialization::FILTER_3x3; | ||
| case Default: return DEFAULT; | ||
| case Filter1x1Pad0: return FILTER_1X1_PAD0; | ||
| case Filter1x1Stride1Pad0: return FILTER_1X1_STRIDE1_PAD0; | ||
| case Filter3x3: return FILTER_3x3; | ||
| } |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. When the specialization doesn't match any case, the function will have undefined behavior. Consider adding a default case to handle unexpected values.
| switch(ck_sched) | ||
| { | ||
| case Default: return DEFAULT; | ||
| case Interwave: return INTERWAVE; |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. If ck_sched doesn't match any of the enumerated cases (Default, Interwave), the function will have undefined behavior. Consider adding a default case or ensuring all enum values are covered.
| case Interwave: return INTERWAVE; | |
| case Interwave: return INTERWAVE; | |
| default: | |
| static_assert(ck_sched == Default || ck_sched == Interwave, | |
| "Unsupported ck::LoopScheduler value in convert_pipeline_scheduler"); | |
| // To satisfy all code paths, return a value (will never be reached) | |
| return DEFAULT; |
| switch(InstTraits::kConvBwdDataSpecialization) | ||
| { | ||
| return builder::ConvBwdDataSpecialization::DEFAULT; | ||
| } | ||
| else if constexpr(InstTraits::kConvBwdDataSpecialization == Filter1x1Stride1Pad0) | ||
| { | ||
| return builder::ConvBwdDataSpecialization::FILTER_1X1_STRIDE1_PAD0; | ||
| case Default: return DEFAULT; | ||
| case Filter1x1Stride1Pad0: return FILTER_1X1_STRIDE1_PAD0; | ||
| } |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. When the specialization doesn't match any case, the function will have undefined behavior. Consider adding a default case to handle unexpected values.
| switch(InstTraits::kConvBwdWeightSpecialization) | ||
| { | ||
| return builder::ConvBwdWeightSpecialization::DEFAULT; | ||
| } | ||
| else if constexpr(InstTraits::kConvBwdWeightSpecialization == Filter1x1Stride1Pad0) | ||
| { | ||
| return builder::ConvBwdWeightSpecialization::FILTER_1X1_STRIDE1_PAD0; | ||
| } | ||
| else if constexpr(InstTraits::kConvBwdWeightSpecialization == Filter1x1Pad0) | ||
| { | ||
| return builder::ConvBwdWeightSpecialization::FILTER_1X1_PAD0; | ||
| } | ||
| else if constexpr(InstTraits::kConvBwdWeightSpecialization == OddC) | ||
| { | ||
| return builder::ConvBwdWeightSpecialization::ODD_C; | ||
| case Default: return DEFAULT; | ||
| case Filter1x1Stride1Pad0: return FILTER_1X1_STRIDE1_PAD0; | ||
| case Filter1x1Pad0: return FILTER_1X1_PAD0; | ||
| case OddC: return ODD_C; | ||
| } |
There was a problem hiding this comment.
The switch statement lacks a default case and may not return a value in all code paths. When the specialization doesn't match any case, the function will have undefined behavior. Consider adding a default case to handle unexpected values.
We were getting ugly error messages when we tried to call
describe<Instance>on an instance that is not a supported (forward XDL) kernel.Changes in this PR:
When I asked for a description of operators that didn't have ConvTraits, I was getting very long confusing errors about ConvTraits not being defined. Now we get specific errors explaining which concepts are violated, making it easier to know which code to generalize or update.