mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-11 17:00:18 +00:00
6a6163a3d1
* Sync the order of type string with template parameter * Add more instances * Check the vector size and remove redundant var * Extract var to static, prepare to separate sweep once kernel * Separate sweeponce flow and optimize the flow * 1. Rename AccDatatype in normalization to computeData 2. Rename AccElementwiseOperation to YElementwiseOperation in normalization * Remove useless code * Update naive variance kernel * Refine string * Fix typo * Support naive variance for device_normalization * Check the blocksize * Share the VGPR of x and y * Share the VGPR of gamma and beta * Add more instances * Support fp16 sqrt for experiment * Add CHANGELOG * Fix typo * clang-format
46 lines
1.6 KiB
C++
46 lines
1.6 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#include "gtest/gtest.h"
|
|
#include "profiler/profile_layernorm_impl.hpp"
|
|
|
|
using F16 = ck::half_t;
|
|
using F32 = float;
|
|
using ck::index_t;
|
|
|
|
template <typename Tuple>
|
|
class TestLayernorm2d : public ::testing::Test
|
|
{
|
|
protected:
|
|
using XDataType = std::tuple_element_t<0, Tuple>;
|
|
using GammaDataType = std::tuple_element_t<1, Tuple>;
|
|
using BetaDataType = std::tuple_element_t<2, Tuple>;
|
|
using ComputeDataType = std::tuple_element_t<3, Tuple>;
|
|
using YDataType = std::tuple_element_t<4, Tuple>;
|
|
|
|
void Run()
|
|
{
|
|
// [N, D], reduce D
|
|
std::vector<std::vector<ck::index_t>> lengths = {
|
|
{4, 256}, {8, 511}, {9, 1032}, {4, 2048}, {1, 8192}, {4000, 2000}};
|
|
|
|
for(auto length : lengths)
|
|
{
|
|
bool success = ck::profiler::profile_layernorm_impl<XDataType,
|
|
GammaDataType,
|
|
BetaDataType,
|
|
ComputeDataType,
|
|
YDataType,
|
|
2>(true, 2, false, false, length);
|
|
EXPECT_TRUE(success);
|
|
}
|
|
}
|
|
};
|
|
|
|
using KernelTypes = ::testing::Types<
|
|
// XDataType, GammaDataType, BetaDataType, ComputeDataType, YDataType>
|
|
std::tuple<F32, F32, F32, F32, F32>>;
|
|
|
|
TYPED_TEST_SUITE(TestLayernorm2d, KernelTypes);
|
|
TYPED_TEST(TestLayernorm2d, Test_FP32) { this->Run(); }
|