Skip to content

Commit 20f66c1

Browse files
adressed review comments from PR3459 (#3526)
Co-authored-by: Kevin Abraham <kevin.abraham@streamhpc.com>
1 parent b352a68 commit 20f66c1

2 files changed

Lines changed: 20 additions & 11 deletions

File tree

experimental/builder/include/ck_tile/builder/README.md

Lines changed: 7 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -85,21 +85,23 @@ The top-level signature contains global properties that apply to the entire conv
8585
template <typename T>
8686
concept ConvSignatureDescriptor = requires(T t) {
8787
{ t.spatial_dim } -> std::convertible_to<unsigned int>; // 1, 2, or 3
88-
{ t.data_type } -> std::convertible_to<DataType>; // Default data type
8988
{ t.input } -> ConvTensorDescriptor;
9089
{ t.weight } -> ConvTensorDescriptor;
9190
{ t.output } -> ConvTensorDescriptor;
9291
requires ConvolutionDirectionWellDefinedIfProvided<T>; // Optional direction
92+
requires detail::DataTypeWellDefinedIfProvided<T>; // Optional default data type
93+
requires detail::ElementwiseOpWellDefinedIfProvided<T>; // Optional default elementwise operation
9394
};
9495
```
9596

9697
**Properties:**
9798
- **`spatial_dim`**: Dimensionality of the convolution (1D, 2D, or 3D)
98-
- **`direction`**: Operation type (optional, defaults to FORWARD)
99+
- **`direction`**: Operation type (Optional, defaults to FORWARD)
99100
- `FORWARD`: Standard forward convolution
100101
- `BACKWARD_DATA`: Gradient computation w.r.t. input
101102
- `BACKWARD_WEIGHT`: Gradient computation w.r.t. weights
102-
- **`data_type`**: Default data type for all tensors (FP32, FP16, BF16, FP8, I8, U8)
103+
- **`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)
104+
- **`operation`**: Default Operation (Optional, defaults to PASS_THROUGH, may be overridden by tensors)
103105
- **`accumulation_data_type`**: Type used for internal accumulation
104106

105107
#### 2. Tensor Level
@@ -116,7 +118,7 @@ concept ConvTensorDescriptor = requires(T t) {
116118

117119
A tensor descriptor encapsulates:
118120
- **Configuration**: Layout and data type information
119-
- **Operation** (optional): Fused elementwise operations on this tensor
121+
- **operation** Fused elementwise operations on this tensor (Optional, default provided by ConvSignatureDescriptor)
120122

121123
#### 3. Tensor Configuration
122124

@@ -126,7 +128,7 @@ Describes the memory layout and data types:
126128
template <typename T>
127129
concept TensorConfigDescriptor = requires(T t) {
128130
{ t.layout } -> std::convertible_to<ConvLayout>;
129-
{ t.data_type } -> std::convertible_to<DataType>; // Optional override
131+
requires detail::DataTypeWellDefinedIfProvided<T>; // Override data type (Optional, default provided by ConvSignatureDescriptor)
130132
};
131133
```
132134

experimental/builder/include/ck_tile/builder/conv_signature_concepts.hpp

Lines changed: 13 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -80,6 +80,7 @@ concept ConvOutputLayout3D =
8080
(L == TensorLayout::GNKDHW) || (L == TensorLayout::GNDHWK) || (L == TensorLayout::NDHWGK) ||
8181
(L == TensorLayout::NGKDHW) || (L == TensorLayout::G_NDHW_K_strided);
8282

83+
namespace detail {
8384
template <typename T>
8485
concept HasDataType = requires(T t) {
8586
{ t.data_type };
@@ -94,10 +95,11 @@ concept DataTypeWellDefinedIfProvided = requires(T t) {
9495
};
9596
};
9697

98+
} // namespace detail
9799
template <typename T>
98100
concept TensorConfigDescriptor = requires(T t) {
99101
{ t.layout } -> std::convertible_to<TensorLayout>;
100-
requires DataTypeWellDefinedIfProvided<T>;
102+
requires detail::DataTypeWellDefinedIfProvided<T>;
101103
};
102104

103105
template <typename T>
@@ -116,7 +118,6 @@ template <typename T, std::size_t N>
116118
struct IsArrayOfTensorConfigDescriptors<std::array<T, N>> : std::true_type
117119
{
118120
};
119-
} // namespace detail
120121

121122
template <typename T>
122123
concept ConvertibleToArrayOfTensorConfigs =
@@ -128,18 +129,21 @@ concept AuxiliaryOperandConfigsWellDefinedIfProvided = requires(T t) {
128129
{ t.auxiliary_operand_configs } -> ConvertibleToArrayOfTensorConfigs;
129130
};
130131
};
132+
} // namespace detail
131133

132134
template <typename T>
133135
concept TensorOperatorDescriptor = requires(T t) {
134136
{ t.elementwise_operation } -> std::convertible_to<ElementwiseOperation>;
135-
requires AuxiliaryOperandConfigsWellDefinedIfProvided<T>;
137+
requires detail::AuxiliaryOperandConfigsWellDefinedIfProvided<T>;
136138
};
137139

138140
template <typename T>
139141
concept HasTensorOp = requires(T t) {
140142
{ t.operation };
141143
};
142144

145+
namespace detail {
146+
143147
template <typename T>
144148
concept HasConvolutionDirection = requires(T t) {
145149
{ t.direction };
@@ -159,11 +163,13 @@ concept ConvolutionDirectionWellDefinedIfProvided = requires(T t) {
159163
};
160164
};
161165

166+
} // namespace detail
167+
162168
// Concept for the convolution tensor
163169
template <typename T>
164170
concept ConvTensorDescriptor = requires(T t) {
165171
{ t.config } -> TensorConfigDescriptor;
166-
requires ElementwiseOpWellDefinedIfProvided<T>;
172+
requires detail::ElementwiseOpWellDefinedIfProvided<T>;
167173
};
168174

169175
template <typename T>
@@ -179,8 +185,9 @@ concept ConvSignatureDescriptor = requires(T t) {
179185
{ t.input } -> ConvTensorDescriptor;
180186
{ t.weight } -> ConvTensorDescriptor;
181187
{ t.output } -> ConvTensorDescriptor;
182-
requires ConvolutionDirectionWellDefinedIfProvided<T>;
183-
requires DataTypeWellDefinedIfProvided<T>;
188+
requires detail::ConvolutionDirectionWellDefinedIfProvided<T>;
189+
requires detail::DataTypeWellDefinedIfProvided<T>;
190+
requires detail::ElementwiseOpWellDefinedIfProvided<T>;
184191
};
185192

186193
// Concept to validate a convolution signature's values.

0 commit comments

Comments
 (0)