#pragma once #include "common.hip.hpp" #include "ConstantMatrixDescriptor.hip.hpp" template __device__ void threadwise_matrix_set_zero(Matrix, Float* __restrict__ p_thread) { for(index_t i = 0; i < Matrix::NRow(); ++i) { for(index_t j = 0; j < Matrix::NCol(); ++j) { const index_t id = Matrix::Get1dIndex(i, j); p_thread[id] = 0; } } } template __device__ void threadwise_matrix_copy(SrcMatrix, const Float* __restrict__ p_src, DstMatrix, Float* __restrict__ p_dst, Sequence, Number) { static_assert(NCol % DataPerRead == 0, "wrong! should be NCol % == DataPerRead == 0"); using vector_t = typename vector_type::MemoryType; constexpr auto src_mtx = SrcMatrix{}; constexpr auto dst_mtx = DstMatrix{}; for(index_t i = 0; i < NRow; ++i) { for(index_t j = 0; j < NCol; j += DataPerRead) { const index_t src_index = src_mtx.Get1dIndex(i, j); const index_t dst_index = dst_mtx.Get1dIndex(i, j); *reinterpret_cast(&p_dst[dst_index]) = *reinterpret_cast(&p_src[src_index]); } } } template __device__ void threadwise_gemm(MatrixA, integral_constant, const FloatA* __restrict__ p_a_thread, MatrixB, integral_constant, const FloatB* __restrict__ p_b_thread, MatrixC, integral_constant, FloatC* __restrict__ p_c_thread) { if(TransA && (!TransB) && (!TransC)) { constexpr auto a_mtx = MatrixA{}; constexpr auto b_mtx = MatrixB{}; constexpr auto c_mtx = MatrixC{}; constexpr index_t M = c_mtx.NRow(); constexpr index_t N = c_mtx.NCol(); constexpr index_t K = a_mtx.NRow(); // A is transposed for(index_t k = 0; k < K; ++k) { for(index_t i = 0; i < M; ++i) { for(index_t j = 0; j < N; ++j) { const index_t aindex = a_mtx.Get1dIndex(k, i); // A is transposed const index_t bindex = b_mtx.Get1dIndex(k, j); const index_t cindex = c_mtx.Get1dIndex(i, j); p_c_thread[cindex] += p_a_thread[aindex] * p_b_thread[bindex]; } } } } else { // not implemented assert(false); } }