mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-20 12:59:49 +00:00
75bd80b1be
* Add cshuffle epilogue test
* add the poc implementation to the epilogue and tests
* refactor cshuffle epilogue
* WIP: adding tensor/tile usage to scale_tile
* fix usage of tile_elementwise_inout
* add gemm_quant_kernel for generalizing gemm quant kernel
* Add problem specific to different quants, add QuantType to Traits
* Add quant_type to quant_kernel template parameters
* Create aq/bq_block_windows and views depending on QuantType
* Use tile windows as inputs in cshuffle epilogue
* Fix some issues in epilogue
* initial new example code for new general gemm quant kernel test
* Fix issues in kernel
* Add verification check for rowcol Quantmode
* use AccDataType instead of AQ in pipeline
* fix aquant preshuffle
* fix formatting
* some cleanup
* remove gemm_aquant_basic.cpp
* remove gemm_aquant_kernel.hpp
* fix tests for the renamed quant kernel
* fix formatting
* clean example files
* fix some merge conflicts
* fix preshufflequant rename issue
* fix some templates after merging with develop
* fix test preshuffle parameter
* fix formatting
* Unify bquant kernel to the common quant kernel
* remove bquant kernel also from common header
* fix formatting
* clean up commented code
* fix formatting config hpp
* fix merge mistake
* Non-const for movable windows
* fix formatting
* Fix grammar in README
Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>
* Remove #include<bit> and clean up example
* fix strides
* Add some descriptions for move_windows
---------
Co-authored-by: Mohsen Saffari <mohsen.saffari@amd.com>
Co-authored-by: spolifroni-amd <Sandra.Polifroni@amd.com>
[ROCm/composable_kernel commit: c6010f2953]
85 lines
3.2 KiB
C++
85 lines
3.2 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#include "test_cshuffle_epilogue_util.hpp"
|
|
#include <gtest/gtest.h>
|
|
#include <hip/hip_runtime.h>
|
|
|
|
using namespace ck_tile;
|
|
|
|
class CShuffleEpilogueTest : public ::testing::Test
|
|
{
|
|
protected:
|
|
void SetUp() override {}
|
|
};
|
|
|
|
TEST_F(CShuffleEpilogueTest, BasicHalfTest)
|
|
{
|
|
// Basic test configuration with half_t data types
|
|
using ADataType = ck_tile::half_t;
|
|
using BDataType = ck_tile::half_t;
|
|
using AccDataType = float;
|
|
using ODataType = ck_tile::half_t;
|
|
|
|
constexpr index_t kMPerBlock = 256;
|
|
constexpr index_t kNPerBlock = 256;
|
|
constexpr index_t MWave = 2;
|
|
constexpr index_t NWave = 2;
|
|
constexpr index_t MPerXdl = 32;
|
|
constexpr index_t NPerXdl = 32;
|
|
constexpr index_t KPerXdl = 8;
|
|
|
|
using TestProblem = SimpleCShuffleEpilogueProblem<ADataType,
|
|
BDataType,
|
|
AccDataType,
|
|
ODataType,
|
|
kMPerBlock,
|
|
kNPerBlock,
|
|
MWave,
|
|
NWave,
|
|
MPerXdl,
|
|
NPerXdl,
|
|
KPerXdl>;
|
|
|
|
bool result = run_cshuffle_epilogue_test<TestProblem, kMPerBlock, kNPerBlock>();
|
|
EXPECT_TRUE(result) << "Basic CShuffleEpilogue test failed";
|
|
}
|
|
|
|
TEST_F(CShuffleEpilogueTest, BasicHalfTestWithScale)
|
|
{
|
|
// Basic test configuration with half_t data types
|
|
using ADataType = ck_tile::half_t;
|
|
using BDataType = ck_tile::half_t;
|
|
using AccDataType = float;
|
|
using ODataType = ck_tile::half_t;
|
|
|
|
constexpr index_t kMPerBlock = 256;
|
|
constexpr index_t kNPerBlock = 256;
|
|
constexpr index_t MWave = 2;
|
|
constexpr index_t NWave = 2;
|
|
constexpr index_t MPerXdl = 32;
|
|
constexpr index_t NPerXdl = 32;
|
|
constexpr index_t KPerXdl = 8;
|
|
|
|
using TestProblem = SimpleCShuffleEpilogueProblem<ADataType,
|
|
BDataType,
|
|
AccDataType,
|
|
ODataType,
|
|
kMPerBlock,
|
|
kNPerBlock,
|
|
MWave,
|
|
NWave,
|
|
MPerXdl,
|
|
NPerXdl,
|
|
KPerXdl>;
|
|
|
|
bool result = run_cshuffle_epilogue_test<TestProblem, kMPerBlock, kNPerBlock>(true);
|
|
EXPECT_TRUE(result) << "Scale CShuffleEpilogue test failed";
|
|
}
|
|
|
|
int main(int argc, char** argv)
|
|
{
|
|
::testing::InitGoogleTest(&argc, argv);
|
|
return RUN_ALL_TESTS();
|
|
}
|