Skip to content

GitLab

Commit 734c2e74 authored by Alan Turner's avatar Alan Turner
Browse files

Clean up ck_gemm.hpp

parent 7c636a51
Hide whitespace changes
Inline Side-by-side
Showing with 0 additions and 68 deletions
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
View file @ 734c2e74
......@@ -62,74 +62,6 @@ __device__ void ck_gemm_matrix(E e, A a, B b, Ds... ds)
to_ck_const_pointer(b.data()),
ck::make_tuple(to_ck_const_pointer(ds.data())...),
to_ck_pointer(e.data()));
// constexpr const auto M = a.get_shape().lens[0];
// constexpr const auto N = b.get_shape().lens[1];
// constexpr const auto K = a.get_shape().lens[1];
// constexpr const auto K1Number = ck::Number<4>{};
// constexpr const auto K0 = K / 4;
// using GridwiseGemm = typename G::GridwiseGemm;
// constexpr auto a_grid_desc_k0_m_k1 = ck::transform_tensor_descriptor(
// to_ck_tensor<A>(),
// ck::make_tuple(ck::make_unmerge_transform(ck::make_tuple(K0, K1Number)),
// ck::make_pass_through_transform(M)),
// ck::make_tuple(ck::Sequence<1>{}, ck::Sequence<0>{}),
// ck::make_tuple(ck::Sequence<0, 2>{}, ck::Sequence<1>{}));
// constexpr const auto a_grid_desc_k0_m0_m1_k1 =
// GridwiseGemm::MakeAGridDescriptor_K0_M0_M1_K1(a_grid_desc_k0_m_k1);
// noop(a, b, e, M, N, K, a_grid_desc_k0_m_k1, a_grid_desc_k0_m0_m1_k1, ds...);
//////////////////////////
// constexpr const G gemm{};
// constexpr const auto a_grid_desc_m_k =
// gemm.matrix_padder.PadADescriptor_M_K(to_ck_tensor<A>()); constexpr const auto
// b_grid_desc_n_k =
// gemm.matrix_padder.PadBDescriptor_N_K(to_ck_tensor<ck_transposeb<B>>());
// constexpr const auto e_grid_desc_m_n =
// gemm.matrix_padder.PadCDescriptor_M_N(to_ck_tensor<E>()); constexpr const auto
// ds_grid_desc_m_n =
// ck::make_tuple(gemm.matrix_padder.PadCDescriptor_M_N(to_ck_tensor<Ds>())...);
// constexpr const auto block_2_etile_map = gemm.MakeDefaultBlock2ETileMap(e_grid_desc_m_n);
// using GridwiseGemm = typename G::GridwiseGemm;
// // tensor descriptors for block/thread-wise copy
// constexpr auto a_grid_desc_ak0_m_ak1 =
// GridwiseGemm::MakeDefaultAGridDescriptor_AK0_M_AK1(a_grid_desc_m_k);
// constexpr auto b_grid_desc_bk0_n_bk1 =
// GridwiseGemm::MakeDefaultBGridDescriptor_BK0_N_BK1(b_grid_desc_n_k);
// constexpr auto ds_grid_desc_mblock_mperblock_nblock_nperblock =
// GridwiseGemm::MakeDsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(ds_grid_desc_m_n);
// constexpr auto e_grid_desc_mblock_mperblock_nblock_nperblock =
// GridwiseGemm::MakeEGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock(e_grid_desc_m_n);
// static_assert(GridwiseGemm::CheckValidity(
// a_grid_desc_m_k, b_grid_desc_n_k, ds_grid_desc_m_n, e_grid_desc_m_n, block_2_etile_map));
// __shared__ char p_shared_block[GridwiseGemm::GetSharedMemoryNumberOfByte()];
// constexpr const bool HasMainKBlockLoop =
// GridwiseGemm::CalculateHasMainKBlockLoop(a_grid_desc_ak0_m_ak1.GetLength(ck::Number<0>{})
// *
// a_grid_desc_ak0_m_ak1.GetLength(ck::Number<2>{}));
// GridwiseGemm::template Run<HasMainKBlockLoop>(to_ck_const_pointer(a.data()),
// to_ck_const_pointer(b.data()),
// ck::make_tuple(to_ck_const_pointer(ds.data())...),
// to_ck_pointer(e.data()),
// p_shared_block,
// gemm.a_element_op,
// gemm.b_element_op,
// gemm.cde_element_op,
// a_grid_desc_ak0_m_ak1,
// b_grid_desc_bk0_n_bk1,
// ds_grid_desc_mblock_mperblock_nblock_nperblock,
// e_grid_desc_mblock_mperblock_nblock_nperblock,
// block_2_etile_map);
}
template <class G, index_int BlocksPerBatch, class... Ts>
......
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