mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-12 01:10:17 +00:00
109 lines
2.4 KiB
C++
109 lines
2.4 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#include <tuple>
|
|
|
|
#include "gtest/gtest.h"
|
|
|
|
#include "test_gemm_mx_util.hpp"
|
|
|
|
using E8M0 = ck::e8m0_bexp_t;
|
|
using F8 = ck::f8_t;
|
|
using BF8 = ck::bf8_t;
|
|
using F6 = ck::f6_t;
|
|
using BF6 = ck::bf6_t;
|
|
using F4 = ck::f4_t;
|
|
using F16 = ck::half_t;
|
|
using BF16 = ck::bhalf_t;
|
|
using F32 = float;
|
|
|
|
using Row = ck::tensor_layout::gemm::RowMajor;
|
|
using Col = ck::tensor_layout::gemm::ColumnMajor;
|
|
|
|
namespace {
|
|
|
|
template <typename X, typename Y>
|
|
struct tuple_concat;
|
|
|
|
template <typename... Xs, typename... Ys>
|
|
struct tuple_concat<std::tuple<Xs...>, std::tuple<Ys...>>
|
|
{
|
|
using type = std::tuple<Xs..., Ys...>;
|
|
};
|
|
|
|
} // namespace
|
|
|
|
template <typename Tuple>
|
|
class TestGemmMX_MK_NK
|
|
: public ck::test::TestGemmMX<typename tuple_concat<std::tuple<Row, Col>, Tuple>::type>
|
|
{
|
|
};
|
|
|
|
// clang-format off
|
|
using KernelTypes_MK_NK = ::testing::Types<
|
|
#if defined(CK_ENABLE_FP8)
|
|
// ADataType, BDataType, CDataType, ScaleBlockSize
|
|
std::tuple< F8, F8, F16, ck::Number<32> >,
|
|
std::tuple< F8, F8, BF16, ck::Number<32> >
|
|
#endif
|
|
>;
|
|
// clang-format on
|
|
|
|
TYPED_TEST_SUITE(TestGemmMX_MK_NK, KernelTypes_MK_NK);
|
|
|
|
TYPED_TEST(TestGemmMX_MK_NK, SmallM)
|
|
{
|
|
std::vector<int> Ms{1, 2, 3, 4, 5, 6};
|
|
constexpr int N = 256;
|
|
constexpr int K = 512;
|
|
|
|
constexpr int StrideA = K;
|
|
constexpr int StrideB = K;
|
|
constexpr int StrideC = N;
|
|
|
|
for(int M : Ms)
|
|
this->Run(M, N, K, StrideA, StrideB, StrideC);
|
|
}
|
|
|
|
TYPED_TEST(TestGemmMX_MK_NK, MidLargeM)
|
|
{
|
|
std::vector<int> Ms{127, 255, 312, 799, 1573};
|
|
constexpr int N = 256;
|
|
constexpr int K = 512;
|
|
|
|
constexpr int StrideA = K;
|
|
constexpr int StrideB = K;
|
|
constexpr int StrideC = N;
|
|
|
|
for(int M : Ms)
|
|
this->Run(M, N, K, StrideA, StrideB, StrideC);
|
|
}
|
|
|
|
TYPED_TEST(TestGemmMX_MK_NK, Regular)
|
|
{
|
|
std::vector<int> Ms{3840};
|
|
constexpr int N = 512;
|
|
constexpr int K = 1024;
|
|
|
|
constexpr int StrideA = K;
|
|
constexpr int StrideB = K;
|
|
constexpr int StrideC = N;
|
|
|
|
for(int M : Ms)
|
|
this->Run(M, N, K, StrideA, StrideB, StrideC);
|
|
}
|
|
|
|
TYPED_TEST(TestGemmMX_MK_NK, Large)
|
|
{
|
|
std::vector<int> Ms{4096};
|
|
constexpr int N = 3840;
|
|
constexpr int K = 4096;
|
|
|
|
constexpr int StrideA = K;
|
|
constexpr int StrideB = K;
|
|
constexpr int StrideC = N;
|
|
|
|
for(int M : Ms)
|
|
this->Run(M, N, K, StrideA, StrideB, StrideC);
|
|
}
|