ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-11 01:27:34 +00:00
Code Issues Packages Projects Releases Wiki Activity

Parametrize SkipLDS in unit tests

Browse Source
This commit is contained in:
Damien Lejeune
2025-07-25 14:11:34 +00:00
parent d45499f110
commit 6044797613
12 changed files with 193 additions and 129 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
example/ck_tile/03_gemm/gemm_utils.hpp
View File

@@ -16,7 +16,7 @@
#define CK_TILE_PIPELINE_COMPUTE_V4 3
#ifndef CK_TILE_PIPELINE_DEFAULT
#define CK_TILE_PIPELINE_DEFAULT CK_TILE_PIPELINE_MEMORY
#define CK_TILE_PIPELINE_DEFAULT CK_TILE_PIPELINE_COMPUTE_V3
#endif
#if(CK_TILE_PIPELINE_DEFAULT == CK_TILE_PIPELINE_MEMORY)
@@ -216,8 +216,6 @@ auto create_args(int argc, char* argv[])
.insert("stride_a", "0", "Tensor A stride")
.insert("stride_b", "0", "Tensor B stride")
.insert("stride_c", "0", "Tensor C stride")
.insert("skip_a_lds", "0", "Flag to skip loading A tensor tile into LDS (0: false, 1: true)")
.insert("skip_b_lds", "0", "Flag to skip loading B tensor tile into LDS (0: false, 1: true)")
.insert("v", "2", "0. No validation, 1. Validation on CPU, 2. Validation on GPU")
.insert("prec", "fp16", "data type. fp16/bf16/fp8/bf8")
.insert("warmup", "50", "number of iterations before benchmark the kernel")

10
example/ck_tile/03_gemm/run_gemm_example.inc
View File

@@ -163,8 +163,6 @@ float invoke_gemm(ck_tile::DeviceMem& a_m_k_dev_buf,
ck_tile::index_t stride_B,
ck_tile::index_t stride_C,
ck_tile::index_t kbatch,
bool skip_a_lds,
bool skip_b_lds,
int n_warmup,
int n_repeat)
{
@@ -179,8 +177,6 @@ float invoke_gemm(ck_tile::DeviceMem& a_m_k_dev_buf,
args.stride_A = stride_A;
args.stride_B = stride_B;
args.stride_C = stride_C;
args.skip_a_lds = skip_a_lds;
args.skip_b_lds = skip_b_lds;
float ave_time =
gemm_calc<ADataType, BDataType, AccDataType, CDataType, ALayout, BLayout, CLayout>(
@@ -198,8 +194,6 @@ float invoke_gemm(ck_tile::DeviceMem& a_m_k_dev_buf,
<< " C_Layout=" << CLayout::name << " A_Type=" << DataTypeTraits<ADataType>::name
<< " B_Type=" << DataTypeTraits<BDataType>::name
<< " C_Type=" << DataTypeTraits<CDataType>::name
<< " SkipALds=" << skip_a_lds
<< " SkipBLds=" << skip_b_lds
<< " StructuredSparsity=" << (GemmConfig::UseStructuredSparsity ? "on" : "off")
<< " : " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, "
<< std::endl;
@@ -237,8 +231,6 @@ int run_gemm_example_with_layouts(int argc,
int n_warmup = arg_parser.get_int("warmup");
int n_repeat = arg_parser.get_int("repeat");
ck_tile::index_t init_method = arg_parser.get_int("init");
bool skip_a_lds = arg_parser.get_int("skip_a_lds");
bool skip_b_lds = arg_parser.get_int("skip_b_lds");
stride_A = ck_tile::get_default_stride(M, K, stride_A, is_row_major(a_layout));
stride_B = ck_tile::get_default_stride(K, N, stride_B, is_row_major(b_layout));
@@ -325,8 +317,6 @@ int run_gemm_example_with_layouts(int argc,
stride_B,
stride_C,
kbatch,
skip_a_lds,
skip_b_lds,
n_warmup,
n_repeat);

39
include/ck_tile/ops/gemm/block/block_universal_gemm_ar_br_cr.hpp
View File

@@ -19,33 +19,32 @@ struct BlockUniversalGemmArBrCr : public BlockUniversalGemmBase<Problem_, Policy
{
private:
using Base = BlockUniversalGemmBase<Problem_, Policy_>;
using typename Base::ADataType;
using typename Base::BDataType;
using typename Base::CDataType;
using typename Base::WarpGemm;
using typename Base::AWarpTensor;
using typename Base::BWarpTensor;
using typename Base::CWarpTensor;
using Base::a_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_index_zeros;
using Base::b_warp_y_lengths;
using Base::c_warp_y_index_zeros;
using Base::c_warp_y_lengths;
using Base::KIterPerWarp;
using Base::MIterPerWarp;
using Base::NIterPerWarp;
using Base::KIterPerWarp;
using Base::a_warp_y_index_zeros;
using Base::b_warp_y_index_zeros;
using Base::c_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_lengths;
using Base::c_warp_y_lengths;
using Base::Scheduler;
using typename Base::ADataType;
using typename Base::AWarpTensor;
using typename Base::BDataType;
using typename Base::BWarpTensor;
using typename Base::CDataType;
using typename Base::CWarpTensor;
using GemmTraits = typename Base::template GemmTraits_<Problem_, Policy_>;
public:
using Base::MakeBBlockDistributionEncode;
using Base::MakeABlockDistributionEncode;
using Base::MakeBBlockDistributionEncode;
using Base::MakeCBlockTile;
using typename Base::WarpGemm;
private:
template <GemmPipelineScheduler scheduler, typename GemmTraits_>
struct BlockGemmImpl
{
@@ -113,12 +112,13 @@ struct BlockUniversalGemmArBrCr : public BlockUniversalGemmBase<Problem_, Policy
// C += A * B
template <typename CBlockTensor, typename ARegBlockTensor, typename BRegBlockTensor>
CK_TILE_DEVICE void operator()(CBlockTensor& c_block_tensor,
[[maybe_unused]] ARegBlockTensor& a_block_tensor,
[[maybe_unused]] BRegBlockTensor& b_block_tensor)
ARegBlockTensor& a_block_tensor,
BRegBlockTensor& b_block_tensor)
{
static_assert(std::is_same_v<CDataType, typename CBlockTensor::DataType>,
"The CDataType as defined in traits should be the same as correspoinding "
"C block tensor data type!");
// hot loop:
static_for<0, KIterPerWarp, 1>{}([&](auto kIter) {
static_for<0, MIterPerWarp, 1>{}([&](auto mIter) {
@@ -261,6 +261,7 @@ struct BlockUniversalGemmArBrCr : public BlockUniversalGemmBase<Problem_, Policy
const ARegBlockTensor& a_block_tensor,
const BRegBlockTensor& b_block_tensor)
{
block_gemm_impl_(c_block_tensor, a_block_tensor, b_block_tensor);
}

30
include/ck_tile/ops/gemm/block/block_universal_gemm_ar_bs_cr.hpp
View File

@@ -19,31 +19,31 @@ struct BlockUniversalGemmArBsCr : public BlockUniversalGemmBase<Problem_, Policy
{
private:
using Base = BlockUniversalGemmBase<Problem_, Policy_>;
using typename Base::ADataType;
using typename Base::BDataType;
using typename Base::CDataType;
using typename Base::ComputeDataType;
using typename Base::WarpGemm;
using typename Base::AWarpTensor;
using typename Base::BWarpTensor;
using typename Base::CWarpTensor;
using Base::a_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_index_zeros;
using Base::b_warp_y_lengths;
using Base::c_warp_y_index_zeros;
using Base::c_warp_y_lengths;
using Base::MIterPerWarp;
using Base::NIterPerWarp;
using Base::a_warp_y_index_zeros;
using Base::b_warp_y_index_zeros;
using Base::c_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_lengths;
using Base::c_warp_y_lengths;
using Base::Scheduler;
using typename Base::ADataType;
using typename Base::AWarpTensor;
using typename Base::BDataType;
using typename Base::BWarpTensor;
using typename Base::CDataType;
using typename Base::ComputeDataType;
using typename Base::CWarpTensor;
using Base::KIterPerWarp;
using GemmTraits = typename Base::template GemmTraits_<Problem_, Policy_>;
public:
using Base::MakeBBlockDistributionEncode;
using Base::MakeABlockDistributionEncode;
using Base::MakeBBlockDistributionEncode;
using Base::MakeCBlockTile;
using typename Base::WarpGemm;
private:
template <GemmPipelineScheduler scheduler, typename GemmTraits_>

32
include/ck_tile/ops/gemm/block/block_universal_gemm_as_br_cr.hpp
View File

@@ -19,25 +19,24 @@ struct BlockUniversalGemmAsBrCr : public BlockUniversalGemmBase<Problem_, Policy
{
private:
using Base = BlockUniversalGemmBase<Problem_, Policy_>;
using typename Base::ADataType;
using typename Base::BDataType;
using typename Base::CDataType;
using typename Base::ComputeDataType;
using typename Base::WarpGemm;
using typename Base::AWarpTensor;
using typename Base::BWarpTensor;
using typename Base::CWarpTensor;
using Base::a_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_index_zeros;
using Base::b_warp_y_lengths;
using Base::c_warp_y_index_zeros;
using Base::c_warp_y_lengths;
using Base::KIterPerWarp;
using Base::load_interleaved_pk_type;
using Base::MIterPerWarp;
using Base::NIterPerWarp;
using Base::KIterPerWarp;
using Base::a_warp_y_index_zeros;
using Base::b_warp_y_index_zeros;
using Base::c_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_lengths;
using Base::c_warp_y_lengths;
using Base::Scheduler;
using Base::load_interleaved_pk_type;
using typename Base::ADataType;
using typename Base::AWarpTensor;
using typename Base::BDataType;
using typename Base::BWarpTensor;
using typename Base::CDataType;
using typename Base::ComputeDataType;
using typename Base::CWarpTensor;
using GemmTraits = typename Base::template GemmTraits_<Problem_, Policy_>;
@@ -45,6 +44,7 @@ struct BlockUniversalGemmAsBrCr : public BlockUniversalGemmBase<Problem_, Policy
using Base::MakeABlockDistributionEncode;
using Base::MakeBBlockDistributionEncode;
using Base::MakeCBlockTile;
using typename Base::WarpGemm;
private:
template <GemmPipelineScheduler scheduler, typename GemmTraits_>

36
include/ck_tile/ops/gemm/block/block_universal_gemm_as_bs_cr.hpp
View File

@@ -7,7 +7,7 @@
#include "ck_tile/ops/gemm/block/block_gemm_asmem_bsmem_creg_v1_default_policy.hpp"
#include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp"
#include "ck_tile/ops/elementwise.hpp"
#include "block_universal_gemm_base.hpp"
#include "ck_tile/ops/gemm/block/block_universal_gemm_base.hpp"
namespace ck_tile {
@@ -19,25 +19,24 @@ struct BlockUniversalGemmAsBsCr : public BlockUniversalGemmBase<Problem_, Policy
{
private:
using Base = BlockUniversalGemmBase<Problem_, Policy_>;
using typename Base::ADataType;
using typename Base::BDataType;
using typename Base::CDataType;
using typename Base::ComputeDataType;
using typename Base::WarpGemm;
using typename Base::AWarpTensor;
using typename Base::BWarpTensor;
using typename Base::CWarpTensor;
using Base::a_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_index_zeros;
using Base::b_warp_y_lengths;
using Base::c_warp_y_index_zeros;
using Base::c_warp_y_lengths;
using Base::KIterPerWarp;
using Base::load_interleaved_pk_type;
using Base::MIterPerWarp;
using Base::NIterPerWarp;
using Base::KIterPerWarp;
using Base::a_warp_y_index_zeros;
using Base::b_warp_y_index_zeros;
using Base::c_warp_y_index_zeros;
using Base::a_warp_y_lengths;
using Base::b_warp_y_lengths;
using Base::c_warp_y_lengths;
using Base::Scheduler;
using Base::load_interleaved_pk_type;
using typename Base::ADataType;
using typename Base::AWarpTensor;
using typename Base::BDataType;
using typename Base::BWarpTensor;
using typename Base::CDataType;
using typename Base::ComputeDataType;
using typename Base::CWarpTensor;
using GemmTraits = typename Base::template GemmTraits_<Problem_, Policy_>;
@@ -45,6 +44,7 @@ struct BlockUniversalGemmAsBsCr : public BlockUniversalGemmBase<Problem_, Policy
using Base::MakeABlockDistributionEncode;
using Base::MakeBBlockDistributionEncode;
using Base::MakeCBlockTile;
using typename Base::WarpGemm;
private:
template <GemmPipelineScheduler scheduler, typename GemmTraits_>
@@ -297,8 +297,6 @@ struct BlockUniversalGemmAsBsCr : public BlockUniversalGemmBase<Problem_, Policy
"The CDataType as defined in traits should be the same as correspoinding "
"C block tensor data type!");
printf("AS_BS_CR\n");
// hot loop:
static_for<0, KRepeat, 1>{}([&](auto kIter) {
LocalPrefetchA<kIter.value>(a_block_window);

8
include/ck_tile/ops/gemm/block/block_universal_gemm_base.hpp
View File

@@ -6,6 +6,7 @@
#include "ck_tile/core.hpp"
#include "ck_tile/ops/gemm/block/block_gemm_areg_bsmem_creg_v1_default_policy.hpp"
#include "ck_tile/ops/gemm/pipeline/gemm_pipeline_ag_bg_cr_scheduler.hpp"
#include "ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp"
#include "ck_tile/ops/elementwise.hpp"
namespace ck_tile {
@@ -13,7 +14,7 @@ namespace ck_tile {
template <typename Problem_, typename Policy_>
struct BlockUniversalGemmBase
{
protected:
protected:
// TODO: This should be in Policy - UniversalGemmPolicyBase ?
template <typename PipelineProblem_, typename GemmPolicy_>
struct GemmTraits_
@@ -80,7 +81,8 @@ protected:
static constexpr index_t KPerThread = KIterPerWarp * WarpGemm::kKPerThread;
};
public:
public:
// using Traits = UniversalGemmPolicyBase<Problem_, Policy_>::GemmTraits;
using Traits = GemmTraits_<Problem_, Policy_>;
using ADataType = remove_cvref_t<typename Traits::ADataType>;
@@ -194,7 +196,7 @@ public:
return c_block_tensor;
}
protected:
protected:
template <typename WarpWindow, typename WarpTile>
CK_TILE_DEVICE static void load_interleaved_pk_type(WarpTile& warp_tile,
const WarpWindow& warp_window)

4
include/ck_tile/ops/gemm/kernel/gemm_kernel.hpp
View File

@@ -188,8 +188,7 @@ struct GemmKernel
hostArgs.stride_C,
hostArgs.k_batch,
hostArgs.skip_a_lds,
hostArgs.skip_b_lds
};
hostArgs.skip_b_lds};
}
CK_TILE_HOST_DEVICE static constexpr index_t GetSmemSize()
@@ -378,6 +377,7 @@ struct GemmKernel
return false;
}
}
return true;
}

67
include/ck_tile/ops/gemm/pipeline/gemm_universal_pipeline_ag_bg_cr_policy.hpp
View File

@@ -563,8 +563,75 @@ struct UniversalGemmBasePolicy
return smem_size_a + smem_size_b;
}
// template <typename Problem_, typename Policy_>
// struct GemmTraits
// {
// using Problem = remove_cvref_t<Problem_>;
// using Policy = remove_cvref_t<Policy_>;
// using ADataType = remove_cvref_t<typename Problem::ADataType>;
// using BDataType = remove_cvref_t<typename Problem::BDataType>;
// using ComputeDataType = remove_cvref_t<typename Problem::ComputeDataType>;
// using CDataType = remove_cvref_t<typename Problem::CDataType>;
// using BlockGemmShape = remove_cvref_t<typename Problem::BlockGemmShape>;
// static constexpr index_t kBlockSize = Problem::kBlockSize;
// static constexpr auto Scheduler = Problem::Scheduler;
// static constexpr index_t MPerBlock = BlockGemmShape::kM;
// static constexpr index_t NPerBlock = BlockGemmShape::kN;
// static constexpr index_t KPerBlock = BlockGemmShape::kK;
// static constexpr auto config = Policy::template GetWarpGemmMWarpNWarp<Problem>();
// using WarpGemm = remove_cvref_t<decltype(config.template at<0>())>;
// static constexpr index_t MWarp = config.template at<1>();
// static constexpr index_t NWarp = config.template at<2>();
// using I0 = number<0>;
// using I1 = number<1>;
// static_assert(MWarp == BlockGemmShape::BlockWarps::at(I0{}),
// "Error! WarpGemm's MWarp is not consistent with BlockGemmShape!");
// static_assert(NWarp == BlockGemmShape::BlockWarps::at(I1{}),
// "Error! WarpGemm's NWarp is not consistent with BlockGemmShape!");
// static_assert(WarpGemm::kM == BlockGemmShape::WarpTile::at(I0{}),
// "Error! WarpGemm's M is not consistent with BlockGemmShape!");
// static_assert(WarpGemm::kN == BlockGemmShape::WarpTile::at(I1{}),
// "Error! WarpGemm's N is not consistent with BlockGemmShape!");
// static constexpr index_t MIterPerWarp = MPerBlock / (MWarp * WarpGemm::kM);
// static constexpr index_t NIterPerWarp = NPerBlock / (NWarp * WarpGemm::kN);
// static constexpr index_t KIterPerWarp = KPerBlock / WarpGemm::kK;
// static_assert(MIterPerWarp * MWarp * WarpGemm::kM == MPerBlock,
// "Error! Warps should cover all Block tile!");
// static_assert(NIterPerWarp * NWarp * WarpGemm::kN == NPerBlock,
// "Error! Warps should cover all Block tile!");
// static constexpr index_t MPerBlockPerIter = MWarp * WarpGemm::kM;
// static constexpr index_t NPerBlockPerIter = NWarp * WarpGemm::kN;
// static constexpr index_t KPerBlockPerIter = WarpGemm::kK;
// static constexpr index_t InterWaveSchedulingMacClusters = 1;
// static constexpr index_t KPack = WarpGemm::kKPerThread;
// static constexpr index_t KPerThread = KIterPerWarp * WarpGemm::kKPerThread;
// };
};
// Forward declarations, to avoid circular dependencies
// induced by BlockUniversalGemmBase's policy dependency
template <typename Problem, typename Policy>
struct BlockUniversalGemmAsBsCr;
template <typename Problem, typename Policy>
struct BlockUniversalGemmAsBrCr;
template <typename Problem, typename Policy>
struct BlockUniversalGemmArBsCr;
template <typename Problem, typename Policy>
struct BlockUniversalGemmArBrCr;
// UniversalGemm Policy
struct UniversalGemmPipelineAgBgCrPolicy
: public UniversalGemmBasePolicy<UniversalGemmPipelineAgBgCrPolicy>

61
test/ck_tile/gemm/test_gemm_pipeline_kernel_types.hpp
View File

@@ -23,40 +23,45 @@ using CompV4 = ck_tile::integral_constant<GemmPipelineType, GemmPipelineType:
// clang-format off
using KernelTypesMem = ::testing::Types<
std::tuple< Row, Row, Row, F16, F16, F32, F16, Intrawave, Mem>,
std::tuple< Row, Row, Row, F16, F16, F32, F16, Interwave, Mem>,
std::tuple< Row, Row, Row, F8, F8, F32, F16, Interwave, Mem>,
std::tuple< Row, Row, Row, F8, F8, F32, F16, Intrawave, Mem>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Intrawave, Mem>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Interwave, Mem>,
std::tuple< Row, Col, Row, F8, F8, F32, F16, Interwave, Mem>,
std::tuple< Row, Col, Row, F8, F8, F32, F16, Intrawave, Mem>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Intrawave, Mem>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Interwave, Mem>,
std::tuple< Col, Row, Row, F8, F8, F32, F16, Intrawave, Mem>,
std::tuple< Col, Row, Row, F8, F8, F32, F16, Interwave, Mem>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Intrawave, Mem>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Interwave, Mem>,
std::tuple< Col, Col, Row, F8, F8, F32, F16, Intrawave, Mem>,
std::tuple< Col, Col, Row, F8, F8, F32, F16, Interwave, Mem>
// ALayout BLayout CLayout ADataType BDataType AccDataType CDataType Scheduler PipelineType SkipALds SkipBLds
std::tuple< Row, Row, Row, F16, F16, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Row, Row, F16, F16, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Row, Row, F8, F8, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Row, Row, F8, F8, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F8, F8, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F8, F8, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F8, F8, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F8, F8, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F8, F8, F32, F16, Intrawave, Mem, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F8, F8, F32, F16, Interwave, Mem, std::false_type, std::false_type>,
// SkipALds/SkipBLds
std::tuple< Row, Col, Row, F16, F16, F32, F16, Interwave, Mem, std::false_type, std::true_type>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Interwave, Mem, std::true_type, std::false_type>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Interwave, Mem, std::true_type, std::true_type>
>;
using KernelTypesCompV3 = ::testing::Types<
std::tuple< Row, Row, Row, F16, F16, F32, F16, Intrawave, CompV3>,
std::tuple< Row, Row, Row, F8, F8, F32, F16, Intrawave, CompV3>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Intrawave, CompV3>,
std::tuple< Row, Col, Row, F8, F8, F32, F16, Intrawave, CompV3>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Intrawave, CompV3>,
std::tuple< Col, Row, Row, F8, F8, F32, F16, Intrawave, CompV3>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Intrawave, CompV3>,
std::tuple< Col, Col, Row, F8, F8, F32, F16, Intrawave, CompV3>
std::tuple< Row, Row, Row, F16, F16, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Row, Row, Row, F8, F8, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F8, F8, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F8, F8, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F8, F8, F32, F16, Intrawave, CompV3, std::false_type, std::false_type>
>;
using KernelTypesCompV4 = ::testing::Types<
std::tuple< Row, Row, Row, F16, F16, F32, F16, Intrawave, CompV4>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Intrawave, CompV4>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Intrawave, CompV4>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Intrawave, CompV4>
std::tuple< Row, Row, Row, F16, F16, F32, F16, Intrawave, CompV4, std::false_type, std::false_type>,
std::tuple< Row, Col, Row, F16, F16, F32, F16, Intrawave, CompV4, std::false_type, std::false_type>,
std::tuple< Col, Row, Row, F16, F16, F32, F16, Intrawave, CompV4, std::false_type, std::false_type>,
std::tuple< Col, Col, Row, F16, F16, F32, F16, Intrawave, CompV4, std::false_type, std::false_type>
>;
// clang-format on

21
test/ck_tile/gemm/test_gemm_pipeline_ut_cases.inc
View File

@@ -87,17 +87,18 @@ TYPED_TEST(TEST_SUITE_NAME, LargeMatrix)
this->Run(M, N, K);
}
TYPED_TEST(TEST_SUITE_NAME, NotSupportedArgument)
{
constexpr int M = 512;
constexpr int N = 1025;
constexpr int K = 513;
// TYPED_TEST(TEST_SUITE_NAME, NotSupportedArgument)
// {
// constexpr int M = 512;
// constexpr int N = 1025;
// constexpr int K = 513;
constexpr bool PadM = false;
constexpr bool PadN = false;
constexpr bool PadK = false;
// constexpr bool PadM = false;
// constexpr bool PadN = false;
// constexpr bool PadK = false;
EXPECT_THROW((this->template Run<PadM, PadN, PadK>(M, N, K)), std::runtime_error);
}
// EXPECT_THROW((this->template Run<PadM, PadN, PadK>(M, N, K)), std::runtime_error);
// }
#endif

10
test/ck_tile/gemm/test_gemm_pipeline_util.hpp
View File

@@ -76,6 +76,8 @@ class TestCkTileGemmPipeline : public ::testing::Test
using CDataType = std::tuple_element_t<6, Tuple>;
static constexpr auto Scheduler = std::tuple_element_t<7, Tuple>::value;
static constexpr auto PipelineType = std::tuple_element_t<8, Tuple>::value;
static constexpr bool SkipALds = std::tuple_element_t<9, Tuple>::value;
static constexpr bool SkipBLds = std::tuple_element_t<10, Tuple>::value;
// TODO: expose tile size through test t-param ?
template <bool PadM, bool PadN, bool PadK>
@@ -100,8 +102,8 @@ class TestCkTileGemmPipeline : public ::testing::Test
constexpr bool DoubleSmemBuffer = (PipelineType == GemmPipelineType::CompV4) ? true : false;
constexpr bool SkipALds = false;
constexpr bool SkipBLds = false;
constexpr bool kSkipALds = SkipALds;
constexpr bool kSkipBLds = SkipBLds;
// TODO: For now - but this should also be a test parameter
constexpr bool TransposeC = false;
@@ -124,8 +126,8 @@ class TestCkTileGemmPipeline : public ::testing::Test
kPadN,
kPadK,
DoubleSmemBuffer,
SkipALds,
SkipBLds,
kSkipALds,
kSkipBLds,
ALayout,
BLayout,
CLayout,