Skip to content
Original file line number Diff line number Diff line change
Expand Up @@ -580,6 +580,18 @@ struct UniversalGemmKernel
using DiLayout = remove_cvref_t<std::tuple_element_t<index.value, DsLayout>>;
if(std::is_same_v<DiLayout, CLayout> == false)
{
// Every other rejection in this function reports itself under
// CK_TILE_LOGGING; this one used to fail silently, so a kernel
// simply never ran and the reason had to be found by reading
// the source. A broadcast scale D (an M-vector with a column
// layout, paired with a row-major C) hits exactly this path.
if(ck_tile::EnvIsEnabled(CK_TILE_ENV(CK_TILE_LOGGING)))
{
CK_TILE_ERROR("D tensor layout must match the C layout: D",
index.value,
" differs, so this kernel cannot run. Give D the C layout, "
"or use a kernel/epilogue that supports a transposed D.");
}
DTensorIsValid = false;
}
if constexpr(std::is_same_v<DiLayout, tensor_layout::gemm::RowMajor>)
Expand Down
7 changes: 7 additions & 0 deletions projects/composablekernel/test/ck_tile/gemm/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,13 @@ if(GPU_TARGETS MATCHES "gfx94|gfx95|gfx90a|gfx11|gfx12")
set(CK_TILE_GEMM_TEST_TARGETS)

if(GPU_TARGETS MATCHES "gfx94|gfx95")
# This synthetic argument-validation fixture models the XDL pipeline
# interface; keep it out of the gfx11/gfx12 WMMA build.
add_gtest_executable(test_ck_tile_universal_gemm_argument_validation
test_universal_gemm_argument_validation.cpp)
list(APPEND CK_TILE_GEMM_TEST_TARGETS
test_ck_tile_universal_gemm_argument_validation)

add_gtest_executable(test_ck_tile_gemm_pipeline_mem test_gemm_pipeline_mem.cpp)
add_gtest_executable(test_ck_tile_gemm_pipeline_compv3 test_gemm_pipeline_compv3.cpp)
add_gtest_executable(test_ck_tile_gemm_pipeline_compv4 test_gemm_pipeline_compv4.cpp)
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,107 @@
// Copyright (c) Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT

#include <gtest/gtest.h>

#include "ck_tile/core.hpp"
#include "ck_tile/ops/elementwise.hpp"
#include "ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp"

namespace {

using Row = ck_tile::tensor_layout::gemm::RowMajor;
using Col = ck_tile::tensor_layout::gemm::ColumnMajor;

struct MockBlockGemmShape
{
static constexpr ck_tile::index_t kclusterM = 1;
static constexpr ck_tile::index_t kclusterN = 1;
static constexpr ck_tile::index_t kclusterK = 1;
static constexpr bool PermuteA = false;

struct WarpTile
{
template <typename Index>
static constexpr ck_tile::index_t at(Index)
{
return 16;
}
};
};

struct MockTilePartitioner
{
using BlockGemmShape = MockBlockGemmShape;

static constexpr ck_tile::index_t MPerBlock = 64;
static constexpr ck_tile::index_t NPerBlock = 64;
static constexpr ck_tile::index_t KPerBlock = 64;
};

struct MockGemmPipeline
{
using AsDataType = ck_tile::tuple<float>;
using BsDataType = ck_tile::tuple<float>;
using AsLayout = ck_tile::tuple<Row>;
using BsLayout = ck_tile::tuple<Col>;
using CLayout = Row;
using AElementWise = ck_tile::element_wise::PassThrough;
using BElementWise = ck_tile::element_wise::PassThrough;
using BlockGemmShape = MockBlockGemmShape;

static constexpr ck_tile::index_t BlockSize = 64;
static constexpr bool kPadM = false;
static constexpr bool kPadN = false;
static constexpr bool kPadK = false;

template <bool>
static constexpr ck_tile::index_t GetVectorSizeA()
{
return 1;
}

template <bool>
static constexpr ck_tile::index_t GetVectorSizeB()
{
return 1;
}
};

struct MockEpiloguePipeline
{
using DsDataType = ck_tile::tuple<float>;
using DsLayout = ck_tile::tuple<Col>;
using ODataType = float;

static constexpr ck_tile::index_t GetVectorSizeC() { return 1; }

template <typename Index>
static constexpr ck_tile::index_t GetVectorSizeD(Index)
{
return 1;
}
};

using Kernel =
ck_tile::UniversalGemmKernel<MockTilePartitioner, MockGemmPipeline, MockEpiloguePipeline>;

TEST(UniversalGemmArgumentValidation, ReportsMismatchedDTensorLayout)
{
Kernel::KernelArgs args{};
args.M = 64;
args.N = 64;
args.K = 64;
args.k_batch = 1;

ck_tile::UpdateEnvVar(CK_TILE_ENV(CK_TILE_LOGGING), true);
testing::internal::CaptureStderr();
const bool supported = Kernel::IsSupportedArgument(args);
const std::string diagnostic = testing::internal::GetCapturedStderr();
ck_tile::EnvUnset(CK_TILE_ENV(CK_TILE_LOGGING));

EXPECT_FALSE(supported);
EXPECT_NE(diagnostic.find("D tensor layout must match the C layout: D0"), std::string::npos)
<< diagnostic;
}

} // namespace
Loading