mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 10:09:41 +00:00
* added working example for 5D input using 1D kernel
* example with 5D input tensor and 2d kernel - not working: issues with arguments
* added updated version of 3d device op - changed descriptors/dims
* added example file to check kernel
* fixed descriptor and isSupportedArgument stride problem
* added and modified kernel for 3d - updated tids/loop
* adding some more 5d example files
* fixed some issues
* changes made for testing
* working version: fixed error in stride for A, still a bit inefficient
* cleaned up formatting/comments
* updating formatting
* more formatting fixes
* fixing cmake, adding back gpu targets in cmake script
* adding client example
* added instances for client example
* fixed errors in client example
* implemented client ex with device_elementwise.hpp and device_elementwise_3d_impl.hpp
* removed extra files
* minor formatting and naming fixes
* adding test files and profiler
* fixing minor error
* minor fix
* removed unneccesary comments, renamed files
* updated instance list for client example, added different layout example
* removing instances
* fixed error in instance generation
* remove comments
* update profiler and client example tensor layouts
* fixed errors in test/profiler
* updated vector dim access to enable vector load
* updated test/profiler files
* updated example with 1d kernel
* updating profiler
* renamed files
* disabled device op for MI300
* skip elementwise_permute_2d on gfx94x
* Update CMakeLists.txt
* fixing CMake - disabling some GPU targets
* added transpose profiler to CMake
* fixed transpose profiler errors
* fixed instances for tests/profiler
* cleaned up code in transpose profiler source code
* added some comments, updated copyright
* made function arguments const where possible
---------
Co-authored-by: Jing Zhang <jizha@amd.com>
Co-authored-by: Jing Zhang <jizhan@amd.com>
Co-authored-by: zjing14 <zhangjing14@gmail.com>
[ROCm/composable_kernel commit: aa3e2d7967]
36 lines
1.1 KiB
C++
36 lines
1.1 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2018-2022, Advanced Micro Devices, Inc. All rights reserved.
|
|
#include "gtest/gtest.h"
|
|
#include "profiler/profile_transpose_impl.hpp"
|
|
|
|
using F16 = ck::half_t;
|
|
using F32 = float;
|
|
using ck::index_t;
|
|
|
|
template <typename Tuple>
|
|
class TestTranspose : public ::testing::Test
|
|
{
|
|
protected:
|
|
using ADataType = std::tuple_element_t<0, Tuple>;
|
|
using BDataType = std::tuple_element_t<1, Tuple>;
|
|
|
|
void Run()
|
|
{
|
|
std::vector<std::vector<ck::index_t>> lengths = {
|
|
{4, 16, 16, 32, 5}, {8, 16, 16, 32, 8} /**{32, 16, 16, 32, 8},**/};
|
|
|
|
for(auto length : lengths)
|
|
{
|
|
bool success = ck::profiler::profile_transpose_impl<ADataType, BDataType, 5>(
|
|
true, 2, false, false, length);
|
|
EXPECT_TRUE(success);
|
|
}
|
|
}
|
|
};
|
|
|
|
using KernelTypes = ::testing::Types<std::tuple<F16, F16>, std::tuple<F32, F32>>;
|
|
|
|
TYPED_TEST_SUITE(TestTranspose, KernelTypes);
|
|
TYPED_TEST(TestTranspose, Test_FP16) { this->Run(); }
|
|
TYPED_TEST(TestTranspose, Test_FP32) { this->Run(); }
|