transform_conv_bwd_weight_to_gemm_v2.hpp Source File

transform_conv_bwd_weight_to_gemm_v2.hpp Source File#

Composable Kernel: transform_conv_bwd_weight_to_gemm_v2.hpp Source File
transform_conv_bwd_weight_to_gemm_v2.hpp
Go to the documentation of this file.
1// SPDX-License-Identifier: MIT
2// Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved.
3
4#pragma once
5
6#include "ck/library/utility/numeric.hpp"
7#include "ck/utility/common_header.hpp"
8#include "ck/tensor_description/tensor_descriptor.hpp"
9#include "ck/tensor_description/tensor_descriptor_helper.hpp"
10#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
11#include "ck/tensor_operation/gpu/device/convolution_backward_weight_specialization.hpp"
12
13namespace ck {
14namespace tensor_operation {
15
25template <index_t NDimSpatial,
26 index_t MPerBlock,
27 index_t NPerBlock,
28 index_t GemmK1Number,
29 index_t K0PerBlock,
30 index_t NumGroupsToMerge,
31 device::ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization>
32struct TransformConvBwdWeightToGemmV2
33{
34 static constexpr auto I0 = Number<0>{};
35 static constexpr auto I1 = Number<1>{};
36
37 template <index_t NDim, typename enable_if<NDim == 1, bool>::type = false>
38 constexpr static auto
39 make_out_grid_desc(const index_t N,
40 const index_t Wo,
41 const index_t K,
42 const std::array<index_t, NDimSpatial + 3>& output_strides)
43 {
44 const index_t BatchStride = output_strides[0];
45 const index_t WoStride = output_strides[3];
46 const auto KStride = Number<1>{};
47 return make_naive_tensor_descriptor(make_tuple(N * Wo, NumGroupsToMerge, K),
48 make_tuple(WoStride, BatchStride, KStride));
49 }
50
51 template <index_t NDim, typename enable_if<NDim == 1, bool>::type = false>
52 constexpr static auto
53 make_in_grid_desc(const index_t N,
54 const index_t Wi,
55 const index_t C,
56 const std::array<index_t, NDimSpatial + 3>& input_strides)
57 {
58 const index_t BatchStride = input_strides[0];
59 const index_t NStride = input_strides[1];
60 const index_t WiStride = input_strides[3];
61 const auto CStride = input_strides[2];
62 if constexpr(ConvBackwardWeightSpecialization ==
63 device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)
64 {
65 return make_naive_tensor_descriptor(make_tuple(N * Wi, NumGroupsToMerge, C),
66 make_tuple(WiStride, BatchStride, CStride));
67 }
68 else
69 {
70 return make_naive_tensor_descriptor(
71 make_tuple(N, Wi, NumGroupsToMerge, C),
72 make_tuple(NStride, WiStride, BatchStride, CStride));
73 }
74 }
75
76 template <index_t NDim, typename enable_if<NDim == 1, bool>::type = false>
77 constexpr static auto
78 make_wei_grid_desc(const index_t K,
79 const index_t X,
80 const index_t C,
81 const std::array<index_t, NDimSpatial + 3>& weights_strides)
82 {
83 const auto CStride = Number<1>{};
84 const auto KStride = weights_strides[1];
85 const auto XStride = weights_strides[3];
86 const auto BatchStride = weights_strides[0];
87 // Add NumGroupsToMerge for Batch+M dimension and, 1 as a placehorder
88 // for Batch+N dimension
89 const auto desc = make_naive_tensor_descriptor(
90 make_tuple(NumGroupsToMerge, K, X, 1, C),
91 make_tuple(BatchStride, KStride, XStride, BatchStride, CStride));
92 // Padd 1 to NumGroupsToMerge
93 const auto padded_desc = transform_tensor_descriptor(
94 desc,
95 make_tuple(make_pass_through_transform(NumGroupsToMerge),
96 make_pass_through_transform(K),
97 make_pass_through_transform(X),
98 make_pad_transform(1, 0, NumGroupsToMerge - 1),
99 make_pass_through_transform(C)),
100 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}),
101 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}));
102 // We need only matrices from diagonal. Xor returns 0 for the same
103 // values. So if matrices is not on diagonal then it will be stored in padding.
104 // To avoid use of modulo after xor we assume that NumBatch to merge is power of 2.
105 static_assert(NumGroupsToMerge == 1 || NumGroupsToMerge == 2 || NumGroupsToMerge == 4 ||
106 NumGroupsToMerge == 8 || NumGroupsToMerge == 16 || NumGroupsToMerge == 32 ||
107 NumGroupsToMerge == 64);
108 const auto unmerged_padded_desc = transform_tensor_descriptor(
109 padded_desc,
110 make_tuple(make_xor_transform(make_tuple(NumGroupsToMerge, NumGroupsToMerge)),
111 make_pass_through_transform(K),
112 make_pass_through_transform(X),
113 make_pass_through_transform(C)),
114 make_tuple(Sequence<0, 3>{}, Sequence<1>{}, Sequence<2>{}, Sequence<4>{}),
115 make_tuple(Sequence<0, 3>{}, Sequence<1>{}, Sequence<2>{}, Sequence<4>{}));
116 // Merge To M, N
117 return transform_tensor_descriptor(
118 unmerged_padded_desc,
119 make_tuple(make_merge_transform(make_tuple(NumGroupsToMerge, K)),
120 make_merge_transform(make_tuple(X, NumGroupsToMerge, C))),
121 make_tuple(Sequence<0, 1>{}, Sequence<2, 3, 4>{}),
122 make_tuple(Sequence<0>{}, Sequence<1>{}));
123 }
124
125 template <index_t NDim, typename enable_if<NDim == 2, bool>::type = false>
126 constexpr static auto
127 make_out_grid_desc(const index_t N,
128 const index_t Ho,
129 const index_t Wo,
130 const index_t K,
131 const std::array<index_t, NDimSpatial + 3>& output_strides)
132 {
133 const index_t BatchStride = output_strides[0];
134 const index_t WoStride = output_strides[4];
135 const auto KStride = Number<1>{};
136 return make_naive_tensor_descriptor(make_tuple(N * Ho * Wo, NumGroupsToMerge, K),
137 make_tuple(WoStride, BatchStride, KStride));
138 }
139
140 template <index_t NDim, typename enable_if<NDim == 2, bool>::type = false>
141 constexpr static auto
142 make_in_grid_desc(const index_t N,
143 const index_t Hi,
144 const index_t Wi,
145 const index_t C,
146 const std::array<index_t, NDimSpatial + 3>& input_strides)
147 {
148 const index_t BatchStride = input_strides[0];
149 const index_t NStride = input_strides[1];
150 const index_t HiStride = input_strides[3];
151 const index_t WiStride = input_strides[4];
152 const auto CStride = input_strides[2];
153 if constexpr(ConvBackwardWeightSpecialization ==
154 device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)
155 {
156 return make_naive_tensor_descriptor(make_tuple(N * Hi * Wi, NumGroupsToMerge, C),
157 make_tuple(WiStride, BatchStride, CStride));
158 }
159 else
160 {
161 return make_naive_tensor_descriptor(
162 make_tuple(N, Hi, Wi, NumGroupsToMerge, C),
163 make_tuple(NStride, HiStride, WiStride, BatchStride, CStride));
164 }
165 }
166
167 template <index_t NDim, typename enable_if<NDim == 2, bool>::type = false>
168 constexpr static auto
169 make_wei_grid_desc(const index_t K,
170 const index_t Y,
171 const index_t X,
172 const index_t C,
173 const std::array<index_t, NDimSpatial + 3>& weights_strides)
174 {
175 const auto CStride = Number<1>{};
176 const auto KStride = weights_strides[1];
177 const auto XStride = weights_strides[4];
178 const auto BatchStride = weights_strides[0];
179 // Add NumGroupsToMerge for Batch+M dimension and, 1 as a placehorder
180 // for Batch+N dimension
181 const auto desc = make_naive_tensor_descriptor(
182 make_tuple(NumGroupsToMerge, K, Y * X, 1, C),
183 make_tuple(BatchStride, KStride, XStride, BatchStride, CStride));
184 // Padd 1 to NumGroupsToMerge
185 const auto padded_desc = transform_tensor_descriptor(
186 desc,
187 make_tuple(make_pass_through_transform(NumGroupsToMerge),
188 make_pass_through_transform(K),
189 make_pass_through_transform(Y * X),
190 make_pad_transform(1, 0, NumGroupsToMerge - 1),
191 make_pass_through_transform(C)),
192 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}),
193 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}));
194 // We need only matrices from diagonal. Xor returns 0 for the same
195 // values. So if matrices is not on diagonal then it will be stored in padding.
196 // To avoid use of modulo after xor we assume that NumBatch to merge is power of 2.
197 static_assert(NumGroupsToMerge == 1 || NumGroupsToMerge == 2 || NumGroupsToMerge == 4 ||
198 NumGroupsToMerge == 8 || NumGroupsToMerge == 16 || NumGroupsToMerge == 32 ||
199 NumGroupsToMerge == 64);
200 const auto unmerged_padded_desc = transform_tensor_descriptor(
201 padded_desc,
202 make_tuple(make_xor_transform(make_tuple(NumGroupsToMerge, NumGroupsToMerge)),
203 make_pass_through_transform(K),
204 make_pass_through_transform(Y * X),
205 make_pass_through_transform(C)),
206 make_tuple(Sequence<0, 3>{}, Sequence<1>{}, Sequence<2>{}, Sequence<4>{}),
207 make_tuple(Sequence<0, 3>{}, Sequence<1>{}, Sequence<2>{}, Sequence<4>{}));
208 // Merge To M, N
209 return transform_tensor_descriptor(
210 unmerged_padded_desc,
211 make_tuple(make_merge_transform(make_tuple(NumGroupsToMerge, K)),
212 make_merge_transform(make_tuple(Y * X, NumGroupsToMerge, C))),
213 make_tuple(Sequence<0, 1>{}, Sequence<2, 3, 4>{}),
214 make_tuple(Sequence<0>{}, Sequence<1>{}));
215 }
216
217 template <index_t NDim, typename enable_if<NDim == 3, bool>::type = false>
218 constexpr static auto
219 make_out_grid_desc(const index_t N,
220 const index_t Do,
221 const index_t Ho,
222 const index_t Wo,
223 const index_t K,
224 const std::array<index_t, NDimSpatial + 3>& output_strides)
225 {
226 const index_t BatchStride = output_strides[0];
227 const index_t WoStride = output_strides[5];
228 const auto KStride = Number<1>{};
229 return make_naive_tensor_descriptor(make_tuple(N * Do * Ho * Wo, NumGroupsToMerge, K),
230 make_tuple(WoStride, BatchStride, KStride));
231 }
232
233 template <index_t NDim, typename enable_if<NDim == 3, bool>::type = false>
234 constexpr static auto
235 make_in_grid_desc(const index_t N,
236 const index_t Di,
237 const index_t Hi,
238 const index_t Wi,
239 const index_t C,
240 const std::array<index_t, NDimSpatial + 3>& input_strides)
241 {
242 const index_t BatchStride = input_strides[0];
243 const index_t NStride = input_strides[1];
244 const index_t DiStride = input_strides[3];
245 const index_t HiStride = input_strides[4];
246 const index_t WiStride = input_strides[5];
247 const auto CStride = input_strides[2];
248 if constexpr(ConvBackwardWeightSpecialization ==
249 device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)
250 {
251 return make_naive_tensor_descriptor(make_tuple(N * Di * Hi * Wi, NumGroupsToMerge, C),
252 make_tuple(WiStride, BatchStride, CStride));
253 }
254 else
255 {
256 return make_naive_tensor_descriptor(
257 make_tuple(N, Di, Hi, Wi, NumGroupsToMerge, C),
258 make_tuple(NStride, DiStride, HiStride, WiStride, BatchStride, CStride));
259 }
260 }
261
262 template <index_t NDim, typename enable_if<NDim == 3, bool>::type = false>
263 constexpr static auto
264 make_wei_grid_desc(const index_t K,
265 const index_t Z,
266 const index_t Y,
267 const index_t X,
268 const index_t C,
269 const std::array<index_t, NDimSpatial + 3>& weights_strides)
270 {
271 const auto CStride = Number<1>{};
272 const auto KStride = weights_strides[1];
273 const auto XStride = weights_strides[5];
274 const auto BatchStride = weights_strides[0];
275 // Add NumGroupsToMerge for Batch+M dimension and, 1 for placehord for Batch+N dimension
276 const auto desc = make_naive_tensor_descriptor(
277 make_tuple(NumGroupsToMerge, K, Z * Y * X, 1, C),
278 make_tuple(BatchStride, KStride, XStride, BatchStride, CStride));
279 // Padd 1 to NumGroupsToMerge
280 const auto padded_desc = transform_tensor_descriptor(
281 desc,
282 make_tuple(make_pass_through_transform(NumGroupsToMerge),
283 make_pass_through_transform(K),
284 make_pass_through_transform(Z * Y * X),
285 make_pad_transform(1, 0, NumGroupsToMerge - 1),
286 make_pass_through_transform(C)),
287 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}),
288 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}));
289 // We need only matrices from diagonal. Xor returns 0 for the same
290 // values. So if matrices is not on diagonal then it will be stored in padding.
291 // To avoid use of modulo after xor we assume that NumBatch to merge is power of 2.
292 static_assert(NumGroupsToMerge == 1 || NumGroupsToMerge == 2 || NumGroupsToMerge == 4 ||
293 NumGroupsToMerge == 8 || NumGroupsToMerge == 16 || NumGroupsToMerge == 32 ||
294 NumGroupsToMerge == 64);
295 const auto unmerged_padded_desc = transform_tensor_descriptor(
296 padded_desc,
297 make_tuple(make_xor_transform(make_tuple(NumGroupsToMerge, NumGroupsToMerge)),
298 make_pass_through_transform(K),
299 make_pass_through_transform(Z * Y * X),
300 make_pass_through_transform(C)),
301 make_tuple(Sequence<0, 3>{}, Sequence<1>{}, Sequence<2>{}, Sequence<4>{}),
302 make_tuple(Sequence<0, 3>{}, Sequence<1>{}, Sequence<2>{}, Sequence<4>{}));
303 // Merge To M, N
304 return transform_tensor_descriptor(
305 unmerged_padded_desc,
306 make_tuple(make_merge_transform(make_tuple(NumGroupsToMerge, K)),
307 make_merge_transform(make_tuple(Z * Y * X, NumGroupsToMerge, C))),
308 make_tuple(Sequence<0, 1>{}, Sequence<2, 3, 4>{}),
309 make_tuple(Sequence<0>{}, Sequence<1>{}));
310 }
311
312 template <index_t NDim, typename enable_if<NDim == 1, bool>::type = false>
313 static auto MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(
314 const index_t N,
315 const index_t K,
316 const index_t C,
317 const std::array<index_t, NDimSpatial>& input_spatial_lengths,
318 const std::array<index_t, NDimSpatial>& filter_spatial_lengths,
319 const std::array<index_t, NDimSpatial>& output_spatial_lengths,
320 const std::array<index_t, NDimSpatial + 3>& input_strides,
321 const std::array<index_t, NDimSpatial + 3>& weights_strides,
322 const std::array<index_t, NDimSpatial + 3>& output_strides,
323 const std::array<index_t, NDimSpatial>& conv_filter_strides,
324 const std::array<index_t, NDimSpatial>& conv_filter_dilations,
325 const std::array<index_t, NDimSpatial>& input_left_pads,
326 const std::array<index_t, NDimSpatial>& input_right_pads,
327 const index_t batch_k)
328 {
329 using namespace ck;
330
331 const index_t Wi = input_spatial_lengths[0];
332
333 const index_t Wo = output_spatial_lengths[0];
334
335 const index_t X = filter_spatial_lengths[0];
336
337 const index_t ConvStrideW = conv_filter_strides[0];
338
339 const index_t ConvDilationW = conv_filter_dilations[0];
340
341 const index_t InLeftPadW = input_left_pads[0];
342
343 const index_t InRightPadW = input_right_pads[0];
344
345 const index_t GemmKTotal = N * Wo;
346 const index_t GemmM = K * NumGroupsToMerge;
347 const index_t GemmN = C * X * NumGroupsToMerge;
348
349 const auto PadGemmM = GemmM % MPerBlock == 0 ? 0 : MPerBlock - GemmM % MPerBlock;
350 const auto PadGemmN = GemmN % NPerBlock == 0 ? 0 : NPerBlock - GemmN % NPerBlock;
351
352 const index_t GemmKBatch = batch_k;
353 const index_t GemmK0 =
354 math::integer_divide_ceil(GemmKTotal, GemmK1Number * K0PerBlock * GemmKBatch) *
355 K0PerBlock;
356 const index_t GemmKPad = GemmKBatch * GemmK0 * GemmK1Number;
357
358 const auto out_grid_desc = make_out_grid_desc<NDim>(N, Wo, K, output_strides);
359 const auto in_grid_desc = make_in_grid_desc<NDim>(N, Wi, C, input_strides);
360 const auto wei_grid_desc = make_wei_grid_desc<NDim>(K, X, C, weights_strides);
361
362 if constexpr(ConvBackwardWeightSpecialization ==
363 device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)
364 {
365 // A: output tensor
366 const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
367 out_grid_desc,
368 make_tuple(
369 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
370 make_merge_transform(make_tuple(NumGroupsToMerge, GemmM / NumGroupsToMerge))),
371 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
372 make_tuple(Sequence<0>{}, Sequence<1>{}));
373
374 const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
375 out_gemmkpad_gemmm_grid_desc,
376 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
377 make_right_pad_transform(GemmM, PadGemmM)),
378 make_tuple(Sequence<0>{}, Sequence<1>{}),
379 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
380
381 // B: input tensor
382 const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
383 in_grid_desc,
384 make_tuple(
385 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
386 make_merge_transform(make_tuple(NumGroupsToMerge, GemmN / NumGroupsToMerge))),
387 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
388 make_tuple(Sequence<0>{}, Sequence<1>{}));
389
390 const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
391 in_gemmkpad_gemmn_grid_desc,
392 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
393 make_right_pad_transform(GemmN, PadGemmN)),
394 make_tuple(Sequence<0>{}, Sequence<1>{}),
395 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
396
397 // Padd
398 const auto wei_gemmm_gemmn_pad_grid_desc =
399 transform_tensor_descriptor(wei_grid_desc,
400 make_tuple(make_right_pad_transform(GemmM, PadGemmM),
401 make_right_pad_transform(GemmN, PadGemmN)),
402 make_tuple(Sequence<0>{}, Sequence<1>{}),
403 make_tuple(Sequence<0>{}, Sequence<1>{}));
404
405 return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
406 in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
407 wei_gemmm_gemmn_pad_grid_desc);
408 }
409 else
410 {
411 // A: output tensor
412 const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
413 out_grid_desc,
414 make_tuple(
415 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
416 make_merge_transform(make_tuple(NumGroupsToMerge, GemmM / NumGroupsToMerge))),
417 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
418 make_tuple(Sequence<0>{}, Sequence<1>{}));
419
420 const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
421 out_gemmkpad_gemmm_grid_desc,
422 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
423 make_right_pad_transform(GemmM, PadGemmM)),
424 make_tuple(Sequence<0>{}, Sequence<1>{}),
425 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
426
427 // B: input tensor
428 const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
429 in_grid_desc,
430 make_tuple(make_pass_through_transform(N),
431 make_pad_transform(Wi, InLeftPadW, InRightPadW),
432 make_pass_through_transform(NumGroupsToMerge),
433 make_pass_through_transform(C)),
434 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
435 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}));
436
437 const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
438 in_n_hip_wip_c_grid_desc,
439 make_tuple(
440 make_pass_through_transform(N),
441 make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
442 make_pass_through_transform(NumGroupsToMerge),
443 make_pass_through_transform(C)),
444 make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}),
445 make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3>{}, Sequence<4>{}));
446
447 const auto in_gemmktotal_gemmn_grid_desc = transform_tensor_descriptor(
448 in_n_y_ho_x_wo_c_grid_desc,
449 make_tuple(make_merge_transform(make_tuple(X, NumGroupsToMerge, C)),
450 make_merge_transform(make_tuple(N, Wo))),
451 make_tuple(Sequence<1, 3, 4>{}, Sequence<0, 2>{}),
452 make_tuple(Sequence<1>{}, Sequence<0>{}));
453
454 const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
455 in_gemmktotal_gemmn_grid_desc,
456 make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
457 make_pass_through_transform(GemmN)),
458 make_tuple(Sequence<0>{}, Sequence<1>{}),
459 make_tuple(Sequence<0>{}, Sequence<1>{}));
460
461 const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
462 in_gemmkpad_gemmn_grid_desc,
463 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
464 make_right_pad_transform(GemmN, PadGemmN)),
465 make_tuple(Sequence<0>{}, Sequence<1>{}),
466 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
467
468 // Padd
469 const auto wei_gemmm_gemmn_pad_grid_desc =
470 transform_tensor_descriptor(wei_grid_desc,
471 make_tuple(make_right_pad_transform(GemmM, PadGemmM),
472 make_right_pad_transform(GemmN, PadGemmN)),
473 make_tuple(Sequence<0>{}, Sequence<1>{}),
474 make_tuple(Sequence<0>{}, Sequence<1>{}));
475
476 return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
477 in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
478 wei_gemmm_gemmn_pad_grid_desc);
479 }
480
481 } // function end
482
483 template <index_t NDim, typename enable_if<NDim == 2, bool>::type = false>
484 static auto MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(
485 const index_t N,
486 const index_t K,
487 const index_t C,
488 const std::array<index_t, NDimSpatial>& input_spatial_lengths,
489 const std::array<index_t, NDimSpatial>& filter_spatial_lengths,
490 const std::array<index_t, NDimSpatial>& output_spatial_lengths,
491 const std::array<index_t, NDimSpatial + 3>& input_strides,
492 const std::array<index_t, NDimSpatial + 3>& weights_strides,
493 const std::array<index_t, NDimSpatial + 3>& output_strides,
494 const std::array<index_t, NDimSpatial>& conv_filter_strides,
495 const std::array<index_t, NDimSpatial>& conv_filter_dilations,
496 const std::array<index_t, NDimSpatial>& input_left_pads,
497 const std::array<index_t, NDimSpatial>& input_right_pads,
498 const index_t batch_k)
499 {
500 using namespace ck;
501
502 const index_t Hi = input_spatial_lengths[0];
503 const index_t Wi = input_spatial_lengths[1];
504
505 const index_t Ho = output_spatial_lengths[0];
506 const index_t Wo = output_spatial_lengths[1];
507
508 const index_t Y = filter_spatial_lengths[0];
509 const index_t X = filter_spatial_lengths[1];
510
511 const index_t ConvStrideH = conv_filter_strides[0];
512 const index_t ConvStrideW = conv_filter_strides[1];
513
514 const index_t ConvDilationH = conv_filter_dilations[0];
515 const index_t ConvDilationW = conv_filter_dilations[1];
516
517 const index_t InLeftPadH = input_left_pads[0];
518 const index_t InLeftPadW = input_left_pads[1];
519
520 const index_t InRightPadH = input_right_pads[0];
521 const index_t InRightPadW = input_right_pads[1];
522
523 const index_t GemmKTotal = N * Ho * Wo;
524 const index_t GemmM = K * NumGroupsToMerge;
525 const index_t GemmN = C * X * Y * NumGroupsToMerge;
526
527 const auto PadGemmM = GemmM % MPerBlock == 0 ? 0 : MPerBlock - GemmM % MPerBlock;
528 const auto PadGemmN = GemmN % NPerBlock == 0 ? 0 : NPerBlock - GemmN % NPerBlock;
529
530 const index_t GemmKBatch = batch_k;
531 const index_t GemmK0 =
532 math::integer_divide_ceil(GemmKTotal, GemmK1Number * K0PerBlock * GemmKBatch) *
533 K0PerBlock;
534 const index_t GemmKPad = GemmKBatch * GemmK0 * GemmK1Number;
535
536 const auto out_grid_desc = make_out_grid_desc<NDim>(N, Ho, Wo, K, output_strides);
537 const auto in_grid_desc = make_in_grid_desc<NDim>(N, Hi, Wi, C, input_strides);
538 const auto wei_grid_desc = make_wei_grid_desc<NDim>(K, Y, X, C, weights_strides);
539
540 if constexpr(ConvBackwardWeightSpecialization ==
541 device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)
542 {
543 // A: output tensor
544 const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
545 out_grid_desc,
546 make_tuple(
547 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
548 make_merge_transform(make_tuple(NumGroupsToMerge, GemmM / NumGroupsToMerge))),
549 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
550 make_tuple(Sequence<0>{}, Sequence<1>{}));
551
552 const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
553 out_gemmkpad_gemmm_grid_desc,
554 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
555 make_right_pad_transform(GemmM, PadGemmM)),
556 make_tuple(Sequence<0>{}, Sequence<1>{}),
557 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
558
559 // B: input tensor
560 const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
561 in_grid_desc,
562 make_tuple(
563 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
564 make_merge_transform(make_tuple(NumGroupsToMerge, GemmN / NumGroupsToMerge))),
565 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
566 make_tuple(Sequence<0>{}, Sequence<1>{}));
567
568 const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
569 in_gemmkpad_gemmn_grid_desc,
570 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
571 make_right_pad_transform(GemmN, PadGemmN)),
572 make_tuple(Sequence<0>{}, Sequence<1>{}),
573 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
574
575 // Padd
576 const auto wei_gemmm_gemmn_pad_grid_desc =
577 transform_tensor_descriptor(wei_grid_desc,
578 make_tuple(make_right_pad_transform(GemmM, PadGemmM),
579 make_right_pad_transform(GemmN, PadGemmN)),
580 make_tuple(Sequence<0>{}, Sequence<1>{}),
581 make_tuple(Sequence<0>{}, Sequence<1>{}));
582
583 return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
584 in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
585 wei_gemmm_gemmn_pad_grid_desc);
586 }
587 else
588 {
589 // A: output tensor
590 const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
591 out_grid_desc,
592 make_tuple(
593 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
594 make_merge_transform(make_tuple(NumGroupsToMerge, GemmM / NumGroupsToMerge))),
595 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
596 make_tuple(Sequence<0>{}, Sequence<1>{}));
597
598 const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
599 out_gemmkpad_gemmm_grid_desc,
600 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
601 make_right_pad_transform(GemmM, PadGemmM)),
602 make_tuple(Sequence<0>{}, Sequence<1>{}),
603 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
604
605 // B: input tensor
606 const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor(
607 in_grid_desc,
608 make_tuple(make_pass_through_transform(N),
609 make_pad_transform(Hi, InLeftPadH, InRightPadH),
610 make_pad_transform(Wi, InLeftPadW, InRightPadW),
611 make_pass_through_transform(NumGroupsToMerge),
612 make_pass_through_transform(C)),
613 make_tuple(
614 Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}),
615 make_tuple(
616 Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}));
617
618 const auto in_n_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
619 in_n_hip_wip_c_grid_desc,
620 make_tuple(
621 make_pass_through_transform(N),
622 make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
623 make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
624 make_pass_through_transform(NumGroupsToMerge),
625 make_pass_through_transform(C)),
626 make_tuple(
627 Sequence<0>{}, Sequence<1>{}, Sequence<2>{}, Sequence<3>{}, Sequence<4>{}),
628 make_tuple(Sequence<0>{},
629 Sequence<1, 2>{},
630 Sequence<3, 4>{},
631 Sequence<5>{},
632 Sequence<6>{}));
633
634 const auto in_gemmktotal_gemmn_grid_desc = transform_tensor_descriptor(
635 in_n_y_ho_x_wo_c_grid_desc,
636 make_tuple(make_merge_transform(make_tuple(Y, X, NumGroupsToMerge, C)),
637 make_merge_transform(make_tuple(N, Ho, Wo))),
638 make_tuple(Sequence<1, 3, 5, 6>{}, Sequence<0, 2, 4>{}),
639 make_tuple(Sequence<1>{}, Sequence<0>{}));
640
641 const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
642 in_gemmktotal_gemmn_grid_desc,
643 make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
644 make_pass_through_transform(GemmN)),
645 make_tuple(Sequence<0>{}, Sequence<1>{}),
646 make_tuple(Sequence<0>{}, Sequence<1>{}));
647
648 const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
649 in_gemmkpad_gemmn_grid_desc,
650 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
651 make_right_pad_transform(GemmN, PadGemmN)),
652 make_tuple(Sequence<0>{}, Sequence<1>{}),
653 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
654
655 // Padd
656 const auto wei_gemmm_gemmn_pad_grid_desc =
657 transform_tensor_descriptor(wei_grid_desc,
658 make_tuple(make_right_pad_transform(GemmM, PadGemmM),
659 make_right_pad_transform(GemmN, PadGemmN)),
660 make_tuple(Sequence<0>{}, Sequence<1>{}),
661 make_tuple(Sequence<0>{}, Sequence<1>{}));
662
663 return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
664 in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
665 wei_gemmm_gemmn_pad_grid_desc);
666 }
667 }
668
669 template <index_t NDim, typename enable_if<NDim == 3, bool>::type = false>
670 static auto MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(
671 const index_t N,
672 const index_t K,
673 const index_t C,
674 const std::array<index_t, NDimSpatial>& input_spatial_lengths,
675 const std::array<index_t, NDimSpatial>& filter_spatial_lengths,
676 const std::array<index_t, NDimSpatial>& output_spatial_lengths,
677 const std::array<index_t, NDimSpatial + 3>& input_strides,
678 const std::array<index_t, NDimSpatial + 3>& weights_strides,
679 const std::array<index_t, NDimSpatial + 3>& output_strides,
680 const std::array<index_t, NDimSpatial>& conv_filter_strides,
681 const std::array<index_t, NDimSpatial>& conv_filter_dilations,
682 const std::array<index_t, NDimSpatial>& input_left_pads,
683 const std::array<index_t, NDimSpatial>& input_right_pads,
684 const index_t batch_k)
685 {
686 using namespace ck;
687
688 const index_t Di = input_spatial_lengths[0];
689 const index_t Hi = input_spatial_lengths[1];
690 const index_t Wi = input_spatial_lengths[2];
691
692 const index_t Do = output_spatial_lengths[0];
693 const index_t Ho = output_spatial_lengths[1];
694 const index_t Wo = output_spatial_lengths[2];
695
696 const index_t Z = filter_spatial_lengths[0];
697 const index_t Y = filter_spatial_lengths[1];
698 const index_t X = filter_spatial_lengths[2];
699
700 const index_t ConvStrideD = conv_filter_strides[0];
701 const index_t ConvStrideH = conv_filter_strides[1];
702 const index_t ConvStrideW = conv_filter_strides[2];
703
704 const index_t ConvDilationD = conv_filter_dilations[0];
705 const index_t ConvDilationH = conv_filter_dilations[1];
706 const index_t ConvDilationW = conv_filter_dilations[2];
707
708 const index_t InLeftPadD = input_left_pads[0];
709 const index_t InLeftPadH = input_left_pads[1];
710 const index_t InLeftPadW = input_left_pads[2];
711
712 const index_t InRightPadD = input_right_pads[0];
713 const index_t InRightPadH = input_right_pads[1];
714 const index_t InRightPadW = input_right_pads[2];
715
716 const index_t GemmKTotal = N * Do * Ho * Wo;
717 const index_t GemmM = K * NumGroupsToMerge;
718 const index_t GemmN = C * Z * X * Y * NumGroupsToMerge;
719
720 const auto PadGemmM = GemmM % MPerBlock == 0 ? 0 : MPerBlock - GemmM % MPerBlock;
721 const auto PadGemmN = GemmN % NPerBlock == 0 ? 0 : NPerBlock - GemmN % NPerBlock;
722
723 const index_t GemmKBatch = batch_k;
724 const index_t GemmK0 =
725 math::integer_divide_ceil(GemmKTotal, GemmK1Number * K0PerBlock * GemmKBatch) *
726 K0PerBlock;
727 const index_t GemmKPad = GemmKBatch * GemmK0 * GemmK1Number;
728
729 const auto out_grid_desc = make_out_grid_desc<NDim>(N, Do, Ho, Wo, K, output_strides);
730 const auto in_grid_desc = make_in_grid_desc<NDim>(N, Di, Hi, Wi, C, input_strides);
731 const auto wei_grid_desc = make_wei_grid_desc<NDim>(K, Z, Y, X, C, weights_strides);
732
733 if constexpr(ConvBackwardWeightSpecialization ==
734 device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)
735 {
736 // A: output tensor
737 const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
738 out_grid_desc,
739 make_tuple(
740 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
741 make_merge_transform(make_tuple(NumGroupsToMerge, GemmM / NumGroupsToMerge))),
742 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
743 make_tuple(Sequence<0>{}, Sequence<1>{}));
744
745 const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
746 out_gemmkpad_gemmm_grid_desc,
747 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
748 make_right_pad_transform(GemmM, PadGemmM)),
749 make_tuple(Sequence<0>{}, Sequence<1>{}),
750 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
751
752 // B: input tensor
753 const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
754 in_grid_desc,
755 make_tuple(
756 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
757 make_merge_transform(make_tuple(NumGroupsToMerge, GemmN / NumGroupsToMerge))),
758 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
759 make_tuple(Sequence<0>{}, Sequence<1>{}));
760
761 const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
762 in_gemmkpad_gemmn_grid_desc,
763 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
764 make_right_pad_transform(GemmN, PadGemmN)),
765 make_tuple(Sequence<0>{}, Sequence<1>{}),
766 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
767
768 // Padd
769 const auto wei_gemmm_gemmn_pad_grid_desc =
770 transform_tensor_descriptor(wei_grid_desc,
771 make_tuple(make_right_pad_transform(GemmM, PadGemmM),
772 make_right_pad_transform(GemmN, PadGemmN)),
773 make_tuple(Sequence<0>{}, Sequence<1>{}),
774 make_tuple(Sequence<0>{}, Sequence<1>{}));
775
776 return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
777 in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
778 wei_gemmm_gemmn_pad_grid_desc);
779 }
780 else
781 {
782 // A: output tensor
783 const auto out_gemmkpad_gemmm_grid_desc = transform_tensor_descriptor(
784 out_grid_desc,
785 make_tuple(
786 make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
787 make_merge_transform(make_tuple(NumGroupsToMerge, GemmM / NumGroupsToMerge))),
788 make_tuple(Sequence<0>{}, Sequence<1, 2>{}),
789 make_tuple(Sequence<0>{}, Sequence<1>{}));
790
791 const auto out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc = transform_tensor_descriptor(
792 out_gemmkpad_gemmm_grid_desc,
793 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
794 make_right_pad_transform(GemmM, PadGemmM)),
795 make_tuple(Sequence<0>{}, Sequence<1>{}),
796 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
797
798 // B: input tensor
799 const auto in_n_dip_hip_wip_c_grid_desc = transform_tensor_descriptor(
800 in_grid_desc,
801 make_tuple(make_pass_through_transform(N),
802 make_pad_transform(Di, InLeftPadD, InRightPadD),
803 make_pad_transform(Hi, InLeftPadH, InRightPadH),
804 make_pad_transform(Wi, InLeftPadW, InRightPadW),
805 make_pass_through_transform(NumGroupsToMerge),
806 make_pass_through_transform(C)),
807 make_tuple(Sequence<0>{},
808 Sequence<1>{},
809 Sequence<2>{},
810 Sequence<3>{},
811 Sequence<4>{},
812 Sequence<5>{}),
813 make_tuple(Sequence<0>{},
814 Sequence<1>{},
815 Sequence<2>{},
816 Sequence<3>{},
817 Sequence<4>{},
818 Sequence<5>{}));
819
820 const auto in_n_z_do_y_ho_x_wo_c_grid_desc = transform_tensor_descriptor(
821 in_n_dip_hip_wip_c_grid_desc,
822 make_tuple(
823 make_pass_through_transform(N),
824 make_embed_transform(make_tuple(Z, Do), make_tuple(ConvDilationD, ConvStrideD)),
825 make_embed_transform(make_tuple(Y, Ho), make_tuple(ConvDilationH, ConvStrideH)),
826 make_embed_transform(make_tuple(X, Wo), make_tuple(ConvDilationW, ConvStrideW)),
827 make_pass_through_transform(NumGroupsToMerge),
828 make_pass_through_transform(C)),
829 make_tuple(Sequence<0>{},
830 Sequence<1>{},
831 Sequence<2>{},
832 Sequence<3>{},
833 Sequence<4>{},
834 Sequence<5>{}),
835 make_tuple(Sequence<0>{},
836 Sequence<1, 2>{},
837 Sequence<3, 4>{},
838 Sequence<5, 6>{},
839 Sequence<7>{},
840 Sequence<8>{}));
841
842 const auto in_gemmktotal_gemmn_grid_desc = transform_tensor_descriptor(
843 in_n_z_do_y_ho_x_wo_c_grid_desc,
844 make_tuple(make_merge_transform(make_tuple(Z, Y, X, NumGroupsToMerge, C)),
845 make_merge_transform(make_tuple(N, Do, Ho, Wo))),
846 make_tuple(Sequence<1, 3, 5, 7, 8>{}, Sequence<0, 2, 4, 6>{}),
847 make_tuple(Sequence<1>{}, Sequence<0>{}));
848
849 const auto in_gemmkpad_gemmn_grid_desc = transform_tensor_descriptor(
850 in_gemmktotal_gemmn_grid_desc,
851 make_tuple(make_right_pad_transform(GemmKTotal, GemmKPad - GemmKTotal),
852 make_pass_through_transform(GemmN)),
853 make_tuple(Sequence<0>{}, Sequence<1>{}),
854 make_tuple(Sequence<0>{}, Sequence<1>{}));
855
856 const auto in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc = transform_tensor_descriptor(
857 in_gemmkpad_gemmn_grid_desc,
858 make_tuple(make_unmerge_transform(make_tuple(GemmKBatch * GemmK0, GemmK1Number)),
859 make_right_pad_transform(GemmN, PadGemmN)),
860 make_tuple(Sequence<0>{}, Sequence<1>{}),
861 make_tuple(Sequence<0, 2>{}, Sequence<1>{}));
862
863 // Padd
864 const auto wei_gemmm_gemmn_pad_grid_desc =
865 transform_tensor_descriptor(wei_grid_desc,
866 make_tuple(make_right_pad_transform(GemmM, PadGemmM),
867 make_right_pad_transform(GemmN, PadGemmN)),
868 make_tuple(Sequence<0>{}, Sequence<1>{}),
869 make_tuple(Sequence<0>{}, Sequence<1>{}));
870
871 return make_tuple(out_gemmkbatch_gemmk0_gemmm_gemmk1_grid_desc,
872 in_gemmkbatch_gemmk0_gemmn_gemmk1_grid_desc,
873 wei_gemmm_gemmn_pad_grid_desc);
874 }
875 } // function end
876};
877
878} // namespace tensor_operation
879} // namespace ck
common_header.hpp
convolution_backward_weight_specialization.hpp
ck::math::integer_divide_ceil
__host__ __device__ constexpr auto integer_divide_ceil(X x, Y y)
Definition utility/math.hpp:72
ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization
ConvolutionBackwardWeightSpecialization
Definition convolution_backward_weight_specialization.hpp:13
ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0
@ Filter1x1Stride1Pad0
Definition convolution_backward_weight_specialization.hpp:15
ck::tensor_operation
Definition convolution_backward_data_specialization.hpp:7
ck
Definition ck.hpp:268
ck::make_pass_through_transform
__host__ __device__ constexpr auto make_pass_through_transform(const LowLength &low_length)
Definition multi_index_transform_helper.hpp:12
ck::index_t
int32_t index_t
Definition ck.hpp:299
ck::make_naive_tensor_descriptor
__host__ __device__ constexpr auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)
Definition tensor_descriptor_helper.hpp:49
ck::make_right_pad_transform
__host__ __device__ constexpr auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:37
ck::Number
integral_constant< index_t, N > Number
Definition number.hpp:12
ck::make_pad_transform
__host__ __device__ constexpr auto make_pad_transform(const LowLength &low_length, const LeftPad &left_pad, const RightPad &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:19
ck::make_embed_transform
__host__ __device__ constexpr auto make_embed_transform(const UpLengths &up_lengths, const Coefficients &coefficients)
Definition multi_index_transform_helper.hpp:48
ck::make_merge_transform
__host__ __device__ constexpr auto make_merge_transform(const LowLengths &low_lengths)
Definition multi_index_transform_helper.hpp:55
ck::make_tuple
__host__ __device__ constexpr auto make_tuple(Xs &&... xs)
Definition utility/tuple.hpp:211
ck::transform_tensor_descriptor
__host__ __device__ constexpr auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)
Definition tensor_description/tensor_descriptor.hpp:319
ck::make_unmerge_transform
__host__ __device__ constexpr auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})
Definition multi_index_transform_helper.hpp:90
ck::make_xor_transform
__host__ __device__ constexpr auto make_xor_transform(const LowLengths &low_lengths)
Definition multi_index_transform_helper.hpp:191
ck::Sequence
Definition utility/sequence.hpp:43
ck::tensor_operation::TransformConvBwdWeightToGemmV2
Transform conv bwd weight to gemm v2.
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:33
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_wei_grid_desc
static constexpr auto make_wei_grid_desc(const index_t K, const index_t X, const index_t C, const std::array< index_t, NDimSpatial+3 > &weights_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:78
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_in_grid_desc
static constexpr auto make_in_grid_desc(const index_t N, const index_t Hi, const index_t Wi, const index_t C, const std::array< index_t, NDimSpatial+3 > &input_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:142
ck::tensor_operation::TransformConvBwdWeightToGemmV2::I1
static constexpr auto I1
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:35
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_out_grid_desc
static constexpr auto make_out_grid_desc(const index_t N, const index_t Do, const index_t Ho, const index_t Wo, const index_t K, const std::array< index_t, NDimSpatial+3 > &output_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:219
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_wei_grid_desc
static constexpr auto make_wei_grid_desc(const index_t K, const index_t Y, const index_t X, const index_t C, const std::array< index_t, NDimSpatial+3 > &weights_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:169
ck::tensor_operation::TransformConvBwdWeightToGemmV2::I0
static constexpr auto I0
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:34
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_out_grid_desc
static constexpr auto make_out_grid_desc(const index_t N, const index_t Ho, const index_t Wo, const index_t K, const std::array< index_t, NDimSpatial+3 > &output_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:127
ck::tensor_operation::TransformConvBwdWeightToGemmV2::MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N
static auto MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(const index_t N, const index_t K, const index_t C, const std::array< index_t, NDimSpatial > &input_spatial_lengths, const std::array< index_t, NDimSpatial > &filter_spatial_lengths, const std::array< index_t, NDimSpatial > &output_spatial_lengths, const std::array< index_t, NDimSpatial+3 > &input_strides, const std::array< index_t, NDimSpatial+3 > &weights_strides, const std::array< index_t, NDimSpatial+3 > &output_strides, const std::array< index_t, NDimSpatial > &conv_filter_strides, const std::array< index_t, NDimSpatial > &conv_filter_dilations, const std::array< index_t, NDimSpatial > &input_left_pads, const std::array< index_t, NDimSpatial > &input_right_pads, const index_t batch_k)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:313
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_wei_grid_desc
static constexpr auto make_wei_grid_desc(const index_t K, const index_t Z, const index_t Y, const index_t X, const index_t C, const std::array< index_t, NDimSpatial+3 > &weights_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:264
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_out_grid_desc
static constexpr auto make_out_grid_desc(const index_t N, const index_t Wo, const index_t K, const std::array< index_t, NDimSpatial+3 > &output_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:39
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_in_grid_desc
static constexpr auto make_in_grid_desc(const index_t N, const index_t Di, const index_t Hi, const index_t Wi, const index_t C, const std::array< index_t, NDimSpatial+3 > &input_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:235
ck::tensor_operation::TransformConvBwdWeightToGemmV2::make_in_grid_desc
static constexpr auto make_in_grid_desc(const index_t N, const index_t Wi, const index_t C, const std::array< index_t, NDimSpatial+3 > &input_strides)
Definition transform_conv_bwd_weight_to_gemm_v2.hpp:53
tensor_descriptor.hpp
tensor_descriptor_helper.hpp
tensor_layout.hpp