Skip to content

GitLab

Commit 792b62dd authored by Bartlomiej Wroblewski's avatar Bartlomiej Wroblewski
Browse files

Review: Remove construction of dpp_datatypes object

parent 77e63b34
Hide whitespace changes
Inline Side-by-side
Showing with 5 additions and 5 deletions
include/ck/utility/amd_gemm_dpp.hpp
View file @ 792b62dd
...@@ -35,10 +35,10 @@ template <index_t MPerThread, ...@@ -35,10 +35,10 @@ template <index_t MPerThread,
bool ShareA> bool ShareA>
struct DppInstrRunner struct DppInstrRunner
{ {
static constexpr auto datatypes_conf = dpp_datatypes<BaseInputType>{}; using datatypes_conf = dpp_datatypes<BaseInputType>;
using ADataType = typename decltype(datatypes_conf)::a_dtype; using ADataType = typename datatypes_conf::a_dtype;
using BDataType = typename decltype(datatypes_conf)::b_dtype; using BDataType = typename datatypes_conf::b_dtype;
using CDataType = typename decltype(datatypes_conf)::c_dtype; using CDataType = typename datatypes_conf::c_dtype;
__device__ void Run(const AVecDataType& a_vec, const BVecDataType& b_vec, CVecDataType& c_vec) __device__ void Run(const AVecDataType& a_vec, const BVecDataType& b_vec, CVecDataType& c_vec)
{ {
...@@ -51,7 +51,7 @@ struct DppInstrRunner ...@@ -51,7 +51,7 @@ struct DppInstrRunner
float c = c_vec.template AsType<CDataType>()(c_idx); float c = c_vec.template AsType<CDataType>()(c_idx);
// Next `c_idx` implies that we need to pull data from the next lane. // Next `c_idx` implies that we need to pull data from the next lane.
constexpr index_t source_lane = c_idx; constexpr index_t source_lane = c_idx;
static_for<0, KPerThread / datatypes_conf.k_per_instr, 1>{}([&](auto k_chunk) { static_for<0, KPerThread / datatypes_conf::k_per_instr, 1>{}([&](auto k_chunk) {
const auto a_k_vec = a_vector.template AsType<AVecDataType>()[k_chunk]; const auto a_k_vec = a_vector.template AsType<AVecDataType>()[k_chunk];
const auto b_k_vec = b_vector.template AsType<BVecDataType>()[k_chunk]; const auto b_k_vec = b_vector.template AsType<BVecDataType>()[k_chunk];
ck::dpp8:: ck::dpp8::
......
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