Skip to content

GitLab

Commit f24bb5a0 authored by jefyang1's avatar jefyang1
Browse files

Fix clang format

parent e95cb82b
Hide whitespace changes
Inline Side-by-side
Showing with 15 additions and 11 deletions
include/ck/tensor_operation/gpu/grid/gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp
View file @ f24bb5a0
...@@ -866,8 +866,8 @@ struct GridwiseBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle ...@@ -866,8 +866,8 @@ struct GridwiseBatchedGemmMultipleDGemmMultipleD_Xdl_CShuffle
// cause mismatch in summation index for example c[0:7] = a1[[0:3, 8:11]] * b1[0:7]. // cause mismatch in summation index for example c[0:7] = a1[[0:3, 8:11]] * b1[0:7].
// therefore we may just as well assign Gemm1KPack = group_size // therefore we may just as well assign Gemm1KPack = group_size
constexpr index_t Gemm1KPack = constexpr index_t Gemm1KPack =
MfmaSelector<A0B0B1DataType, Gemm0MPerXdl, Gemm0NPerXdl>::selected_mfma.group_size; MfmaSelector<A0B0B1DataType, Gemm0MPerXdl, Gemm0NPerXdl>::selected_mfma.group_size;
auto blockwise_gemm1 = BlockwiseGemmXdlops_v2< auto blockwise_gemm1 = BlockwiseGemmXdlops_v2<
BlockSize, BlockSize,
......
include/ck/tensor_operation/gpu/warp/xdlops_gemm.hpp
View file @ f24bb5a0
...@@ -881,14 +881,14 @@ template <typename base_type, ...@@ -881,14 +881,14 @@ template <typename base_type,
index_t MPerXdlops, index_t MPerXdlops,
index_t NPerXdlops, index_t NPerXdlops,
typename additional_type = base_type, typename additional_type = base_type,
bool is_single_rate_mfma = false> bool is_single_rate_mfma = false>
struct MfmaSelector struct MfmaSelector
{ {
template <typename base_type_, template <typename base_type_,
index_t MPerXdlops_, index_t MPerXdlops_,
index_t NPerXdlops_, index_t NPerXdlops_,
typename additional_type_ = base_type_, typename additional_type_ = base_type_,
bool is_single_rate_mfma_ = false> bool is_single_rate_mfma_ = false>
static constexpr auto GetMfma(); static constexpr auto GetMfma();
template <> template <>
...@@ -1127,8 +1127,8 @@ struct MfmaSelector ...@@ -1127,8 +1127,8 @@ struct MfmaSelector
return MfmaInstr::mfma_f32_16x16x32bf8f8; return MfmaInstr::mfma_f32_16x16x32bf8f8;
} }
static constexpr auto selected_mfma = static constexpr auto selected_mfma = mfma_type<
mfma_type<GetMfma<base_type, MPerXdlops, NPerXdlops, additional_type, is_single_rate_mfma>()>{}; GetMfma<base_type, MPerXdlops, NPerXdlops, additional_type, is_single_rate_mfma>()>{};
__host__ __device__ constexpr MfmaSelector() __host__ __device__ constexpr MfmaSelector()
{ {
...@@ -1431,8 +1431,12 @@ struct XdlopsGemm ...@@ -1431,8 +1431,12 @@ struct XdlopsGemm
} }
// Falls back to single rate instruction on gfx950 if KPack <= 4; no change on gfx942- // Falls back to single rate instruction on gfx950 if KPack <= 4; no change on gfx942-
static constexpr auto mfma = MfmaSelector<base_type, MPerXdlops, NPerXdlops, additional_type, static constexpr auto
((is_same<base_type, half_t>::value || is_same<base_type, bhalf_t>::value) && KPack <= 4) ? true : false>{}; mfma = MfmaSelector < base_type,
MPerXdlops, NPerXdlops, additional_type,
((is_same<base_type, half_t>::value || is_same<base_type, bhalf_t>::value) && KPack <= 4)
? true
: false > {};
static constexpr auto mfma_instr = mfma.selected_mfma; static constexpr auto mfma_instr = mfma.selected_mfma;
......
include/ck_tile/ops/flatmm/block/uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16.inc
View file @ f24bb5a0
...@@ -824,4 +824,4 @@ ...@@ -824,4 +824,4 @@
#undef _UK_PK_CVT_ #undef _UK_PK_CVT_
#undef _UK_ATOMIC_ADD_ #undef _UK_ATOMIC_ADD_
#undef CK_TILE_FLATMM_UK_MFMA #undef CK_TILE_FLATMM_UK_MFMA
// clang-format on // clang-format on
include/ck_tile/ops/flatmm/block/uk/flatmm_sn_uk_gfx9_32x128x512_1x4x1_16x16x16_itl.inc
View file @ f24bb5a0
...@@ -722,4 +722,4 @@ ...@@ -722,4 +722,4 @@
#undef _UK_PK_CVT_ #undef _UK_PK_CVT_
#undef _UK_ATOMIC_ADD_ #undef _UK_ATOMIC_ADD_
#undef CK_TILE_FLATMM_UK_MFMA #undef CK_TILE_FLATMM_UK_MFMA
// clang-format on // clang-format on
include/ck_tile/ops/flatmm/block/uk/flatmm_uk_gfx9_32x512x128_1x1x1_16x16x16.inc
View file @ f24bb5a0
...@@ -771,4 +771,4 @@ ...@@ -771,4 +771,4 @@
#undef _UK_MFMA_ #undef _UK_MFMA_
#undef CK_TILE_FLATMM_UK_2B #undef CK_TILE_FLATMM_UK_2B
#undef CK_TILE_FLATMM_UK_MFMA #undef CK_TILE_FLATMM_UK_MFMA
// clang-format on // clang-format on
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment