ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 09:08:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge branch 'develop' into amd-develop

Browse Source
This commit is contained in:
Jun Liu
2024-08-12 15:13:41 -07:00
parent f2ba10b826 cbb6f2ab8c
commit e2eb041891
22 changed files with 430 additions and 644 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
CMakeLists.txt
View File

@@ -541,6 +541,9 @@ if(NOT DEFINED INSTANCES_ONLY)
PACKAGE_NAME examples
)
add_subdirectory(example)
if(GPU_TARGETS MATCHES "gfx9" AND NOT INSTANCES_ONLY)
add_subdirectory(codegen)
endif()
if(BUILD_TESTING)
add_subdirectory(test)
endif()

31
Jenkinsfile vendored
View File

@@ -746,10 +746,6 @@ pipeline {
name: "RUN_PERFORMANCE_TESTS",
defaultValue: true,
description: "Run the performance tests (default: ON)")
booleanParam(
name: "RUN_CODEGEN_TESTS",
defaultValue: true,
description: "Run the codegen tests (default: ON)")
booleanParam(
name: "RUN_CK_TILE_TESTS",
defaultValue: false,
@@ -841,33 +837,6 @@ pipeline {
}
}
}
stage("Run Codegen Tests")
{
parallel
{
stage("Run Codegen Tests on gfx90a")
{
when {
beforeAgent true
expression { params.RUN_CODEGEN_TESTS.toBoolean() }
}
agent{ label rocmnode("gfx90a")}
environment{
setup_args = "NO_CK_BUILD"
execute_args = """ cd ../codegen && rm -rf build && mkdir build && cd build && \
cmake -D CMAKE_PREFIX_PATH=/opt/rocm \
-D CMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ \
-D CMAKE_BUILD_TYPE=Release \
-D GPU_TARGETS="gfx90a" \
-DCMAKE_CXX_FLAGS=" -O3 " .. && make -j check"""
}
steps{
buildHipClangJobAndReboot(setup_args:setup_args, no_reboot:true, build_type: 'Release', execute_cmd: execute_args)
cleanWs()
}
}
}
}
stage("Run CK_TILE Tests")
{
parallel

18
codegen/CMakeLists.txt
View File

@@ -1,6 +1,3 @@
cmake_minimum_required(VERSION 3.16)
project(composable_kernel_host LANGUAGES CXX HIP)
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
@@ -8,17 +5,9 @@ set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin)
set(CK_ROOT ${CMAKE_CURRENT_SOURCE_DIR}/..)
find_package(ROCM)
include(ROCMInstallTargets)
include(ROCMTest)
add_compile_options(-std=c++17)
find_package(hip)
## HIP
set(CMAKE_HIP_PLATFORM amd)
set(CMAKE_HIP_COMPILER ${CMAKE_CXX_COMPILER})
set(CMAKE_HIP_EXTENSIONS ON)
message("CMAKE_HIP_COMPILER: ${CMAKE_HIP_COMPILER}")
add_custom_target(codegen)
# add include directories
include_directories(BEFORE
@@ -32,8 +21,9 @@ list(APPEND CMAKE_MODULE_PATH ${CK_ROOT}/cmake)
include(Embed)
file(GLOB_RECURSE KERNEL_FILES CONFIGURE_DEPENDS
${CK_ROOT}/include/ck/*.hpp)
message(STATUS "KERNEL_FILES: ${KERNEL_FILES}")
message(STATUS "RELATIVE: ${CK_ROOT}/include")
#printouts fot debug purposes
#message(STATUS "KERNEL_FILES: ${KERNEL_FILES}")
#message(STATUS "RELATIVE: ${CK_ROOT}/include")
add_embed_library(ck_headers ${KERNEL_FILES} RELATIVE ${CK_ROOT}/include)
file(GLOB SOURCES CONFIGURE_DEPENDS src/*.cpp)

3
codegen/include/ck/host/types.hpp
View File

@@ -76,8 +76,11 @@ std::string SequenceStr(const std::vector<int>& v);
std::string MakeTuple(const std::vector<std::string>& v);
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
template <int... xs>
const std::string S = SequenceStr({xs...});
#pragma clang diagnostic pop
constexpr const char* PassThrough = "ck::tensor_operation::element_wise::PassThrough";
constexpr const char* Bilinear = "ck::tensor_operation::element_wise::Bilinear";

13
codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp
View File

@@ -3,6 +3,7 @@
#include "ck/host/device_gemm_multiple_d/operation.hpp"
#include "ck/host/stringutils.hpp"
#include "ck/host/types.hpp"
#include "ck/host/utils.hpp"
#include <cassert>
@@ -32,11 +33,11 @@ static std::string GetGemmSpec(const std::size_t m,
}
// function to update prologue/epilogue with user provided operation
void Operation_Xdl_CShuffle::update_prologue(const std::string& prologue)
void Operation_Xdl_CShuffle::update_prologue(const std::string& pro)
{
if(!prologue.empty())
if(!pro.empty())
{
this->prologue = prologue;
this->prologue = pro;
this->cde_elem_op = "CDEElementOp";
}
else
@@ -45,11 +46,11 @@ void Operation_Xdl_CShuffle::update_prologue(const std::string& prologue)
}
}
void Operation_Xdl_CShuffle::update_epilogue(const std::string& epilogue)
void Operation_Xdl_CShuffle::update_epilogue(const std::string& epi)
{
if(!epilogue.empty())
if(!epi.empty())
{
this->epilogue = epilogue;
this->epilogue = epi;
this->cde_elem_op = "CDEElementOp";
}
else

36
codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp
View File

@@ -4,6 +4,7 @@
#include "ck/host/device_grouped_conv_fwd_multiple_d/conv_fwd_op.hpp"
#include <iostream>
#include "ck/host/stringutils.hpp"
#include "ck/host/types.hpp"
#include "ck/host/utils.hpp"
#include <cassert>
@@ -11,34 +12,15 @@ namespace ck {
namespace host {
namespace conv {
// calculate appropriate Gemm Specification based on input tensor dimensions
// NOTE: in CK, MNKPadding is always used for forward convolution
static std::string GetGemmSpec(const std::size_t m,
const std::size_t n,
const std::size_t k,
const std::size_t m_per_block,
const std::size_t n_per_block,
const std::size_t k_per_block)
{
std::string spec = "";
if(integer_divide_ceil(m, m_per_block) * m_per_block - m != 0)
spec += "M";
if(integer_divide_ceil(n, n_per_block) * n_per_block - n != 0)
spec += "N";
if(integer_divide_ceil(k, k_per_block) * k_per_block - k != 0)
spec += "K";
if(spec == "")
return "ck::tensor_operation::device::GemmSpecialization::Default";
return "ck::tensor_operation::device::GemmSpecialization::" + spec + "Padding";
}
// NOTE: in CK, MNKPadding is always used for forward convolution, so didn't
// add GemmSpec function here
// function to update prologue/epilogue with user provided operation
void Operation_Conv_Fwd_Xdl_Cshuffle::update_prologue(const std::string& prologue)
void Operation_Conv_Fwd_Xdl_Cshuffle::update_prologue(const std::string& pro)
{
if(!prologue.empty())
if(!pro.empty())
{
this->prologue = prologue;
this->prologue = pro;
this->cde_elem_op = "CDEElementOp";
}
else
@@ -47,11 +29,11 @@ void Operation_Conv_Fwd_Xdl_Cshuffle::update_prologue(const std::string& prologu
}
}
void Operation_Conv_Fwd_Xdl_Cshuffle::update_epilogue(const std::string& epilogue)
void Operation_Conv_Fwd_Xdl_Cshuffle::update_epilogue(const std::string& epi)
{
if(!epilogue.empty())
if(!epi.empty())
{
this->epilogue = epilogue;
this->epilogue = epi;
this->cde_elem_op = "CDEElementOp";
}
else

3
codegen/src/headers.cpp
View File

@@ -4,7 +4,10 @@
namespace ck {
namespace host {
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
const std::string config_header = "";
#pragma clang diagnostic pop
std::unordered_map<std::string_view, std::string_view> GetHeaders()
{

4
codegen/test/CMakeLists.txt
View File

@@ -4,7 +4,9 @@ file(GLOB TEST_SRCS CONFIGURE_DEPENDS *.cpp)
foreach(TEST_SRC ${TEST_SRCS})
set_source_files_properties(${TEST_SRC} PROPERTIES LANGUAGE HIP)
get_filename_component(BASE_NAME ${TEST_SRC} NAME_WE)
rocm_add_test_executable(test_host_${BASE_NAME} ${TEST_SRC})
add_executable(test_host_${BASE_NAME} ${TEST_SRC})
add_dependencies(codegen test_host_${BASE_NAME})
add_test(NAME codegen_test_${BASE_NAME} COMMAND test_host_${BASE_NAME})
target_link_libraries(test_host_${BASE_NAME} ck_rtc ck_host)
# target_link_libraries(test_host_${BASE_NAME} ${CK_ROOT}/build/lib/libutility.a)
target_include_directories(test_host_${BASE_NAME} PUBLIC include())

2
codegen/test/grouped_conv_fwd_multiple_d_v1.cpp
View File

@@ -92,7 +92,6 @@ struct Epilogue
static_cast<int>(prob.C),
static_cast<int>(prob.Y),
static_cast<int>(prob.X)};
ck::Array<ck::index_t, 5> d_lengths = {};
ck::Array<ck::index_t, 5> in_strides{static_cast<int>(prob.C),
static_cast<int>(prob.Hi * prob.Wi * prob.G * prob.C),
@@ -109,7 +108,6 @@ struct Epilogue
1,
static_cast<int>(prob.X * prob.C),
static_cast<int>(prob.C)};
ck::Array<ck::index_t, 5> d_strides = {};
ck::Array<ck::index_t, 2> conv_filter_strides = {2, 2};
ck::Array<ck::index_t, 2> conv_filter_dilations = {1, 1};

2
codegen/test/grouped_conv_fwd_multiple_d_v2.cpp
View File

@@ -92,7 +92,6 @@ struct Epilogue
static_cast<int>(prob.C),
static_cast<int>(prob.Y),
static_cast<int>(prob.X)};
ck::Array<ck::index_t, 5> d_lengths = {};
ck::Array<ck::index_t, 5> in_strides{static_cast<int>(prob.C),
static_cast<int>(prob.Hi * prob.Wi * prob.G * prob.C),
@@ -109,7 +108,6 @@ struct Epilogue
1,
static_cast<int>(prob.X * prob.C),
static_cast<int>(prob.C)};
ck::Array<ck::index_t, 5> d_strides = {};
ck::Array<ck::index_t, 2> conv_filter_strides = {1, 1};
ck::Array<ck::index_t, 2> conv_filter_dilations = {1, 1};

2
codegen/test/grouped_conv_fwd_multiple_d_v3.cpp
View File

@@ -92,7 +92,6 @@ struct Epilogue
static_cast<int>(prob.C),
static_cast<int>(prob.Y),
static_cast<int>(prob.X)};
ck::Array<ck::index_t, 5> d_lengths = {};
ck::Array<ck::index_t, 5> in_strides{static_cast<int>(prob.C),
static_cast<int>(prob.Hi * prob.Wi * prob.G * prob.C),
@@ -109,7 +108,6 @@ struct Epilogue
1,
static_cast<int>(prob.X * prob.C),
static_cast<int>(prob.C)};
ck::Array<ck::index_t, 5> d_strides = {};
ck::Array<ck::index_t, 2> conv_filter_strides = {2, 2};
ck::Array<ck::index_t, 2> conv_filter_dilations = {1, 1};

2
codegen/test/grouped_conv_fwd_multiple_d_v4.cpp
View File

@@ -92,7 +92,6 @@ struct Epilogue
static_cast<int>(prob.C),
static_cast<int>(prob.Y),
static_cast<int>(prob.X)};
ck::Array<ck::index_t, 5> d_lengths = {};
ck::Array<ck::index_t, 5> in_strides{static_cast<int>(prob.C),
static_cast<int>(prob.Hi * prob.Wi * prob.G * prob.C),
@@ -109,7 +108,6 @@ struct Epilogue
1,
static_cast<int>(prob.X * prob.C),
static_cast<int>(prob.C)};
ck::Array<ck::index_t, 5> d_strides = {};
ck::Array<ck::index_t, 2> conv_filter_strides = {1, 1};
ck::Array<ck::index_t, 2> conv_filter_dilations = {1, 1};

2
codegen/test/rtc/src/kernel.cpp
View File

@@ -118,4 +118,4 @@ void kernel::launch(hipStream_t stream,
launch_kernel(impl->fun, stream, global, local, kernargs.data(), size);
}
} // namespace rtc
} // namespace rtc

2
codegen/test/rtc/src/tmp_dir.cpp
View File

@@ -45,4 +45,4 @@ void tmp_dir::execute(const std::string& cmd) const
tmp_dir::~tmp_dir() { std::filesystem::remove_all(this->path); }
} // namespace rtc
} // namespace rtc

4
example/CMakeLists.txt
View File

@@ -75,7 +75,7 @@ function(add_example_executable EXAMPLE_NAME FILE_NAME)
#only continue if there are some source files left on the list
if(FILE_NAME)
if(FILE_NAME MATCHES "_xdl")
list(REMOVE_ITEM EX_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103)
list(REMOVE_ITEM EX_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201)
elseif(FILE_NAME MATCHES "_wmma")
list(REMOVE_ITEM EX_TARGETS gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030)
endif()
@@ -162,7 +162,7 @@ function(add_example_executable_no_testing EXAMPLE_NAME FILE_NAME)
#only continue if there are some source files left on the list
if(FILE_NAME)
if(FILE_NAME MATCHES "_xdl")
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103)
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201)
elseif(FILE_NAME MATCHES "_wmma")
list(REMOVE_ITEM EX_TARGETS gfx900 gfx906 gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030)
endif()

18
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp
View File

@@ -86,7 +86,6 @@ __global__ void
const AElementwiseOperation a_element_op,
const BElementwiseOperation b_element_op,
const CDEElementwiseOperation cde_element_op,
const index_t groups_count,
const AGridDesc_AK0_M_AK1 a_grid_desc_k0_m_k1,
const BGridDesc_BK0_N_BK1 b_grid_desc_k0_n_k1,
const DsGridDescriptor_MBlock_MPerBlock_NBlock_NPerBlock
@@ -101,10 +100,8 @@ __global__ void
defined(__gfx94__))
// offset base pointer for each work-group
const index_t num_blocks_per_batch = __builtin_amdgcn_readfirstlane(gridDim.y / groups_count);
const index_t& num_blocks_per_n = groups_count;
const index_t g_idx = __builtin_amdgcn_readfirstlane(blockIdx.y / num_blocks_per_batch);
const index_t n_idx = __builtin_amdgcn_readfirstlane(blockIdx.y / num_blocks_per_n);
const index_t g_idx = __builtin_amdgcn_readfirstlane(blockIdx.y);
const index_t n_idx = __builtin_amdgcn_readfirstlane(blockIdx.z);
const long_index_t e_batch_offset =
amd_wave_read_first_lane(compute_ptr_offset_of_groups.GetEPtrOffset(g_idx));
@@ -200,7 +197,6 @@ __global__ void
ignore = p_bs_grid;
ignore = p_ds_grid;
ignore = p_e_grid;
ignore = groups_count;
ignore = a_grid_desc_k0_m_k1;
ignore = b_grid_desc_k0_n_k1;
ignore = ds_grid_desc_mblock_mperblock_nblock_nperblock;
@@ -321,8 +317,8 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
using ConvToGemmFwdTransformer = TransformConvFwdToGemm<NDimSpatial,
ConvForwardSpecialization,
true /*SplitN*/,
ALayout,
ELayout>;
ADataType,
EDataType>;
static constexpr auto matrix_padder =
MatrixPadder<GemmSpec, index_t, index_t, index_t>{MPerBlock, NPerBlock, KPerBlock};
@@ -730,8 +726,8 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
arg.a_g_n_c_wis_lengths_[I1] / arg.conv_N_per_block_;
const index_t gdx = arg.block_2_etile_map_.CalculateGridSize(arg.e_grid_desc_m_n_);
const index_t gdy = arg.num_group_ * num_workgroups_per_Conv_N;
const index_t gdz = 1;
const index_t gdy = arg.num_group_;
const index_t gdz = num_workgroups_per_Conv_N;
const auto K =
arg.a_grid_desc_ak0_m_ak1_.GetLength(I0) * arg.a_grid_desc_ak0_m_ak1_.GetLength(I2);
@@ -780,7 +776,6 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
arg.a_element_op_,
arg.b_element_op_,
arg.cde_element_op_,
arg.a_g_n_c_wis_lengths_[0], // Group count
as_grid_desc_ak0_m_ak1,
bs_grid_desc_bk0_n_bk1,
arg.ds_grid_desc_mblock_mperblock_nblock_nperblock_,
@@ -824,7 +819,6 @@ struct DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle
arg.a_element_op_,
arg.b_element_op_,
arg.cde_element_op_,
arg.a_g_n_c_wis_lengths_[0], // Group count
arg.a_grid_desc_ak0_m_ak1_,
arg.b_grid_desc_bk0_n_bk1_,
arg.ds_grid_desc_mblock_mperblock_nblock_nperblock_,

4
library/src/tensor_operation_instance/gpu/CMakeLists.txt
View File

@@ -81,11 +81,11 @@ function(add_instance_library INSTANCE_NAME)
set(INST_TARGETS ${GPU_TARGETS})
endif()
if(source MATCHES "_xdl")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103)
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201)
elseif(ARGN MATCHES "_wmma")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030)
elseif(ARGN MATCHES "mha")
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103)
list(REMOVE_ITEM INST_TARGETS gfx900 gfx906 gfx908 gfx90a gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201)
endif()
set(offload_targets)
foreach(target IN LISTS INST_TARGETS)

63
script/test_reduce_with_index.sh
View File

@@ -1,63 +0,0 @@
#!/bin/bash
## The following will be used for CI
set -x
## for float
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2,3 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,3 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,2,3 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1,2,3 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 2 0 2
bin/test_reduce_with_index -D 64,4,280,82 -R 3 0 2
## for float64
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2,3 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,3 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,2,3 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1,2,3 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 2 6 2
bin/test_reduce_with_index -D 64,4,280,82 -R 3 6 2
## for float16
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2,3 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,3 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,2,3 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1,2,3 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 2 1 2
bin/test_reduce_with_index -D 64,4,280,82 -R 3 1 2
## for int8_t
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2,3 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,3 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,2,3 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1,2,3 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 2 3 2
bin/test_reduce_with_index -D 64,4,280,82 -R 3 3 2
## for bfloat16
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2,3 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,2 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,1,3 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0,2,3 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1,2,3 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 0 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 1 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 2 5 2
bin/test_reduce_with_index -D 64,4,280,82 -R 3 5 2
set +x

8
test/CMakeLists.txt
View File

@@ -68,11 +68,11 @@ function(add_test_executable TEST_NAME)
#only continue if there are some source files left on the list
if(ARGN)
if(ARGN MATCHES "_xdl")
list(REMOVE_ITEM TEST_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103)
list(REMOVE_ITEM TEST_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201)
elseif(ARGN MATCHES "_wmma")
list(REMOVE_ITEM TEST_TARGETS gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030)
elseif(ARGN MATCHES "_smfmac")
list(REMOVE_ITEM TEST_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx908 gfx90a)
list(REMOVE_ITEM TEST_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx908 gfx90a gfx1200 gfx1201)
endif()
set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
add_executable(${TEST_NAME} ${ARGN})
@@ -149,11 +149,11 @@ function(add_gtest_executable TEST_NAME)
#only continue if there are some source files left on the list
if(ARGN)
if(ARGN MATCHES "_xdl")
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103)
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx1200 gfx1201)
elseif(ARGN MATCHES "_wmma")
list(REMOVE_ITEM TEST_TARGETS gfx900 gfx906 gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030)
elseif(ARGN MATCHES "_smfmac")
list(REMOVE_ITEM TEST_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx908 gfx90a)
list(REMOVE_ITEM TEST_TARGETS gfx1030 gfx1100 gfx1101 gfx1102 gfx1103 gfx908 gfx90a gfx1200 gfx1201)
endif()
set_source_files_properties(${ARGN} PROPERTIES LANGUAGE HIP)
add_executable(${TEST_NAME} ${ARGN})

4
test/reduce/CMakeLists.txt
View File

@@ -1,5 +1,5 @@
add_test_executable(test_reduce_no_index reduce_no_index.cpp)
add_test_executable(test_reduce_with_index reduce_with_index.cpp)
add_gtest_executable(test_reduce_no_index reduce_no_index.cpp)
add_gtest_executable(test_reduce_with_index reduce_with_index.cpp)
target_link_libraries(test_reduce_no_index PRIVATE utility device_reduce_instance)
target_link_libraries(test_reduce_with_index PRIVATE utility device_reduce_instance)

425
test/reduce/reduce_no_index.cpp
View File

@@ -1,248 +1,203 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <getopt.h>
#include "ck/library/utility/host_common_util.hpp"
#include "profiler/profile_reduce_impl.hpp"
#include <gtest/gtest.h>
using namespace ck;
static struct option long_options[] = {{"inLengths", required_argument, nullptr, 'D'},
{"reduceDimensions", required_argument, nullptr, 'R'},
{"scales", required_argument, nullptr, 'S'},
{"help", no_argument, nullptr, '?'},
{nullptr, 0, nullptr, 0}};
class SimpleAppArgs
struct ReduceParam
{
private:
int option_index = 0;
public:
std::vector<size_t> inLengths;
std::vector<int> reduceDims;
std::vector<float> scales;
int data_type;
int init_method = 1;
public:
void show_usage(const char* cmd)
{
std::cout << "Usage of " << cmd << std::endl;
std::cout << "--inLengths or -D, comma separated list of input tensor dimension lengths "
"(only 4-d tensor supported)"
<< std::endl;
std::cout << "--reduceDimensions or -R comma seperated list of dimension indexes to reduce "
"(only 1 or 3 or 4 dimensions supported)"
<< std::endl;
std::cout << "--scales or -S, comma separated two float values for alpha and beta"
<< std::endl;
std::cout << "Arg1 -- data type (0: fp16, 1: fp32, 3: int8, 5: bp16, 6: fp64)" << std::endl;
std::cout << "Arg2 -- init method(0=no init, 1=single integer value, 2=scope integer "
"value, 3=decimal value)"
<< std::endl;
};
int processArgs(int argc, char* argv[])
{
using ck::host_common::getTypeValuesFromString;
int ch;
while(1)
{
ch = getopt_long(argc, argv, "D:R:S:", long_options, &option_index);
if(ch == -1)
break;
switch(ch)
{
case 'D':
if(!optarg)
throw std::runtime_error("Invalid option format!");
inLengths = getTypeValuesFromString<size_t>(optarg);
break;
case 'R':
if(!optarg)
throw std::runtime_error("Invalid option format!");
reduceDims = getTypeValuesFromString<int>(optarg);
break;
case 'S':
if(!optarg)
throw std::runtime_error("Invalid option format!");
scales = getTypeValuesFromString<float>(optarg);
break;
case '?':
if(std::string(long_options[option_index].name) == "help")
{
show_usage(argv[0]);
return (-1);
};
break;
default: show_usage(argv[0]); return (-1);
};
};
if(optind + 2 > argc)
throw std::runtime_error("Invalid cmd-line arguments, more argumetns are needed!");
data_type = std::atoi(argv[optind++]);
init_method = std::atoi(argv[optind]);
if(scales.empty())
{
scales.push_back(1.0f);
scales.push_back(0.0f);
};
if(inLengths.size() != 4 ||
(reduceDims.size() != 1 && reduceDims.size() != 3 && reduceDims.size() != 4))
return (-1);
if(data_type != 0 && data_type != 1 && data_type != 3 && data_type != 5 && data_type != 6)
return (-1);
return (0);
};
bool do_verification{true};
bool propagateNan{false};
bool useIndex{false};
bool time_kernel{false};
bool do_dumpout{false};
int init_method{2};
float alpha{1.0f};
float beta{0.0f};
std::vector<size_t> inLengths{64, 4, 280, 82};
std::vector<int> reduceDims{0, 1, 2, 3};
};
bool test_reduce_no_index(int data_type,
int init_method,
std::vector<int> reduceDims,
std::vector<size_t> inLengths,
ReduceTensorOp reduceOpId,
bool propagateNan,
float alpha,
float beta)
std::vector<std::vector<int>> SetGenericReduceDim()
{
using ck::profiler::profile_reduce_impl;
bool result = true;
if(data_type == 0)
{
result = profile_reduce_impl<float, float, float>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
false,
alpha,
beta);
}
else if(data_type == 1)
{
result = profile_reduce_impl<ck::half_t, float, ck::half_t>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
false,
alpha,
beta);
}
else if(data_type == 3)
{
result = profile_reduce_impl<int8_t, int32_t, int8_t>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
false,
alpha,
beta);
}
else if(data_type == 5)
{
result = profile_reduce_impl<ck::bhalf_t, float, ck::bhalf_t>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
false,
alpha,
beta);
}
else if(data_type == 6)
{
result = profile_reduce_impl<double, double, double>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
false,
alpha,
beta);
}
return (result);
};
constexpr ReduceTensorOp reduceOpId = ReduceTensorOp::AVG;
constexpr bool propagateNan = false;
int main(int argc, char* argv[])
{
SimpleAppArgs args;
bool result = true;
if(argc == 1)
{
int data_type = 1;
int init_method = 2;
std::vector<size_t> inLengths{64, 4, 280, 80};
std::vector<std::vector<int>> v_reduceDims{
{0, 1, 2, 3}, {0, 1, 2}, {1, 2, 3}, {0, 1, 3}, {0, 2, 3}, {0}, {1}, {2}, {3}};
for(auto& reduceDims : v_reduceDims)
result = result && test_reduce_no_index(data_type,
init_method,
reduceDims,
inLengths,
reduceOpId,
propagateNan,
1.0f,
0.0f);
}
else
{
if(args.processArgs(argc, argv) < 0)
{
throw std::runtime_error(
"Invalid input arguments, test_reduce_no_index could not be executed!");
};
result = test_reduce_no_index(args.data_type,
args.init_method,
args.reduceDims,
args.inLengths,
reduceOpId,
propagateNan,
args.scales[0],
args.scales[1]);
}
std::cout << "test_reduce_no_index ..... " << (result ? "SUCCESS" : "FAILURE") << std::endl;
return (result ? 0 : -1);
return {{0, 1, 2, 3}, {0, 1, 2}, {0, 1, 3}, {0, 2, 3}, {1, 2, 3}, {0}, {1}, {2}, {3}};
}
template <typename T>
class ReduceWithIndexTest : public ::testing::Test
{
protected:
using InDataType = std::tuple_element_t<0, T>;
using AccDataType = std::tuple_element_t<1, T>;
using OutDataType = std::tuple_element_t<2, T>;
static std::vector<ReduceParam> params;
static void SetUpTestSuite()
{
// set testcase variables
ReduceParam set;
const auto setReduceDim = SetGenericReduceDim();
for(std::size_t i(0); i < setReduceDim.size(); ++i)
{
set.reduceDims = setReduceDim[i];
params.emplace_back(set);
}
}
template <ReduceTensorOp ReduceOpIdType>
void Run()
{
for(auto param : this->params)
{
bool success = ck::profiler::profile_reduce_impl<InDataType, AccDataType, OutDataType>(
param.do_verification,
param.init_method,
param.do_dumpout,
param.time_kernel,
param.inLengths,
param.reduceDims,
ReduceOpIdType,
param.propagateNan,
param.useIndex,
param.alpha,
param.beta);
EXPECT_TRUE(success);
}
}
};
template <typename T>
std::vector<ReduceParam> ReduceWithIndexTest<T>::params = {};
using Reduce_float_types = ::testing::Types<std::tuple<float, float, float>>;
using Reduce_double_types = ::testing::Types<std::tuple<double, double, double>>;
using Reduce_int8t_types = ::testing::Types<std::tuple<int8_t, int8_t, int8_t>>;
using Reduce_half_types = ::testing::Types<std::tuple<ck::half_t, ck::half_t, ck::half_t>>;
using Reduce_bhalf_float_Types = ::testing::Types<std::tuple<ck::bhalf_t, float, ck::bhalf_t>>;
template <typename TType>
class ReduceWithNoIndexFloat : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithNoIndexDouble : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithNoIndexInt8 : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithNoIndexHalf : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithNoIndexBHalfFloat : public ReduceWithIndexTest<TType>
{
};
TYPED_TEST_SUITE(ReduceWithNoIndexFloat, Reduce_float_types);
TYPED_TEST_SUITE(ReduceWithNoIndexDouble, Reduce_double_types);
TYPED_TEST_SUITE(ReduceWithNoIndexInt8, Reduce_int8t_types);
TYPED_TEST_SUITE(ReduceWithNoIndexHalf, Reduce_half_types);
TYPED_TEST_SUITE(ReduceWithNoIndexBHalfFloat, Reduce_bhalf_float_Types);
TYPED_TEST(ReduceWithNoIndexFloat, ReduceWithNoIndexTestFloat_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithNoIndexFloat, ReduceWithNoIndexTestFloat_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithNoIndexFloat, ReduceWithNoIndexTestFloat_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithNoIndexDouble, ReduceWithNoIndexTestDouble_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithNoIndexDouble, ReduceWithNoIndexTestDouble_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithNoIndexDouble, ReduceWithNoIndexTestDouble_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithNoIndexInt8, ReduceWithNoIndexTestInt8_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithNoIndexInt8, ReduceWithNoIndexTestInt8_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithNoIndexInt8, ReduceWithNoIndexTestInt8_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithNoIndexHalf, ReduceWithNoIndexTestHalf_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithNoIndexHalf, ReduceWithNoIndexTestHalf_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithNoIndexHalf, ReduceWithNoIndexTestHalf_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithNoIndexBHalfFloat, ReduceWithNoIndexTesBtHalfFloat_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithNoIndexBHalfFloat, ReduceWithNoIndexTestBHalfFloat_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithNoIndexBHalfFloat, ReduceWithNoIndexTestBHalfFloat_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}

425
test/reduce/reduce_with_index.cpp
View File

@@ -1,248 +1,203 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
// Copyright (c) 2018-2024, Advanced Micro Devices, Inc. All rights reserved.
#include <getopt.h>
#include "ck/library/utility/host_common_util.hpp"
#include "profiler/profile_reduce_impl.hpp"
#include <gtest/gtest.h>
using namespace ck;
static struct option long_options[] = {{"inLengths", required_argument, nullptr, 'D'},
{"reduceDimensions", required_argument, nullptr, 'R'},
{"scales", required_argument, nullptr, 'S'},
{"help", no_argument, nullptr, '?'},
{nullptr, 0, nullptr, 0}};
class SimpleAppArgs
struct ReduceParam
{
private:
int option_index = 0;
public:
std::vector<size_t> inLengths;
std::vector<int> reduceDims;
std::vector<float> scales;
int data_type;
int init_method = 1;
public:
void show_usage(const char* cmd)
{
std::cout << "Usage of " << cmd << std::endl;
std::cout << "--inLengths or -D, comma separated list of input tensor dimension lengths "
"(only 4-d tensor supported)"
<< std::endl;
std::cout << "--reduceDimensions or -R comma seperated list of dimension indexes to reduce "
"(only 1 or 3 or 4 dimensions supported)"
<< std::endl;
std::cout << "--scales or -S, comma separated two float values for alpha and beta"
<< std::endl;
std::cout << "Arg1 -- data type (1: fp32, 3: int8, 5: bp16, 6: fp64)" << std::endl;
std::cout << "Arg2 -- init method(0=no init, 1=single integer value, 2=scope integer "
"value, 3=decimal value)"
<< std::endl;
};
int processArgs(int argc, char* argv[])
{
using ck::host_common::getTypeValuesFromString;
int ch;
while(1)
{
ch = getopt_long(argc, argv, "D:R:S:", long_options, &option_index);
if(ch == -1)
break;
switch(ch)
{
case 'D':
if(!optarg)
throw std::runtime_error("Invalid option format!");
inLengths = getTypeValuesFromString<size_t>(optarg);
break;
case 'R':
if(!optarg)
throw std::runtime_error("Invalid option format!");
reduceDims = getTypeValuesFromString<int>(optarg);
break;
case 'S':
if(!optarg)
throw std::runtime_error("Invalid option format!");
scales = getTypeValuesFromString<float>(optarg);
break;
case '?':
if(std::string(long_options[option_index].name) == "help")
{
show_usage(argv[0]);
return (-1);
};
break;
default: show_usage(argv[0]); return (-1);
};
};
if(optind + 2 > argc)
throw std::runtime_error("Invalid cmd-line arguments, more argumetns are needed!");
data_type = std::atoi(argv[optind++]);
init_method = std::atoi(argv[optind]);
if(scales.empty())
{
scales.push_back(1.0f);
scales.push_back(0.0f);
};
if(inLengths.size() != 4 ||
(reduceDims.size() != 1 && reduceDims.size() != 3 && reduceDims.size() != 4))
return (-1);
if(data_type != 0 && data_type != 1 && data_type != 3 && data_type != 5 && data_type != 6)
return (-1);
return (0);
};
bool do_verification{true};
bool propagateNan{false};
bool useIndex{false};
bool time_kernel{false};
bool do_dumpout{false};
int init_method{2};
float alpha{1.0f};
float beta{0.0f};
std::vector<size_t> inLengths{64, 4, 280, 82};
std::vector<int> reduceDims{0, 1, 2, 3};
};
bool test_reduce_with_index(int data_type,
int init_method,
std::vector<int> reduceDims,
std::vector<size_t> inLengths,
ReduceTensorOp reduceOpId,
bool propagateNan,
float alpha,
float beta)
std::vector<std::vector<int>> SetGenericReduceDim()
{
using ck::profiler::profile_reduce_impl;
bool result = true;
if(data_type == 0)
{
result = profile_reduce_impl<float, float, float>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
true,
alpha,
beta);
}
else if(data_type == 1)
{
result = profile_reduce_impl<ck::half_t, ck::half_t, ck::half_t>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
true,
alpha,
beta);
}
else if(data_type == 3)
{
result = profile_reduce_impl<int8_t, int8_t, int8_t>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
true,
alpha,
beta);
}
else if(data_type == 5)
{
result = profile_reduce_impl<ck::bhalf_t, float, ck::bhalf_t>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
true,
alpha,
beta);
}
else if(data_type == 6)
{
result = profile_reduce_impl<double, double, double>(true,
init_method,
false,
false,
inLengths,
reduceDims,
reduceOpId,
propagateNan,
true,
alpha,
beta);
}
return (result);
};
constexpr ReduceTensorOp reduceOpId = ReduceTensorOp::AMAX;
constexpr bool propagateNan = false;
int main(int argc, char* argv[])
{
SimpleAppArgs args;
bool result = true;
if(argc == 1)
{
int data_type = 1;
int init_method = 2;
std::vector<size_t> inLengths{64, 4, 280, 80};
std::vector<std::vector<int>> v_reduceDims{
{0, 1, 2, 3}, {0, 1, 2}, {1, 2, 3}, {0, 1, 3}, {0, 2, 3}, {0}, {1}, {2}, {3}};
for(auto& reduceDims : v_reduceDims)
result = result && test_reduce_with_index(data_type,
init_method,
reduceDims,
inLengths,
reduceOpId,
propagateNan,
1.0f,
0.0f);
}
else
{
if(args.processArgs(argc, argv) < 0)
{
throw std::runtime_error(
"Invalid input arguments, test_reduce_with_index could not be executed!");
};
result = test_reduce_with_index(args.data_type,
args.init_method,
args.reduceDims,
args.inLengths,
reduceOpId,
propagateNan,
args.scales[0],
args.scales[1]);
}
std::cout << "test_reduce_with_index ..... " << (result ? "SUCCESS" : "FAILURE") << std::endl;
return (result ? 0 : -1);
return {{0, 1, 2, 3}, {0, 1, 2}, {0, 1, 3}, {0, 2, 3}, {1, 2, 3}, {0}, {1}, {2}, {3}};
}
template <typename T>
class ReduceWithIndexTest : public ::testing::Test
{
protected:
using InDataType = std::tuple_element_t<0, T>;
using AccDataType = std::tuple_element_t<1, T>;
using OutDataType = std::tuple_element_t<2, T>;
static std::vector<ReduceParam> params;
static void SetUpTestSuite()
{
// set testcase variables
ReduceParam set;
const auto setReduceDim = SetGenericReduceDim();
for(std::size_t i(0); i < setReduceDim.size(); ++i)
{
set.reduceDims = setReduceDim[i];
params.emplace_back(set);
}
}
template <ReduceTensorOp ReduceOpIdType>
void Run()
{
for(auto param : this->params)
{
bool success = ck::profiler::profile_reduce_impl<InDataType, AccDataType, OutDataType>(
param.do_verification,
param.init_method,
param.do_dumpout,
param.time_kernel,
param.inLengths,
param.reduceDims,
ReduceOpIdType,
param.propagateNan,
param.useIndex,
param.alpha,
param.beta);
EXPECT_TRUE(success);
}
}
};
template <typename T>
std::vector<ReduceParam> ReduceWithIndexTest<T>::params = {};
using Reduce_float_types = ::testing::Types<std::tuple<float, float, float>>;
using Reduce_double_types = ::testing::Types<std::tuple<double, double, double>>;
using Reduce_int8t_types = ::testing::Types<std::tuple<int8_t, int8_t, int8_t>>;
using Reduce_half_types = ::testing::Types<std::tuple<ck::half_t, ck::half_t, ck::half_t>>;
using Reduce_bhalf_float_Types = ::testing::Types<std::tuple<ck::bhalf_t, float, ck::bhalf_t>>;
template <typename TType>
class ReduceWithIndexFloat : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithIndexDouble : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithIndexInt8 : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithIndexHalf : public ReduceWithIndexTest<TType>
{
};
template <typename TType>
class ReduceWithIndexBHalfFloat : public ReduceWithIndexTest<TType>
{
};
TYPED_TEST_SUITE(ReduceWithIndexFloat, Reduce_float_types);
TYPED_TEST_SUITE(ReduceWithIndexDouble, Reduce_double_types);
TYPED_TEST_SUITE(ReduceWithIndexInt8, Reduce_int8t_types);
TYPED_TEST_SUITE(ReduceWithIndexHalf, Reduce_half_types);
TYPED_TEST_SUITE(ReduceWithIndexBHalfFloat, Reduce_bhalf_float_Types);
TYPED_TEST(ReduceWithIndexFloat, ReduceWithIndexTestFloat_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithIndexFloat, ReduceWithIndexTestFloat_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithIndexFloat, ReduceWithIndexTestFloat_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithIndexDouble, ReduceWithIndexTestDouble_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithIndexDouble, ReduceWithIndexTestDouble_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithIndexDouble, ReduceWithIndexTestDouble_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithIndexInt8, ReduceWithIndexTestInt8_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithIndexInt8, ReduceWithIndexTestInt8_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithIndexInt8, ReduceWithIndexTestInt8_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithIndexHalf, ReduceWithIndexTestHalf_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithIndexHalf, ReduceWithIndexTestHalf_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithIndexHalf, ReduceWithIndexTestHalf_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}
TYPED_TEST(ReduceWithIndexBHalfFloat, ReduceWithIndexTesBtHalfFloat_AMAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::AMAX>();
}
TYPED_TEST(ReduceWithIndexBHalfFloat, ReduceWithIndexTestBHalfFloat_MIN)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MIN>();
}
TYPED_TEST(ReduceWithIndexBHalfFloat, ReduceWithIndexTestBHalfFloat_MAX)
{
// trigger Run() -> Generic
this->template Run<ReduceTensorOp::MAX>();
}