/home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-composable-kernel/checkouts/develop/include/ck_tile/ops/gemm/warp/warp_gemm_attribute_mfma_impl.hpp Source File

• develop
• include
• ck_tile
• ops
• gemm
• warp

warp_gemm_attribute_mfma_impl.hpp

Go to the documentation of this file.

1 // SPDX-License-Identifier: MIT

2 // Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.

3 

4 #pragma once

5 

6 #include "ck_tile/core.hpp"

7 

8 namespace ck_tile {

9 

10 // TODO: refactor warp-gemm

11 // currently there is a discrepency for vav/vva if we need transpose C/D

12 // e.g. if we want A:agpr, B:vgpr, we have to use vva in WGAttrEnum

13 // because we swap the A/B pointer in _impl code (but not known this info here)

14 enum class WGAttrCtlEnum

15 {

16  Default_ = 0,

17  Raw_vvv = 1, // c-vgpr, a-vgpr, b-vgpr

18  Raw_vaa = 2, // c-vgpr, a-agpr, b-agpr

19  Raw_vav = 3, // c-vgpr, a-agpr, b-vgpr

20  Raw_vva = 4, // c-vgpr, a-vgpr, b-agpr

21  Raw_avv = 5, // c-agpr, a-vgpr, b-vgpr

22  // raw_a_a_a = 3, // c-agpr, a-agpr, b-agpr

23 };

24 

25 #define DISPATCH_MFMA_(mfma_, dmod_, amod_, bmod_, cmod_) \

26  if constexpr(post_nop_) \

27  { \

28  asm volatile(mfma_ " %0, %1, %2, %3 ; yyy\n" \

29  "s_nop 3" \

30  : dmod_(c_vec) \

31  : amod_(a_vec), bmod_(b_vec), cmod_(c_vec) \

32  :); \

33  } \

34  else \

35  { \

36  asm volatile(mfma_ " %0, %1, %2, %3\n" \

37  : dmod_(c_vec) \

38  : amod_(a_vec), bmod_(b_vec), cmod_(c_vec) \

39  :); \

40  }

41 

42 #define DISPATCH_MFMA_CTRL_(mfma_, ctrl_) \

43  if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vvv) \

44  { \

45  DISPATCH_MFMA_(mfma_, "+v", "v", "v", "v") \

46  } \

47  else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vaa) \

48  { \

49  DISPATCH_MFMA_(mfma_, "+v", "a", "a", "v") \

50  } \

51  else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vav) \

52  { \

53  DISPATCH_MFMA_(mfma_, "+v", "a", "v", "v") \

54  } \

55  else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_vva) \

56  { \

57  DISPATCH_MFMA_(mfma_, "+v", "v", "a", "v") \

58  } \

59  else if constexpr(ctrl_ == WGAttrCtlEnum::Raw_avv) \

60  { \

61  DISPATCH_MFMA_(mfma_, "+a", "v", "v", "a") \

62  }

63 

64 // F32

65 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

66 struct WarpGemmAttributeMfmaImplF32F32F32M16N16K4

67 {

68  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

69 

70  using ADataType = float;

71  using BDataType = float;

72  using CDataType = float;

73 

74  using AVecType = ext_vector_t<ADataType, 1>;

75  using BVecType = ext_vector_t<BDataType, 1>;

76  using CVecType = ext_vector_t<CDataType, 4>;

77 

78  static constexpr index_t kM = 16;

79  static constexpr index_t kN = 16;

80  static constexpr index_t kK = 4;

81 

82  static constexpr index_t kAMBlock = 1;

83  static constexpr index_t kBNBlock = 1;

84 

85  static constexpr index_t kAMLane = 16;

86  static constexpr index_t kBNLane = 16;

87  static constexpr index_t kABKLane = 4;

88  static constexpr index_t kABKPerLane = 1;

89 

90  static constexpr index_t kCMLane = 4;

91  static constexpr index_t kCNLane = 16;

92  static constexpr index_t kCM0PerLane = 1;

93  static constexpr index_t kCM1PerLane = 4;

94 

95  // c_vec += a_vec * b_vec

96  template <bool post_nop_ = false>

97  CK_TILE_DEVICE void operator()(CVecType& c_vec,

98  const AVecType& a_vec,

99  const BVecType& b_vec,

100  bool_constant<post_nop_> = {}) const

101  {

102  DISPATCH_MFMA_CTRL_("v_mfma_f32_16x16x4f32", Ctrl)

103  else

104  {

105 #if defined(__gfx9__)

106  c_vec = __builtin_amdgcn_mfma_f32_16x16x4f32(a_vec[0], b_vec[0], c_vec, 0, 0, 0);

107 #else

108  ck_tile::ignore = c_vec;

109  ck_tile::ignore = a_vec;

110  ck_tile::ignore = b_vec;

111 #endif

112  }

113  }

114 

115  // c_vec = a_vec * b_vec

116  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

117  {

118 #if defined(__gfx9__)

119  return bit_cast<CVecType>(

120  __builtin_amdgcn_mfma_f32_16x16x4f32(a_vec[0], b_vec[0], CVecType{0.f}, 0, 0, 0));

121 #else

122  ck_tile::ignore = a_vec;

123  ck_tile::ignore = b_vec;

124  return CVecType{0.f};

125 #endif

126  }

127 };

128 

129 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

130 struct WarpGemmAttributeMfmaImplF32F32F32M32N32K2

131 {

132  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

133 

134  using ADataType = float;

135  using BDataType = float;

136  using CDataType = float;

137 

138  using AVecType = ext_vector_t<ADataType, 1>;

139  using BVecType = ext_vector_t<BDataType, 1>;

140  using CVecType = ext_vector_t<CDataType, 16>;

141 

142  static constexpr index_t kM = 32;

143  static constexpr index_t kN = 32;

144  static constexpr index_t kK = 2;

145 

146  static constexpr index_t kAMBlock = 1;

147  static constexpr index_t kBNBlock = 1;

148 

149  static constexpr index_t kAMLane = 32;

150  static constexpr index_t kBNLane = 32;

151  static constexpr index_t kABKLane = 2;

152  static constexpr index_t kABKPerLane = 1;

153 

154  static constexpr index_t kCMLane = 2;

155  static constexpr index_t kCNLane = 32;

156  static constexpr index_t kCM0PerLane = 4;

157  static constexpr index_t kCM1PerLane = 4;

158 

159  // c_vec += a_vec * b_vec

160  template <bool post_nop_ = false>

161  CK_TILE_DEVICE void operator()(CVecType& c_vec,

162  const AVecType& a_vec,

163  const BVecType& b_vec,

164  bool_constant<post_nop_> = {}) const

165  {

166  DISPATCH_MFMA_CTRL_("v_mfma_f32_32x32x2f32", Ctrl)

167  else

168  {

169 #if defined(__gfx9__)

170  c_vec = __builtin_amdgcn_mfma_f32_32x32x2f32(a_vec[0], b_vec[0], c_vec, 0, 0, 0);

171 #else

172  ck_tile::ignore = c_vec;

173  ck_tile::ignore = a_vec;

174  ck_tile::ignore = b_vec;

175 #endif

176  }

177  }

178 

179  // c_vec = a_vec * b_vec

180  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

181  {

182 #if defined(__gfx9__)

183  return bit_cast<CVecType>(

184  __builtin_amdgcn_mfma_f32_32x32x2f32(a_vec[0], b_vec[0], CVecType{0.f}, 0, 0, 0));

185 #else

186  ck_tile::ignore = a_vec;

187  ck_tile::ignore = b_vec;

188  return CVecType{0.f};

189 #endif

190  }

191 };

192 

193 // V_MFMA_F32_16x16x32_BF16

194 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

195 struct WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32

196 {

197  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

198  using ADataType = bf16_t;

199  using BDataType = bf16_t;

200  using CDataType = float;

201 

202  using AVecType = ext_vector_t<bf16_t, 8>;

203  using BVecType = ext_vector_t<bf16_t, 8>;

204  using CVecType = ext_vector_t<float, 4>;

205 

206  static constexpr index_t kM = 16;

207  static constexpr index_t kN = 16;

208  static constexpr index_t kK = 32;

209 

210  static constexpr index_t kAMBlock = 1;

211  static constexpr index_t kBNBlock = 1;

212 

213  static constexpr index_t kAMLane = 16;

214  static constexpr index_t kBNLane = 16;

215  static constexpr index_t kABKLane = 4;

216  static constexpr index_t kABKPerLane = 8;

217 

218  static constexpr index_t kCMLane = 4;

219  static constexpr index_t kCNLane = 16;

220  static constexpr index_t kCM0PerLane = 1;

221  static constexpr index_t kCM1PerLane = 4;

222 

223  // c_vec += a_vec * b_vec

224  template <bool post_nop_ = false>

225  CK_TILE_DEVICE void operator()(CVecType& c_vec,

226  const AVecType& a_vec,

227  const BVecType& b_vec,

228  bool_constant<post_nop_> = {}) const

229  {

230  DISPATCH_MFMA_CTRL_("v_mfma_f32_16x16x32_bf16", Ctrl)

231  else

232  {

233 #if defined(__gfx950__)

234  c_vec = __builtin_amdgcn_mfma_f32_16x16x32_bf16(a_vec, b_vec, c_vec, 0, 0, 0);

235 #else

236  ck_tile::ignore = c_vec;

237  ck_tile::ignore = a_vec;

238  ck_tile::ignore = b_vec;

239 #endif

240  }

241  }

242 

243  // c_vec = a_vec * b_vec

244  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

245  {

246 #if defined(__gfx950__)

247  return bit_cast<CVecType>(

248  __builtin_amdgcn_mfma_f32_16x16x32_bf16(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

249 #else

250  ck_tile::ignore = a_vec;

251  ck_tile::ignore = b_vec;

252  return CVecType{0.f};

253 #endif

254  }

255 };

256 // FP16

257 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

258 struct WarpGemmAttributeMfmaImplF16F16F32M32N32K8

259 {

260  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

261  using ADataType = fp16_t;

262  using BDataType = fp16_t;

263  using CDataType = float;

264 

265  using AVecType = ext_vector_t<fp16_t, 4>;

266  using BVecType = ext_vector_t<fp16_t, 4>;

267  using CVecType = ext_vector_t<float, 16>;

268 

269  static constexpr index_t kM = 32;

270  static constexpr index_t kN = 32;

271  static constexpr index_t kK = 8;

272 

273  static constexpr index_t kAMBlock = 1;

274  static constexpr index_t kBNBlock = 1;

275 

276  static constexpr index_t kAMLane = 32;

277  static constexpr index_t kBNLane = 32;

278  static constexpr index_t kABKLane = 2;

279  static constexpr index_t kABKPerLane = 4;

280 

281  static constexpr index_t kCMLane = 2;

282  static constexpr index_t kCNLane = 32;

283  static constexpr index_t kCM0PerLane = 4;

284  static constexpr index_t kCM1PerLane = 4;

285 

286  // c_vec += a_vec * b_vec

287  template <bool post_nop_ = false>

288  CK_TILE_DEVICE void operator()(CVecType& c_vec,

289  const AVecType& a_vec,

290  const BVecType& b_vec,

291  bool_constant<post_nop_> = {}) const

292  {

293  DISPATCH_MFMA_CTRL_("v_mfma_f32_32x32x8f16", Ctrl)

294  else

295  {

296 #if defined(__gfx9__)

297  c_vec = __builtin_amdgcn_mfma_f32_32x32x8f16(a_vec, b_vec, c_vec, 0, 0, 0);

298 #else

299  ck_tile::ignore = c_vec;

300  ck_tile::ignore = a_vec;

301  ck_tile::ignore = b_vec;

302 #endif

303  }

304  }

305 

306  // c_vec = a_vec * b_vec

307  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

308  {

309 #if defined(__gfx9__)

310  return bit_cast<CVecType>(

311  __builtin_amdgcn_mfma_f32_32x32x8f16(a_vec, b_vec, fp32x16_t{0.f}, 0, 0, 0));

312 #else

313  ck_tile::ignore = a_vec;

314  ck_tile::ignore = b_vec;

315  return CVecType{0.f};

316 #endif

317  }

318 };

319 

320 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

321 struct WarpGemmAttributeMfmaImplF16F16F32M16N16K16

322 {

323  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

324  using ADataType = fp16_t;

325  using BDataType = fp16_t;

326  using CDataType = float;

327 

328  using AVecType = ext_vector_t<fp16_t, 4>;

329  using BVecType = ext_vector_t<fp16_t, 4>;

330  using CVecType = ext_vector_t<float, 4>;

331 

332  static constexpr index_t kM = 16;

333  static constexpr index_t kN = 16;

334  static constexpr index_t kK = 16;

335 

336  static constexpr index_t kAMBlock = 1;

337  static constexpr index_t kBNBlock = 1;

338 

339  static constexpr index_t kAMLane = 16;

340  static constexpr index_t kBNLane = 16;

341  static constexpr index_t kABKLane = 4;

342  static constexpr index_t kABKPerLane = 4;

343 

344  static constexpr index_t kCMLane = 4;

345  static constexpr index_t kCNLane = 16;

346  static constexpr index_t kCM0PerLane = 1;

347  static constexpr index_t kCM1PerLane = 4;

348 

349  // c_vec += a_vec * b_vec

350  template <bool post_nop_ = false>

351  CK_TILE_DEVICE void operator()(CVecType& c_vec,

352  const AVecType& a_vec,

353  const BVecType& b_vec,

354  bool_constant<post_nop_> = {}) const

355  {

356  DISPATCH_MFMA_CTRL_("v_mfma_f32_16x16x16f16", Ctrl)

357  else

358  {

359 #if defined(__gfx9__)

360  c_vec = __builtin_amdgcn_mfma_f32_16x16x16f16(a_vec, b_vec, c_vec, 0, 0, 0);

361 #else

362  ck_tile::ignore = c_vec;

363  ck_tile::ignore = a_vec;

364  ck_tile::ignore = b_vec;

365 #endif

366  }

367  }

368 

369  // c_vec = a_vec * b_vec

370  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

371  {

372 #if defined(__gfx9__)

373  return bit_cast<CVecType>(

374  __builtin_amdgcn_mfma_f32_16x16x16f16(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

375 #else

376  ck_tile::ignore = a_vec;

377  ck_tile::ignore = b_vec;

378  return CVecType{0.f};

379 #endif

380  }

381 };

382 

383 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

384 struct WarpGemmAttributeMfmaImplF16F16F32M16N16K32

385 {

386  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

387  using ADataType = fp16_t;

388  using BDataType = fp16_t;

389  using CDataType = float;

390 

391  using AVecType = ext_vector_t<fp16_t, 8>;

392  using BVecType = ext_vector_t<fp16_t, 8>;

393  using CVecType = ext_vector_t<float, 4>;

394 

395  static constexpr index_t kM = 16;

396  static constexpr index_t kN = 16;

397  static constexpr index_t kK = 32;

398 

399  static constexpr index_t kAMBlock = 1;

400  static constexpr index_t kBNBlock = 1;

401 

402  static constexpr index_t kAMLane = 16;

403  static constexpr index_t kBNLane = 16;

404  static constexpr index_t kABKLane = 4;

405  static constexpr index_t kABKPerLane = 8;

406 

407  static constexpr index_t kCMLane = 4;

408  static constexpr index_t kCNLane = 16;

409  static constexpr index_t kCM0PerLane = 1;

410  static constexpr index_t kCM1PerLane = 4;

411 

412  // c_vec += a_vec * b_vec

413  template <bool post_nop_ = false>

414  CK_TILE_DEVICE void operator()(CVecType& c_vec,

415  const AVecType& a_vec,

416  const BVecType& b_vec,

417  bool_constant<post_nop_> = {}) const

418  {

419  DISPATCH_MFMA_CTRL_("v_mfma_f32_16x16x32f16", Ctrl)

420  else

421  {

422 #if defined(__gfx950__)

423  c_vec = __builtin_amdgcn_mfma_f32_16x16x32_f16(a_vec, b_vec, c_vec, 0, 0, 0);

424 #else

425  ck_tile::ignore = c_vec;

426  ck_tile::ignore = a_vec;

427  ck_tile::ignore = b_vec;

428 #endif

429  }

430  }

431 

432  // c_vec = a_vec * b_vec

433  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

434  {

435 #if defined(__gfx950__)

436  return bit_cast<CVecType>(

437  __builtin_amdgcn_mfma_f32_16x16x32_f16(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

438 #else

439  ck_tile::ignore = a_vec;

440  ck_tile::ignore = b_vec;

441  return CVecType{0.f};

442 #endif

443  }

444 };

445 

446 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

447 struct WarpGemmAttributeMfmaImplF16F16F32M4N64K4

448 {

449  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

450  using ADataType = fp16_t;

451  using BDataType = fp16_t;

452  using CDataType = float;

453 

454  using AVecType = ext_vector_t<fp16_t, 4>;

455  using BVecType = ext_vector_t<fp16_t, 4>;

456  using CVecType = ext_vector_t<float, 4>;

457 

458  static constexpr index_t kM = 4;

459  static constexpr index_t kN = 64;

460  static constexpr index_t kK = 4;

461 

462  static constexpr index_t kAMBlock = 1;

463  static constexpr index_t kBNBlock = 16;

464 

465  // we only write down single block (4 threads) thread mapping here

466  static constexpr index_t kAMLane = 4;

467  static constexpr index_t kBNLane = 4;

468  static constexpr index_t kABKLane = 1;

469  static constexpr index_t kABKPerLane = 4;

470 

471  static constexpr index_t kCMLane = 1;

472  static constexpr index_t kCNLane = 4;

473  static constexpr index_t kCM0PerLane = 1;

474  static constexpr index_t kCM1PerLane = 4;

475 

476  // c_vec += a_vec * b_vec

477  template <bool post_nop_ = false>

478  CK_TILE_DEVICE void operator()(CVecType& c_vec,

479  const AVecType& a_vec,

480  const BVecType& b_vec,

481  bool_constant<post_nop_> = {}) const

482  {

483  DISPATCH_MFMA_CTRL_("v_mfma_f32_4x4x4f16", Ctrl)

484  else

485  {

486 #if defined(__gfx9__)

487  c_vec = __builtin_amdgcn_mfma_f32_4x4x4f16(a_vec, b_vec, c_vec, 0, 0, 0);

488 #else

489  ignore = c_vec;

490  ignore = a_vec;

491  ignore = b_vec;

492 #endif

493  }

494  }

495 

496  // c_vec = a_vec * b_vec

497  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

498  {

499 #if defined(__gfx9__)

500  return bit_cast<CVecType>(

501  __builtin_amdgcn_mfma_f32_4x4x4f16(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

502 #else

503  ignore = a_vec;

504  ignore = b_vec;

505  return CVecType{0.f};

506 #endif

507  }

508 };

509 

510 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

511 struct WarpGemmAttributeMfmaImplF16F16F32M64N4K4

512 {

513  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

514  using ADataType = fp16_t;

515  using BDataType = fp16_t;

516  using CDataType = float;

517 

518  using AVecType = ext_vector_t<fp16_t, 4>;

519  using BVecType = ext_vector_t<fp16_t, 4>;

520  using CVecType = ext_vector_t<float, 4>;

521 

522  static constexpr index_t kM = 64;

523  static constexpr index_t kN = 4;

524  static constexpr index_t kK = 4;

525 

526  static constexpr index_t kAMBlock = 16;

527  static constexpr index_t kBNBlock = 1;

528 

529  // we only write down single block (4 threads) thread mapping here

530  static constexpr index_t kAMLane = 4;

531  static constexpr index_t kBNLane = 4;

532  static constexpr index_t kABKLane = 1;

533  static constexpr index_t kABKPerLane = 4;

534 

535  static constexpr index_t kCMLane = 1;

536  static constexpr index_t kCNLane = 4;

537  static constexpr index_t kCM0PerLane = 1;

538  static constexpr index_t kCM1PerLane = 4;

539 

540  // c_vec += a_vec * b_vec

541  template <bool post_nop_ = false>

542  CK_TILE_DEVICE void operator()(CVecType& c_vec,

543  const AVecType& a_vec,

544  const BVecType& b_vec,

545  bool_constant<post_nop_> = {}) const

546  {

547  DISPATCH_MFMA_CTRL_("v_mfma_f32_4x4x4f16", Ctrl)

548  else

549  {

550 #if defined(__gfx9__)

551  c_vec = __builtin_amdgcn_mfma_f32_4x4x4f16(a_vec, b_vec, c_vec, 0, 0, 0);

552 #else

553  ignore = c_vec;

554  ignore = a_vec;

555  ignore = b_vec;

556 #endif

557  }

558  }

559 

560  // c_vec = a_vec * b_vec

561  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

562  {

563 #if defined(__gfx9__)

564  return bit_cast<CVecType>(

565  __builtin_amdgcn_mfma_f32_4x4x4f16(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

566 #else

567  ignore = a_vec;

568  ignore = b_vec;

569  return CVecType{0.f};

570 #endif

571  }

572 };

573 

574 // Bf16

575 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

576 struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8

577 {

578  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

579  using ADataType = bf16_t;

580  using BDataType = bf16_t;

581  using CDataType = float;

582 

583  using AVecType = ext_vector_t<bf16_t, 4>;

584  using BVecType = ext_vector_t<bf16_t, 4>;

585  using CVecType = ext_vector_t<float, 16>;

586 

587  static constexpr index_t kM = 32;

588  static constexpr index_t kN = 32;

589  static constexpr index_t kK = 8;

590 

591  static constexpr index_t kAMBlock = 1;

592  static constexpr index_t kBNBlock = 1;

593 

594  static constexpr index_t kAMLane = 32;

595  static constexpr index_t kBNLane = 32;

596  static constexpr index_t kABKLane = 2;

597  static constexpr index_t kABKPerLane = 4;

598 

599  static constexpr index_t kCMLane = 2;

600  static constexpr index_t kCNLane = 32;

601  static constexpr index_t kCM0PerLane = 4;

602  static constexpr index_t kCM1PerLane = 4;

603 

604  // c_vec += a_vec * b_vec

605  template <bool post_nop_ = false>

606  CK_TILE_DEVICE void operator()(CVecType& c_vec,

607  const AVecType& a_vec,

608  const BVecType& b_vec,

609  bool_constant<post_nop_> = {}) const

610  {

611  DISPATCH_MFMA_CTRL_("v_mfma_f32_32x32x8bf16_1k", Ctrl)

612  else

613  {

614 #if defined(__gfx90a__) || defined(__gfx94__)

615  c_vec = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a_vec, b_vec, c_vec, 0, 0, 0);

616 #elif defined(__gfx908__)

617  static_for<0, 2, 1>{}([&](auto k) {

618  c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16(

619  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

620  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

621  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

622  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

623  c_vec,

624  0,

625  0,

626  0);

627  });

628 #else

629  ck_tile::ignore = c_vec;

630  ck_tile::ignore = a_vec;

631  ck_tile::ignore = b_vec;

632 #endif

633  }

634  }

635 

636  // c_vec = a_vec * b_vec

637  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

638  {

639 #if defined(__gfx90a__) || defined(__gfx94__)

640  return bit_cast<CVecType>(

641  __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a_vec, b_vec, fp32x16_t{0.f}, 0, 0, 0));

642 #elif defined(__gfx908__)

643  CVecType c_vec{0.f};

644  static_for<0, 2, 1>{}([&](auto k) {

645  c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16(

646  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

647  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

648  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

649  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

650  c_vec,

651  0,

652  0,

653  0);

654  });

655  return c_vec;

656 #else

657  ck_tile::ignore = a_vec;

658  ck_tile::ignore = b_vec;

659  return CVecType{0.f};

660 #endif

661  }

662 };

663 

664 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

665 struct WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16

666 {

667  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

668  using ADataType = bf16_t;

669  using BDataType = bf16_t;

670  using CDataType = float;

671 

672  using AVecType = ext_vector_t<bf16_t, 4>;

673  using BVecType = ext_vector_t<bf16_t, 4>;

674  using CVecType = ext_vector_t<float, 4>;

675 

676  static constexpr index_t kM = 16;

677  static constexpr index_t kN = 16;

678  static constexpr index_t kK = 16;

679 

680  static constexpr index_t kAMBlock = 1;

681  static constexpr index_t kBNBlock = 1;

682 

683  static constexpr index_t kAMLane = 16;

684  static constexpr index_t kBNLane = 16;

685  static constexpr index_t kABKLane = 4;

686  static constexpr index_t kABKPerLane = 4;

687 

688  static constexpr index_t kCMLane = 4;

689  static constexpr index_t kCNLane = 16;

690  static constexpr index_t kCM0PerLane = 1;

691  static constexpr index_t kCM1PerLane = 4;

692 

693  // c_vec += a_vec * b_vec

694  template <bool post_nop_ = false>

695  CK_TILE_DEVICE void operator()(CVecType& c_vec,

696  const AVecType& a_vec,

697  const BVecType& b_vec,

698  bool_constant<post_nop_> = {}) const

699  {

700  DISPATCH_MFMA_CTRL_("v_mfma_f32_16x16x16bf16_1k", Ctrl)

701  {

702 #if defined(__gfx90a__) || defined(__gfx94__)

703  c_vec = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a_vec, b_vec, c_vec, 0, 0, 0);

704 #elif defined(__gfx908__)

705  static_for<0, 2, 1>{}([&](auto k) {

706  c_vec = __builtin_amdgcn_mfma_f32_16x16x8bf16(

707  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

708  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

709  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

710  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

711  c_vec,

712  0,

713  0,

714  0);

715  });

716 #else

717  ck_tile::ignore = c_vec;

718  ck_tile::ignore = a_vec;

719  ck_tile::ignore = b_vec;

720 #endif

721  }

722  }

723 

724  // c_vec = a_vec * b_vec

725  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

726  {

727 #if defined(__gfx90a__) || defined(__gfx94__)

728  return bit_cast<CVecType>(

729  __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

730 #elif defined(__gfx908__)

731  CVecType c_vec{0.f};

732  static_for<0, 2, 1>{}([&](auto k) {

733  c_vec = __builtin_amdgcn_mfma_f32_16x16x8bf16(

734  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

735  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

736  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

737  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

738  c_vec,

739  0,

740  0,

741  0);

742  });

743  return c_vec;

744 #else

745  ck_tile::ignore = a_vec;

746  ck_tile::ignore = b_vec;

747  return CVecType{0.f};

748 #endif

749  }

750 };

751 

752 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

753 struct WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4

754 {

755  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

756  using ADataType = bf16_t;

757  using BDataType = bf16_t;

758  using CDataType = float;

759 

760  using AVecType = ext_vector_t<bf16_t, 4>;

761  using BVecType = ext_vector_t<bf16_t, 4>;

762  using CVecType = ext_vector_t<float, 4>;

763 

764  static constexpr index_t kM = 4;

765  static constexpr index_t kN = 64;

766  static constexpr index_t kK = 4;

767 

768  static constexpr index_t kAMBlock = 1;

769  static constexpr index_t kBNBlock = 16;

770 

771  // we only write down single block (4 threads) thread mapping here

772  static constexpr index_t kAMLane = 4;

773  static constexpr index_t kBNLane = 4;

774  static constexpr index_t kABKLane = 1;

775  static constexpr index_t kABKPerLane = 4;

776 

777  static constexpr index_t kCMLane = 1;

778  static constexpr index_t kCNLane = 4;

779  static constexpr index_t kCM0PerLane = 1;

780  static constexpr index_t kCM1PerLane = 4;

781 

782  // c_vec += a_vec * b_vec

783  template <bool post_nop_ = false>

784  CK_TILE_DEVICE void operator()(CVecType& c_vec,

785  const AVecType& a_vec,

786  const BVecType& b_vec,

787  bool_constant<post_nop_> = {}) const

788  {

789  DISPATCH_MFMA_CTRL_("v_mfma_f32_4x4x4bf16_1k", Ctrl)

790  else

791  {

792 #if defined(__gfx90a__) || defined(__gfx94__)

793  c_vec = __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a_vec, b_vec, c_vec, 0, 0, 0);

794 #elif defined(__gfx908__)

795  static_for<0, 2, 1>{}([&](auto k) {

796  c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16(

797  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

798  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

799  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

800  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

801  c_vec,

802  0,

803  0,

804  0);

805  });

806 #else

807  ignore = c_vec;

808  ignore = a_vec;

809  ignore = b_vec;

810 #endif

811  }

812  }

813 

814  // c_vec = a_vec * b_vec

815  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

816  {

817 #if defined(__gfx90a__) || defined(__gfx94__)

818  return bit_cast<CVecType>(

819  __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

820 #elif defined(__gfx908__)

821  CVecType c_vec{0.f};

822  static_for<0, 2, 1>{}([&](auto k) {

823  c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16(

824  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

825  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

826  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

827  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

828  c_vec,

829  0,

830  0,

831  0);

832  });

833  return c_vec;

834 #else

835  ignore = a_vec;

836  ignore = b_vec;

837  return CVecType{0.f};

838 #endif

839  }

840 };

841 

842 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

843 struct WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4

844 {

845  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

846  using ADataType = bf16_t;

847  using BDataType = bf16_t;

848  using CDataType = float;

849 

850  using AVecType = ext_vector_t<bf16_t, 4>;

851  using BVecType = ext_vector_t<bf16_t, 4>;

852  using CVecType = ext_vector_t<float, 4>;

853 

854  static constexpr index_t kM = 64;

855  static constexpr index_t kN = 4;

856  static constexpr index_t kK = 4;

857 

858  static constexpr index_t kAMBlock = 16;

859  static constexpr index_t kBNBlock = 1;

860 

861  // we only write down single block (4 threads) thread mapping here

862  static constexpr index_t kAMLane = 4;

863  static constexpr index_t kBNLane = 4;

864  static constexpr index_t kABKLane = 1;

865  static constexpr index_t kABKPerLane = 4;

866 

867  static constexpr index_t kCMLane = 1;

868  static constexpr index_t kCNLane = 4;

869  static constexpr index_t kCM0PerLane = 1;

870  static constexpr index_t kCM1PerLane = 4;

871 

872  // c_vec += a_vec * b_vec

873  template <bool post_nop_ = false>

874  CK_TILE_DEVICE void operator()(CVecType& c_vec,

875  const AVecType& a_vec,

876  const BVecType& b_vec,

877  bool_constant<post_nop_> = {}) const

878  {

879  DISPATCH_MFMA_CTRL_("v_mfma_f32_4x4x4bf16_1k", Ctrl)

880  else

881  {

882 #if defined(__gfx90a__) || defined(__gfx94__)

883  c_vec = __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a_vec, b_vec, c_vec, 0, 0, 0);

884 #elif defined(__gfx908__)

885  static_for<0, 2, 1>{}([&](auto k) {

886  c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16(

887  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

888  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

889  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

890  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

891  c_vec,

892  0,

893  0,

894  0);

895  });

896 #else

897  ignore = c_vec;

898  ignore = a_vec;

899  ignore = b_vec;

900 #endif

901  }

902  }

903 

904  // c_vec = a_vec * b_vec

905  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

906  {

907 #if defined(__gfx90a__) || defined(__gfx94__)

908  return bit_cast<CVecType>(

909  __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a_vec, b_vec, fp32x4_t{0.f}, 0, 0, 0));

910 #elif defined(__gfx908__)

911  CVecType c_vec{0.f};

912  static_for<0, 2, 1>{}([&](auto k) {

913  c_vec = __builtin_amdgcn_mfma_f32_4x4x2bf16(

914  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

915  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

916  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

917  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

918  c_vec,

919  0,

920  0,

921  0);

922  });

923  return c_vec;

924 #else

925  ignore = a_vec;

926  ignore = b_vec;

927  return CVecType{0.f};

928 #endif

929  }

930 };

931 

932 // gfx950

933 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

934 struct WarpGemmAttributeMfmaImplF16F16F32M32N32K16

935 {

936  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

937  using ADataType = fp16_t;

938  using BDataType = fp16_t;

939  using CDataType = float;

940 

941  using AVecType = ext_vector_t<fp16_t, 8>;

942  using BVecType = ext_vector_t<fp16_t, 8>;

943  using CVecType = ext_vector_t<float, 16>;

944 

945  static constexpr index_t kM = 32;

946  static constexpr index_t kN = 32;

947  static constexpr index_t kK = 16;

948 

949  static constexpr index_t kAMBlock = 1;

950  static constexpr index_t kBNBlock = 1;

951 

952  static constexpr index_t kAMLane = 32;

953  static constexpr index_t kBNLane = 32;

954  static constexpr index_t kABKLane = 2;

955  static constexpr index_t kABKPerLane = 8;

956 

957  static constexpr index_t kCMLane = 2;

958  static constexpr index_t kCNLane = 32;

959  static constexpr index_t kCM0PerLane = 4;

960  static constexpr index_t kCM1PerLane = 4;

961 

962  // c_vec += a_vec * b_vec

963  template <bool post_nop_ = false>

964  CK_TILE_DEVICE void operator()(CVecType& c_vec,

965  const AVecType& a_vec,

966  const BVecType& b_vec,

967  bool_constant<post_nop_> = {}) const

968  {

969  DISPATCH_MFMA_CTRL_("v_mfma_f32_32x32x16_f16", Ctrl)

970  else

971  {

972 #if defined(__gfx950__)

973  c_vec = __builtin_amdgcn_mfma_f32_32x32x16_f16(a_vec, b_vec, c_vec, 0, 0, 0);

974 #elif defined(__gfx90a__) || defined(__gfx94__)

975  static_for<0, 2, 1>{}([&](auto k) {

976  c_vec = __builtin_amdgcn_mfma_f32_32x32x8f16(

977  reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

978  .template get_as<ext_vector_t<fp16_t, 4>>()[number<k>{}],

979  reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

980  .template get_as<ext_vector_t<fp16_t, 4>>()[number<k>{}],

981  c_vec,

982  0,

983  0,

984  0);

985  });

986 #elif defined(__gfx908__)

987  static_for<0, 4, 1>{}([&](auto k) {

988  c_vec = __builtin_amdgcn_mfma_f32_32x32x4f16(

989  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

990  .template get_as<ext_vector_t<fp16_t, 2>>()[number<k>{}],

991  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

992  .template get_as<ext_vector_t<fp16_t, 2>>()[number<k>{}],

993  c_vec,

994  0,

995  0,

996  0);

997  });

998 #else

999  ck_tile::ignore = c_vec;

1000  ck_tile::ignore = a_vec;

1001  ck_tile::ignore = b_vec;

1002 #endif

1003  }

1004  }

1005 

1006  // c_vec = a_vec * b_vec

1007  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1008  {

1009 #if defined(__gfx950__)

1010  return __builtin_amdgcn_mfma_f32_32x32x16_f16(a_vec, b_vec, fp32x16_t{0.f}, 0, 0, 0);

1011 #elif defined(__gfx90a__) || defined(__gfx94__)

1012  CVecType c_vec{0.f};

1013  static_for<0, 2, 1>{}([&](auto k) {

1014  c_vec = __builtin_amdgcn_mfma_f32_32x32x8f16(

1015  reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

1016  .template get_as<ext_vector_t<fp16_t, 4>>()[number<k>{}],

1017  reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

1018  .template get_as<ext_vector_t<fp16_t, 4>>()[number<k>{}],

1019  c_vec,

1020  0,

1021  0,

1022  0);

1023  });

1024  return c_vec;

1025 #elif defined(__gfx908__)

1026  CVecType c_vec{0.f};

1027  static_for<0, 4, 1>{}([&](auto k) {

1028  c_vec = __builtin_amdgcn_mfma_f32_32x32x4f16(

1029  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

1030  .template get_as<ext_vector_t<fp16_t, 2>>()[number<k>{}],

1031  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

1032  .template get_as<ext_vector_t<fp16_t, 2>>()[number<k>{}],

1033  c_vec,

1034  0,

1035  0,

1036  0);

1037  });

1038  return c_vec;

1039 #else

1040  ck_tile::ignore = a_vec;

1041  ck_tile::ignore = b_vec;

1042  return CVecType{0.f};

1043 #endif

1044  }

1045 };

1046 

1047 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1048 struct WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16

1049 {

1050  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1051  using ADataType = bf16_t;

1052  using BDataType = bf16_t;

1053  using CDataType = float;

1054 

1055  using AVecType = ext_vector_t<bf16_t, 8>;

1056  using BVecType = ext_vector_t<bf16_t, 8>;

1057  using CVecType = ext_vector_t<float, 16>;

1058 

1059  static constexpr index_t kM = 32;

1060  static constexpr index_t kN = 32;

1061  static constexpr index_t kK = 16;

1062 

1063  static constexpr index_t kAMBlock = 1;

1064  static constexpr index_t kBNBlock = 1;

1065 

1066  static constexpr index_t kAMLane = 32;

1067  static constexpr index_t kBNLane = 32;

1068  static constexpr index_t kABKLane = 2;

1069  static constexpr index_t kABKPerLane = 8;

1070 

1071  static constexpr index_t kCMLane = 2;

1072  static constexpr index_t kCNLane = 32;

1073  static constexpr index_t kCM0PerLane = 4;

1074  static constexpr index_t kCM1PerLane = 4;

1075 

1076  // c_vec += a_vec * b_vec

1077  template <bool post_nop_ = false>

1078  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1079  const AVecType& a_vec,

1080  const BVecType& b_vec,

1081  bool_constant<post_nop_> = {}) const

1082  {

1083  DISPATCH_MFMA_CTRL_("v_mfma_f32_32x32x16_bf16", Ctrl)

1084  else

1085  {

1086 #if defined(__gfx950__)

1087  c_vec = __builtin_amdgcn_mfma_f32_32x32x16_bf16(a_vec, b_vec, c_vec, 0, 0, 0);

1088 #elif defined(__gfx90a__) || defined(__gfx94__)

1089  static_for<0, 2, 1>{}([&](auto k) {

1090  c_vec = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(

1091  reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

1092  .template get_as<ext_vector_t<bf16_t, 4>>()[number<k>{}],

1093  reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

1094  .template get_as<ext_vector_t<bf16_t, 4>>()[number<k>{}],

1095  c_vec,

1096  0,

1097  0,

1098  0);

1099  });

1100 #elif defined(__gfx908__)

1101  static_for<0, 4, 1>{}([&](auto k) {

1102  c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16(

1103  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

1104  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

1105  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

1106  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

1107  c_vec,

1108  0,

1109  0,

1110  0);

1111  });

1112 #else

1113  ck_tile::ignore = c_vec;

1114  ck_tile::ignore = a_vec;

1115  ck_tile::ignore = b_vec;

1116 #endif

1117  }

1118  }

1119 

1120  // c_vec = a_vec * b_vec

1121  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1122  {

1123 #if defined(__gfx950__)

1124  return __builtin_amdgcn_mfma_f32_32x32x16_bf16(a_vec, b_vec, fp32x16_t{0.f}, 0, 0, 0);

1125 #elif defined(__gfx90a__) || defined(__gfx94__)

1126  CVecType c_vec{0.f};

1127  static_for<0, 2, 1>{}([&](auto k) {

1128  c_vec = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(

1129  reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

1130  .template get_as<ext_vector_t<bf16_t, 4>>()[number<k>{}],

1131  reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

1132  .template get_as<ext_vector_t<bf16_t, 4>>()[number<k>{}],

1133  c_vec,

1134  0,

1135  0,

1136  0);

1137  });

1138  return c_vec;

1139 #elif defined(__gfx908__)

1140  CVecType c_vec{0.f};

1141  static_for<0, 4, 1>{}([&](auto k) {

1142  c_vec = __builtin_amdgcn_mfma_f32_32x32x4bf16(

1143  reinterpret_cast<const thread_buffer<ADataType, 4>&>(a_vec)

1144  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

1145  reinterpret_cast<const thread_buffer<BDataType, 4>&>(b_vec)

1146  .template get_as<ext_vector_t<bf16_t, 2>>()[number<k>{}],

1147  c_vec,

1148  0,

1149  0,

1150  0);

1151  });

1152  return c_vec;

1153 #else

1154  ck_tile::ignore = a_vec;

1155  ck_tile::ignore = b_vec;

1156  return CVecType{0.f};

1157 #endif

1158  }

1159 };

1160 

1161 // FP8

1162 template <typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1163 struct WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base

1164 {

1165  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1166  using ADataType = AType_;

1167  using BDataType = BType_;

1168  using CDataType = float;

1169 

1170  using AVecType = ext_vector_t<ADataType, 8>;

1171  using BVecType = ext_vector_t<BDataType, 8>;

1172  using CVecType = ext_vector_t<CDataType, 4>;

1173 

1174  static constexpr index_t kM = 16;

1175  static constexpr index_t kN = 16;

1176  static constexpr index_t kK = 32;

1177 

1178  static constexpr index_t kAMBlock = 1;

1179  static constexpr index_t kBNBlock = 1;

1180 

1181  static constexpr index_t kAMLane = 16;

1182  static constexpr index_t kBNLane = 16;

1183  static constexpr index_t kABKLane = 4;

1184  static constexpr index_t kABKPerLane = 8;

1185 

1186  static constexpr index_t kCMLane = 4;

1187  static constexpr index_t kCNLane = 16;

1188  static constexpr index_t kCM0PerLane = 1;

1189  static constexpr index_t kCM1PerLane = 4;

1190 

1191  // c_vec += a_vec * b_vec

1192  template <bool post_nop_ = false>

1193  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1194  const AVecType& a_vec,

1195  const BVecType& b_vec,

1196  bool_constant<post_nop_> = {}) const

1197  {

1198  if constexpr(Ctrl == WGAttrCtlEnum::Raw_vvv)

1199  {

1200  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1201  {

1202  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_fp8", "+v", "v", "v", "v")

1203  }

1204  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1205  {

1206  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_bf8", "+v", "v", "v", "v")

1207  }

1208  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1209  {

1210  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_fp8", "+v", "v", "v", "v")

1211  }

1212  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1213  {

1214  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_bf8", "+v", "v", "v", "v")

1215  }

1216  }

1217  else if constexpr(Ctrl == WGAttrCtlEnum::Raw_vaa)

1218  {

1219  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1220  {

1221  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_fp8", "+v", "a", "a", "v")

1222  }

1223  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1224  {

1225  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_bf8", "+v", "a", "a", "v")

1226  }

1227  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1228  {

1229  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_fp8", "+v", "a", "a", "v")

1230  }

1231  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1232  {

1233  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_bf8", "+v", "a", "a", "v")

1234  }

1235  }

1236  else if constexpr(Ctrl == WGAttrCtlEnum::Raw_vav)

1237  {

1238  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1239  {

1240  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_fp8", "+v", "a", "v", "v")

1241  }

1242  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1243  {

1244  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_bf8", "+v", "a", "v", "v")

1245  }

1246  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1247  {

1248  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_fp8", "+v", "a", "v", "v")

1249  }

1250  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1251  {

1252  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_bf8", "+v", "a", "v", "v")

1253  }

1254  }

1255  else if constexpr(Ctrl == WGAttrCtlEnum::Raw_vva)

1256  {

1257  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1258  {

1259  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_fp8", "+v", "v", "a", "v")

1260  }

1261  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1262  {

1263  DISPATCH_MFMA_("mfma_f32_16x16x32_fp8_bf8", "+v", "v", "a", "v")

1264  }

1265  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1266  {

1267  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_fp8", "+v", "v", "a", "v")

1268  }

1269  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1270  {

1271  DISPATCH_MFMA_("mfma_f32_16x16x32_bf8_bf8", "+v", "v", "a", "v")

1272  }

1273  }

1274  else

1275  {

1276 #if defined(__gfx94__) or defined(__gfx95__)

1277  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1278  c_vec = __builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(

1279  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1280  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1281  c_vec = __builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8(

1282  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1283  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1284  c_vec = __builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8(

1285  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1286  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1287  c_vec = __builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8(

1288  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1289 #else

1290  ck_tile::ignore = c_vec;

1291  ck_tile::ignore = a_vec;

1292  ck_tile::ignore = b_vec;

1293 #endif

1294  }

1295  }

1296 

1297  // c_vec = a_vec * b_vec

1298  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1299  {

1300 #if defined(__gfx94__) or defined(__gfx95__)

1301  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1302  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_16x16x32_fp8_fp8(

1303  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1304  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1305  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_16x16x32_fp8_bf8(

1306  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1307  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1308  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_16x16x32_bf8_fp8(

1309  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1310  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1311  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_16x16x32_bf8_bf8(

1312  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1313 #else

1314  ck_tile::ignore = a_vec;

1315  ck_tile::ignore = b_vec;

1316  return CVecType{0.f};

1317 #endif

1318  }

1319 };

1320 

1321 template <typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1322 struct WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base

1323 {

1324  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1325  using ADataType = AType_;

1326  using BDataType = BType_;

1327  using CDataType = float;

1328 

1329  using AVecType = ext_vector_t<ADataType, 8>;

1330  using BVecType = ext_vector_t<BDataType, 8>;

1331  using CVecType = ext_vector_t<CDataType, 16>;

1332 

1333  static constexpr index_t kM = 32;

1334  static constexpr index_t kN = 32;

1335  static constexpr index_t kK = 16;

1336 

1337  static constexpr index_t kAMBlock = 1;

1338  static constexpr index_t kBNBlock = 1;

1339 

1340  static constexpr index_t kAMLane = 32;

1341  static constexpr index_t kBNLane = 32;

1342  static constexpr index_t kABKLane = 2;

1343  static constexpr index_t kABKPerLane = 8;

1344 

1345  static constexpr index_t kCMLane = 2;

1346  static constexpr index_t kCNLane = 32;

1347  static constexpr index_t kCM0PerLane = 4;

1348  static constexpr index_t kCM1PerLane = 4;

1349 

1350  // c_vec += a_vec * b_vec

1351  template <bool post_nop_ = false>

1352  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1353  const AVecType& a_vec,

1354  const BVecType& b_vec,

1355  bool_constant<post_nop_> = {}) const

1356  {

1357  if constexpr(Ctrl == WGAttrCtlEnum::Raw_vvv)

1358  {

1359  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1360  {

1361  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_fp8", "+v", "v", "v", "v")

1362  }

1363  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1364  {

1365  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_bf8", "+v", "v", "v", "v")

1366  }

1367  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1368  {

1369  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_fp8", "+v", "v", "v", "v")

1370  }

1371  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1372  {

1373  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_bf8", "+v", "v", "v", "v")

1374  }

1375  }

1376  else if constexpr(Ctrl == WGAttrCtlEnum::Raw_vaa)

1377  {

1378  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1379  {

1380  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_fp8", "+v", "a", "a", "v")

1381  }

1382  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1383  {

1384  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_bf8", "+v", "a", "a", "v")

1385  }

1386  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1387  {

1388  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_fp8", "+v", "a", "a", "v")

1389  }

1390  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1391  {

1392  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_bf8", "+v", "a", "a", "v")

1393  }

1394  }

1395  else if constexpr(Ctrl == WGAttrCtlEnum::Raw_vav)

1396  {

1397  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1398  {

1399  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_fp8", "+v", "a", "v", "v")

1400  }

1401  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1402  {

1403  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_bf8", "+v", "a", "v", "v")

1404  }

1405  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1406  {

1407  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_fp8", "+v", "a", "v", "v")

1408  }

1409  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1410  {

1411  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_bf8", "+v", "a", "v", "v")

1412  }

1413  }

1414  else if constexpr(Ctrl == WGAttrCtlEnum::Raw_vva)

1415  {

1416  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1417  {

1418  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_fp8", "+v", "v", "a", "v")

1419  }

1420  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1421  {

1422  DISPATCH_MFMA_("mfma_f32_32x32x16_fp8_bf8", "+v", "v", "a", "v")

1423  }

1424  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1425  {

1426  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_fp8", "+v", "v", "a", "v")

1427  }

1428  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1429  {

1430  DISPATCH_MFMA_("mfma_f32_32x32x16_bf8_bf8", "+v", "v", "a", "v")

1431  }

1432  }

1433  else

1434  {

1435 #if defined(__gfx94__) or defined(__gfx95__)

1436  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1437  c_vec = __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8(

1438  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1439  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1440  c_vec = __builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8(

1441  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1442  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1443  c_vec = __builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8(

1444  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1445  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1446  c_vec = __builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8(

1447  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1448 #elif defined(__gfx908__) || defined(__gfx90a__)

1449  static_for<0, 8, 1>{}([&](auto k) {

1450  float a_f32 =

1451  type_convert<float>(reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

1452  .template get_as<ADataType>()[number<k>{}]);

1453  float b_f32 =

1454  type_convert<float>(reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

1455  .template get_as<BDataType>()[number<k>{}]);

1456 

1457  c_vec = __builtin_amdgcn_mfma_f32_32x32x2f32(a_f32, b_f32, c_vec, 0, 0, 0);

1458  });

1459 #else

1460  ck_tile::ignore = c_vec;

1461  ck_tile::ignore = a_vec;

1462  ck_tile::ignore = b_vec;

1463 #endif

1464  }

1465  }

1466 

1467  // c_vec = a_vec * b_vec

1468  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1469  {

1470 #if defined(__gfx94__) or defined(__gfx95__)

1471  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1472  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8(

1473  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1474  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1475  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_32x32x16_fp8_bf8(

1476  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1477  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1478  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_32x32x16_bf8_fp8(

1479  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1480  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1481  return bit_cast<CVecType>(__builtin_amdgcn_mfma_f32_32x32x16_bf8_bf8(

1482  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), CVecType{0.f}, 0, 0, 0));

1483 #elif defined(__gfx908__) || defined(__gfx90a__)

1484  CVecType c_vec{0.f};

1485  static_for<0, 8, 1>{}([&](auto k) {

1486  float a_f32 =

1487  type_convert<float>(reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

1488  .template get_as<ADataType>()[number<k>{}]);

1489  float b_f32 =

1490  type_convert<float>(reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

1491  .template get_as<BDataType>()[number<k>{}]);

1492 

1493  c_vec = __builtin_amdgcn_mfma_f32_32x32x2f32(a_f32, b_f32, c_vec, 0, 0, 0);

1494  });

1495  return c_vec;

1496 #else

1497  ck_tile::ignore = a_vec;

1498  ck_tile::ignore = b_vec;

1499  return CVecType{0.f};

1500 #endif

1501  }

1502 };

1503 

1504 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1505 using WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_fp8 =

1506  WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base<fp8_t, fp8_t, Ctrl_>;

1507 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1508 using WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_fp8 =

1509  WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base<fp8_t, fp8_t, Ctrl_>;

1510 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1511 using WarpGemmAttributeMfmaImpl_f32_32x32x16_fp8_bf8 =

1512  WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base<fp8_t, bf8_t, Ctrl_>;

1513 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1514 using WarpGemmAttributeMfmaImpl_f32_16x16x32_fp8_bf8 =

1515  WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base<fp8_t, bf8_t, Ctrl_>;

1516 

1517 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1518 using WarpGemmAttributeMfmaImpl_f32_16x16x32_bf8_bf8 =

1519  WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base<bf8_t, bf8_t, Ctrl_>;

1520 

1521 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1522 using WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_fp8 =

1523  WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base<bf8_t, fp8_t, Ctrl_>;

1524 

1525 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1526 using WarpGemmAttributeMfmaImpl_f32_32x32x16_bf8_bf8 =

1527  WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base<bf8_t, bf8_t, Ctrl_>;

1528 

1529 template <typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1530 struct WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base

1531 {

1532  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1533  using ADataType = AType_;

1534  using BDataType = BType_;

1535  using CDataType = float;

1536 

1537  using AVecType = ext_vector_t<ADataType, 32>;

1538  using BVecType = ext_vector_t<BDataType, 32>;

1539  using CVecType = ext_vector_t<CDataType, 4>;

1540 

1541  static constexpr index_t kM = 16;

1542  static constexpr index_t kN = 16;

1543  static constexpr index_t kK = 128;

1544 

1545  static constexpr index_t kAMBlock = 1;

1546  static constexpr index_t kBNBlock = 1;

1547 

1548  static constexpr index_t kAMLane = 16;

1549  static constexpr index_t kBNLane = 16;

1550  static constexpr index_t kABKLane = 4;

1551  static constexpr index_t kABKPerLane = 32;

1552 

1553  static constexpr index_t kCMLane = 4;

1554  static constexpr index_t kCNLane = 16;

1555  static constexpr index_t kCM0PerLane = 1;

1556  static constexpr index_t kCM1PerLane = 4;

1557 

1558  // c_vec += a_vec * b_vec

1559  template <bool post_nop_ = false>

1560  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1561  const AVecType& a_vec,

1562  const BVecType& b_vec,

1563  bool_constant<post_nop_> = {}) const

1564  {

1565  //__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, cbsz, blgp, opsel, scale_a,

1566  // opsel, scale_b)

1567 #if defined(__gfx950__)

1568  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1569  c_vec = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1570  a_vec, b_vec, c_vec, 0, 0, 0, 0, 0, 0);

1571  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1572  c_vec = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1573  a_vec, b_vec, c_vec, 0, 1, 0, 0, 0, 0);

1574  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1575  c_vec = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1576  a_vec, b_vec, c_vec, 1, 0, 0, 0, 0, 0);

1577  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1578  c_vec = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1579  a_vec, b_vec, c_vec, 1, 1, 0, 0, 0, 0);

1580 #else

1581  ck_tile::ignore = c_vec;

1582  ck_tile::ignore = a_vec;

1583  ck_tile::ignore = b_vec;

1584 #endif

1585  }

1586 

1587  // c_vec = a_vec * b_vec

1588  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1589  {

1590 #if defined(__gfx950__)

1591  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1592  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1593  a_vec, b_vec, CVecType{0.f}, 0, 0, 0, 0, 0, 0));

1594  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1595  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1596  a_vec, b_vec, CVecType{0.f}, 0, 1, 0, 0, 0, 0));

1597  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1598  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1599  a_vec, b_vec, CVecType{0.f}, 1, 0, 0, 0, 0, 0));

1600  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1601  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1602  a_vec, b_vec, CVecType{0.f}, 1, 1, 0, 0, 0, 0));

1603 #else

1604  ck_tile::ignore = a_vec;

1605  ck_tile::ignore = b_vec;

1606  return CVecType{0.f};

1607 #endif

1608  }

1609 };

1610 

1611 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1612 using WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_fp8 =

1613  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base<fp8_t, fp8_t, Ctrl_>;

1614 

1615 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1616 using WarpGemmAttributeMfmaImpl_f32_16x16x128_fp8_bf8 =

1617  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base<fp8_t, bf8_t, Ctrl_>;

1618 

1619 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1620 using WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_fp8 =

1621  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base<bf8_t, fp8_t, Ctrl_>;

1622 

1623 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1624 using WarpGemmAttributeMfmaImpl_f32_16x16x128_bf8_bf8 =

1625  WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base<bf8_t, bf8_t, Ctrl_>;

1626 

1627 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1628 struct WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4

1629 {

1630  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1631  using ADataType = pk_fp4_t;

1632  using BDataType = pk_fp4_t;

1633  using CDataType = float;

1634 

1635  using AVecType = ext_vector_t<ADataType, 16>;

1636  using BVecType = ext_vector_t<BDataType, 16>;

1637  using CVecType = ext_vector_t<CDataType, 4>;

1638 

1639  static constexpr index_t kM = 16;

1640  static constexpr index_t kN = 16;

1641  static constexpr index_t kK = 128;

1642 

1643  static constexpr index_t kAMBlock = 1;

1644  static constexpr index_t kBNBlock = 1;

1645 

1646  static constexpr index_t kAMLane = 16;

1647  static constexpr index_t kBNLane = 16;

1648  static constexpr index_t kABKLane = 4;

1649  static constexpr index_t kABKPerLane = 32;

1650 

1651  static constexpr index_t kCMLane = 4;

1652  static constexpr index_t kCNLane = 16;

1653  static constexpr index_t kCM0PerLane = 1;

1654  static constexpr index_t kCM1PerLane = 4;

1655 

1656  // c_vec += a_vec * b_vec

1657  template <index_t opselA, index_t opselB, bool post_nop_ = false>

1658  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1659  const AVecType& a_vec,

1660  const int32_t& a_scale,

1661  const BVecType& b_vec,

1662  const int32_t& b_scale,

1663  bool_constant<post_nop_> = {}) const

1664  {

1665  //__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(a, b, c, cbsz, blgp, opsel, scale_a,

1666  // opsel, scale_b)

1667 #if defined(__gfx950__)

1668  auto arg_a = bit_cast<int32x4_t>(a_vec);

1669  auto arg_b = bit_cast<int32x4_t>(b_vec);

1670  c_vec = __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1671  int32x8_t{arg_a[0], arg_a[1], arg_a[2], arg_a[3], 0, 0, 0, 0},

1672  int32x8_t{arg_b[0], arg_b[1], arg_b[2], arg_b[3], 0, 0, 0, 0},

1673  c_vec,

1674  4,

1675  4,

1676  opselA,

1677  a_scale,

1678  opselB,

1679  b_scale);

1680 #else

1681  ck_tile::ignore = c_vec;

1682  ck_tile::ignore = a_vec;

1683  ck_tile::ignore = b_vec;

1684  ck_tile::ignore = a_scale;

1685  ck_tile::ignore = b_scale;

1686 #endif

1687  }

1688 

1689  // c_vec = a_vec * b_vec

1690  template <index_t opselA, index_t opselB>

1691  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec,

1692  const int32_t& a_scale,

1693  const BVecType& b_vec,

1694  const int32_t& b_scale) const

1695  {

1696 #if defined(__gfx950__)

1697  auto arg_a = bit_cast<int32x4_t>(a_vec);

1698  auto arg_b = bit_cast<int32x4_t>(b_vec);

1699  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(

1700  int32x8_t{arg_a[0], arg_a[1], arg_a[2], arg_a[3], 0, 0, 0, 0},

1701  int32x8_t{arg_b[0], arg_b[1], arg_b[2], arg_b[3], 0, 0, 0, 0},

1702  CVecType{0.f},

1703  4,

1704  4,

1705  opselA,

1706  a_scale,

1707  opselB,

1708  b_scale));

1709 #else

1710  ck_tile::ignore = a_vec;

1711  ck_tile::ignore = b_vec;

1712  ck_tile::ignore = a_scale;

1713  ck_tile::ignore = b_scale;

1714  return CVecType{0.f};

1715 #endif

1716  }

1717 };

1718 

1719 template <typename AType_, typename BType_, WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1720 struct WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base

1721 {

1722  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1723  using ADataType = AType_;

1724  using BDataType = BType_;

1725  using CDataType = float;

1726 

1727  using AVecType = ext_vector_t<ADataType, 32>;

1728  using BVecType = ext_vector_t<BDataType, 32>;

1729  using CVecType = ext_vector_t<CDataType, 16>;

1730 

1731  static constexpr index_t kM = 32;

1732  static constexpr index_t kN = 32;

1733  static constexpr index_t kK = 64;

1734 

1735  static constexpr index_t kAMBlock = 1;

1736  static constexpr index_t kBNBlock = 1;

1737 

1738  static constexpr index_t kAMLane = 32;

1739  static constexpr index_t kBNLane = 32;

1740  static constexpr index_t kABKLane = 2;

1741  static constexpr index_t kABKPerLane = 32;

1742 

1743  static constexpr index_t kCMLane = 2;

1744  static constexpr index_t kCNLane = 32;

1745  static constexpr index_t kCM0PerLane = 4;

1746  static constexpr index_t kCM1PerLane = 4;

1747 

1748  // c_vec += a_vec * b_vec

1749  template <bool post_nop_ = false>

1750  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1751  const AVecType& a_vec,

1752  const BVecType& b_vec,

1753  bool_constant<post_nop_> = {}) const

1754  {

1755  //__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, cbsz, blgp, opsel, scale_a,

1756  // opsel, scale_b)

1757 #if defined(__gfx950__)

1758  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1759  c_vec = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1760  a_vec, b_vec, c_vec, 0, 0, 0, 0, 0, 0);

1761  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1762  c_vec = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1763  a_vec, b_vec, c_vec, 0, 1, 0, 0, 0, 0);

1764  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1765  c_vec = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1766  a_vec, b_vec, c_vec, 1, 0, 0, 0, 0, 0);

1767  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1768  c_vec = __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1769  a_vec, b_vec, c_vec, 1, 1, 0, 0, 0, 0);

1770 #else

1771  ck_tile::ignore = c_vec;

1772  ck_tile::ignore = a_vec;

1773  ck_tile::ignore = b_vec;

1774 #endif

1775  }

1776 

1777  // c_vec = a_vec * b_vec

1778  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1779  {

1780 #if defined(__gfx950__)

1781  if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, fp8_t>)

1782  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1783  a_vec, b_vec, CVecType{0.f}, 0, 0, 0, 0, 0, 0));

1784  else if constexpr(std::is_same_v<ADataType, fp8_t> && std::is_same_v<BDataType, bf8_t>)

1785  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1786  a_vec, b_vec, CVecType{0.f}, 0, 1, 0, 0, 0, 0));

1787  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, fp8_t>)

1788  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1789  a_vec, b_vec, CVecType{0.f}, 1, 0, 0, 0, 0, 0));

1790  else if constexpr(std::is_same_v<ADataType, bf8_t> && std::is_same_v<BDataType, bf8_t>)

1791  return bit_cast<CVecType>(__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(

1792  a_vec, b_vec, CVecType{0.f}, 1, 1, 0, 0, 0, 0));

1793 #else

1794  ck_tile::ignore = a_vec;

1795  ck_tile::ignore = b_vec;

1796  return CVecType{0.f};

1797 #endif

1798  }

1799 };

1800 

1801 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1802 using WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_fp8 =

1803  WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base<fp8_t, fp8_t, Ctrl_>;

1804 

1805 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1806 using WarpGemmAttributeMfmaImpl_f32_32x32x64_fp8_bf8 =

1807  WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base<fp8_t, bf8_t, Ctrl_>;

1808 

1809 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1810 using WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_fp8 =

1811  WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base<bf8_t, fp8_t, Ctrl_>;

1812 

1813 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1814 using WarpGemmAttributeMfmaImpl_f32_32x32x64_bf8_bf8 =

1815  WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base<bf8_t, bf8_t, Ctrl_>;

1816 

1817 // int8

1818 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1819 struct WarpGemmAttributeMfmaImpl_i32_32x32x16_i8

1820 {

1821  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1822  using ADataType = int8_t;

1823  using BDataType = int8_t;

1824  using CDataType = int32_t;

1825 

1826  using AVecType = ext_vector_t<ADataType, 8>;

1827  using BVecType = ext_vector_t<BDataType, 8>;

1828  using CVecType = ext_vector_t<CDataType, 16>;

1829 

1830  static constexpr index_t kM = 32;

1831  static constexpr index_t kN = 32;

1832  static constexpr index_t kK = 16;

1833 

1834  static constexpr index_t kAMBlock = 1;

1835  static constexpr index_t kBNBlock = 1;

1836 

1837  static constexpr index_t kAMLane = 32;

1838  static constexpr index_t kBNLane = 32;

1839  static constexpr index_t kABKLane = 2;

1840  static constexpr index_t kABKPerLane = 8;

1841 

1842  static constexpr index_t kCMLane = 2;

1843  static constexpr index_t kCNLane = 32;

1844  static constexpr index_t kCM0PerLane = 4;

1845  static constexpr index_t kCM1PerLane = 4;

1846 

1847  // c_vec += a_vec * b_vec

1848  template <bool post_nop_ = false>

1849  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1850  const AVecType& a_vec,

1851  const BVecType& b_vec,

1852  bool_constant<post_nop_> = {}) const

1853  {

1854  DISPATCH_MFMA_CTRL_("v_mfma_i32_32x32x16_i8", Ctrl)

1855  else

1856  {

1857 #if defined(__gfx94__) or defined(__gfx95__)

1858  c_vec = __builtin_amdgcn_mfma_i32_32x32x16_i8(

1859  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1860 #elif defined(__gfx908__) || defined(__gfx90a__)

1861  static_for<0, 8, 1>{}([&](auto k) {

1862  float a_f32 =

1863  type_convert<float>(reinterpret_cast<const thread_buffer<ADataType, 8>&>(a_vec)

1864  .template get_as<ADataType>()[number<k>{}]);

1865  float b_f32 =

1866  type_convert<float>(reinterpret_cast<const thread_buffer<BDataType, 8>&>(b_vec)

1867  .template get_as<BDataType>()[number<k>{}]);

1868 

1869  c_vec = __builtin_amdgcn_mfma_f32_32x32x2f32(a_f32, b_f32, c_vec, 0, 0, 0);

1870  });

1871 #else

1872  ck_tile::ignore = c_vec;

1873  ck_tile::ignore = a_vec;

1874  ck_tile::ignore = b_vec;

1875 #endif

1876  }

1877  }

1878 

1879  // c_vec = a_vec * b_vec

1880  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1881  {

1882  CVecType c_vec{0};

1883  operator()(c_vec, a_vec, b_vec);

1884  return c_vec;

1885  }

1886 };

1887 

1888 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1889 struct WarpGemmAttributeMfmaImpl_i32_16x16x32_i8

1890 {

1891  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1892  using ADataType = int8_t;

1893  using BDataType = int8_t;

1894  using CDataType = int32_t;

1895 

1896  using AVecType = ext_vector_t<ADataType, 8>;

1897  using BVecType = ext_vector_t<BDataType, 8>;

1898  using CVecType = ext_vector_t<CDataType, 4>;

1899 

1900  static constexpr index_t kM = 16;

1901  static constexpr index_t kN = 16;

1902  static constexpr index_t kK = 32;

1903 

1904  static constexpr index_t kAMBlock = 1;

1905  static constexpr index_t kBNBlock = 1;

1906 

1907  static constexpr index_t kAMLane = 16;

1908  static constexpr index_t kBNLane = 16;

1909  static constexpr index_t kABKLane = 4;

1910  static constexpr index_t kABKPerLane = 8;

1911 

1912  static constexpr index_t kCMLane = 4;

1913  static constexpr index_t kCNLane = 16;

1914  static constexpr index_t kCM0PerLane = 1;

1915  static constexpr index_t kCM1PerLane = 4; // write to 4x AccVGPRs

1916 

1917  // c_vec += a_vec * b_vec

1918  template <bool post_nop_ = false>

1919  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1920  const AVecType& a_vec,

1921  const BVecType& b_vec,

1922  bool_constant<post_nop_> = {}) const

1923  {

1924  DISPATCH_MFMA_CTRL_("v_mfma_i32_16x16x32_i8", Ctrl)

1925  else

1926  {

1927 #if defined(__gfx94__) or defined(__gfx95__)

1928  c_vec = __builtin_amdgcn_mfma_i32_16x16x32_i8(

1929  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1930 #else

1931  ck_tile::ignore = c_vec;

1932  ck_tile::ignore = a_vec;

1933  ck_tile::ignore = b_vec;

1934 #endif

1935  }

1936  }

1937 

1938  // c_vec = a_vec * b_vec

1939  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1940  {

1941  CVecType c_vec{0};

1942  operator()(c_vec, a_vec, b_vec);

1943  return c_vec;

1944  }

1945 };

1946 

1947 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

1948 struct WarpGemmAttributeMfmaImpl_i32_16x16x64_i8

1949 {

1950  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

1951  using ADataType = int8_t;

1952  using BDataType = int8_t;

1953  using CDataType = int32_t;

1954 

1955  using AVecType = ext_vector_t<ADataType, 16>;

1956  using BVecType = ext_vector_t<BDataType, 16>;

1957  using CVecType = ext_vector_t<CDataType, 4>;

1958 

1959  static constexpr index_t kM = 16;

1960  static constexpr index_t kN = 16;

1961  static constexpr index_t kK = 64;

1962 

1963  static constexpr index_t kAMBlock = 1;

1964  static constexpr index_t kBNBlock = 1;

1965 

1966  static constexpr index_t kAMLane = 16;

1967  static constexpr index_t kBNLane = 16;

1968  static constexpr index_t kABKLane = 4;

1969  static constexpr index_t kABKPerLane = 16;

1970 

1971  static constexpr index_t kCMLane = 4;

1972  static constexpr index_t kCNLane = 16;

1973  static constexpr index_t kCM0PerLane = 1;

1974  static constexpr index_t kCM1PerLane = 4; // write to 4x AccVGPRs

1975 

1976  // c_vec += a_vec * b_vec

1977  template <bool post_nop_ = false>

1978  CK_TILE_DEVICE void operator()(CVecType& c_vec,

1979  const AVecType& a_vec,

1980  const BVecType& b_vec,

1981  bool_constant<post_nop_> = {}) const

1982  {

1983  DISPATCH_MFMA_CTRL_("v_mfma_i32_16x16x64_i8", Ctrl)

1984  else

1985  {

1986 #if defined(__gfx95__)

1987  c_vec = __builtin_amdgcn_mfma_i32_16x16x64_i8(

1988  bit_cast<int64_t>(a_vec), bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

1989 #else

1990  ck_tile::ignore = c_vec;

1991  ck_tile::ignore = a_vec;

1992  ck_tile::ignore = b_vec;

1993 #endif

1994  }

1995  }

1996 

1997  // c_vec = a_vec * b_vec

1998  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

1999  {

2000  CVecType c_vec{0};

2001  operator()(c_vec, a_vec, b_vec);

2002  return c_vec;

2003  }

2004 };

2005 

2006 template <WGAttrCtlEnum Ctrl_ = WGAttrCtlEnum::Default_>

2007 struct WarpGemmAttributeMfmaImpl_i32_32x32x32_i8

2008 {

2009  static constexpr WGAttrCtlEnum Ctrl = Ctrl_;

2010  using ADataType = int8_t;

2011  using BDataType = int8_t;

2012  using CDataType = int32_t;

2013 

2014  using AVecType = ext_vector_t<ADataType, 16>;

2015  using BVecType = ext_vector_t<BDataType, 16>;

2016  using CVecType = ext_vector_t<CDataType, 16>;

2017 

2018  static constexpr index_t kM = 32;

2019  static constexpr index_t kN = 32;

2020  static constexpr index_t kK = 32;

2021 

2022  static constexpr index_t kAMBlock = 1;

2023  static constexpr index_t kBNBlock = 1;

2024 

2025  static constexpr index_t kAMLane = 32;

2026  static constexpr index_t kBNLane = 32;

2027  static constexpr index_t kABKLane = 2;

2028  static constexpr index_t kABKPerLane = 16;

2029 

2030  static constexpr index_t kCMLane = 2;

2031  static constexpr index_t kCNLane = 32;

2032  static constexpr index_t kCM0PerLane = 4;

2033  static constexpr index_t kCM1PerLane = 4;

2034 

2035  // c_vec += a_vec * b_vec

2036  template <bool post_nop_ = false>

2037  CK_TILE_DEVICE void operator()(CVecType& c_vec,

2038  const AVecType& a_vec,

2039  const BVecType& b_vec,

2040  bool_constant<post_nop_> = {}) const

2041  {

2042  DISPATCH_MFMA_CTRL_("v_mfma_i32_32x32x32_i8", Ctrl)

2043  else

2044  {

2045 #if defined(__gfx95__)

2046  c_vec = __builtin_amdgcn_mfma_i32_32x32x32_i8(

2047  a_vec, bit_cast<int64_t>(b_vec), c_vec, 0, 0, 0);

2048 #else

2049  ck_tile::ignore = c_vec;

2050  ck_tile::ignore = a_vec;

2051  ck_tile::ignore = b_vec;

2052 #endif

2053  }

2054  }

2055 

2056  // c_vec = a_vec * b_vec

2057  CK_TILE_DEVICE CVecType operator()(const AVecType& a_vec, const BVecType& b_vec) const

2058  {

2059  CVecType c_vec{0};

2060  operator()(c_vec, a_vec, b_vec);

2061  return c_vec;

2062  }

2063 };

2064 

2065 #undef DISPATCH_MFMA_

2066 

2067 } // namespace ck_tile

CK_TILE_DEVICE

#define CK_TILE_DEVICE

Definition: config.hpp:41

core.hpp

ck_tile

Definition: cluster_descriptor.hpp:13

ck_tile::WGAttrCtlEnum

WGAttrCtlEnum

Definition: warp_gemm_attribute_mfma_impl.hpp:15

ck_tile::WGAttrCtlEnum::Raw_vav

@ Raw_vav

ck_tile::WGAttrCtlEnum::Raw_avv

@ Raw_avv

ck_tile::WGAttrCtlEnum::Raw_vva

@ Raw_vva

ck_tile::WGAttrCtlEnum::Default_

@ Default_

ck_tile::WGAttrCtlEnum::Raw_vvv

@ Raw_vvv

ck_tile::WGAttrCtlEnum::Raw_vaa

@ Raw_vaa

ck_tile::int32x8_t

int32_t int32x8_t

Definition: vector_type.hpp:156

ck_tile::fp16_t

_Float16 fp16_t

Definition: half.hpp:110

ck_tile::int8_t

int8_t int8_t

Definition: int8.hpp:20

ck_tile::bf16_t

bfloat16_t bf16_t

Definition: bfloat16.hpp:113

ck_tile::index_t

int32_t index_t

Definition: integer.hpp:9

ck_tile::pk_fp4_t

pk_float4_e2m1_t pk_fp4_t

Definition: pk_fp4.hpp:151

ck_tile::ignore

constexpr detail::ignore_t ignore

Definition: ignore.hpp:20

ck_tile::ext_vector_t

typename impl::ext_vector< T, N >::type ext_vector_t

Definition: vector_type.hpp:84

ck_tile::int32_t

int32_t int32_t

Definition: integer.hpp:10

ck_tile::fp32x16_t

float fp32x16_t

Definition: vector_type.hpp:130

ck_tile::fp32x4_t

float fp32x4_t

Definition: vector_type.hpp:128

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1531

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1541

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1555

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1546

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1532

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::CVecType

ext_vector_t< CDataType, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1539

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1550

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::AVecType

ext_vector_t< ADataType, 32 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1537

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1553

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1535

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1548

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::BDataType

BType_ BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1534

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1556

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1543

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1551

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1545

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1554

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1549

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1542

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1560

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1588

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::BVecType

ext_vector_t< BDataType, 32 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1538

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x128_f8_bf8_base::ADataType

AType_ ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1533

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1164

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1178

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1179

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1188

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1174

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::ADataType

AType_ ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1166

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::CVecType

ext_vector_t< CDataType, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1172

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::AVecType

ext_vector_t< ADataType, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1170

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1165

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1186

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1183

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1193

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1189

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1168

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1298

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1182

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::BDataType

BType_ BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1167

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1187

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1184

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1176

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1181

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1175

ck_tile::WarpGemmAttributeMfmaImpl_f32_16x16x32_f8_base::BVecType

ext_vector_t< BDataType, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1171

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1323

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1345

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1324

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::AVecType

ext_vector_t< ADataType, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1329

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::BDataType

BType_ BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1326

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1348

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::CVecType

ext_vector_t< CDataType, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1331

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1343

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1327

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1337

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1346

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1352

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1468

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1340

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1333

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::ADataType

AType_ ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1325

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::BVecType

ext_vector_t< BDataType, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1330

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1335

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1334

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1347

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1338

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1342

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x16_f8_base::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1341

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base

Definition: warp_gemm_attribute_mfma_impl.hpp:1721

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1778

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1732

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1743

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::BDataType

BType_ BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1724

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1741

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1731

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1739

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1738

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::CVecType

ext_vector_t< CDataType, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1729

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1725

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1750

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1745

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1736

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1735

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1744

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::BVecType

ext_vector_t< BDataType, 32 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1728

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1722

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::ADataType

AType_ ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1723

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1740

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1746

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::AVecType

ext_vector_t< ADataType, 32 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1727

ck_tile::WarpGemmAttributeMfmaImpl_f32_32x32x64_f8_bf8_base::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1733

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1890

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1939

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::CVecType

ext_vector_t< CDataType, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1898

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::BVecType

ext_vector_t< BDataType, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1897

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::AVecType

ext_vector_t< ADataType, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1896

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1907

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1919

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1891

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1901

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1905

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::ADataType

int8_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1892

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1908

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1910

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1909

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1904

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1912

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1914

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1900

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1915

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1913

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::CDataType

int32_t CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1894

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1902

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x32_i8::BDataType

int8_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1893

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1949

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::CDataType

int32_t CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1953

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1974

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1961

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1960

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1967

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::AVecType

ext_vector_t< ADataType, 16 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1955

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1969

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1973

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1959

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1966

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1978

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1972

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1964

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1968

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1950

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::BDataType

int8_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1952

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::BVecType

ext_vector_t< BDataType, 16 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1956

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1998

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::ADataType

int8_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1951

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1971

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::CVecType

ext_vector_t< CDataType, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1957

ck_tile::WarpGemmAttributeMfmaImpl_i32_16x16x64_i8::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1963

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:1820

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::ADataType

int8_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1822

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1849

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1842

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::AVecType

ext_vector_t< ADataType, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1826

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::BVecType

ext_vector_t< BDataType, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1827

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1844

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::CDataType

int32_t CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1824

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1821

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1880

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::BDataType

int8_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1823

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1837

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1831

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1835

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1832

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1845

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::CVecType

ext_vector_t< CDataType, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1828

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1843

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1840

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1834

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1838

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1839

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x16_i8::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1830

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8

Definition: warp_gemm_attribute_mfma_impl.hpp:2008

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:2018

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::ADataType

int8_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:2010

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2032

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2027

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:2037

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::BDataType

int8_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:2011

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2031

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::CDataType

int32_t CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:2012

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2026

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::BVecType

ext_vector_t< BDataType, 16 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:2015

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:2023

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2025

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2030

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:2020

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::AVecType

ext_vector_t< ADataType, 16 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:2014

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::CVecType

ext_vector_t< CDataType, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:2016

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:2057

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:2019

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2033

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:2028

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:2009

ck_tile::WarpGemmAttributeMfmaImpl_i32_32x32x32_i8::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:2022

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16

Definition: warp_gemm_attribute_mfma_impl.hpp:666

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:674

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:684

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::BVecType

ext_vector_t< bf16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:673

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:695

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:688

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:725

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:678

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:683

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:667

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:689

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:690

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:685

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:680

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:670

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:691

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:677

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::BDataType

bf16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:669

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:681

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::ADataType

bf16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:668

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::AVecType

ext_vector_t< bf16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:672

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:686

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K16::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:676

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32

Definition: warp_gemm_attribute_mfma_impl.hpp:196

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:219

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::BDataType

bf16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:199

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:204

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:207

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:210

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:220

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:215

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:206

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:208

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::AVecType

ext_vector_t< bf16_t, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:202

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:197

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:216

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:200

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:214

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::BVecType

ext_vector_t< bf16_t, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:203

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:218

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::ADataType

bf16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:198

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:225

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:221

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:244

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:211

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M16N16K32::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:213

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16

Definition: warp_gemm_attribute_mfma_impl.hpp:1049

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1072

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1078

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1053

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1050

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::BDataType

bf16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1052

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1063

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1064

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1067

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1074

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1121

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1068

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1073

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1066

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1059

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::AVecType

ext_vector_t< bf16_t, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1055

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1060

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1071

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::BVecType

ext_vector_t< bf16_t, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1056

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::ADataType

bf16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1051

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1069

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::CVecType

ext_vector_t< float, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1057

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K16::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1061

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8

Definition: warp_gemm_attribute_mfma_impl.hpp:577

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::BVecType

ext_vector_t< bf16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:584

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::BDataType

bf16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:580

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:581

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:592

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::CVecType

ext_vector_t< float, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:585

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::AVecType

ext_vector_t< bf16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:583

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:637

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:588

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:601

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:587

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:600

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:591

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:596

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:595

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:594

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:589

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:599

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::ADataType

bf16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:579

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:578

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:606

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:602

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M32N32K8::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:597

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4

Definition: warp_gemm_attribute_mfma_impl.hpp:754

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:775

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::ADataType

bf16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:756

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:778

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:777

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:779

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:815

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::BDataType

bf16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:757

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:766

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:755

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:773

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:784

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::AVecType

ext_vector_t< bf16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:760

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:765

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:764

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:774

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:769

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::BVecType

ext_vector_t< bf16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:761

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:780

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:762

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:772

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:768

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M4N64K4::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:758

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4

Definition: warp_gemm_attribute_mfma_impl.hpp:844

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::BDataType

bf16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:847

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:867

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:862

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:905

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:845

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:868

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:855

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::ADataType

bf16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:846

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:858

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:869

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:859

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:863

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:856

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:848

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:864

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:870

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:865

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::AVecType

ext_vector_t< bf16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:850

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:874

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::BVecType

ext_vector_t< bf16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:851

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:854

ck_tile::WarpGemmAttributeMfmaImplBf16Bf16F32M64N4K4::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:852

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16

Definition: warp_gemm_attribute_mfma_impl.hpp:322

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:323

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:339

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:326

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:345

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:336

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::BDataType

fp16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:325

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:344

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:347

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::AVecType

ext_vector_t< fp16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:328

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:337

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:333

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::ADataType

fp16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:324

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:341

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:340

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:330

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:346

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:342

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:370

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:351

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:332

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:334

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K16::BVecType

ext_vector_t< fp16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:329

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32

Definition: warp_gemm_attribute_mfma_impl.hpp:385

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:396

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:408

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:410

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::ADataType

fp16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:387

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:407

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::AVecType

ext_vector_t< fp16_t, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:391

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:404

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:405

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::BVecType

ext_vector_t< fp16_t, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:392

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:395

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:393

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:433

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::BDataType

fp16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:388

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:389

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:402

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:409

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:400

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:399

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:386

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:403

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:397

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M16N16K32::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:414

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16

Definition: warp_gemm_attribute_mfma_impl.hpp:935

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:936

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::ADataType

fp16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:937

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:949

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:957

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::AVecType

ext_vector_t< fp16_t, 8 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:941

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:958

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::BVecType

ext_vector_t< fp16_t, 8 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:942

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:945

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:939

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:964

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:946

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:955

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1007

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:954

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:953

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:960

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:950

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:947

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:952

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:959

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::CVecType

ext_vector_t< float, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:943

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K16::BDataType

fp16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:938

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8

Definition: warp_gemm_attribute_mfma_impl.hpp:259

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:282

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:276

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:283

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:281

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:288

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:284

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:263

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:271

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::BDataType

fp16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:262

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:277

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::CVecType

ext_vector_t< float, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:267

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:274

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:260

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:269

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::BVecType

ext_vector_t< fp16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:266

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:270

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:307

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:273

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::AVecType

ext_vector_t< fp16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:265

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::ADataType

fp16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:261

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:278

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M32N32K8::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:279

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4

Definition: warp_gemm_attribute_mfma_impl.hpp:448

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:459

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:471

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::ADataType

fp16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:450

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::AVecType

ext_vector_t< fp16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:454

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:497

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::BVecType

ext_vector_t< fp16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:455

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:463

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:452

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:478

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:456

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:466

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:467

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:472

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::BDataType

fp16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:451

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:460

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:468

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:458

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:473

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:449

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:469

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:474

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M4N64K4::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:462

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4

Definition: warp_gemm_attribute_mfma_impl.hpp:512

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:533

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:516

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:522

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:530

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::BDataType

fp16_t BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:515

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:513

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:527

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:542

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:524

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:526

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::CVecType

ext_vector_t< float, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:520

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:523

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:561

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:535

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::AVecType

ext_vector_t< fp16_t, 4 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:518

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:531

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:536

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:537

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:532

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::ADataType

fp16_t ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:514

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::BVecType

ext_vector_t< fp16_t, 4 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:519

ck_tile::WarpGemmAttributeMfmaImplF16F16F32M64N4K4::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:538

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4

Definition: warp_gemm_attribute_mfma_impl.hpp:67

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:68

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:87

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::AVecType

ext_vector_t< ADataType, 1 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:74

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:80

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:78

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:97

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:90

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:93

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:85

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::ADataType

float ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:70

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::CVecType

ext_vector_t< CDataType, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:76

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:83

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:82

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::BDataType

float BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:71

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:92

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::BVecType

ext_vector_t< BDataType, 1 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:75

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:88

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:72

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:116

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:91

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:79

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M16N16K4::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:86

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2

Definition: warp_gemm_attribute_mfma_impl.hpp:131

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:156

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::BDataType

float BDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:135

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:142

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:157

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:155

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:132

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const BVecType &b_vec, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:161

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:136

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:149

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::CVecType

ext_vector_t< CDataType, 16 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:140

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:154

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:150

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::AVecType

ext_vector_t< ADataType, 1 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:138

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const BVecType &b_vec) const

Definition: warp_gemm_attribute_mfma_impl.hpp:180

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:152

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::BVecType

ext_vector_t< BDataType, 1 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:139

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:143

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:146

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:151

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:147

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:144

ck_tile::WarpGemmAttributeMfmaImplF32F32F32M32N32K2::ADataType

float ADataType

Definition: warp_gemm_attribute_mfma_impl.hpp:134

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4

Definition: warp_gemm_attribute_mfma_impl.hpp:1629

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kABKLane

static constexpr index_t kABKLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1648

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kK

static constexpr index_t kK

Definition: warp_gemm_attribute_mfma_impl.hpp:1641

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::operator()

CK_TILE_DEVICE CVecType operator()(const AVecType &a_vec, const int32_t &a_scale, const BVecType &b_vec, const int32_t &b_scale) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1691

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::BVecType

ext_vector_t< BDataType, 16 > BVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1636

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::CDataType

float CDataType

Definition: warp_gemm_attribute_mfma_impl.hpp:1633

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::CVecType

ext_vector_t< CDataType, 4 > CVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1637

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kN

static constexpr index_t kN

Definition: warp_gemm_attribute_mfma_impl.hpp:1640

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kABKPerLane

static constexpr index_t kABKPerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1649

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kM

static constexpr index_t kM

Definition: warp_gemm_attribute_mfma_impl.hpp:1639

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kCNLane

static constexpr index_t kCNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1652

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kCMLane

static constexpr index_t kCMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1651

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kBNLane

static constexpr index_t kBNLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1647

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::operator()

CK_TILE_DEVICE void operator()(CVecType &c_vec, const AVecType &a_vec, const int32_t &a_scale, const BVecType &b_vec, const int32_t &b_scale, bool_constant< post_nop_ >={}) const

Definition: warp_gemm_attribute_mfma_impl.hpp:1658

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::AVecType

ext_vector_t< ADataType, 16 > AVecType

Definition: warp_gemm_attribute_mfma_impl.hpp:1635

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kCM0PerLane

static constexpr index_t kCM0PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1653

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kBNBlock

static constexpr index_t kBNBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1644

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kAMLane

static constexpr index_t kAMLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1646

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::Ctrl

static constexpr WGAttrCtlEnum Ctrl

Definition: warp_gemm_attribute_mfma_impl.hpp:1630

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kAMBlock

static constexpr index_t kAMBlock

Definition: warp_gemm_attribute_mfma_impl.hpp:1643

ck_tile::WarpGemmAttributeMfmaScaleImpl_f32_16x16x128_fp4::kCM1PerLane

static constexpr index_t kCM1PerLane

Definition: warp_gemm_attribute_mfma_impl.hpp:1654

ck_tile::constant

Definition: integral_constant.hpp:13

ck_tile::pk_float4_e2m1_t

Definition: pk_fp4.hpp:76

ck_tile::static_for

Definition: functional.hpp:43

ck_tile::thread_buffer

Definition: debug.hpp:67

DISPATCH_MFMA_

#define DISPATCH_MFMA_(mfma_, dmod_, amod_, bmod_, cmod_)

Definition: warp_gemm_attribute_mfma_impl.hpp:25

DISPATCH_MFMA_CTRL_

#define DISPATCH_MFMA_CTRL_(mfma_, ctrl_)

Definition: warp_gemm_attribute_mfma_impl.hpp:42