device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp Source File

• include
• ck
• tensor_operation
• gpu
• device
• impl

device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.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 <iostream>

7#include <numeric>

8#include <sstream>

9

10#include "ck/utility/common_header.hpp"

11#include "ck/utility/env.hpp"

12#include "ck/tensor_description/tensor_descriptor.hpp"

13#include "ck/tensor_description/tensor_descriptor_helper.hpp"

14#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"

15#include "ck/tensor_operation/gpu/device/device_grouped_conv_bwd_weight.hpp"

16#include "ck/tensor_operation/operator_transform/transform_conv_bwd_weight_to_gemm.hpp"

17#include "ck/tensor_operation/operator_transform/transform_conv_bwd_weight_to_gemm_v2.hpp"

18#include "ck/tensor_operation/operator_transform/transform_conv_ngchw_to_nhwgc.hpp"

19#include "ck/tensor_operation/gpu/device/convolution_backward_weight_specialization.hpp"

20#include "ck/tensor_operation/gpu/grid/gridwise_elementwise_2d.hpp"

21#include "ck/tensor_operation/gpu/device/gemm_specialization.hpp"

22#include "ck/tensor_operation/gpu/grid/gridwise_gemm_xdl_cshuffle_conv_v3.hpp"

23#include <ck/tensor_operation/gpu/grid/block_to_ctile_map.hpp>

24#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp"

25#include "ck/tensor_operation/gpu/device/impl/split_k_utils.hpp"

26#include "ck/tensor_operation/gpu/device/impl/split_k_arg.hpp"

27#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"

28

29#include "ck/host_utility/device_prop.hpp"

30#include "ck/host_utility/kernel_launch.hpp"

31#include "ck/host_utility/flush_cache.hpp"

32

33namespace ck {

34namespace tensor_operation {

35namespace device {

36

37template <typename GridwiseGemm,

38 typename AGridDesc_AK0_M_K1,

39 typename BGridDesc_BK0_N_K1,

40 typename CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,

41 typename ComputePtrOffsetOfBatch,

42 index_t NumGroupsToMerge,

43 bool HasMainKBlockLoop,

44 InMemoryDataOperationEnum CGlobalMemoryDataOperation,

45 index_t MinimumOccupancy = 1,

46 TailNumber TailNum = TailNumber::Full>

47__global__ void

48#if CK_USE_LAUNCH_BOUNDS

49__launch_bounds__(CK_MAX_THREAD_PER_BLOCK, MinimumOccupancy)

50#endif

51 kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3(

52 typename GridwiseGemm::Argument karg,

53 [[maybe_unused]] const AGridDesc_AK0_M_K1 a_grid_desc_ak0_m_ak1,

54 [[maybe_unused]] const BGridDesc_BK0_N_K1 b_grid_desc_bk0_n_bk1,

55 [[maybe_unused]] const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock

56 c_grid_desc_mblock_mperblock_nblock_nperblock,

57 [[maybe_unused]] const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch,

58 [[maybe_unused]] const index_t num_k_per_block)

59{

60#if defined(__gfx9__) || defined(__gfx11__) || defined(__gfx12__)

61 if constexpr(GridwiseGemm::template IsValidCompilationParameter<CGlobalMemoryDataOperation>())

62 {

63 const index_t g_idx = __builtin_amdgcn_readfirstlane(blockIdx.z * NumGroupsToMerge);

64 const index_t k_idx = __builtin_amdgcn_readfirstlane(blockIdx.y * num_k_per_block);

65

66 const long_index_t a_batch_offset = amd_wave_read_first_lane(

67 static_cast<long_index_t>(compute_ptr_offset_of_batch.GetAPtrOffset(g_idx)));

68 const long_index_t b_batch_offset = amd_wave_read_first_lane(

69 static_cast<long_index_t>(compute_ptr_offset_of_batch.GetBPtrOffset(g_idx)));

70 const long_index_t e_batch_offset = amd_wave_read_first_lane(

71 static_cast<long_index_t>(compute_ptr_offset_of_batch.GetEPtrOffset(g_idx)));

72

73 __shared__ char p_shared[GridwiseGemm::GetSharedMemoryNumberOfByte()];

74

75 GridwiseGemm::template Run<AGridDesc_AK0_M_K1,

76 BGridDesc_BK0_N_K1,

77 CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,

78 HasMainKBlockLoop,

79 CGlobalMemoryDataOperation,

80 TailNum>(karg.p_a_grid + a_batch_offset,

81 karg.p_b_grid + b_batch_offset,

82 karg.p_c_grid + e_batch_offset,

83 p_shared,

84 karg,

85 a_grid_desc_ak0_m_ak1,

86 b_grid_desc_bk0_n_bk1,

87 c_grid_desc_mblock_mperblock_nblock_nperblock,

88 k_idx);

89 }

90#else

91 ignore = karg;

92#endif // end of if (defined(__gfx9__))

93}

94

95template <typename GridwiseGemm,

96 typename AGridDesc_AK0_M_K1,

97 typename BGridDesc_BK0_N_K1,

98 typename CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,

99 typename ComputePtrOffsetOfBatch,

100 index_t NumGroupsToMerge,

101 bool HasMainKBlockLoop,

102 InMemoryDataOperationEnum CGlobalMemoryDataOperation,

103 index_t MinimumOccupancy = 1,

104 TailNumber TailNum = TailNumber::Full>

105__global__ void

106#if CK_USE_LAUNCH_BOUNDS

107__launch_bounds__(CK_MAX_THREAD_PER_BLOCK, MinimumOccupancy)

108#endif

109 kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds(

110 typename GridwiseGemm::Argument karg,

111 [[maybe_unused]] const AGridDesc_AK0_M_K1 a_grid_desc_ak0_m_ak1,

112 [[maybe_unused]] const BGridDesc_BK0_N_K1 b_grid_desc_bk0_n_bk1,

113 [[maybe_unused]] const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock

114 c_grid_desc_mblock_mperblock_nblock_nperblock,

115 [[maybe_unused]] const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch,

116 [[maybe_unused]] const index_t num_k_per_block)

117{

118#if defined(__gfx9__) || defined(__gfx11__) || defined(__gfx12__)

119 if constexpr(GridwiseGemm::template IsValidCompilationParameter<CGlobalMemoryDataOperation>())

120 {

121 // offset base pointer for each work-group

122 const index_t g_idx = __builtin_amdgcn_readfirstlane(blockIdx.z * NumGroupsToMerge);

123 const index_t k_idx = __builtin_amdgcn_readfirstlane(blockIdx.y * num_k_per_block);

124

125 const long_index_t a_batch_offset = amd_wave_read_first_lane(

126 static_cast<long_index_t>(compute_ptr_offset_of_batch.GetAPtrOffset(g_idx)));

127 const long_index_t b_batch_offset = amd_wave_read_first_lane(

128 static_cast<long_index_t>(compute_ptr_offset_of_batch.GetBPtrOffset(g_idx)));

129 const long_index_t e_batch_offset = amd_wave_read_first_lane(

130 static_cast<long_index_t>(compute_ptr_offset_of_batch.GetEPtrOffset(g_idx)));

131

132 // Pass two lds pointer is the key to tell compiler that ds_read/write

133 // operate on different lds chunk at same time without order dependecy

134 __shared__ char p_shared_0[GridwiseGemm::GetSharedMemoryNumberOfByte()];

135 __shared__ char p_shared_1[GridwiseGemm::GetSharedMemoryNumberOfByte()];

136

137 GridwiseGemm::template Run_2Lds<AGridDesc_AK0_M_K1,

138 BGridDesc_BK0_N_K1,

139 CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock,

140 HasMainKBlockLoop,

141 CGlobalMemoryDataOperation,

142 TailNum>(karg.p_a_grid + a_batch_offset,

143 karg.p_b_grid + b_batch_offset,

144 karg.p_c_grid + e_batch_offset,

145 p_shared_0,

146 p_shared_1,

147 karg,

148 a_grid_desc_ak0_m_ak1,

149 b_grid_desc_bk0_n_bk1,

150 c_grid_desc_mblock_mperblock_nblock_nperblock,

151 k_idx);

152 }

153#else

154 ignore = karg;

155#endif // end of if (defined(__gfx9__))

156}

157

158template <ck::index_t NDimSpatial,

159 typename InLayout,

160 typename WeiLayout,

161 typename OutLayout,

162 typename InDataType,

163 typename WeiDataType,

164 typename OutDataType,

165 typename AccDataType,

166 typename InElementwiseOperation,

167 typename WeiElementwiseOperation,

168 typename OutElementwiseOperation,

169 ConvolutionBackwardWeightSpecialization ConvBackwardWeightSpecialization,

170 ck::index_t BlockSize,

171 ck::index_t MPerBlock,

172 ck::index_t NPerBlock,

173 ck::index_t KPerBlock,

174 ck::index_t K1,

175 ck::index_t MPerXDL,

176 ck::index_t NPerXDL,

177 ck::index_t MXdlPerWave,

178 ck::index_t NXdlPerWave,

179 typename ABlockTransferThreadClusterLengths_K0_M_K1,

180 typename ABlockTransferThreadClusterArrangeOrder,

181 typename ABlockTransferSrcAccessOrder,

182 ck::index_t ABlockTransferSrcVectorDim,

183 ck::index_t ABlockTransferSrcScalarPerVector,

184 ck::index_t ABlockTransferDstScalarPerVector_K1,

185 bool ABlockLdsAddExtraM,

186 typename BBlockTransferThreadClusterLengths_K0_N_K1,

187 typename BBlockTransferThreadClusterArrangeOrder,

188 typename BBlockTransferSrcAccessOrder,

189 ck::index_t BBlockTransferSrcVectorDim,

190 ck::index_t BBlockTransferSrcScalarPerVector,

191 ck::index_t BBlockTransferDstScalarPerVector_K1,

192 bool BBlockLdsAddExtraN,

193 index_t CShuffleMXdlPerWavePerShuffle,

194 index_t CShuffleNXdlPerWavePerShuffle,

195 typename CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,

196 index_t CBlockTransferScalarPerVector_NWaveNPerXdl,

197 BlockGemmPipelineScheduler BlkGemmPipeSched = BlockGemmPipelineScheduler::Intrawave,

198 BlockGemmPipelineVersion BlkGemmPipelineVer = BlockGemmPipelineVersion::v1,

199 index_t NumGroupsToMerge = 1,

200 typename ComputeTypeA = InDataType,

201 typename ComputeTypeB = ComputeTypeA,

202 index_t TransposeTransferSrcScalarPerVector = 1,

203 index_t TransposeTransferDstScalarPerVector = 1>

204struct DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle

205 : public DeviceGroupedConvBwdWeight<NDimSpatial,

206 InLayout,

207 WeiLayout,

208 OutLayout,

209 InDataType,

210 WeiDataType,

211 OutDataType,

212 InElementwiseOperation,

213 WeiElementwiseOperation,

214 OutElementwiseOperation,

215 ComputeTypeA,

216 ComputeTypeB>

217{

218 static_assert(is_same_v<InElementwiseOperation, element_wise::PassThrough>);

219 static_assert(is_same_v<WeiElementwiseOperation, element_wise::PassThrough>);

220 static_assert(is_same_v<OutElementwiseOperation, element_wise::PassThrough>);

221

222 using DeviceOp = DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle;

223 GET_NXDL_PER_WAVE_IMPL

224 static constexpr auto NXdlPerWave64 = GetNXdlPerWave<true>();

225 static constexpr auto NXdlPerWave32 = GetNXdlPerWave<false>();

226

227 using ADataType = OutDataType;

228 using BDataType = InDataType;

229 using EDataType = WeiDataType;

230

231 // If NGCHW then ADataType must be equal to BDataType

232 static_assert(!(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

233 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>()) ||

234 is_same_v<ADataType, BDataType>);

235

236 using AElementwiseOperation = OutElementwiseOperation;

237 using BElementwiseOperation = InElementwiseOperation;

238 using CDEElementwiseOperation = WeiElementwiseOperation;

239

240 // TODO make A/B datatype different

241 using ABDataType = InDataType;

242

243 static constexpr auto I0 = Number<0>{};

244 static constexpr auto I1 = Number<1>{};

245 static constexpr auto I2 = Number<2>{};

246 static constexpr auto I3 = Number<3>{};

247 static constexpr auto I4 = Number<4>{};

248 static constexpr auto I5 = Number<5>{};

249

250 static constexpr auto K1Number = Number<K1>{};

251

252 static constexpr auto conv_to_gemm_transformer_v2 =

253 TransformConvBwdWeightToGemmV2<NDimSpatial,

254 MPerBlock,

255 NPerBlock,

256 K1Number,

257 KPerBlock / K1Number,

258 NumGroupsToMerge,

259 ConvBackwardWeightSpecialization>{};

260

261 static constexpr auto conv_to_gemm_transformer_v1 =

262 TransformConvBwdWeightToGemm<NDimSpatial,

263 MPerBlock,

264 NPerBlock,

265 K1Number,

266 KPerBlock / K1Number,

267 ConvBackwardWeightSpecialization>{};

268

269 static constexpr index_t ClusterLengthMPerBlock =

270 CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock::At(1);

271 static constexpr index_t ClusterLengthNPerBlock =

272 CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock::At(3);

273

274 static constexpr auto conv_ngchw_to_nhwgc_transformer =

275 TransformConvNGCHWToNHWGC<InLayout,

276 WeiLayout,

277 OutLayout,

278 NDimSpatial,

279 MPerBlock / ClusterLengthMPerBlock,

280 NPerBlock / ClusterLengthNPerBlock>{};

281

282 static constexpr GemmSpecialization GemmSpec = GemmSpecialization::Default;

283

284 template <ck::index_t NDim, typename ck::enable_if<NDim == 2, bool>::type = false>

285 static auto GetABCGridDesc()

286 {

287 const ck::index_t dim = 1;

288 const ck::index_t batch = 1;

289 const std::array<ck::index_t, NDimSpatial> lengths{1, 1};

290 const std::array<ck::index_t, NDimSpatial + 3> strides{1, 1, 1, 1, 1};

291 const std::array<ck::index_t, NDimSpatial> params{1, 1};

292 return conv_to_gemm_transformer_v2

293 .template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<2>(dim,

294 dim,

295 dim,

296 lengths,

297 lengths,

298 lengths,

299 strides,

300 strides,

301 strides,

302 params,

303 params,

304 params,

305 params,

306 batch);

307 }

308

309 template <ck::index_t NDim, typename ck::enable_if<NDim == 3, bool>::type = false>

310 static auto GetABCGridDesc()

311 {

312 const ck::index_t dim = 1;

313 const ck::index_t batch = 1;

314 const std::array<ck::index_t, NDimSpatial> lengths{1, 1, 1};

315 const std::array<ck::index_t, NDimSpatial + 3> strides{1, 1, 1, 1, 1, 1};

316 const std::array<ck::index_t, NDimSpatial> params{1, 1, 1};

317 return conv_to_gemm_transformer_v2

318 .template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<3>(dim,

319 dim,

320 dim,

321 lengths,

322 lengths,

323 lengths,

324 strides,

325 strides,

326 strides,

327 params,

328 params,

329 params,

330 params,

331 batch);

332 }

333

334 template <ck::index_t NDim, typename ck::enable_if<NDim == 2, bool>::type = false>

335 static auto GetElementwiseCGridDesc()

336 {

337 const ck::index_t dim = 1;

338 const ck::index_t batch = 1;

339 const std::array<ck::index_t, NDimSpatial> lengths{1, 1};

340 const std::array<ck::index_t, NDimSpatial + 3> strides{1, 1, 1, 1, 1};

341 const std::array<ck::index_t, NDimSpatial> params{1, 1};

342 return conv_to_gemm_transformer_v1

343 .template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<2>(dim,

344 dim,

345 dim,

346 lengths,

347 lengths,

348 lengths,

349 strides,

350 strides,

351 strides,

352 params,

353 params,

354 params,

355 params,

356 batch)[I2];

357 }

358

359 template <ck::index_t NDim, typename ck::enable_if<NDim == 3, bool>::type = false>

360 static auto GetElementwiseCGridDesc()

361 {

362 const ck::index_t dim = 1;

363 const ck::index_t batch = 1;

364 const std::array<ck::index_t, NDimSpatial> lengths{1, 1, 1};

365 const std::array<ck::index_t, NDimSpatial + 3> strides{1, 1, 1, 1, 1, 1};

366 const std::array<ck::index_t, NDimSpatial> params{1, 1, 1};

367 return conv_to_gemm_transformer_v1

368 .template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<3>(dim,

369 dim,

370 dim,

371 lengths,

372 lengths,

373 lengths,

374 strides,

375 strides,

376 strides,

377 params,

378 params,

379 params,

380 params,

381 batch)[I2];

382 }

383

384 using NGCHWTransposeDescType =

385 remove_cvref_t<decltype(conv_ngchw_to_nhwgc_transformer

386 .template MakeNGCHWTransposeDesc<NDimSpatial>({}, {}))>;

387 using NHWGCTransposeDescType =

388 remove_cvref_t<decltype(conv_ngchw_to_nhwgc_transformer

389 .template MakeNHWGCTransposeDesc<NDimSpatial>({}, {}))>;

390 using GKCYXTransposeDescType =

391 remove_cvref_t<decltype(conv_ngchw_to_nhwgc_transformer

392 .template MakeGKCYXTransposeDesc<NDimSpatial>({}, {}))>;

393 using GKYXCTransposeDescType =

394 remove_cvref_t<decltype(conv_ngchw_to_nhwgc_transformer

395 .template MakeGKYXCTransposeDesc<NDimSpatial>({}, {}))>;

396

397 using ABCGridDescs = decltype(GetABCGridDesc<NDimSpatial>());

398

399 using AGridDesc_K0_M_K1 = remove_cvref_t<decltype(ABCGridDescs{}[I0])>;

400 using BGridDesc_K0_N_K1 = remove_cvref_t<decltype(ABCGridDescs{}[I1])>;

401 using CGridDesc_M_N = remove_cvref_t<decltype(ABCGridDescs{}[I2])>;

402 using CElementwiseGridDesc_M_N =

403 remove_cvref_t<decltype(GetElementwiseCGridDesc<NDimSpatial>())>;

404

405 template <index_t NXdlPerWave_>

406 using GridwiseGemmBase = GridwiseGemm_xdl_cshuffle_conv_v3<

407 tensor_layout::gemm::RowMajor,

408 tensor_layout::gemm::ColumnMajor,

409 tensor_layout::gemm::RowMajor,

410 ADataType,

411 BDataType,

412 AccDataType,

413 AccDataType,

414 AccDataType,

415 AElementwiseOperation,

416 BElementwiseOperation,

417 CDEElementwiseOperation,

418 GemmSpec,

419 BlockSize,

420 MPerBlock,

421 NPerBlock,

422 KPerBlock,

423 K1,

424 K1,

425 MPerXDL,

426 NPerXDL,

427 MXdlPerWave,

428 NXdlPerWave_,

429 ABlockTransferThreadClusterLengths_K0_M_K1,

430 ABlockTransferThreadClusterArrangeOrder,

431 ABlockTransferSrcAccessOrder,

432 ABlockTransferSrcVectorDim,

433 ABlockTransferSrcScalarPerVector,

434 ABlockTransferDstScalarPerVector_K1,

435 false,

436 ABlockLdsAddExtraM,

437 BBlockTransferThreadClusterLengths_K0_N_K1,

438 BBlockTransferThreadClusterArrangeOrder,

439 BBlockTransferSrcAccessOrder,

440 BBlockTransferSrcVectorDim,

441 BBlockTransferSrcScalarPerVector,

442 BBlockTransferDstScalarPerVector_K1,

443 false,

444 BBlockLdsAddExtraN,

445 CShuffleMXdlPerWavePerShuffle,

446 CShuffleNXdlPerWavePerShuffle,

447 CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock,

448 CBlockTransferScalarPerVector_NWaveNPerXdl,

449 BlkGemmPipeSched,

450 BlkGemmPipelineVer,

451 ComputeTypeA,

452 ComputeTypeB>;

453 using GridwiseGemm64 = GridwiseGemmBase<math::max(NXdlPerWave64, 1)>;

454 using GridwiseGemm32 = GridwiseGemmBase<NXdlPerWave32>;

455

456 using Block2TileMapElementwise = BlockToCTileMap_M00_N0_M01Adapt<MPerBlock, NPerBlock>;

457

458 using GridwiseElementwiseCast =

459 GridwiseElementwise<Tuple<CElementwiseGridDesc_M_N>,

460 Tuple<CElementwiseGridDesc_M_N>,

461 Tuple<const AccDataType*>,

462 Tuple<EDataType*>,

463 Block2TileMapElementwise,

464 CDEElementwiseOperation,

465 BlockSize,

466 MPerBlock,

467 NPerBlock,

468 MPerBlock / ClusterLengthMPerBlock,

469 NPerBlock / ClusterLengthNPerBlock,

470 Sequence<0, 1>,

471 Sequence<CBlockTransferScalarPerVector_NWaveNPerXdl>,

472 Sequence<CBlockTransferScalarPerVector_NWaveNPerXdl>,

473 I1,

474 I1>;

475 // NPerBlock is used for the first dim which is store dimension

476 // (with CBlockTransferScalarPerVector_NWaveNPerXdl scalar per vector).

477 // CBlockTransferScalarPerVector_NWaveNPerXdl is aligned to NPerBlock so

478 // it is more flexible to use this dim for store dimension with such scalar

479 // per vector.

480 using GridwiseElementwiseWeightTransposeCast =

481 GridwiseElementwise<Tuple<GKYXCTransposeDescType>,

482 Tuple<GKCYXTransposeDescType>,

483 Tuple<const AccDataType*>,

484 Tuple<EDataType*>,

485 Block2TileMapElementwise,

486 CDEElementwiseOperation,

487 BlockSize,

488 MPerBlock,

489 NPerBlock,

490 MPerBlock / ClusterLengthMPerBlock,

491 NPerBlock / ClusterLengthNPerBlock,

492 Sequence<0, 1>,

493 Sequence<CBlockTransferScalarPerVector_NWaveNPerXdl>,

494 Sequence<1>,

495 I1,

496 I0>;

497

498 using GridwiseElementwiseTranspose =

499 GridwiseElementwise<Tuple<NGCHWTransposeDescType>,

500 Tuple<NHWGCTransposeDescType>,

501 Tuple<const ADataType*>,

502 Tuple<ADataType*>,

503 Block2TileMapElementwise,

504 element_wise::PassThrough,

505 BlockSize,

506 MPerBlock,

507 NPerBlock,

508 MPerBlock / ClusterLengthMPerBlock,

509 NPerBlock / ClusterLengthNPerBlock,

510 Sequence<1, 0>,

511 Sequence<TransposeTransferSrcScalarPerVector>,

512 Sequence<TransposeTransferDstScalarPerVector>,

513 I1,

514 I0>;

515

516 // Argument

517 using CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock =

518 decltype(GridwiseGemm64::MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(

519 CGridDesc_M_N{}, 1, 1));

520

521 struct ActiveWorkgroupsPerCU

522 {

523 template <typename GridwiseGemm>

524 int GetMaxOccupancy()

525 {

526 constexpr int dynamic_smem_size = 0;

527 constexpr index_t minimum_occupancy =

528 BlkGemmPipeSched == BlockGemmPipelineScheduler::Intrawave ? 1 : 2;

529 int max_occupancy = 0;

530

531 if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v4)

532 {

533 hip_check_error(hipOccupancyMaxActiveBlocksPerMultiprocessor(

534 &max_occupancy,

535 kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds<

536 GridwiseGemm,

537 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

538 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

539 remove_reference_t<DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

540 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

541 NumGroupsToMerge,

542 true,

543 InMemoryDataOperationEnum::AtomicAdd,

544 minimum_occupancy>,

545 BlockSize,

546 dynamic_smem_size));

547 }

548 else

549 {

550 hip_check_error(hipOccupancyMaxActiveBlocksPerMultiprocessor(

551 &max_occupancy,

552 kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

553 GridwiseGemm,

554 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

555 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

556 remove_reference_t<DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

557 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

558 NumGroupsToMerge,

559 true,

560 InMemoryDataOperationEnum::AtomicAdd,

561 minimum_occupancy>,

562 BlockSize,

563 dynamic_smem_size));

564 }

565 return std::max(1, max_occupancy);

566 }

567

568 ActiveWorkgroupsPerCU()

569 {

570 max_occupancy_ = 1;

571 if(get_warp_size() == 64)

572 {

573 if constexpr(NXdlPerWave64 > 0)

574 {

575 max_occupancy_ = GetMaxOccupancy<GridwiseGemm64>();

576 }

577 }

578 else

579 {

580 if constexpr(NXdlPerWave32 > 0)

581 {

582 max_occupancy_ = GetMaxOccupancy<GridwiseGemm32>();

583 }

584 }

585 }

586 int max_occupancy_;

587 };

588

589 struct Argument : public BaseArgument, public ArgumentSplitK

590 {

591 Argument(const InDataType* p_in_grid,

592 WeiDataType* p_wei_grid,

593 const OutDataType* p_out_grid,

594 const std::array<index_t, NDimSpatial + 3>& b_g_n_c_wis_lengths, // input

595 const std::array<index_t, NDimSpatial + 3>& b_g_n_c_wis_strides,

596 const std::array<index_t, NDimSpatial + 3>& e_g_k_c_xs_lengths, // weight

597 const std::array<index_t, NDimSpatial + 3>& e_g_k_c_xs_strides,

598 const std::array<index_t, NDimSpatial + 3>& a_g_n_k_wos_lengths, // output

599 const std::array<index_t, NDimSpatial + 3>& a_g_n_k_wos_strides,

600 const std::array<ck::index_t, NDimSpatial>& conv_filter_strides,

601 const std::array<ck::index_t, NDimSpatial>& conv_filter_dilations,

602 const std::array<ck::index_t, NDimSpatial>& input_left_pads,

603 const std::array<ck::index_t, NDimSpatial>& input_right_pads,

604 const ck::index_t M01,

605 const ck::index_t N01,

606 InElementwiseOperation in_element_op,

607 WeiElementwiseOperation wei_element_op,

608 OutElementwiseOperation out_element_op,

609 ck::index_t split_k)

610 : p_a_grid_{p_out_grid},

611 p_b_grid_{p_in_grid},

612 p_e_grid_{p_wei_grid},

613 a_grid_desc_k0_m_k1_{},

614 b_grid_desc_k0_n_k1_{},

615 ce_grid_desc_m_n_{},

616 c_grid_desc_mblock_mperblock_nblock_nperblock_{},

617 compute_ptr_offset_of_batch_{},

618 M01_{M01},

619 N01_{N01},

620 a_element_op_{out_element_op},

621 b_element_op_{in_element_op},

622 cde_element_op_{wei_element_op},

623 Conv_G_{b_g_n_c_wis_lengths[0]},

624 Conv_N_{b_g_n_c_wis_lengths[1]},

625 Conv_K_{e_g_k_c_xs_lengths[1]},

626 Conv_C_{b_g_n_c_wis_lengths[2]},

627 input_spatial_lengths_{},

628 filter_spatial_lengths_{},

629 output_spatial_lengths_{},

630 conv_filter_strides_{conv_filter_strides},

631 input_left_pads_{input_left_pads},

632 input_right_pads_{input_right_pads}

633 {

634 static ActiveWorkgroupsPerCU active_workgroups_per_cu;

635

636 c_space_size_bytes =

637 ck::accumulate_n<long_index_t>(

638 e_g_k_c_xs_lengths.begin(), NDimSpatial + I3, 1, std::multiplies<>()) *

639 sizeof(AccDataType);

640

641 constexpr index_t spatial_offset = 3;

642 std::copy(begin(b_g_n_c_wis_lengths) + spatial_offset,

643 end(b_g_n_c_wis_lengths),

644 begin(input_spatial_lengths_));

645 std::copy(begin(e_g_k_c_xs_lengths) + spatial_offset,

646 end(e_g_k_c_xs_lengths),

647 begin(filter_spatial_lengths_));

648 std::copy(begin(a_g_n_k_wos_lengths) + spatial_offset,

649 end(a_g_n_k_wos_lengths),

650 begin(output_spatial_lengths_));

651

652 std::array<index_t, NDimSpatial + 3> a_g_n_k_wos_strides_transposed =

653 conv_ngchw_to_nhwgc_transformer.TransposeInOutStrides(a_g_n_k_wos_lengths,

654 a_g_n_k_wos_strides);

655 std::array<index_t, NDimSpatial + 3> b_g_n_c_wis_strides_transposed =

656 conv_ngchw_to_nhwgc_transformer.TransposeInOutStrides(b_g_n_c_wis_lengths,

657 b_g_n_c_wis_strides);

658 std::array<index_t, NDimSpatial + 3> e_g_k_c_xs_strides_transposed =

659 conv_ngchw_to_nhwgc_transformer.TransposeWeiStrides(e_g_k_c_xs_lengths,

660 e_g_k_c_xs_strides);

661

662 if(split_k < 0)

663 {

664 ck::index_t gemmM, gemmN, gemmK;

665 std::tie(gemmM, gemmN, gemmK) =

666 get_bwd_weight_gemm_sizes<NDimSpatial>(a_g_n_k_wos_lengths, e_g_k_c_xs_lengths);

667

668 const auto grid_size = calculate_mn_grid_size<MPerBlock, NPerBlock>(gemmM, gemmN) *

669 Conv_G_ / NumGroupsToMerge;

670 k_batch_ = get_best_occupancy_k_batch_value(active_workgroups_per_cu.max_occupancy_,

671 grid_size);

672

673 // Ensure that k_batch_ does not exceed the maximum value

674 // for the GEMM pipeline.

675 const auto k_batch_max = static_cast<index_t>((gemmK - 1) / KPerBlock);

676 k_batch_ = std::max(std::min(k_batch_, k_batch_max), 1);

677

678 if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))

679 {

680 std::cout << "[SPLIT-K AUTODEDUCE] k_batch max value: " << k_batch_max

681 << std::endl;

682 std::cout << "[SPLIT-K AUTODEDUCE] Final k_batch value: " << k_batch_

683 << std::endl;

684 }

685 }

686 else

687 {

688 k_batch_ = split_k;

689 }

690

691 const auto descs =

692 conv_to_gemm_transformer_v2

693 .template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<NDimSpatial>(

694 Conv_N_,

695 Conv_K_,

696 Conv_C_,

697 input_spatial_lengths_,

698 filter_spatial_lengths_,

699 output_spatial_lengths_,

700 b_g_n_c_wis_strides_transposed,

701 e_g_k_c_xs_strides_transposed,

702 a_g_n_k_wos_strides_transposed,

703 conv_filter_strides,

704 conv_filter_dilations,

705 input_left_pads,

706 input_right_pads,

707 k_batch_);

708

709 a_grid_desc_k0_m_k1_ = descs[I0];

710 b_grid_desc_k0_n_k1_ = descs[I1];

711 ce_grid_desc_m_n_ = descs[I2];

712

713 ce_elementwise_grid_desc_m_n_ =

714 conv_to_gemm_transformer_v1

715 .template MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N<NDimSpatial>(

716 Conv_N_,

717 Conv_K_,

718 Conv_C_,

719 input_spatial_lengths_,

720 filter_spatial_lengths_,

721 output_spatial_lengths_,

722 b_g_n_c_wis_strides,

723 e_g_k_c_xs_strides,

724 a_g_n_k_wos_strides,

725 conv_filter_strides,

726 conv_filter_dilations,

727 input_left_pads,

728 input_right_pads,

729 k_batch_)[I2];

730

731 const index_t GemmM = a_grid_desc_k0_m_k1_.GetLength(I1);

732 const index_t GemmN = b_grid_desc_k0_n_k1_.GetLength(I1);

733

734 // A/B/C Batch Stride

735 compute_ptr_offset_of_batch_.BatchStrideA_ = a_g_n_k_wos_strides_transposed[0];

736 compute_ptr_offset_of_batch_.BatchStrideB_ = b_g_n_c_wis_strides_transposed[0];

737 compute_ptr_offset_of_batch_.BatchStrideC_ = e_g_k_c_xs_strides_transposed[0];

738 c_grid_desc_mblock_mperblock_nblock_nperblock_ =

739 GridwiseGemm64::MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(

740 ce_grid_desc_m_n_,

741 GridwiseGemm64::CalculateMBlock(GemmM),

742 GridwiseGemm64::CalculateNBlock(GemmN));

743

744 if constexpr(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

745 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>())

746 {

747 a_in_transpose_desc_ =

748 conv_ngchw_to_nhwgc_transformer.template MakeNGCHWTransposeDesc<NDimSpatial>(

749 a_g_n_k_wos_lengths, a_g_n_k_wos_strides);

750 a_out_transpose_desc_ =

751 conv_ngchw_to_nhwgc_transformer.template MakeNHWGCTransposeDesc<NDimSpatial>(

752 a_g_n_k_wos_lengths, a_g_n_k_wos_strides);

753

754 b_in_transpose_desc_ =

755 conv_ngchw_to_nhwgc_transformer.template MakeNGCHWTransposeDesc<NDimSpatial>(

756 b_g_n_c_wis_lengths, b_g_n_c_wis_strides);

757 b_out_transpose_desc_ =

758 conv_ngchw_to_nhwgc_transformer.template MakeNHWGCTransposeDesc<NDimSpatial>(

759 b_g_n_c_wis_lengths, b_g_n_c_wis_strides);

760

761 e_in_transpose_desc_ =

762 conv_ngchw_to_nhwgc_transformer.template MakeGKYXCTransposeDesc<NDimSpatial>(

763 e_g_k_c_xs_lengths, e_g_k_c_xs_strides);

764 e_out_transpose_desc_ =

765 conv_ngchw_to_nhwgc_transformer.template MakeGKCYXTransposeDesc<NDimSpatial>(

766 e_g_k_c_xs_lengths, e_g_k_c_xs_strides);

767

768 elementwise_block_2_ctile_map_transpose_a_ = Block2TileMapElementwise{

769 a_in_transpose_desc_.GetLength(I0), a_in_transpose_desc_.GetLength(I1)};

770

771 elementwise_block_2_ctile_map_transpose_b_ = Block2TileMapElementwise{

772 b_in_transpose_desc_.GetLength(I0), b_in_transpose_desc_.GetLength(I1)};

773 }

774

775 elementwise_block_2_ctile_map_ =

776 is_NGCHW_GKCYX_NGKHW<InLayout, WeiLayout, OutLayout>() ||

777 is_NGCDHW_GKCZYX_NGKDHW<InLayout, WeiLayout, OutLayout>()

778 ? Block2TileMapElementwise{e_in_transpose_desc_.GetLength(I0),

779 e_in_transpose_desc_.GetLength(I1)}

780 : Block2TileMapElementwise{ce_grid_desc_m_n_.GetLength(I0),

781 ce_grid_desc_m_n_.GetLength(I1)};

782 }

783

784 std::size_t GetWorkspaceATensorSizeBytes() const

785 {

786 // Align to 128B

787 return math::integer_divide_ceil(

788 sizeof(ADataType) * a_in_transpose_desc_.GetElementSpaceSize(), 128) *

789 128;

790 }

791

792 std::size_t GetWorkspaceBTensorSizeBytes() const

793 {

794 return sizeof(BDataType) * b_in_transpose_desc_.GetElementSpaceSize();

795 }

796

797 std::size_t GetWorkspaceETensorSizeBytes() const

798 {

799 // Align to 128B

800 return math::integer_divide_ceil(sizeof(AccDataType) *

801 ce_grid_desc_m_n_.GetElementSpaceSize() * Conv_G_,

802 128) *

803 128;

804 }

805

806 std::size_t GetWorkspaceSizeBytes() const

807 {

808 // 1. We need to transpose A and B for NGCHW and NGKHW layouts

809 // 2. If C format is GKCYX then tranpose during second stage.

810 // If C format is GKYXC then just perform second stage.

811 // Due to the fact that E workspace is always needed, we

812 // allocate them as the first part of the workspace.

813 // [EWorkspace, AWorkspace, BWorkspace]

814 if constexpr(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

815 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>())

816 {

817 return GetWorkspaceATensorSizeBytes() + GetWorkspaceBTensorSizeBytes() +

818 GetWorkspaceETensorSizeBytes();

819 }

820 else

821 {

822 return GetWorkspaceETensorSizeBytes();

823 }

824 }

825

826 const ADataType* p_a_grid_;

827 const BDataType* p_b_grid_;

828 EDataType* p_e_grid_;

829

830 AGridDesc_K0_M_K1 a_grid_desc_k0_m_k1_;

831 BGridDesc_K0_N_K1 b_grid_desc_k0_n_k1_;

832 CGridDesc_M_N ce_grid_desc_m_n_;

833 CElementwiseGridDesc_M_N ce_elementwise_grid_desc_m_n_;

834 CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock c_grid_desc_mblock_mperblock_nblock_nperblock_;

835

836 Block2TileMapElementwise elementwise_block_2_ctile_map_;

837 Block2TileMapElementwise elementwise_block_2_ctile_map_transpose_a_,

838 elementwise_block_2_ctile_map_transpose_b_;

839

840 NGCHWTransposeDescType a_in_transpose_desc_, b_in_transpose_desc_;

841 NHWGCTransposeDescType a_out_transpose_desc_, b_out_transpose_desc_;

842 GKYXCTransposeDescType e_in_transpose_desc_;

843 GKCYXTransposeDescType e_out_transpose_desc_;

844

845 // for computing batch offset

846 ComputePtrOffsetOfStridedBatch<I1, I1, I0> compute_ptr_offset_of_batch_;

847

848 index_t M01_;

849 index_t N01_;

850

851 OutElementwiseOperation a_element_op_;

852 InElementwiseOperation b_element_op_;

853 WeiElementwiseOperation cde_element_op_;

854

855 // for checking IsSupportedArgument()

856 const index_t Conv_G_;

857 const index_t Conv_N_;

858 const index_t Conv_K_;

859 const index_t Conv_C_;

860 std::array<ck::index_t, NDimSpatial> input_spatial_lengths_;

861 std::array<ck::index_t, NDimSpatial> filter_spatial_lengths_;

862 std::array<ck::index_t, NDimSpatial> output_spatial_lengths_;

863 const std::array<ck::index_t, NDimSpatial>& conv_filter_strides_;

864 const std::array<ck::index_t, NDimSpatial>& input_left_pads_;

865 const std::array<ck::index_t, NDimSpatial>& input_right_pads_;

866 long_index_t c_space_size_bytes;

867 };

868

869 // Invoker

870 struct Invoker : public BaseInvoker

871 {

872 using Argument = DeviceOp::Argument;

873

874 void ShowInfo(const Argument& arg)

875 {

876 std::cout << "arg.a_grid_desc_k0_m_k1_{" << arg.a_grid_desc_k0_m_k1_.GetLength(I0)

877 << ", " << arg.a_grid_desc_k0_m_k1_.GetLength(I1) << ", "

878 << arg.a_grid_desc_k0_m_k1_.GetLength(I2) << "}" << std::endl;

879

880 std::cout << "arg.b_grid_desc_k0_n_k1_{" << arg.b_grid_desc_k0_n_k1_.GetLength(I0)

881 << ", " << arg.b_grid_desc_k0_n_k1_.GetLength(I1) << ", "

882 << arg.b_grid_desc_k0_n_k1_.GetLength(I2) << "}" << std::endl;

883

884 std::cout << "arg.ce_grid_desc_m_n_{" << arg.ce_grid_desc_m_n_.GetLength(I0) << ", "

885 << arg.ce_grid_desc_m_n_.GetLength(I1) << "}" << std::endl;

886 }

887

888 template <typename GridwiseGemm>

889 float RunGemmV3(const Argument& arg, const StreamConfig& stream_config = StreamConfig{})

890 {

891 const index_t GemmM = arg.a_grid_desc_k0_m_k1_.GetLength(I1);

892 const index_t GemmN = arg.b_grid_desc_k0_n_k1_.GetLength(I1);

893 const index_t GemmK =

894 arg.a_grid_desc_k0_m_k1_.GetLength(I0) * arg.a_grid_desc_k0_m_k1_.GetLength(I2);

895

896 AccDataType* p_c_grid = type_convert<AccDataType*>(arg.p_workspace_);

897

898 const ADataType* p_a_grid = arg.p_a_grid_;

899 const BDataType* p_b_grid = arg.p_b_grid_;

900

901 if constexpr(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

902 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>())

903 {

904 p_a_grid = type_convert<const ADataType*>(arg.p_workspace_) +

905 arg.GetWorkspaceETensorSizeBytes() / sizeof(ADataType);

906 p_b_grid =

907 type_convert<const BDataType*>(arg.p_workspace_) +

908 (arg.GetWorkspaceETensorSizeBytes() + arg.GetWorkspaceATensorSizeBytes()) /

909 sizeof(BDataType);

910 }

911

912 // nullptr for output, will be set after workspace set

913 typename GridwiseGemm::Argument gemm_arg{

914 p_a_grid, p_b_grid, p_c_grid, GemmM, GemmN, GemmK, I0, I0, I0, arg.k_batch_};

915

916 index_t gdx, gdy, gdz;

917 std::tie(gdx, gdy, gdz) = GridwiseGemm::CalculateGridSize(

918 gemm_arg.M, gemm_arg.N, gemm_arg.KBatch, arg.Conv_G_ / NumGroupsToMerge);

919

920 float ave_time = 0;

921

922 index_t k_grain = gemm_arg.KBatch * KPerBlock;

923 index_t K_split = (gemm_arg.K + k_grain - 1) / k_grain * (KPerBlock);

924

925 const bool has_main_k_block_loop = GridwiseGemm::CalculateHasMainKBlockLoop(K_split);

926

927 const auto num_k_per_block =

928 arg.a_grid_desc_k0_m_k1_.GetLength(Number<0>{}) / gemm_arg.KBatch;

929

930 const auto clear_workspace = [&]() {

931 if(arg.k_batch_ > 1)

932 {

933 hip_check_error(hipMemsetAsync(

934 gemm_arg.p_c_grid, 0, arg.c_space_size_bytes, stream_config.stream_id_));

935 }

936 };

937

938 const auto Run = [&](const auto& kernel) {

939 if(stream_config.flush_cache)

940 {

941 typename GridwiseGemm::Argument gemm_arg_ = gemm_arg;

942 ck::utility::RotatingMemWrapper<typename GridwiseGemm::Argument> rotating_mem(

943 gemm_arg_,

944 stream_config.rotating_count,

945 gemm_arg_.M * gemm_arg_.K * sizeof(ADataType),

946 gemm_arg_.K * gemm_arg_.N * sizeof(BDataType));

947 rotating_mem.Print();

948

949 auto run_flush_cache = [&]() {

950 // flush icache

951 ck::utility::flush_icache();

952 // rotating mem

953 rotating_mem.Next();

954 clear_workspace();

955 };

956

957 ave_time += ck::utility::launch_and_time_kernel_with_preprocess<false>(

958 stream_config,

959 run_flush_cache,

960 kernel,

961 dim3(gdx, gdy, gdz),

962 dim3(BlockSize),

963 0,

964 gemm_arg_,

965 arg.a_grid_desc_k0_m_k1_,

966 arg.b_grid_desc_k0_n_k1_,

967 arg.c_grid_desc_mblock_mperblock_nblock_nperblock_,

968 arg.compute_ptr_offset_of_batch_,

969 num_k_per_block);

970 }

971 else

972 {

973 ave_time += launch_and_time_kernel_with_preprocess(

974 stream_config,

975 clear_workspace,

976 kernel,

977 dim3(gdx, gdy, gdz),

978 dim3(BlockSize),

979 0,

980 gemm_arg,

981 arg.a_grid_desc_k0_m_k1_,

982 arg.b_grid_desc_k0_n_k1_,

983 arg.c_grid_desc_mblock_mperblock_nblock_nperblock_,

984 arg.compute_ptr_offset_of_batch_,

985 num_k_per_block);

986 }

987 };

988

989 constexpr index_t minimum_occupancy =

990 BlkGemmPipeSched == BlockGemmPipelineScheduler::Intrawave ? 1 : 2;

991

992 if(has_main_k_block_loop)

993 {

994 // Tail number always full

995 if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v1 ||

996 BlkGemmPipelineVer == BlockGemmPipelineVersion::v3)

997 {

998 if(gemm_arg.KBatch > 1)

999 {

1000 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1001 GridwiseGemm,

1002 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1003 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1004 remove_reference_t<

1005 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1006 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1007 NumGroupsToMerge,

1008 true,

1009 InMemoryDataOperationEnum::AtomicAdd,

1010 minimum_occupancy>;

1011 Run(kernel);

1012 }

1013 else

1014 {

1015 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1016 GridwiseGemm,

1017 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1018 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1019 remove_reference_t<

1020 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1021 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1022 NumGroupsToMerge,

1023 true,

1024 InMemoryDataOperationEnum::Set,

1025 minimum_occupancy>;

1026 Run(kernel);

1027 }

1028 }

1029 // Tail number could be One to Seven

1030 else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v2)

1031 {

1032 if(gemm_arg.KBatch > 1)

1033 {

1034 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::One)

1035 {

1036 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1037 GridwiseGemm,

1038 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1039 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1040 remove_reference_t<

1041 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1042 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1043 NumGroupsToMerge,

1044 true,

1045 InMemoryDataOperationEnum::AtomicAdd,

1046 minimum_occupancy,

1047 TailNumber::One>;

1048 Run(kernel);

1049 }

1050 else if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1051 TailNumber::Full)

1052 {

1053 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1054 GridwiseGemm,

1055 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1056 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1057 remove_reference_t<

1058 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1059 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1060 NumGroupsToMerge,

1061 true,

1062 InMemoryDataOperationEnum::AtomicAdd,

1063 minimum_occupancy,

1064 TailNumber::Full>;

1065 Run(kernel);

1066 }

1067

1068 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 2)

1069 {

1070 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Two)

1071 {

1072 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1073 GridwiseGemm,

1074 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1075 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1076 remove_reference_t<

1077 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1078 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1079 NumGroupsToMerge,

1080 true,

1081 InMemoryDataOperationEnum::AtomicAdd,

1082 minimum_occupancy,

1083 TailNumber::Two>;

1084 Run(kernel);

1085 }

1086 }

1087

1088 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 3)

1089 {

1090 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1091 TailNumber::Three)

1092 {

1093 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1094 GridwiseGemm,

1095 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1096 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1097 remove_reference_t<

1098 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1099 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1100 NumGroupsToMerge,

1101 true,

1102 InMemoryDataOperationEnum::AtomicAdd,

1103 minimum_occupancy,

1104 TailNumber::Three>;

1105 Run(kernel);

1106 }

1107 }

1108

1109 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 4)

1110 {

1111 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1112 TailNumber::Four)

1113 {

1114 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1115 GridwiseGemm,

1116 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1117 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1118 remove_reference_t<

1119 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1120 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1121 NumGroupsToMerge,

1122 true,

1123 InMemoryDataOperationEnum::AtomicAdd,

1124 minimum_occupancy,

1125 TailNumber::Four>;

1126 Run(kernel);

1127 }

1128 }

1129

1130 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 5)

1131 {

1132 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1133 TailNumber::Five)

1134 {

1135 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1136 GridwiseGemm,

1137 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1138 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1139 remove_reference_t<

1140 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1141 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1142 NumGroupsToMerge,

1143 true,

1144 InMemoryDataOperationEnum::AtomicAdd,

1145 minimum_occupancy,

1146 TailNumber::Five>;

1147 Run(kernel);

1148 }

1149 }

1150

1151 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 6)

1152 {

1153 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Six)

1154 {

1155 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1156 GridwiseGemm,

1157 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1158 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1159 remove_reference_t<

1160 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1161 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1162 NumGroupsToMerge,

1163 true,

1164 InMemoryDataOperationEnum::AtomicAdd,

1165 minimum_occupancy,

1166 TailNumber::Six>;

1167 Run(kernel);

1168 }

1169 }

1170

1171 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 7)

1172 {

1173 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1174 TailNumber::Seven)

1175 {

1176 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1177 GridwiseGemm,

1178 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1179 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1180 remove_reference_t<

1181 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1182 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1183 NumGroupsToMerge,

1184 true,

1185 InMemoryDataOperationEnum::AtomicAdd,

1186 minimum_occupancy,

1187 TailNumber::Seven>;

1188 Run(kernel);

1189 }

1190 }

1191 }

1192 else

1193 {

1194 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::One)

1195 {

1196 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1197 GridwiseGemm,

1198 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1199 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1200 remove_reference_t<

1201 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1202 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1203 NumGroupsToMerge,

1204 true,

1205 InMemoryDataOperationEnum::Set,

1206 minimum_occupancy,

1207 TailNumber::One>;

1208 Run(kernel);

1209 }

1210 else if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1211 TailNumber::Full)

1212 {

1213 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1214 GridwiseGemm,

1215 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1216 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1217 remove_reference_t<

1218 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1219 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1220 NumGroupsToMerge,

1221 true,

1222 InMemoryDataOperationEnum::Set,

1223 minimum_occupancy,

1224 TailNumber::Full>;

1225 Run(kernel);

1226 }

1227

1228 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 2)

1229 {

1230 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Two)

1231 {

1232 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1233 GridwiseGemm,

1234 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1235 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1236 remove_reference_t<

1237 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1238 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1239 NumGroupsToMerge,

1240 true,

1241 InMemoryDataOperationEnum::Set,

1242 minimum_occupancy,

1243 TailNumber::Two>;

1244 Run(kernel);

1245 }

1246 }

1247

1248 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 3)

1249 {

1250 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1251 TailNumber::Three)

1252 {

1253 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1254 GridwiseGemm,

1255 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1256 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1257 remove_reference_t<

1258 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1259 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1260 NumGroupsToMerge,

1261 true,

1262 InMemoryDataOperationEnum::Set,

1263 minimum_occupancy,

1264 TailNumber::Three>;

1265 Run(kernel);

1266 }

1267 }

1268

1269 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 4)

1270 {

1271 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1272 TailNumber::Four)

1273 {

1274 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1275 GridwiseGemm,

1276 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1277 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1278 remove_reference_t<

1279 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1280 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1281 NumGroupsToMerge,

1282 true,

1283 InMemoryDataOperationEnum::Set,

1284 minimum_occupancy,

1285 TailNumber::Four>;

1286 Run(kernel);

1287 }

1288 }

1289

1290 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 5)

1291 {

1292 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1293 TailNumber::Five)

1294 {

1295 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1296 GridwiseGemm,

1297 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1298 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1299 remove_reference_t<

1300 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1301 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1302 NumGroupsToMerge,

1303 true,

1304 InMemoryDataOperationEnum::Set,

1305 minimum_occupancy,

1306 TailNumber::Five>;

1307 Run(kernel);

1308 }

1309 }

1310

1311 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 6)

1312 {

1313 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Six)

1314 {

1315 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1316 GridwiseGemm,

1317 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1318 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1319 remove_reference_t<

1320 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1321 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1322 NumGroupsToMerge,

1323 true,

1324 InMemoryDataOperationEnum::Set,

1325 minimum_occupancy,

1326 TailNumber::Six>;

1327 Run(kernel);

1328 }

1329 }

1330

1331 if constexpr(GridwiseGemm::BlockwiseGemmPipe::PrefetchStages > 7)

1332 {

1333 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) ==

1334 TailNumber::Seven)

1335 {

1336 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1337 GridwiseGemm,

1338 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1339 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1340 remove_reference_t<

1341 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1342 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1343 NumGroupsToMerge,

1344 true,

1345 InMemoryDataOperationEnum::Set,

1346 minimum_occupancy,

1347 TailNumber::Seven>;

1348 Run(kernel);

1349 }

1350 }

1351 }

1352 }

1353 // Tail number could be Odd or Even

1354 else if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v4)

1355 {

1356 if(gemm_arg.KBatch > 1)

1357 {

1358 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd)

1359 {

1360 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds<

1361 GridwiseGemm,

1362 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1363 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1364 remove_reference_t<

1365 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1366 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1367 NumGroupsToMerge,

1368 true,

1369 InMemoryDataOperationEnum::AtomicAdd,

1370 minimum_occupancy,

1371 TailNumber::Odd>;

1372 Run(kernel);

1373 }

1374 else

1375 {

1376 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds<

1377 GridwiseGemm,

1378 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1379 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1380 remove_reference_t<

1381 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1382 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1383 NumGroupsToMerge,

1384 true,

1385 InMemoryDataOperationEnum::AtomicAdd,

1386 minimum_occupancy,

1387 TailNumber::Even>;

1388 Run(kernel);

1389 }

1390 }

1391 else

1392 {

1393 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd)

1394 {

1395 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds<

1396 GridwiseGemm,

1397 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1398 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1399 remove_reference_t<

1400 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1401 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1402 NumGroupsToMerge,

1403 true,

1404 InMemoryDataOperationEnum::Set,

1405 minimum_occupancy,

1406 TailNumber::Odd>;

1407 Run(kernel);

1408 }

1409 else

1410 {

1411 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds<

1412 GridwiseGemm,

1413 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1414 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1415 remove_reference_t<

1416 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1417 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1418 NumGroupsToMerge,

1419 true,

1420 InMemoryDataOperationEnum::Set,

1421 minimum_occupancy,

1422 TailNumber::Even>;

1423 Run(kernel);

1424 }

1425 }

1426 }

1427 else

1428 {

1429 if(gemm_arg.KBatch > 1)

1430 {

1431 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd)

1432 {

1433 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1434 GridwiseGemm,

1435 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1436 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1437 remove_reference_t<

1438 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1439 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1440 NumGroupsToMerge,

1441 true,

1442 InMemoryDataOperationEnum::AtomicAdd,

1443 minimum_occupancy,

1444 TailNumber::Odd>;

1445 Run(kernel);

1446 }

1447 else

1448 {

1449 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1450 GridwiseGemm,

1451 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1452 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1453 remove_reference_t<

1454 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1455 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1456 NumGroupsToMerge,

1457 true,

1458 InMemoryDataOperationEnum::AtomicAdd,

1459 minimum_occupancy,

1460 TailNumber::Even>;

1461 Run(kernel);

1462 }

1463 }

1464 else

1465 {

1466 if(GridwiseGemm::CalculateKBlockLoopTailNum(K_split) == TailNumber::Odd)

1467 {

1468 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1469 GridwiseGemm,

1470 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1471 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1472 remove_reference_t<

1473 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1474 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1475 NumGroupsToMerge,

1476 true,

1477 InMemoryDataOperationEnum::Set,

1478 minimum_occupancy,

1479 TailNumber::Odd>;

1480 Run(kernel);

1481 }

1482 else

1483 {

1484 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1485 GridwiseGemm,

1486 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1487 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1488 remove_reference_t<

1489 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1490 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1491 NumGroupsToMerge,

1492 true,

1493 InMemoryDataOperationEnum::Set,

1494 minimum_occupancy,

1495 TailNumber::Even>;

1496 Run(kernel);

1497 }

1498 }

1499 }

1500 }

1501 else

1502 {

1503 // Tail number always 1

1504 if constexpr(BlkGemmPipelineVer == BlockGemmPipelineVersion::v1)

1505 {

1506 if(gemm_arg.KBatch > 1)

1507 {

1508 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1509 GridwiseGemm,

1510 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1511 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1512 remove_reference_t<

1513 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1514 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1515 NumGroupsToMerge,

1516 false,

1517 InMemoryDataOperationEnum::AtomicAdd,

1518 minimum_occupancy>;

1519 Run(kernel);

1520 }

1521 else

1522 {

1523 const auto kernel = kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3<

1524 GridwiseGemm,

1525 remove_reference_t<DeviceOp::AGridDesc_K0_M_K1>,

1526 remove_reference_t<DeviceOp::BGridDesc_K0_N_K1>,

1527 remove_reference_t<

1528 DeviceOp::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock>,

1529 ComputePtrOffsetOfStridedBatch<I1, I1, I0>,

1530 NumGroupsToMerge,

1531 false,

1532 InMemoryDataOperationEnum::Set,

1533 minimum_occupancy>;

1534 Run(kernel);

1535 }

1536 }

1537 }

1538

1539 return ave_time;

1540 }

1541

1542 template <typename GridwiseGemm>

1543 float RunImp(const Argument& arg, const StreamConfig& stream_config = StreamConfig{})

1544 {

1545 float avg_time = 0.f;

1546 auto launch_elementwise_kernel = [&]() {

1547 const AccDataType* p_c_grid = type_convert<const AccDataType*>(arg.p_workspace_);

1548

1549 std::array<index_t, I1> in_out_batch_strides = {

1550 static_cast<index_t>(arg.compute_ptr_offset_of_batch_.BatchStrideC_)};

1551

1552 if constexpr(is_NGCHW_GKCYX_NGKHW<InLayout, WeiLayout, OutLayout>() ||

1553 is_NGCDHW_GKCZYX_NGKDHW<InLayout, WeiLayout, OutLayout>())

1554 {

1555 const index_t grid_size = arg.elementwise_block_2_ctile_map_.CalculateGridSize(

1556 arg.e_in_transpose_desc_);

1557

1558 const auto kernel = kernel_elementwise<GridwiseElementwiseWeightTransposeCast,

1559 ck::Tuple<GKYXCTransposeDescType>,

1560 ck::Tuple<GKCYXTransposeDescType>,

1561 ck::Tuple<const AccDataType*>,

1562 ck::Tuple<EDataType*>,

1563 Block2TileMapElementwise,

1564 CDEElementwiseOperation>;

1565

1566 return launch_and_time_kernel(stream_config,

1567 kernel,

1568 dim3(grid_size),

1569 dim3(BlockSize),

1570 0,

1571 make_tuple(arg.e_in_transpose_desc_),

1572 make_tuple(arg.e_out_transpose_desc_),

1573 make_tuple(p_c_grid),

1574 make_tuple(arg.p_e_grid_),

1575 arg.elementwise_block_2_ctile_map_,

1576 arg.cde_element_op_);

1577 }

1578 else

1579 {

1580 const index_t grid_size = arg.elementwise_block_2_ctile_map_.CalculateGridSize(

1581 arg.ce_elementwise_grid_desc_m_n_) *

1582 arg.Conv_G_;

1583

1584 const auto kernel =

1585 kernel_batched_elementwise<GridwiseElementwiseCast,

1586 ck::Tuple<CElementwiseGridDesc_M_N>,

1587 ck::Tuple<CElementwiseGridDesc_M_N>,

1588 ck::Tuple<const AccDataType*>,

1589 ck::Tuple<EDataType*>,

1590 Block2TileMapElementwise,

1591 CDEElementwiseOperation,

1592 I1,

1593 I1>;

1594

1595 return launch_and_time_kernel(stream_config,

1596 kernel,

1597 dim3(grid_size),

1598 dim3(BlockSize),

1599 0,

1600 make_tuple(arg.ce_elementwise_grid_desc_m_n_),

1601 make_tuple(arg.ce_elementwise_grid_desc_m_n_),

1602 make_tuple(p_c_grid),

1603 make_tuple(arg.p_e_grid_),

1604 arg.elementwise_block_2_ctile_map_,

1605 arg.cde_element_op_,

1606 arg.Conv_G_,

1607 in_out_batch_strides,

1608 in_out_batch_strides);

1609 }

1610 };

1611

1612 if constexpr(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

1613 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>())

1614 {

1615 const index_t grid_size_a =

1616 arg.elementwise_block_2_ctile_map_transpose_a_.CalculateGridSize(

1617 arg.a_in_transpose_desc_);

1618 const index_t grid_size_b =

1619 arg.elementwise_block_2_ctile_map_transpose_b_.CalculateGridSize(

1620 arg.b_in_transpose_desc_);

1621

1622 ADataType* p_a_out_grid = type_convert<ADataType*>(arg.p_workspace_) +

1623 arg.GetWorkspaceETensorSizeBytes() / sizeof(ADataType);

1624 BDataType* p_b_out_grid =

1625 type_convert<BDataType*>(arg.p_workspace_) +

1626 (arg.GetWorkspaceETensorSizeBytes() + arg.GetWorkspaceATensorSizeBytes()) /

1627 sizeof(BDataType);

1628

1629 // Different data type for A and B is not supported

1630 auto kernel_transpose = kernel_elementwise_dual<GridwiseElementwiseTranspose,

1631 GridwiseElementwiseTranspose,

1632 ck::Tuple<NGCHWTransposeDescType>,

1633 ck::Tuple<NGCHWTransposeDescType>,

1634 ck::Tuple<NHWGCTransposeDescType>,

1635 ck::Tuple<NHWGCTransposeDescType>,

1636 ck::Tuple<const ADataType*>,

1637 ck::Tuple<const ADataType*>,

1638 ck::Tuple<ADataType*>,

1639 ck::Tuple<ADataType*>,

1640 Block2TileMapElementwise,

1641 Block2TileMapElementwise,

1642 element_wise::PassThrough>;

1643

1644 avg_time += launch_and_time_kernel(stream_config,

1645 kernel_transpose,

1646 dim3(grid_size_a + grid_size_b),

1647 dim3(BlockSize),

1648 0,

1649 make_tuple(arg.a_in_transpose_desc_),

1650 make_tuple(arg.b_in_transpose_desc_),

1651 make_tuple(arg.a_out_transpose_desc_),

1652 make_tuple(arg.b_out_transpose_desc_),

1653 make_tuple(arg.p_a_grid_),

1654 make_tuple(arg.p_b_grid_),

1655 make_tuple(p_a_out_grid),

1656 make_tuple(p_b_out_grid),

1657 arg.elementwise_block_2_ctile_map_transpose_a_,

1658 arg.elementwise_block_2_ctile_map_transpose_b_,

1659 element_wise::PassThrough{},

1660 grid_size_a);

1661 }

1662

1663 avg_time += RunGemmV3<GridwiseGemm>(arg, stream_config);

1664 avg_time += launch_elementwise_kernel();

1665 return avg_time;

1666 }

1667

1668 INVOKER_RUN_IMPL

1669

1670 float Run(const BaseArgument* p_arg,

1671 const StreamConfig& stream_config = StreamConfig{}) override

1672 {

1673 return Run(*dynamic_cast<const Argument*>(p_arg), stream_config);

1674 }

1675 };

1676

1677 static constexpr bool IsValidCompilationParameter()

1678 {

1679 // TODO: properly implement this check

1680 return true;

1681 }

1682

1683 static bool IsSupportedArgument(const Argument& arg)

1684 {

1685 const index_t GemmM = arg.a_grid_desc_k0_m_k1_.GetLength(I1);

1686 const index_t GemmN = arg.b_grid_desc_k0_n_k1_.GetLength(I1);

1687 const index_t GemmK =

1688 arg.a_grid_desc_k0_m_k1_.GetLength(I0) * arg.a_grid_desc_k0_m_k1_.GetLength(I2);

1689

1690 if constexpr(is_same_v<ComputeTypeA, ck::tf32_t> || is_same_v<ComputeTypeB, ck::tf32_t>)

1691 {

1692 if(!is_tf32_supported())

1693 {

1694 return false;

1695 }

1696 if constexpr(!is_same_v<ComputeTypeA, ComputeTypeB>)

1697 {

1698 if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))

1699 {

1700 std::cout << "ComputeDataType for A and B should be same while using TF32"

1701 << std::endl;

1702 }

1703 return false;

1704 }

1705 }

1706

1707 if(get_warp_size() == 64)

1708 {

1709 if constexpr(NXdlPerWave64 > 0)

1710 {

1711 typename GridwiseGemm64::Argument gemm_arg{

1712 nullptr, nullptr, nullptr, GemmM, GemmN, GemmK, I0, I0, I0, arg.k_batch_};

1713

1714 const auto num_k_loop = gemm_arg.AK0 / (KPerBlock / K1);

1715 if constexpr(BlkGemmPipelineVer != BlockGemmPipelineVersion::v1)

1716 {

1717 if(num_k_loop <= GridwiseGemm64::BlockwiseGemmPipe::PrefetchStages)

1718 {

1719 return false;

1720 }

1721 }

1722 }

1723 else

1724 {

1725 return false;

1726 }

1727 }

1728 else

1729 {

1730 if constexpr(NXdlPerWave32 > 0)

1731 {

1732 typename GridwiseGemm32::Argument gemm_arg{

1733 nullptr, nullptr, nullptr, GemmM, GemmN, GemmK, I0, I0, I0, arg.k_batch_};

1734

1735 const auto num_k_loop = gemm_arg.AK0 / (KPerBlock / K1);

1736 if constexpr(BlkGemmPipelineVer != BlockGemmPipelineVersion::v1)

1737 {

1738 if(num_k_loop <= GridwiseGemm32::BlockwiseGemmPipe::PrefetchStages)

1739 {

1740 return false;

1741 }

1742 }

1743 }

1744 else

1745 {

1746 return false;

1747 }

1748 }

1749

1750 // Check this here, it allows to use other instances from factory even

1751 // if workspace is not allocated

1752 if(!arg.p_workspace_)

1753 {

1754 if(ck::EnvIsEnabled(CK_ENV(CK_LOGGING)))

1755 {

1756 std::cout << "Warning: Workspace for "

1757 "DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument is not "

1758 "allocated, use SetWorkSpacePointer."

1759 << std::endl;

1760 }

1761 return false;

1762 }

1763 if(!ck::is_xdl_wmma_supported<ComputeTypeA, ComputeTypeB, MPerXDL, NPerXDL>())

1764 {

1765 return false;

1766 }

1767 if constexpr(NDimSpatial == 2)

1768 {

1769 if constexpr(!(is_NHWGC_GKYXC_NHWGK<InLayout, WeiLayout, OutLayout>() ||

1770 is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>()))

1771 {

1772 return false;

1773 }

1774 }

1775 else if constexpr(NDimSpatial == 3)

1776 {

1777 if constexpr(!(is_NDHWGC_GKZYXC_NDHWGK<InLayout, WeiLayout, OutLayout>() ||

1778 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>()))

1779 {

1780 return false;

1781 }

1782 }

1783 else

1784 {

1785 return false;

1786 }

1787

1788 if constexpr(ConvBackwardWeightSpecialization ==

1789 ConvolutionBackwardWeightSpecialization::Filter1x1Stride1Pad0)

1790 {

1791 // check if it's 1x1, stride=1 pad = 0 conv

1792 for(int i = 0; i < NDimSpatial; i++)

1793 {

1794 if(!(arg.filter_spatial_lengths_[i] == 1 && arg.conv_filter_strides_[i] == 1 &&

1795 arg.input_left_pads_[i] == 0 && arg.input_right_pads_[i] == 0))

1796 {

1797 return false;

1798 }

1799 }

1800 }

1801

1802 if constexpr(NumGroupsToMerge > 1)

1803 {

1804 // support only if whole M and N can be proccessed on one block

1805 if(!(GemmM <= MPerBlock && GemmN <= NPerBlock))

1806 {

1807 return false;

1808 }

1809 if(!(arg.Conv_C_ == 1 && arg.Conv_K_ == 1))

1810 {

1811 return false;

1812 }

1813 if(arg.Conv_G_ % NumGroupsToMerge != 0)

1814 {

1815 return false;

1816 }

1817 }

1818

1819 const bool is_w_pad_zero = arg.input_left_pads_[NDimSpatial - 1] == 0 &&

1820 arg.input_right_pads_[NDimSpatial - 1] == 0;

1821 const auto X = arg.filter_spatial_lengths_[NDimSpatial - 1];

1822 const bool XC_access_allowed = arg.Conv_G_ == 1 &&

1823 (arg.Conv_C_ * X) % BBlockTransferSrcScalarPerVector == 0 &&

1824 is_w_pad_zero;

1825

1826 if(!((arg.Conv_C_ % BBlockTransferSrcScalarPerVector == 0 || XC_access_allowed) &&

1827 arg.Conv_K_ % ABlockTransferSrcScalarPerVector == 0))

1828 {

1829 if(!(arg.Conv_K_ == 1 && arg.compute_ptr_offset_of_batch_.BatchStrideA_ == 1 &&

1830 NumGroupsToMerge > 1))

1831 {

1832 return false;

1833 }

1834 if(!(arg.Conv_C_ == 1 && arg.compute_ptr_offset_of_batch_.BatchStrideB_ == 1 &&

1835 NumGroupsToMerge > 1))

1836 {

1837 return false;

1838 }

1839 }

1840

1841 // vector load A/B matrix from global memory

1842 if(!(ABlockTransferSrcVectorDim == 1 && BBlockTransferSrcVectorDim == 1))

1843 {

1844 return false;

1845 }

1846

1847 // vector store C matrix into global memory

1848 if(!(arg.Conv_C_ % CBlockTransferScalarPerVector_NWaveNPerXdl == 0))

1849 {

1850 return false;

1851 }

1852

1853 if constexpr(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

1854 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>())

1855 {

1856 if((arg.Conv_G_ * arg.Conv_C_) % TransposeTransferDstScalarPerVector != 0)

1857 {

1858 return false;

1859 }

1860

1861 if((arg.Conv_G_ * arg.Conv_K_) % TransposeTransferDstScalarPerVector != 0)

1862 {

1863 return false;

1864 }

1865

1866 const index_t input_spatial_acum = ck::accumulate_n<index_t>(

1867 arg.input_spatial_lengths_.begin(), NDimSpatial, 1, std::multiplies<>());

1868 const index_t output_spatial_acum = ck::accumulate_n<index_t>(

1869 arg.output_spatial_lengths_.begin(), NDimSpatial, 1, std::multiplies<>());

1870

1871 if(input_spatial_acum % TransposeTransferSrcScalarPerVector != 0)

1872 {

1873 return false;

1874 }

1875

1876 if(output_spatial_acum % TransposeTransferSrcScalarPerVector != 0)

1877 {

1878 return false;

1879 }

1880

1881 constexpr long_index_t TwoGB = (long_index_t{1} << 31);

1882 if(!(arg.a_out_transpose_desc_.GetElementSpaceSize() * sizeof(ADataType) <= TwoGB &&

1883 arg.b_out_transpose_desc_.GetElementSpaceSize() * sizeof(BDataType) <= TwoGB))

1884 {

1885 return false;

1886 }

1887 }

1888

1889 constexpr long_index_t TwoGB = (long_index_t{1} << 31);

1890 if(!(arg.a_grid_desc_k0_m_k1_.GetElementSpaceSize() * sizeof(ADataType) <= TwoGB &&

1891 arg.b_grid_desc_k0_n_k1_.GetElementSpaceSize() * sizeof(BDataType) <= TwoGB &&

1892 arg.ce_grid_desc_m_n_.GetElementSpaceSize() * sizeof(EDataType) <= TwoGB))

1893 {

1894 return false;

1895 }

1896

1897 return true;

1898 }

1899

1900 bool IsSupportedArgument(const BaseArgument* p_arg) override

1901 {

1902 return IsSupportedArgument(*dynamic_cast<const Argument*>(p_arg));

1903 }

1904

1905 static auto

1906 MakeArgument(const InDataType* p_in_grid,

1907 WeiDataType* p_wei_grid,

1908 const OutDataType* p_out_grid,

1909 const std::array<index_t, NDimSpatial + 3>& b_g_n_c_wis_lengths, // input

1910 const std::array<index_t, NDimSpatial + 3>& b_g_n_c_wis_strides,

1911 const std::array<index_t, NDimSpatial + 3>& e_g_k_c_xs_lengths, // weight

1912 const std::array<index_t, NDimSpatial + 3>& e_g_k_c_xs_strides,

1913 const std::array<index_t, NDimSpatial + 3>& a_g_n_k_wos_lengths, // output

1914 const std::array<index_t, NDimSpatial + 3>& a_g_n_k_wos_strides,

1915 const std::array<ck::index_t, NDimSpatial>& conv_filter_strides,

1916 const std::array<ck::index_t, NDimSpatial>& conv_filter_dilations,

1917 const std::array<ck::index_t, NDimSpatial>& input_left_pads,

1918 const std::array<ck::index_t, NDimSpatial>& input_right_pads,

1919 InElementwiseOperation in_element_op,

1920 WeiElementwiseOperation wei_element_op,

1921 OutElementwiseOperation out_element_op,

1922 const ck::index_t split_k)

1923 {

1924 return Argument{p_in_grid,

1925 p_wei_grid,

1926 p_out_grid,

1927 b_g_n_c_wis_lengths, // input

1928 b_g_n_c_wis_strides,

1929 e_g_k_c_xs_lengths, // weight

1930 e_g_k_c_xs_strides,

1931 a_g_n_k_wos_lengths, // output

1932 a_g_n_k_wos_strides,

1933 conv_filter_strides,

1934 conv_filter_dilations,

1935 input_left_pads,

1936 input_right_pads,

1937 1,

1938 1,

1939 in_element_op,

1940 wei_element_op,

1941 out_element_op,

1942 split_k};

1943 }

1944

1945 static auto MakeInvoker() { return Invoker{}; }

1946

1947 std::unique_ptr<BaseArgument>

1948 MakeArgumentPointer(const void* p_in_grid,

1949 void* p_wei_grid,

1950 const void* p_out_grid,

1951 const std::array<index_t, NDimSpatial + 3>& b_g_n_c_wis_lengths, // input

1952 const std::array<index_t, NDimSpatial + 3>& b_g_n_c_wis_strides,

1953 const std::array<index_t, NDimSpatial + 3>& e_g_k_c_xs_lengths, // weight

1954 const std::array<index_t, NDimSpatial + 3>& e_g_k_c_xs_strides,

1955 const std::array<index_t, NDimSpatial + 3>& a_g_n_k_wos_lengths, // output

1956 const std::array<index_t, NDimSpatial + 3>& a_g_n_k_wos_strides,

1957 const std::array<ck::index_t, NDimSpatial>& conv_filter_strides,

1958 const std::array<ck::index_t, NDimSpatial>& conv_filter_dilations,

1959 const std::array<ck::index_t, NDimSpatial>& input_left_pads,

1960 const std::array<ck::index_t, NDimSpatial>& input_right_pads,

1961 InElementwiseOperation in_element_op,

1962 WeiElementwiseOperation wei_element_op,

1963 OutElementwiseOperation out_element_op,

1964 const ck::index_t split_k) override

1965 {

1966 return std::make_unique<Argument>(static_cast<const InDataType*>(p_in_grid),

1967 static_cast<WeiDataType*>(p_wei_grid),

1968 static_cast<const OutDataType*>(p_out_grid),

1969 b_g_n_c_wis_lengths, // input

1970 b_g_n_c_wis_strides,

1971 e_g_k_c_xs_lengths, // weight

1972 e_g_k_c_xs_strides,

1973 a_g_n_k_wos_lengths, // output

1974 a_g_n_k_wos_strides,

1975 conv_filter_strides,

1976 conv_filter_dilations,

1977 input_left_pads,

1978 input_right_pads,

1979 1,

1980 1,

1981 in_element_op,

1982 wei_element_op,

1983 out_element_op,

1984 split_k);

1985 }

1986

1987 std::unique_ptr<BaseInvoker> MakeInvokerPointer() override

1988 {

1989 return std::make_unique<Invoker>(Invoker{});

1990 }

1991

1992 std::string GetTypeString() const override

1993 {

1994 auto str = std::stringstream();

1995

1996 std::map<BlockGemmPipelineScheduler, std::string> BlkGemmPipelineSchedulerToString{

1997 {BlockGemmPipelineScheduler::Intrawave, "Intrawave"},

1998 {BlockGemmPipelineScheduler::Interwave, "Interwave"}};

1999

2000 std::map<BlockGemmPipelineVersion, std::string> BlkGemmPipelineVersionToString{

2001 {BlockGemmPipelineVersion::v1, "v1"},

2002 {BlockGemmPipelineVersion::v2, "v2"},

2003 {BlockGemmPipelineVersion::v3, "v3"},

2004 {BlockGemmPipelineVersion::v4, "v4"},

2005 {BlockGemmPipelineVersion::v5, "v5"}};

2006

2007 // clang-format off

2008 str << "DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle"

2009 << "<"

2010 << BlockSize << ", "

2011 << MPerBlock << ", "

2012 << NPerBlock << ", "

2013 << KPerBlock << ", "

2014 << getConvBackwardWeightSpecializationString(ConvBackwardWeightSpecialization) << ", "

2015 << K1 << ", "

2016 << MXdlPerWave << ", "

2017 << NXdlPerWave << ", "

2018 << ABlockTransferSrcScalarPerVector << ", "

2019 << ABlockTransferDstScalarPerVector_K1 << ", "

2020 << BBlockTransferSrcScalarPerVector << ", "

2021 << BBlockTransferDstScalarPerVector_K1 << ", "

2022 << CShuffleMXdlPerWavePerShuffle << ", "

2023 << CShuffleNXdlPerWavePerShuffle << ", "

2024 << CBlockTransferScalarPerVector_NWaveNPerXdl << ", "

2025 << "BlkGemmPipelineScheduler: "

2026 << BlkGemmPipelineSchedulerToString[BlkGemmPipeSched] << ", "

2027 << "BlkGemmPipelineVersion: "

2028 << BlkGemmPipelineVersionToString[BlkGemmPipelineVer] << ", "

2029 << NumGroupsToMerge;

2030

2031 if constexpr(is_NGCHW_NGKHW<InLayout, WeiLayout, OutLayout>() ||

2032 is_NGCDHW_NGKDHW<InLayout, WeiLayout, OutLayout>()) {

2033 str << ", TransposeTransferSrcScalarPerVector: "

2034 << TransposeTransferSrcScalarPerVector <<", "

2035 << "TransposeTransferDstScalarPerVector: " << TransposeTransferDstScalarPerVector;

2036 }

2037

2038

2039 str << ">";

2040 // clang-format on

2041

2042 return str.str();

2043 }

2044

2045 size_t GetWorkSpaceSize(const BaseArgument* p_arg) const override

2046 {

2047 auto arg = dynamic_cast<const Argument*>(p_arg);

2048 if(arg)

2049 {

2050 return arg->GetWorkspaceSizeBytes();

2051 }

2052 else

2053 throw std::runtime_error(

2054 "The argument pointer is not an object of "

2055 "DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument structure!");

2056 }

2057

2058 void SetWorkSpacePointer(BaseArgument* p_arg,

2059 void* p_workspace,

2060 const StreamConfig& = StreamConfig{}) const override

2061 {

2062 auto p_arg_ = dynamic_cast<Argument*>(p_arg);

2063 if(p_arg_)

2064 {

2065 p_arg_->p_workspace_ = p_workspace;

2066 }

2067 else

2068 throw std::runtime_error(

2069 "The argument pointer is not an object of "

2070 "DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument structure!");

2071 }

2072};

2073

2074} // namespace device

2075} // namespace tensor_operation

2076} // namespace ck

block_to_ctile_map.hpp

CK_MAX_THREAD_PER_BLOCK

#define CK_MAX_THREAD_PER_BLOCK

Definition ck.hpp:30

common_header.hpp

convolution_backward_weight_specialization.hpp

GET_NXDL_PER_WAVE_IMPL

#define GET_NXDL_PER_WAVE_IMPL

Definition device_base.hpp:81

INVOKER_RUN_IMPL

#define INVOKER_RUN_IMPL

Definition device_base.hpp:94

device_grouped_conv_bwd_weight.hpp

device_grouped_conv_utils.hpp

element_wise_operation.hpp

flush_cache.hpp

gemm_specialization.hpp

gridwise_elementwise_2d.hpp

gridwise_gemm_xdl_cshuffle_conv_v3.hpp

device_prop.hpp

hip_check_error

void hip_check_error(hipError_t x)

Definition host_utility/hip_check_error.hpp:10

kernel_launch.hpp

launch_and_time_kernel

float launch_and_time_kernel(const StreamConfig &stream_config, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)

Definition host_utility/kernel_launch.hpp:14

launch_and_time_kernel_with_preprocess

float launch_and_time_kernel_with_preprocess(const StreamConfig &stream_config, PreProcessFunc preprocess, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, Args... args)

Definition host_utility/kernel_launch.hpp:91

ck::math::max

__host__ __device__ constexpr T max(T x)

Definition utility/math.hpp:84

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

Definition convolution_backward_data_specialization.hpp:8

ck::tensor_operation::device::kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds

__global__ void kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3_2lds(typename GridwiseGemm::Argument karg, const AGridDesc_AK0_M_K1 a_grid_desc_ak0_m_ak1, const BGridDesc_BK0_N_K1 b_grid_desc_bk0_n_bk1, const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock c_grid_desc_mblock_mperblock_nblock_nperblock, const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch, const index_t num_k_per_block)

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:109

ck::tensor_operation::device::kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3

__global__ void kernel_grouped_conv_bwd_weight_xdl_cshuffle_v3(typename GridwiseGemm::Argument karg, const AGridDesc_AK0_M_K1 a_grid_desc_ak0_m_ak1, const BGridDesc_BK0_N_K1 b_grid_desc_bk0_n_bk1, const CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock c_grid_desc_mblock_mperblock_nblock_nperblock, const ComputePtrOffsetOfBatch compute_ptr_offset_of_batch, const index_t num_k_per_block)

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:51

ck::tensor_operation::device::get_bwd_weight_gemm_sizes

auto get_bwd_weight_gemm_sizes(const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_lengths)

Definition split_k_utils.hpp:55

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::device::is_NHWGC_GKYXC_NHWGK

constexpr bool is_NHWGC_GKYXC_NHWGK()

Definition device_grouped_conv_utils.hpp:40

ck::tensor_operation::device::get_best_occupancy_k_batch_value

ck::index_t get_best_occupancy_k_batch_value(int max_occupancy, ck::index_t grid_size)

Definition split_k_utils.hpp:30

ck::tensor_operation::device::GemmSpecialization

GemmSpecialization

Definition gemm_specialization.hpp:11

ck::tensor_operation::device::GemmSpecialization::Default

@ Default

Definition gemm_specialization.hpp:13

ck::tensor_operation::device::is_NDHWGC_GKZYXC_NDHWGK

constexpr bool is_NDHWGC_GKZYXC_NDHWGK()

Definition device_grouped_conv_utils.hpp:80

ck::tensor_operation::device::is_NGCDHW_NGKDHW

constexpr bool is_NGCDHW_NGKDHW()

Definition device_grouped_conv_utils.hpp:112

ck::tensor_operation::device::is_NGCHW_GKCYX_NGKHW

constexpr bool is_NGCHW_GKCYX_NGKHW()

Definition device_grouped_conv_utils.hpp:64

ck::tensor_operation::device::getConvBackwardWeightSpecializationString

std::string getConvBackwardWeightSpecializationString(const ConvolutionBackwardWeightSpecialization &s)

Definition convolution_backward_weight_specialization.hpp:21

ck::tensor_operation::device::calculate_mn_grid_size

ck::index_t calculate_mn_grid_size(ck::index_t gemmM, ck::index_t gemmN)

Definition split_k_utils.hpp:84

ck::tensor_operation::device::is_NGCDHW_GKCZYX_NGKDHW

constexpr bool is_NGCDHW_GKCZYX_NGKDHW()

Definition device_grouped_conv_utils.hpp:104

ck::tensor_operation::device::is_NGCHW_NGKHW

constexpr bool is_NGCHW_NGKHW()

Definition device_grouped_conv_utils.hpp:72

ck::tensor_operation

Definition convolution_backward_data_specialization.hpp:7

ck::utility::flush_icache

void flush_icache()

Definition flush_cache.hpp:383

ck::utility::launch_and_time_kernel_with_preprocess

float launch_and_time_kernel_with_preprocess(const StreamConfig &stream_config, PreProcessFunc preprocess, F kernel, dim3 grid_dim, dim3 block_dim, std::size_t lds_byte, GemmArgs &gemm_args, Args... args)

Definition flush_cache.hpp:398

ck

Definition ck.hpp:268

ck::kernel_batched_elementwise

__global__ void kernel_batched_elementwise(const InGridDescTuple in_grid_desc_tuple, const OutGridDescTuple out_grid_desc_tuple, const InDataTypePointerTuple p_in_global_tuple, const OutDataTypePointerTuple p_out_global_tuple, const Block2TileMap block_2_tile_map, const ElementwiseOperation elementwise_op, const index_t batch_count, const std::array< index_t, NumInputs > input_batch_strides, const std::array< index_t, NumOutputs > output_batch_strides)

Definition gridwise_elementwise_2d.hpp:221

ck::index_t

int32_t index_t

Definition ck.hpp:299

ck::InMemoryDataOperationEnum

InMemoryDataOperationEnum

Definition ck.hpp:277

ck::InMemoryDataOperationEnum::Set

@ Set

Definition ck.hpp:278

ck::InMemoryDataOperationEnum::AtomicAdd

@ AtomicAdd

Definition ck.hpp:279

ck::remove_cvref_t

remove_cv_t< remove_reference_t< T > > remove_cvref_t

Definition type.hpp:297

ck::BlockGemmPipelineVersion

BlockGemmPipelineVersion

Definition blkgemmpipe_scheduler.hpp:12

ck::BlockGemmPipelineVersion::v2

@ v2

Definition blkgemmpipe_scheduler.hpp:15

ck::BlockGemmPipelineVersion::v3

@ v3

Definition blkgemmpipe_scheduler.hpp:16

ck::BlockGemmPipelineVersion::v5

@ v5

Definition blkgemmpipe_scheduler.hpp:18

ck::BlockGemmPipelineVersion::v4

@ v4

Definition blkgemmpipe_scheduler.hpp:17

ck::BlockGemmPipelineVersion::v1

@ v1

Definition blkgemmpipe_scheduler.hpp:14

ck::Number

integral_constant< index_t, N > Number

Definition number.hpp:12

ck::TailNumber

TailNumber

Definition blkgemmpipe_scheduler.hpp:31

ck::TailNumber::One

@ One

Definition blkgemmpipe_scheduler.hpp:37

ck::TailNumber::Seven

@ Seven

Definition blkgemmpipe_scheduler.hpp:43

ck::TailNumber::Even

@ Even

Definition blkgemmpipe_scheduler.hpp:34

ck::TailNumber::Odd

@ Odd

Definition blkgemmpipe_scheduler.hpp:33

ck::TailNumber::Four

@ Four

Definition blkgemmpipe_scheduler.hpp:40

ck::TailNumber::Two

@ Two

Definition blkgemmpipe_scheduler.hpp:38

ck::TailNumber::Full

@ Full

Definition blkgemmpipe_scheduler.hpp:49

ck::TailNumber::Three

@ Three

Definition blkgemmpipe_scheduler.hpp:39

ck::TailNumber::Five

@ Five

Definition blkgemmpipe_scheduler.hpp:41

ck::TailNumber::Six

@ Six

Definition blkgemmpipe_scheduler.hpp:42

ck::ignore

constexpr detail::ignore_t ignore

Definition utility/ignore.hpp:20

ck::amd_wave_read_first_lane

__device__ uint32_t amd_wave_read_first_lane(uint32_t value)

Definition amd_wave_read_first_lane.hpp:100

ck::type_convert

__host__ __device__ constexpr Y type_convert(X x)

Definition utility/type_convert.hpp:98

ck::is_xdl_wmma_supported

bool is_xdl_wmma_supported()

Definition host_utility/device_prop.hpp:76

ck::get_warp_size

__device__ constexpr index_t get_warp_size()

Definition get_id.hpp:10

ck::EnvIsEnabled

bool EnvIsEnabled(EnvVar)

Definition utility/env.hpp:140

ck::is_tf32_supported

bool is_tf32_supported()

Definition host_utility/device_prop.hpp:132

ck::is_same_v

constexpr bool is_same_v

Definition type.hpp:283

ck::BlockGemmPipelineScheduler

BlockGemmPipelineScheduler

Definition blkgemmpipe_scheduler.hpp:25

ck::BlockGemmPipelineScheduler::Intrawave

@ Intrawave

Definition blkgemmpipe_scheduler.hpp:26

ck::BlockGemmPipelineScheduler::Interwave

@ Interwave

Definition blkgemmpipe_scheduler.hpp:27

ck::remove_reference_t

typename remove_reference< T >::type remove_reference_t

Definition type.hpp:292

ck::accumulate_n

auto accumulate_n(ForwardIterator first, Size count, T init, BinaryOperation op) -> decltype(std::accumulate(first, std::next(first, count), init, op))

Definition library/utility/numeric.hpp:11

ck::make_tuple

__host__ __device__ constexpr auto make_tuple(Xs &&... xs)

Definition utility/tuple.hpp:211

ck::kernel_elementwise_dual

__global__ void kernel_elementwise_dual(const InAGridDescTuple in_grid_desc_tuple_a, const InBGridDescTuple in_grid_desc_tuple_b, const OutAGridDescTuple out_grid_desc_tuple_a, const OutBGridDescTuple out_grid_desc_tuple_b, const InADataTypePointerTuple p_in_global_tuple_a, const InBDataTypePointerTuple p_in_global_tuple_b, const OutADataTypePointerTuple p_out_global_tuple_a, const OutBDataTypePointerTuple p_out_global_tuple_b, const Block2TileMapA block_2_tile_map_a, const Block2TileMapB block_2_tile_map_b, const ElementwiseOperation elementwise_op, const index_t a_grid_size)

Definition gridwise_elementwise_2d.hpp:61

ck::long_index_t

int64_t long_index_t

Definition ck.hpp:300

ck::kernel_elementwise

__global__ void kernel_elementwise(const InGridDescTuple in_grid_desc_tuple, const OutGridDescTuple out_grid_desc_tuple, const InDataTypePointerTuple p_in_global_tuple, const OutDataTypePointerTuple p_out_global_tuple, const Block2TileMap block_2_tile_map, const ElementwiseOperation elementwise_op)

Definition gridwise_elementwise_2d.hpp:29

split_k_arg.hpp

split_k_utils.hpp

StreamConfig

Definition ck/stream_config.hpp:10

ck::BlockToCTileMap_M00_N0_M01Adapt

Definition block_to_ctile_map.hpp:261

ck::GridwiseElementwise

Definition gridwise_elementwise_2d.hpp:278

ck::GridwiseGemm_xdl_cshuffle_conv_v3

Definition gridwise_gemm_xdl_cshuffle_conv_v3.hpp:66

ck::GridwiseGemm_xdl_cshuffle_conv_v3< tensor_layout::gemm::RowMajor, tensor_layout::gemm::ColumnMajor, tensor_layout::gemm::RowMajor, ADataType, BDataType, AccDataType, AccDataType, AccDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, GemmSpec, BlockSize, MPerBlock, NPerBlock, KPerBlock, K1, K1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave_, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, false, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, false, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, BlkGemmPipeSched, BlkGemmPipelineVer, ComputeTypeA, ComputeTypeB >::MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock

__host__ static __device__ constexpr auto MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(const CGridDesc &c_grid_desc_m_n, index_t MBlock, index_t NBlock)

Definition gridwise_gemm_xdl_cshuffle_conv_v3.hpp:644

ck::GridwiseGemm_xdl_cshuffle_conv_v3< tensor_layout::gemm::RowMajor, tensor_layout::gemm::ColumnMajor, tensor_layout::gemm::RowMajor, ADataType, BDataType, AccDataType, AccDataType, AccDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, GemmSpec, BlockSize, MPerBlock, NPerBlock, KPerBlock, K1, K1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave_, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, false, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, false, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, BlkGemmPipeSched, BlkGemmPipelineVer, ComputeTypeA, ComputeTypeB >::CalculateNBlock

static __host__ auto CalculateNBlock(index_t N)

Definition gridwise_gemm_xdl_cshuffle_conv_v3.hpp:153

ck::GridwiseGemm_xdl_cshuffle_conv_v3< tensor_layout::gemm::RowMajor, tensor_layout::gemm::ColumnMajor, tensor_layout::gemm::RowMajor, ADataType, BDataType, AccDataType, AccDataType, AccDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, GemmSpec, BlockSize, MPerBlock, NPerBlock, KPerBlock, K1, K1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave_, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, false, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, false, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, BlkGemmPipeSched, BlkGemmPipelineVer, ComputeTypeA, ComputeTypeB >::CalculateMBlock

static __host__ auto CalculateMBlock(index_t M)

Definition gridwise_gemm_xdl_cshuffle_conv_v3.hpp:148

ck::Sequence

Definition utility/sequence.hpp:43

ck::Tuple

Definition utility/tuple.hpp:117

ck::tensor_layout::gemm::ColumnMajor

Definition tensor_operation/gpu/device/tensor_layout.hpp:31

ck::tensor_layout::gemm::RowMajor

Definition tensor_operation/gpu/device/tensor_layout.hpp:26

ck::tensor_operation::TransformConvBwdWeightToGemm

Definition tensor_operation/operator_transform/transform_conv_bwd_weight_to_gemm.hpp:24

ck::tensor_operation::TransformConvBwdWeightToGemmV2

Transform conv bwd weight to gemm v2.

Definition transform_conv_bwd_weight_to_gemm_v2.hpp:33

ck::tensor_operation::TransformConvNGCHWToNHWGC

Definition transform_conv_ngchw_to_nhwgc.hpp:31

ck::tensor_operation::device::ArgumentSplitK

Definition split_k_arg.hpp:11

ck::tensor_operation::device::ArgumentSplitK::k_batch_

index_t k_batch_

Definition split_k_arg.hpp:12

ck::tensor_operation::device::BaseArgument

Definition device_base.hpp:197

ck::tensor_operation::device::BaseArgument::p_workspace_

void * p_workspace_

Definition device_base.hpp:204

ck::tensor_operation::device::BaseArgument::BaseArgument

BaseArgument()=default

ck::tensor_operation::device::BaseInvoker::BaseInvoker

BaseInvoker()=default

ck::tensor_operation::device::DeviceGroupedConvBwdWeight

Definition device_grouped_conv_bwd_weight.hpp:29

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ActiveWorkgroupsPerCU

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:522

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ActiveWorkgroupsPerCU::GetMaxOccupancy

int GetMaxOccupancy()

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:524

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ActiveWorkgroupsPerCU::max_occupancy_

int max_occupancy_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:586

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ActiveWorkgroupsPerCU::ActiveWorkgroupsPerCU

ActiveWorkgroupsPerCU()

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:568

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:590

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::c_space_size_bytes

long_index_t c_space_size_bytes

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:866

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::b_in_transpose_desc_

NGCHWTransposeDescType b_in_transpose_desc_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:840

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::input_spatial_lengths_

std::array< ck::index_t, NDimSpatial > input_spatial_lengths_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:860

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::Conv_C_

const index_t Conv_C_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:859

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::b_element_op_

InElementwiseOperation b_element_op_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:852

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::a_element_op_

OutElementwiseOperation a_element_op_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:851

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::cde_element_op_

WeiElementwiseOperation cde_element_op_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:853

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::input_left_pads_

const std::array< ck::index_t, NDimSpatial > & input_left_pads_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:864

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::elementwise_block_2_ctile_map_

Block2TileMapElementwise elementwise_block_2_ctile_map_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:836

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::compute_ptr_offset_of_batch_

ComputePtrOffsetOfStridedBatch< I1, I1, I0 > compute_ptr_offset_of_batch_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:846

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::b_out_transpose_desc_

NHWGCTransposeDescType b_out_transpose_desc_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:841

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::GetWorkspaceETensorSizeBytes

std::size_t GetWorkspaceETensorSizeBytes() const

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:797

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::ce_elementwise_grid_desc_m_n_

CElementwiseGridDesc_M_N ce_elementwise_grid_desc_m_n_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:833

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::p_b_grid_

const BDataType * p_b_grid_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:827

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::elementwise_block_2_ctile_map_transpose_a_

Block2TileMapElementwise elementwise_block_2_ctile_map_transpose_a_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:837

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::Conv_N_

const index_t Conv_N_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:857

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::filter_spatial_lengths_

std::array< ck::index_t, NDimSpatial > filter_spatial_lengths_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:861

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::Argument

Argument(const InDataType *p_in_grid, WeiDataType *p_wei_grid, const OutDataType *p_out_grid, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &input_left_pads, const std::array< ck::index_t, NDimSpatial > &input_right_pads, const ck::index_t M01, const ck::index_t N01, InElementwiseOperation in_element_op, WeiElementwiseOperation wei_element_op, OutElementwiseOperation out_element_op, ck::index_t split_k)

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:591

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::GetWorkspaceSizeBytes

std::size_t GetWorkspaceSizeBytes() const

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:806

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::GetWorkspaceATensorSizeBytes

std::size_t GetWorkspaceATensorSizeBytes() const

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:784

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::Conv_K_

const index_t Conv_K_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:858

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::p_a_grid_

const ADataType * p_a_grid_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:826

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::conv_filter_strides_

const std::array< ck::index_t, NDimSpatial > & conv_filter_strides_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:863

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::c_grid_desc_mblock_mperblock_nblock_nperblock_

CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock c_grid_desc_mblock_mperblock_nblock_nperblock_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:834

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::e_in_transpose_desc_

GKYXCTransposeDescType e_in_transpose_desc_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:842

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::output_spatial_lengths_

std::array< ck::index_t, NDimSpatial > output_spatial_lengths_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:862

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::b_grid_desc_k0_n_k1_

BGridDesc_K0_N_K1 b_grid_desc_k0_n_k1_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:831

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::input_right_pads_

const std::array< ck::index_t, NDimSpatial > & input_right_pads_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:865

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::a_out_transpose_desc_

NHWGCTransposeDescType a_out_transpose_desc_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:841

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::p_e_grid_

EDataType * p_e_grid_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:828

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::M01_

index_t M01_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:848

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::GetWorkspaceBTensorSizeBytes

std::size_t GetWorkspaceBTensorSizeBytes() const

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:792

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::elementwise_block_2_ctile_map_transpose_b_

Block2TileMapElementwise elementwise_block_2_ctile_map_transpose_b_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:838

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::a_grid_desc_k0_m_k1_

AGridDesc_K0_M_K1 a_grid_desc_k0_m_k1_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:830

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::e_out_transpose_desc_

GKCYXTransposeDescType e_out_transpose_desc_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:843

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::a_in_transpose_desc_

NGCHWTransposeDescType a_in_transpose_desc_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:840

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::Conv_G_

const index_t Conv_G_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:856

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::ce_grid_desc_m_n_

CGridDesc_M_N ce_grid_desc_m_n_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:832

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Argument::N01_

index_t N01_

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:849

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Invoker

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:871

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Invoker::Run

INVOKER_RUN_IMPL float Run(const BaseArgument *p_arg, const StreamConfig &stream_config=StreamConfig{}) override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1670

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Invoker::ShowInfo

void ShowInfo(const Argument &arg)

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:874

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Invoker::RunGemmV3

float RunGemmV3(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:889

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Invoker::RunImp

float RunImp(const Argument &arg, const StreamConfig &stream_config=StreamConfig{})

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1543

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Invoker::Argument

DeviceOp::Argument Argument

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:872

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:217

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::BElementwiseOperation

InElementwiseOperation BElementwiseOperation

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:237

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::IsSupportedArgument

bool IsSupportedArgument(const BaseArgument *p_arg) override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1900

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::I4

static constexpr auto I4

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:247

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::SetWorkSpacePointer

void SetWorkSpacePointer(BaseArgument *p_arg, void *p_workspace, const StreamConfig &=StreamConfig{}) const override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:2058

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ABDataType

InDataType ABDataType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:241

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::I5

static constexpr auto I5

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:248

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GKCYXTransposeDescType

remove_cvref_t< decltype(conv_ngchw_to_nhwgc_transformer .template MakeGKCYXTransposeDesc< NDimSpatial >({}, {}))> GKCYXTransposeDescType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:390

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::I1

static constexpr auto I1

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:244

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::NXdlPerWave64

static GET_NXDL_PER_WAVE_IMPL constexpr auto NXdlPerWave64

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:224

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GetTypeString

std::string GetTypeString() const override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1992

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GridwiseGemm32

GridwiseGemmBase< NXdlPerWave32 > GridwiseGemm32

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:454

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ADataType

OutDataType ADataType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:227

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GemmSpec

static constexpr GemmSpecialization GemmSpec

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:282

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GKYXCTransposeDescType

remove_cvref_t< decltype(conv_ngchw_to_nhwgc_transformer .template MakeGKYXCTransposeDesc< NDimSpatial >({}, {}))> GKYXCTransposeDescType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:393

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::IsSupportedArgument

static bool IsSupportedArgument(const Argument &arg)

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1683

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::MakeInvoker

static auto MakeInvoker()

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1945

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GetWorkSpaceSize

size_t GetWorkSpaceSize(const BaseArgument *p_arg) const override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:2045

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::IsValidCompilationParameter

static constexpr bool IsValidCompilationParameter()

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1677

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::EDataType

WeiDataType EDataType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:229

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ABCGridDescs

decltype(GetABCGridDesc< NDimSpatial >()) ABCGridDescs

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:397

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::I3

static constexpr auto I3

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:246

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::MakeArgument

static auto MakeArgument(const InDataType *p_in_grid, WeiDataType *p_wei_grid, const OutDataType *p_out_grid, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &input_left_pads, const std::array< ck::index_t, NDimSpatial > &input_right_pads, InElementwiseOperation in_element_op, WeiElementwiseOperation wei_element_op, OutElementwiseOperation out_element_op, const ck::index_t split_k)

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1906

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GridwiseElementwiseCast

GridwiseElementwise< Tuple< CElementwiseGridDesc_M_N >, Tuple< CElementwiseGridDesc_M_N >, Tuple< const AccDataType * >, Tuple< EDataType * >, Block2TileMapElementwise, CDEElementwiseOperation, BlockSize, MPerBlock, NPerBlock, MPerBlock/ClusterLengthMPerBlock, NPerBlock/ClusterLengthNPerBlock, Sequence< 0, 1 >, Sequence< CBlockTransferScalarPerVector_NWaveNPerXdl >, Sequence< CBlockTransferScalarPerVector_NWaveNPerXdl >, I1, I1 > GridwiseElementwiseCast

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:458

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::CDEElementwiseOperation

WeiElementwiseOperation CDEElementwiseOperation

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:238

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::AElementwiseOperation

OutElementwiseOperation AElementwiseOperation

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:236

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::BDataType

InDataType BDataType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:228

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::Block2TileMapElementwise

BlockToCTileMap_M00_N0_M01Adapt< MPerBlock, NPerBlock > Block2TileMapElementwise

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:456

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::CGridDesc_M_N

remove_cvref_t< decltype(ABCGridDescs{}[I2])> CGridDesc_M_N

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:401

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::I0

static constexpr auto I0

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:243

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::K1Number

static constexpr auto K1Number

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:250

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GridwiseElementwiseWeightTransposeCast

GridwiseElementwise< Tuple< GKYXCTransposeDescType >, Tuple< GKCYXTransposeDescType >, Tuple< const AccDataType * >, Tuple< EDataType * >, Block2TileMapElementwise, CDEElementwiseOperation, BlockSize, MPerBlock, NPerBlock, MPerBlock/ClusterLengthMPerBlock, NPerBlock/ClusterLengthNPerBlock, Sequence< 0, 1 >, Sequence< CBlockTransferScalarPerVector_NWaveNPerXdl >, Sequence< 1 >, I1, I0 > GridwiseElementwiseWeightTransposeCast

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:480

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GridwiseGemm64

GridwiseGemmBase< math::max(NXdlPerWave64, 1)> GridwiseGemm64

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:453

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GridwiseGemmBase

GridwiseGemm_xdl_cshuffle_conv_v3< tensor_layout::gemm::RowMajor, tensor_layout::gemm::ColumnMajor, tensor_layout::gemm::RowMajor, ADataType, BDataType, AccDataType, AccDataType, AccDataType, AElementwiseOperation, BElementwiseOperation, CDEElementwiseOperation, GemmSpec, BlockSize, MPerBlock, NPerBlock, KPerBlock, K1, K1, MPerXDL, NPerXDL, MXdlPerWave, NXdlPerWave_, ABlockTransferThreadClusterLengths_K0_M_K1, ABlockTransferThreadClusterArrangeOrder, ABlockTransferSrcAccessOrder, ABlockTransferSrcVectorDim, ABlockTransferSrcScalarPerVector, ABlockTransferDstScalarPerVector_K1, false, ABlockLdsAddExtraM, BBlockTransferThreadClusterLengths_K0_N_K1, BBlockTransferThreadClusterArrangeOrder, BBlockTransferSrcAccessOrder, BBlockTransferSrcVectorDim, BBlockTransferSrcScalarPerVector, BBlockTransferDstScalarPerVector_K1, false, BBlockLdsAddExtraN, CShuffleMXdlPerWavePerShuffle, CShuffleNXdlPerWavePerShuffle, CBlockTransferClusterLengths_MBlock_MPerBlock_NBlock_NPerBlock, CBlockTransferScalarPerVector_NWaveNPerXdl, BlkGemmPipeSched, BlkGemmPipelineVer, ComputeTypeA, ComputeTypeB > GridwiseGemmBase

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:406

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::DeviceOp

DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle DeviceOp

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:222

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::I2

static constexpr auto I2

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:245

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::NHWGCTransposeDescType

remove_cvref_t< decltype(conv_ngchw_to_nhwgc_transformer .template MakeNHWGCTransposeDesc< NDimSpatial >({}, {}))> NHWGCTransposeDescType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:387

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::NGCHWTransposeDescType

remove_cvref_t< decltype(conv_ngchw_to_nhwgc_transformer .template MakeNGCHWTransposeDesc< NDimSpatial >({}, {}))> NGCHWTransposeDescType

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:384

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ClusterLengthMPerBlock

static constexpr index_t ClusterLengthMPerBlock

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:269

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::NXdlPerWave32

static constexpr auto NXdlPerWave32

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:225

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::BGridDesc_K0_N_K1

remove_cvref_t< decltype(ABCGridDescs{}[I1])> BGridDesc_K0_N_K1

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:400

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::ClusterLengthNPerBlock

static constexpr index_t ClusterLengthNPerBlock

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:271

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GetElementwiseCGridDesc

static auto GetElementwiseCGridDesc()

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:335

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::conv_to_gemm_transformer_v1

static constexpr auto conv_to_gemm_transformer_v1

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:261

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock

decltype(GridwiseGemm64::MakeCGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock( CGridDesc_M_N{}, 1, 1)) CGridDesc_MBlock_MPerBlock_NBlock_NPerBlock

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:517

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GetABCGridDesc

static auto GetABCGridDesc()

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:285

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::conv_ngchw_to_nhwgc_transformer

static constexpr auto conv_ngchw_to_nhwgc_transformer

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:274

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::AGridDesc_K0_M_K1

remove_cvref_t< decltype(ABCGridDescs{}[I0])> AGridDesc_K0_M_K1

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:399

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::MakeArgumentPointer

std::unique_ptr< BaseArgument > MakeArgumentPointer(const void *p_in_grid, void *p_wei_grid, const void *p_out_grid, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_lengths, const std::array< index_t, NDimSpatial+3 > &b_g_n_c_wis_strides, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_lengths, const std::array< index_t, NDimSpatial+3 > &e_g_k_c_xs_strides, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_lengths, const std::array< index_t, NDimSpatial+3 > &a_g_n_k_wos_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_strides, const std::array< ck::index_t, NDimSpatial > &conv_filter_dilations, const std::array< ck::index_t, NDimSpatial > &input_left_pads, const std::array< ck::index_t, NDimSpatial > &input_right_pads, InElementwiseOperation in_element_op, WeiElementwiseOperation wei_element_op, OutElementwiseOperation out_element_op, const ck::index_t split_k) override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1948

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::GridwiseElementwiseTranspose

GridwiseElementwise< Tuple< NGCHWTransposeDescType >, Tuple< NHWGCTransposeDescType >, Tuple< const ADataType * >, Tuple< ADataType * >, Block2TileMapElementwise, element_wise::PassThrough, BlockSize, MPerBlock, NPerBlock, MPerBlock/ClusterLengthMPerBlock, NPerBlock/ClusterLengthNPerBlock, Sequence< 1, 0 >, Sequence< TransposeTransferSrcScalarPerVector >, Sequence< TransposeTransferDstScalarPerVector >, I1, I0 > GridwiseElementwiseTranspose

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:498

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::MakeInvokerPointer

std::unique_ptr< BaseInvoker > MakeInvokerPointer() override

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:1987

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::conv_to_gemm_transformer_v2

static constexpr auto conv_to_gemm_transformer_v2

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:252

ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle::CElementwiseGridDesc_M_N

remove_cvref_t< decltype(GetElementwiseCGridDesc< NDimSpatial >())> CElementwiseGridDesc_M_N

Definition device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp:402

ck::tensor_operation::element_wise::PassThrough

Definition tensor_operation/gpu/element/unary_element_wise_operation.hpp:340

tensor_descriptor.hpp

tensor_descriptor_helper.hpp

tensor_layout.hpp

transform_conv_bwd_weight_to_gemm.hpp

transform_conv_bwd_weight_to_gemm_v2.hpp

transform_conv_ngchw_to_nhwgc.hpp

env.hpp

CK_ENV

#define CK_ENV(name)

Definition utility/env.hpp:129