adding implicit gemm

[ROCm/composable_kernel commit: 913afaeb5d]
This commit is contained in:
Chao Liu
2019-01-16 02:01:56 -06:00
parent 8b3c613be1
commit ba06d1ee1a
6 changed files with 325 additions and 136 deletions

View File

@@ -1,34 +1,36 @@
#pragma once
#include "common.cuh"
template <unsigned NRow, unsigned NCol, unsigned RowStride>
template <unsigned NRow_, unsigned NCol_, unsigned RowStride_>
struct ConstantMatrixDescriptor
{
__host__ __device__ ConstantMatrixDescriptor()
{
static_assert(NCol <= RowStride, "wrong! NCol > RowStride!");
static_assert(NCol_ <= RowStride_, "wrong! NCol > RowStride!");
}
__host__ __device__ constexpr unsigned GetNumberOfRow() const { return NRow; }
__host__ __device__ constexpr unsigned NRow() const { return NRow_; }
__host__ __device__ constexpr unsigned GetNumberOfColumn() const { return NCol; }
__host__ __device__ constexpr unsigned NCol() const { return NCol_; }
__host__ __device__ constexpr unsigned GetRowStride() const { return RowStride; }
__host__ __device__ constexpr unsigned RowStride() const { return RowStride_; }
__host__ __device__ constexpr unsigned GetElementSize() const { return NRow * NCol; }
__host__ __device__ constexpr auto GetLengths() const { return Sequence<NRow_, NCol_>{}; }
__host__ __device__ constexpr unsigned GetElementSpace() const { return NRow * RowStride; }
__host__ __device__ constexpr unsigned GetElementSize() const { return NRow_ * NCol_; }
__host__ __device__ constexpr unsigned GetElementSpace() const { return NRow_ * RowStride_; }
__host__ __device__ unsigned Get1dIndex(unsigned irow, unsigned icol) const
{
return irow * RowStride + icol;
return irow * RowStride_ + icol;
}
template <unsigned SubNRow, unsigned SubNCol>
__host__ __device__ constexpr auto MakeSubMatrixDescriptor(Number<SubNRow>,
Number<SubNCol>) const
{
return ConstantMatrixDescriptor<SubNRow, SubNCol, RowStride>{};
return ConstantMatrixDescriptor<SubNRow, SubNCol, RowStride_>{};
}
};