/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck/tensor_operation/gpu/grid/gridwise_ab_transfer_thread_tiles.hpp Source File

• develop
• include
• ck
• tensor_operation
• gpu
• grid

gridwise_ab_transfer_thread_tiles.hpp

Go to the documentation of this file.

1 // Copyright (c) Advanced Micro Devices, Inc., or its affiliates.

2 // SPDX-License-Identifier: MIT

3 

4 #pragma once

5 

6 #include "ck/utility/amd_address_space.hpp"

7 #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v4r1.hpp"

8 #include "ck/tensor_operation/gpu/block/thread_group_tensor_slice_transfer_v7r2.hpp"

9 

10 namespace ck {

11 

12 template <typename ABLayout,

13  typename ABMajorLayout,

14  typename LDSTypeAB,

15  index_t BlockSize,

16  index_t MNPerBlock,

17  index_t KPerBlock,

18  index_t MNPerWmma,

19  index_t ABK1Value,

20  index_t KPack,

21  index_t KInner,

22  index_t KPerWmmaBlk,

23  bool UseBlockPaddingAB,

24  bool PermuteAB,

25  typename ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1,

26  typename ABBlockTransferThreadClusterArrangeOrder,

27  typename ABBlockTransferSrcAccessOrder,

28  index_t ABBlockTransferSrcVectorDim,

29  index_t ABBlockTransferSrcScalarPerVector,

30  index_t ABBlockTransferDstScalarPerVector_ABK1,

31  bool ABThreadTransferSrcResetCoordinateAfterRun>

32 struct ABTransferThreadTiles

33 {

34  static constexpr auto ABK0Number = Number<KPerBlock / ABK1Value>{};

35  static constexpr auto ABK1Number = Number<ABK1Value>{};

36 

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

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

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

40 

41  static constexpr index_t ABPackedSize = []() {

42  if constexpr(is_same_v<remove_cvref_t<LDSTypeAB>, pk_i4_t>)

43  return 2;

44  else

45  return 1;

46  }();

47 

48  using ThisThreadBlock = ThisThreadBlock<BlockSize>;

49 

50  template <bool PadMN, bool PadK, typename GridDescriptorBase>

51  __host__ __device__ static auto MakeGridDescriptor(const GridDescriptorBase& ab_grid_desc,

52  index_t MN,

53  index_t MNPad,

54  index_t K,

55  index_t KPad,

56  index_t StrideAB,

57  index_t ABK0)

58  {

59 

60  if constexpr(PadMN && PadK)

61  {

62  // pad both MN and K

63  const auto ab_grid_desc_n_k =

64  transform_tensor_descriptor(ab_grid_desc,

65  make_tuple(make_right_pad_transform(MN, MNPad - MN),

66  make_right_pad_transform(K, KPad - K)),

67  make_tuple(Sequence<0>{}, Sequence<1>{}),

68  make_tuple(Sequence<0>{}, Sequence<1>{}));

69 

70  const auto ab_grid_desc_bk0_n_bk1 = transform_tensor_descriptor(

71  ab_grid_desc_n_k,

72  make_tuple(make_unmerge_transform(make_tuple(ABK0, ABK1Value)),

73  make_pass_through_transform(MNPad)),

74  make_tuple(Sequence<1>{}, Sequence<0>{}),

75  make_tuple(Sequence<0, 2>{}, Sequence<1>{}));

76 

77  return ab_grid_desc_bk0_n_bk1;

78  }

79  else if constexpr(PadMN && !PadK)

80  {

81  // pad MN, but not K

82  const auto ab_grid_desc_bk0_n_bk1 = transform_tensor_descriptor(

83  ab_grid_desc,

84  make_tuple(make_unmerge_transform(make_tuple(ABK0, ABK1Value)),

85  make_right_pad_transform(MN, MNPad - MN)),

86  make_tuple(Sequence<1>{}, Sequence<0>{}),

87  make_tuple(Sequence<0, 2>{}, Sequence<1>{}));

88 

89  return ab_grid_desc_bk0_n_bk1;

90  }

91  else if constexpr(!PadMN && PadK)

92  {

93  // pad K, but not MN

94  const auto ab_grid_desc_n_k = transform_tensor_descriptor(

95  ab_grid_desc,

96  make_tuple(make_pass_through_transform(MN), make_right_pad_transform(K, KPad - K)),

97  make_tuple(Sequence<0>{}, Sequence<1>{}),

98  make_tuple(Sequence<0>{}, Sequence<1>{}));

99 

100  const auto ab_grid_desc_bk0_n_bk1 = transform_tensor_descriptor(

101  ab_grid_desc_n_k,

102  make_tuple(make_unmerge_transform(make_tuple(ABK0, ABK1Value)),

103  make_pass_through_transform(MN)),

104  make_tuple(Sequence<1>{}, Sequence<0>{}),

105  make_tuple(Sequence<0, 2>{}, Sequence<1>{}));

106 

107  return ab_grid_desc_bk0_n_bk1;

108  }

109  else

110  {

111  if constexpr(!PermuteAB)

112  {

113  // not pad MN or K

114  const auto ab_grid_desc_bk0_n_bk1 = transform_tensor_descriptor(

115  ab_grid_desc,

116  make_tuple(make_unmerge_transform(make_tuple(ABK0, ABK1Value)),

117  make_pass_through_transform(MN)),

118  make_tuple(Sequence<1>{}, Sequence<0>{}),

119  make_tuple(Sequence<0, 2>{}, Sequence<1>{}));

120 

121  return ab_grid_desc_bk0_n_bk1;

122  }

123  else

124  {

125  // Pre-shuffled Weight

126  // BGlobal[K / KPerBlock, MN, KPerBlock / K1, K1] -> BTile[K / K1, MN, K1]

127  constexpr index_t ABK01 = KPerBlock / ABK1Value;

128  const index_t ABK0_ = StrideAB / ABK1Value;

129  const index_t ABK00 = ABK0_ / ABK01;

130 

131  const auto ab_grid_desc_abk00_mn_abk01_abk1_permute =

132  make_naive_tensor_descriptor_packed(make_tuple(ABK00, MN, ABK01, ABK1Value));

133 

134  const auto ab_grid_desc_abk0_mn_abk1_permute = transform_tensor_descriptor(

135  ab_grid_desc_abk00_mn_abk01_abk1_permute,

136  make_tuple(make_merge_transform(make_tuple(ABK00, ABK01)),

137  make_pass_through_transform(make_tuple(MN)),

138  make_pass_through_transform(ABK1Value)),

139  make_tuple(Sequence<0, 2>{}, Sequence<1>{}, Sequence<3>{}),

140  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));

141 

142  return ab_grid_desc_abk0_mn_abk1_permute;

143  }

144  }

145  }

146 

147  __device__ static constexpr auto GetBlockDescriptor()

148  {

149  // A matrix in LDS memory, dst of blockwise copy

150  if constexpr(UseBlockPaddingAB)

151  {

152  // bank conflict when writting the data into LDS, but don't worry, we have whole entire

153  // loop to hide it in v4. it may give you some benefit from less valu in compute address

154  return make_naive_tensor_descriptor(

155  make_tuple(ABK0Number, Number<MNPerBlock>{}, ABK1Number),

156  make_tuple(Number<MNPerBlock + 1>{} * ABK1Number, ABK1Number, I1));

157  }

158  // xor tensor transformation request more unnecessary vgpr usage, would cause register spill

159  // in some cases.

160  else if constexpr(is_same<ABMajorLayout, ABLayout>::value)

161  {

162  constexpr index_t LdsSize = 32 * 4 / KPerBlock / sizeof(LDSTypeAB) / ABPackedSize;

163  constexpr auto MNLdsLayer = LdsSize < 1 ? 1 : LdsSize;

164  constexpr auto ab_lds_block_desc = make_naive_tensor_descriptor(

165  make_tuple(ABK0Number * Number<MNLdsLayer>{},

166  Number<MNPerBlock / MNLdsLayer>{},

167  ABK1Number),

168  make_tuple(ABK1Number, Number<KPerBlock * MNLdsLayer>{}, I1));

169 

170  constexpr auto ab_lds_block_desc_permuted = transform_tensor_descriptor(

171  ab_lds_block_desc,

172  make_tuple(

173  make_xor_with_modulo_transform(make_tuple(Number<MNPerBlock / MNLdsLayer>{},

174  Number<ABK0Number * MNLdsLayer>{})),

175  make_pass_through_transform(ABK1Number)),

176  make_tuple(Sequence<1, 0>{}, Sequence<2>{}),

177  make_tuple(Sequence<1, 0>{}, Sequence<2>{}));

178 

179  constexpr auto ab_lds_block_desc_abk0_mnldslayer_mn_abk1 = transform_tensor_descriptor(

180  ab_lds_block_desc_permuted,

181  make_tuple(make_unmerge_transform(make_tuple(ABK0Number, Number<MNLdsLayer>{})),

182  make_pass_through_transform(Number<MNPerBlock / MNLdsLayer>{}),

183  make_pass_through_transform(ABK1Number)),

184  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}),

185  make_tuple(Sequence<0, 2>{}, Sequence<1>{}, Sequence<3>{}));

186 

187  constexpr auto ab_lds_block_desc_abk0_mn_abk1 = transform_tensor_descriptor(

188  ab_lds_block_desc_abk0_mnldslayer_mn_abk1,

189  make_tuple(make_pass_through_transform(ABK0Number),

190  make_merge_transform_v3_division_mod(

191  make_tuple(Number<MNPerBlock / MNLdsLayer>{}, Number<MNLdsLayer>{})),

192  make_pass_through_transform(ABK1Number)),

193  make_tuple(Sequence<0>{}, Sequence<1, 2>{}, Sequence<3>{}),

194  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));

195 

196  return ab_lds_block_desc_abk0_mn_abk1;

197  }

198  else

199  {

200  // kfold and mpair dimension is not always required.

201  // more dimension in merge_transform increase the difficulty of generating immarg offset

202  // for compiler.

203  constexpr auto MN0 = ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1{}.At(I1);

204  constexpr auto MN1 = MNPerBlock / MN0;

205 

206  constexpr auto KThreadWrite = ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1{}.At(I0);

207  constexpr auto K0PerThreadWrite = ABK0Number / KThreadWrite;

208  constexpr auto KThreadRead = 64 / MNPerWmma;

209  constexpr auto K0PerThreadRead = ABK0Number / KThreadRead;

210 

211  constexpr auto kfold = (ABK1Number * MN0 * sizeof(LDSTypeAB) > 128)

212  ? 1

213  : 128 / (ABK1Number * MN0 * sizeof(LDSTypeAB));

214  constexpr auto KThreadReadPerm =

215  (kfold * K0PerThreadWrite / K0PerThreadRead) > 1

216  ? KThreadRead / (kfold * K0PerThreadWrite / K0PerThreadRead)

217  : KThreadRead;

218 

219  // 1<=mpair<=n0

220  constexpr auto mpair = (ABK1Number * MNPerWmma * sizeof(LDSTypeAB) > 128)

221  ? 1

222  : ((128 / (ABK1Number * MNPerWmma * sizeof(LDSTypeAB))) > MN0

223  ? MN0

224  : 128 / (ABK1Number * MNPerWmma * sizeof(LDSTypeAB)));

225 

226  constexpr auto ab_lds_block_desc = make_naive_tensor_descriptor_packed(

227  make_tuple(Number<KThreadWrite / kfold / KThreadReadPerm>{},

228  Number<K0PerThreadWrite>{},

229  Number<KThreadReadPerm * MN1>{},

230  Number<kfold * MN0 / mpair>{},

231  Number<mpair>{},

232  ABK1Number));

233 

234  constexpr auto ab_lds_block_desc_permuted = transform_tensor_descriptor(

235  ab_lds_block_desc,

236  make_tuple(

237  make_pass_through_transform(Number<KThreadWrite / kfold / KThreadReadPerm>{}),

238  make_pass_through_transform(Number<K0PerThreadWrite>{}),

239  make_xor_with_modulo_transform(

240  make_tuple(Number<KThreadReadPerm * MN1>{}, Number<kfold * MN0 / mpair>{})),

241  make_pass_through_transform(Number<mpair>{}),

242  make_pass_through_transform(ABK1Number)),

243  make_tuple(

244  Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4>{}, Sequence<5>{}),

245  make_tuple(

246  Sequence<0>{}, Sequence<1>{}, Sequence<2, 3>{}, Sequence<4>{}, Sequence<5>{}));

247 

248  constexpr auto ab_lds_block_desc_unmerged = transform_tensor_descriptor(

249  ab_lds_block_desc_permuted,

250  make_tuple(

251  make_pass_through_transform(Number<KThreadWrite / kfold / KThreadReadPerm>{}),

252  make_pass_through_transform(Number<K0PerThreadWrite>{}),

253  make_unmerge_transform(make_tuple(Number<KThreadReadPerm>{}, Number<MN1>{})),

254  make_unmerge_transform(make_tuple(Number<kfold>{}, Number<MN0 / mpair>{})),

255  make_pass_through_transform(Number<mpair>{}),

256  make_pass_through_transform(ABK1Number)),

257  make_tuple(Sequence<0>{},

258  Sequence<1>{},

259  Sequence<2>{},

260  Sequence<3>{},

261  Sequence<4>{},

262  Sequence<5>{}),

263  make_tuple(Sequence<1>{},

264  Sequence<2>{},

265  Sequence<0, 3>{},

266  Sequence<4, 5>{},

267  Sequence<6>{},

268  Sequence<7>{}));

269 

270  constexpr auto ab_lds_block_desc_abk0_mn_abk1 = transform_tensor_descriptor(

271  ab_lds_block_desc_unmerged,

272  make_tuple(make_merge_transform_v3_division_mod(

273  make_tuple(Number<KThreadReadPerm>{},

274  Number<KThreadWrite / kfold / KThreadReadPerm>{},

275  Number<kfold>{},

276  Number<K0PerThreadWrite>{})),

277  make_merge_transform_v3_division_mod(

278  make_tuple(Number<MN0 / mpair>{}, Number<mpair>{}, Number<MN1>{})),

279  make_pass_through_transform(ABK1Number)),

280  make_tuple(Sequence<0, 1, 4, 2>{}, Sequence<5, 6, 3>{}, Sequence<7>{}),

281  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}));

282 

283  return ab_lds_block_desc_abk0_mn_abk1;

284  }

285  }

286 

287  template <typename GridDescriptor,

288  typename BlockDescriptor,

289  typename ABsDataType,

290  typename ABElementwiseOperation,

291  index_t GlobalBufferNum>

292  __device__ static auto GetBlockTransfer(GridDescriptor& grid_descriptor,

293  BlockDescriptor& block_descriptor,

294  ABElementwiseOperation& ab_element_op,

295  const index_t block_mn_id)

296  {

297  constexpr index_t NumABTensor = ABsDataType::Size();

298  const index_t mn_block_data_idx_on_grid =

299  __builtin_amdgcn_readfirstlane(block_mn_id * MNPerBlock);

300  // workaround because v7r2 is not as general as v4r1

301  if constexpr(NumABTensor > 1)

302  {

303  const auto idx_as_block_begin = generate_tuple(

304  [&](auto) { return make_multi_index(0, mn_block_data_idx_on_grid, 0); },

305  Number<NumABTensor>{});

306 

307  return ThreadGroupTensorSliceTransfer_v7r2<

308  ThisThreadBlock,

309  ABsDataType,

310  Tuple<LDSTypeAB>,

311  GridDescriptor,

312  decltype(tie(block_descriptor)),

313  ABElementwiseOperation,

314  Sequence<static_cast<index_t>(InMemoryDataOperationEnum::Set)>,

315  Sequence<ABK0Number, MNPerBlock, ABK1Number>,

316  ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1,

317  ABBlockTransferThreadClusterArrangeOrder,

318  ABBlockTransferSrcAccessOrder,

319  Sequence<1, 0, 2>,

320  ABBlockTransferSrcVectorDim,

321  2,

322  ABBlockTransferSrcScalarPerVector,

323  ABBlockTransferDstScalarPerVector_ABK1,

324  uniform_sequence_gen_t<NumABTensor, ABThreadTransferSrcResetCoordinateAfterRun>,

325  Sequence<true>,

326  GlobalBufferNum>{grid_descriptor,

327  idx_as_block_begin,

328  tie(block_descriptor),

329  make_tuple(make_multi_index(0, 0, 0)),

330  ab_element_op};

331  }

332  else

333  {

334  return ThreadGroupTensorSliceTransfer_v4r1<

335  ThisThreadBlock,

336  ABElementwiseOperation,

337  ck::tensor_operation::element_wise::PassThrough,

338  InMemoryDataOperationEnum::Set,

339  Sequence<ABK0Number, MNPerBlock, ABK1Number>,

340  ABBlockTransferThreadClusterLengths_ABK0_MN_ABK1,

341  ABBlockTransferThreadClusterArrangeOrder,

342  remove_cvref_t<tuple_element_t<0, ABsDataType>>,

343  remove_cvref_t<tuple_element_t<0, ABsDataType>>,

344  decltype(grid_descriptor[I0]),

345  decltype(block_descriptor),

346  ABBlockTransferSrcAccessOrder,

347  Sequence<0, 1, 2>,

348  ABBlockTransferSrcVectorDim,

349  2,

350  ABBlockTransferSrcScalarPerVector,

351  ABBlockTransferDstScalarPerVector_ABK1,

352  1,

353  1,

354  ABThreadTransferSrcResetCoordinateAfterRun,

355  true,

356  GlobalBufferNum>(grid_descriptor[I0],

357  make_multi_index(0, mn_block_data_idx_on_grid, 0),

358  ab_element_op,

359  block_descriptor,

360  make_multi_index(0, 0, 0),

361  ck::tensor_operation::element_wise::PassThrough{});

362  }

363  }

364 

365  template <index_t MNRepeat, index_t MNWaves>

366  __host__ __device__ static constexpr auto MakeWmmaTileDescriptor()

367  {

368  // This is a block descriptor used to read LDS memory into register

369  // It's defined in a way consistent with the existing implementation to

370  // avoid changes in the pipelines

371  using BlockDesc = decltype(GetBlockDescriptor());

372  // ABK0_MN_ABK1 -> ABK0_MNRepeat_MNWaves_KRow_MNPerWmma_ABK1

373  constexpr auto ABK0 = BlockDesc{}.GetLength(I0);

374  constexpr auto ABK1 = BlockDesc{}.GetLength(I2);

375 #ifdef __gfx12__

376  constexpr auto KRow = I2;

377 #else

378  constexpr auto KRow = I1;

379 #endif

380  if constexpr(KInner > 1)

381  {

382  // KPack = KInner * KPerWmma

383  // K1 = KInner * KPerWmmaBlk

384  // Each thread loads multiple tiles with one instruction

385  // 1 - MNRepeat - K0 / KRow - MNWaves - KRow - MNPerWmma - K1

386  return transform_tensor_descriptor(

387  BlockDesc{},

388  make_tuple(

389  make_unmerge_transform(make_tuple(Number<ABK0 / KRow>{}, KRow, Number<1>{})),

390  make_unmerge_transform(

391  make_tuple(Number<MNRepeat>{}, Number<MNWaves>{}, Number<MNPerWmma>{})),

392  make_pass_through_transform(Number<ABK1>{})),

393  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}),

394  make_tuple(Sequence<2, 4, 0>{}, Sequence<1, 3, 5>{}, Sequence<6>{}));

395  }

396  else

397  {

398  // KPack = KPerWmma (KInner == 1)

399  if constexpr(ABK1 <= KPerWmmaBlk)

400  {

401  // K1 <= single tile (KPerWmmaBlk)

402  // Each thread will load KPerWmmaBlk for the WMMA instruction

403  // Since K1 <= single tile, K0 is unmerged first over KPack / KRow / K1

404  // (rest of the single WMMA tile for single thread) and then over KRow

405  // (rest of the single WMMA tile for single wave)

406  // KPack / KRow / K1 - MNRepeat - K0 / KRow - MNWaves - KRow - MNPerWmma - K1

407  return transform_tensor_descriptor(

408  BlockDesc{},

409  make_tuple(

410  make_unmerge_transform(make_tuple(

411  Number<ABK0 / (KPack / ABK1)>{}, KRow, Number<KPack / KRow / ABK1>{})),

412  make_unmerge_transform(

413  make_tuple(Number<MNRepeat>{}, Number<MNWaves>{}, Number<MNPerWmma>{})),

414  make_pass_through_transform(Number<ABK1>{})),

415  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}),

416  make_tuple(Sequence<2, 4, 0>{}, Sequence<1, 3, 5>{}, Sequence<6>{}));

417  }

418  else

419  {

420  // K1 > single tile (KPerWmmaBlk)

421  // Each thread will load KPerWmmaBlk for the WMMA instruction

422  // Since K1 > single tile, each thread loads KPerWmmaBlk and the next

423  // KPerWmmaBlk chunk is loaded by a different thread in the same wave (WMMA layout).

424  // This layout is needed to support for example AK1 > single tile and

425  // BK1 <= single tile in the same gemm

426  // KPack / KPerWmmaBlk / KRow - MNRepeat - K0 / KRow - MNWaves - KRow - MNPerWmma -

427  // K1

428  constexpr auto desc1 = transform_tensor_descriptor(

429  BlockDesc{},

430  make_tuple(

431  make_pass_through_transform(Number<ABK0>{}),

432  make_unmerge_transform(

433  make_tuple(Number<MNRepeat>{}, Number<MNWaves>{}, Number<MNPerWmma>{})),

434  make_unmerge_transform(make_tuple(Number<ABK1 / KPack>{},

435  Number<KPack / KPerWmmaBlk / KRow>{},

436  Number<KRow>{},

437  Number<KPerWmmaBlk>{}))),

438  make_tuple(Sequence<0>{}, Sequence<1>{}, Sequence<2>{}),

439  make_tuple(Sequence<2>{}, Sequence<1, 4, 6>{}, Sequence<3, 0, 5, 7>{}));

440 

441  return transform_tensor_descriptor(

442  desc1,

443  make_tuple(

444  make_pass_through_transform(Number<KPack / KPerWmmaBlk / KRow>{}),

445  make_pass_through_transform(Number<MNRepeat>{}),

446  make_merge_transform(make_tuple(Number<ABK0>{}, Number<ABK1 / KPack>{})),

447  make_pass_through_transform(Number<MNWaves>{}),

448  make_pass_through_transform(Number<KRow>{}),

449  make_pass_through_transform(Number<MNPerWmma>{}),

450  make_pass_through_transform(Number<KPerWmmaBlk>{})),

451  make_tuple(Sequence<0>{},

452  Sequence<1>{},

453  Sequence<2, 3>{},

454  Sequence<4>{},

455  Sequence<5>{},

456  Sequence<6>{},

457  Sequence<7>{}),

458  make_tuple(Sequence<0>{},

459  Sequence<1>{},

460  Sequence<2>{},

461  Sequence<3>{},

462  Sequence<4>{},

463  Sequence<5>{},

464  Sequence<6>{}));

465  }

466  }

467  }

468 

469  __device__ static constexpr auto GetBlockStep()

470  {

471  // Grid descriptor step (MoveSrcSliceWindow)

472  return make_multi_index(KPerBlock / ABK1Number, 0, 0);

473  }

474 

475  template <typename GridDescriptor>

476  __device__ static constexpr index_t GetKDimension(const GridDescriptor& grid_desc)

477  {

478  // K dimension size. This should always be called with the A matrix grid descriptor

479  // because it doesn't work for B matrix when packed int4 is used

480  return grid_desc.GetLength(I0) * grid_desc.GetLength(I2);

481  }

482 };

483 

484 } // namespace ck

amd_address_space.hpp

ck::tensor_operation::device::grid_desc

auto grid_desc(MatrixPadder< GemmSpec, MPerTileType, NPerTileType, KPerTileType > matrix_padder, CDesc_MRaw_NRaw conv_desc)

Definition: matrix_padder.hpp:190

ck_tile::PassThrough

ck_tile::element_wise::PassThrough PassThrough

Definition: grouped_convolution_utils.hpp:47

ck

Definition: ck.hpp:270

ck::make_multi_index

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

Definition: array_multi_index.hpp:15

ck::make_naive_tensor_descriptor

__host__ constexpr __device__ auto make_naive_tensor_descriptor(const Tuple< Lengths... > &lengths, const Tuple< Strides... > &strides)

Definition: tensor_descriptor_helper.hpp:49

ck::uniform_sequence_gen_t

typename uniform_sequence_gen< NSize, I >::type uniform_sequence_gen_t

Definition: sequence.hpp:928

ck::generate_tuple

__host__ constexpr __device__ auto generate_tuple(F &&f, Number< N >)

Definition: tuple_helper.hpp:21

ck::InMemoryDataOperationEnum::Set

@ Set

ck::make_naive_tensor_descriptor_packed

__host__ constexpr __device__ auto make_naive_tensor_descriptor_packed(const Tuple< Lengths... > &lengths)

Definition: tensor_descriptor_helper.hpp:101

ck::make_merge_transform

__host__ constexpr __device__ auto make_merge_transform(const LowLengths &low_lengths)

Definition: multi_index_transform_helper.hpp:55

ck::make_merge_transform_v3_division_mod

__host__ constexpr __device__ auto make_merge_transform_v3_division_mod(const LowLengths &low_lengths)

Definition: multi_index_transform_helper.hpp:84

ck::tie

constexpr Tuple< Args &... > tie(Args &... args) noexcept

Definition: tuple.hpp:218

ck::make_xor_with_modulo_transform

__host__ constexpr __device__ auto make_xor_with_modulo_transform(const LowLengths &low_lengths)

Definition: multi_index_transform_helper.hpp:185

ck::make_pass_through_transform

__host__ constexpr __device__ auto make_pass_through_transform(const LowLength &low_length)

Definition: multi_index_transform_helper.hpp:12

ck::is_same_v

constexpr bool is_same_v

Definition: type.hpp:283

ck::make_tuple

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

Definition: tuple.hpp:211

ck::remove_cvref_t

remove_cv_t< remove_reference_t< T > > remove_cvref_t

Definition: type.hpp:297

ck::make_unmerge_transform

__host__ constexpr __device__ auto make_unmerge_transform(const UpLengths &up_lengths, integral_constant< bool, Use24BitIntegerCalculation >=integral_constant< bool, false >{})

Definition: multi_index_transform_helper.hpp:90

ck::index_t

int32_t index_t

Definition: ck.hpp:301

ck::transform_tensor_descriptor

__host__ constexpr __device__ auto transform_tensor_descriptor(const OldTensorDescriptor &old_tensor_desc, const NewTransforms &new_transforms, NewLowerDimensionOldVisibleIdss, NewUpperDimensionNewVisibleIdss)

Definition: tensor_descriptor.hpp:319

ck::make_right_pad_transform

__host__ constexpr __device__ auto make_right_pad_transform(const LowLength &low_length, const RightPadLength &right_pad, integral_constant< bool, SkipIsValidCheck >=integral_constant< bool, false >{})

Definition: multi_index_transform_helper.hpp:37

ck::ABTransferThreadTiles

Definition: gridwise_ab_transfer_thread_tiles.hpp:33

ck::ABTransferThreadTiles::ThisThreadBlock

ThisThreadBlock< BlockSize > ThisThreadBlock

Definition: gridwise_ab_transfer_thread_tiles.hpp:48

ck::ABTransferThreadTiles::I1

static constexpr auto I1

Definition: gridwise_ab_transfer_thread_tiles.hpp:38

ck::ABTransferThreadTiles::ABK0Number

static constexpr auto ABK0Number

Definition: gridwise_ab_transfer_thread_tiles.hpp:34

ck::ABTransferThreadTiles::GetBlockStep

static constexpr __device__ auto GetBlockStep()

Definition: gridwise_ab_transfer_thread_tiles.hpp:469

ck::ABTransferThreadTiles::MakeGridDescriptor

__host__ static __device__ auto MakeGridDescriptor(const GridDescriptorBase &ab_grid_desc, index_t MN, index_t MNPad, index_t K, index_t KPad, index_t StrideAB, index_t ABK0)

Definition: gridwise_ab_transfer_thread_tiles.hpp:51

ck::ABTransferThreadTiles::I0

static constexpr auto I0

Definition: gridwise_ab_transfer_thread_tiles.hpp:37

ck::ABTransferThreadTiles::GetBlockTransfer

static __device__ auto GetBlockTransfer(GridDescriptor &grid_descriptor, BlockDescriptor &block_descriptor, ABElementwiseOperation &ab_element_op, const index_t block_mn_id)

Definition: gridwise_ab_transfer_thread_tiles.hpp:292

ck::ABTransferThreadTiles::ABPackedSize

static constexpr index_t ABPackedSize

Definition: gridwise_ab_transfer_thread_tiles.hpp:41

ck::ABTransferThreadTiles::GetBlockDescriptor

static constexpr __device__ auto GetBlockDescriptor()

Definition: gridwise_ab_transfer_thread_tiles.hpp:147

ck::ABTransferThreadTiles::ABK1Number

static constexpr auto ABK1Number

Definition: gridwise_ab_transfer_thread_tiles.hpp:35

ck::ABTransferThreadTiles::MakeWmmaTileDescriptor

__host__ static constexpr __device__ auto MakeWmmaTileDescriptor()

Definition: gridwise_ab_transfer_thread_tiles.hpp:366

ck::ABTransferThreadTiles::GetKDimension

static constexpr __device__ index_t GetKDimension(const GridDescriptor &grid_desc)

Definition: gridwise_ab_transfer_thread_tiles.hpp:476

ck::ABTransferThreadTiles::I2

static constexpr auto I2

Definition: gridwise_ab_transfer_thread_tiles.hpp:39

ck::Sequence

Definition: sequence.hpp:43

ck::ThreadGroupTensorSliceTransfer_v4r1

Blockwise data transfer.

Definition: thread_group_tensor_slice_transfer_v4r1.hpp:46

ck::ThreadGroupTensorSliceTransfer_v7r2

Definition: thread_group_tensor_slice_transfer_v7r2.hpp:47

ck::Tuple

Definition: tuple.hpp:117

ck::integral_constant

Definition: integral_constant.hpp:20

ck::is_same

Definition: type.hpp:177

ck::pk_i4_t

Definition: data_type.hpp:187

ck::tensor_operation::element_wise::PassThrough

Definition: unary_element_wise_operation.hpp:340

thread_group_tensor_slice_transfer_v4r1.hpp

thread_group_tensor_slice_transfer_v7r2.hpp