diff --git a/CMakeLists.txt b/CMakeLists.txt index c00db26f3a..96a49b1c00 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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() diff --git a/Jenkinsfile b/Jenkinsfile index 139f928037..3fccb2881b 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -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 diff --git a/codegen/CMakeLists.txt b/codegen/CMakeLists.txt index d8b22fc943..d08fe2380b 100644 --- a/codegen/CMakeLists.txt +++ b/codegen/CMakeLists.txt @@ -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) diff --git a/codegen/include/ck/host/types.hpp b/codegen/include/ck/host/types.hpp index 812c073678..8bad7bf89c 100644 --- a/codegen/include/ck/host/types.hpp +++ b/codegen/include/ck/host/types.hpp @@ -76,8 +76,11 @@ std::string SequenceStr(const std::vector& v); std::string MakeTuple(const std::vector& v); +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wglobal-constructors" template 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"; diff --git a/codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp b/codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp index a2e8eccbf1..fff75c1962 100644 --- a/codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp +++ b/codegen/src/device_gemm_multiple_d_operation_xdl_cshuffle.cpp @@ -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 @@ -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 diff --git a/codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp b/codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp index 5ed59dd56b..36c9a13b4c 100644 --- a/codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp +++ b/codegen/src/device_grouped_conv_fwd_multiple_abd_operation_xdl_cshuffle.cpp @@ -4,6 +4,7 @@ #include "ck/host/device_grouped_conv_fwd_multiple_d/conv_fwd_op.hpp" #include #include "ck/host/stringutils.hpp" +#include "ck/host/types.hpp" #include "ck/host/utils.hpp" #include @@ -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 diff --git a/codegen/src/headers.cpp b/codegen/src/headers.cpp index f685aca044..5b0c929db3 100644 --- a/codegen/src/headers.cpp +++ b/codegen/src/headers.cpp @@ -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 GetHeaders() { diff --git a/codegen/test/CMakeLists.txt b/codegen/test/CMakeLists.txt index f891286019..5aad1ef877 100644 --- a/codegen/test/CMakeLists.txt +++ b/codegen/test/CMakeLists.txt @@ -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()) diff --git a/codegen/test/grouped_conv_fwd_multiple_d_v1.cpp b/codegen/test/grouped_conv_fwd_multiple_d_v1.cpp index 3c477692e5..50290fa25a 100644 --- a/codegen/test/grouped_conv_fwd_multiple_d_v1.cpp +++ b/codegen/test/grouped_conv_fwd_multiple_d_v1.cpp @@ -92,7 +92,6 @@ struct Epilogue static_cast(prob.C), static_cast(prob.Y), static_cast(prob.X)}; - ck::Array d_lengths = {}; ck::Array in_strides{static_cast(prob.C), static_cast(prob.Hi * prob.Wi * prob.G * prob.C), @@ -109,7 +108,6 @@ struct Epilogue 1, static_cast(prob.X * prob.C), static_cast(prob.C)}; - ck::Array d_strides = {}; ck::Array conv_filter_strides = {2, 2}; ck::Array conv_filter_dilations = {1, 1}; diff --git a/codegen/test/grouped_conv_fwd_multiple_d_v2.cpp b/codegen/test/grouped_conv_fwd_multiple_d_v2.cpp index ec9bd2b781..b558d97c78 100644 --- a/codegen/test/grouped_conv_fwd_multiple_d_v2.cpp +++ b/codegen/test/grouped_conv_fwd_multiple_d_v2.cpp @@ -92,7 +92,6 @@ struct Epilogue static_cast(prob.C), static_cast(prob.Y), static_cast(prob.X)}; - ck::Array d_lengths = {}; ck::Array in_strides{static_cast(prob.C), static_cast(prob.Hi * prob.Wi * prob.G * prob.C), @@ -109,7 +108,6 @@ struct Epilogue 1, static_cast(prob.X * prob.C), static_cast(prob.C)}; - ck::Array d_strides = {}; ck::Array conv_filter_strides = {1, 1}; ck::Array conv_filter_dilations = {1, 1}; diff --git a/codegen/test/grouped_conv_fwd_multiple_d_v3.cpp b/codegen/test/grouped_conv_fwd_multiple_d_v3.cpp index 9850184c5e..e2972a93d2 100644 --- a/codegen/test/grouped_conv_fwd_multiple_d_v3.cpp +++ b/codegen/test/grouped_conv_fwd_multiple_d_v3.cpp @@ -92,7 +92,6 @@ struct Epilogue static_cast(prob.C), static_cast(prob.Y), static_cast(prob.X)}; - ck::Array d_lengths = {}; ck::Array in_strides{static_cast(prob.C), static_cast(prob.Hi * prob.Wi * prob.G * prob.C), @@ -109,7 +108,6 @@ struct Epilogue 1, static_cast(prob.X * prob.C), static_cast(prob.C)}; - ck::Array d_strides = {}; ck::Array conv_filter_strides = {2, 2}; ck::Array conv_filter_dilations = {1, 1}; diff --git a/codegen/test/grouped_conv_fwd_multiple_d_v4.cpp b/codegen/test/grouped_conv_fwd_multiple_d_v4.cpp index 907f744db4..b728096c51 100644 --- a/codegen/test/grouped_conv_fwd_multiple_d_v4.cpp +++ b/codegen/test/grouped_conv_fwd_multiple_d_v4.cpp @@ -92,7 +92,6 @@ struct Epilogue static_cast(prob.C), static_cast(prob.Y), static_cast(prob.X)}; - ck::Array d_lengths = {}; ck::Array in_strides{static_cast(prob.C), static_cast(prob.Hi * prob.Wi * prob.G * prob.C), @@ -109,7 +108,6 @@ struct Epilogue 1, static_cast(prob.X * prob.C), static_cast(prob.C)}; - ck::Array d_strides = {}; ck::Array conv_filter_strides = {1, 1}; ck::Array conv_filter_dilations = {1, 1}; diff --git a/codegen/test/rtc/src/kernel.cpp b/codegen/test/rtc/src/kernel.cpp index f4fb19130c..9fe38e84ad 100644 --- a/codegen/test/rtc/src/kernel.cpp +++ b/codegen/test/rtc/src/kernel.cpp @@ -118,4 +118,4 @@ void kernel::launch(hipStream_t stream, launch_kernel(impl->fun, stream, global, local, kernargs.data(), size); } -} // namespace rtc \ No newline at end of file +} // namespace rtc diff --git a/codegen/test/rtc/src/tmp_dir.cpp b/codegen/test/rtc/src/tmp_dir.cpp index 3b0f0170e8..1cc8f75b29 100644 --- a/codegen/test/rtc/src/tmp_dir.cpp +++ b/codegen/test/rtc/src/tmp_dir.cpp @@ -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 \ No newline at end of file +} // namespace rtc diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 45cfee4de9..cfcb845421 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -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() diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp index f5d9a789fa..bb97d29534 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp @@ -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; + ADataType, + EDataType>; static constexpr auto matrix_padder = MatrixPadder{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_, diff --git a/library/src/tensor_operation_instance/gpu/CMakeLists.txt b/library/src/tensor_operation_instance/gpu/CMakeLists.txt index 0fff0dfa31..72026dc185 100644 --- a/library/src/tensor_operation_instance/gpu/CMakeLists.txt +++ b/library/src/tensor_operation_instance/gpu/CMakeLists.txt @@ -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) diff --git a/script/test_reduce_with_index.sh b/script/test_reduce_with_index.sh deleted file mode 100755 index b0843ba6c1..0000000000 --- a/script/test_reduce_with_index.sh +++ /dev/null @@ -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 - diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index fc1bcfdb27..71bde7e267 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -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}) diff --git a/test/reduce/CMakeLists.txt b/test/reduce/CMakeLists.txt index fabaf7cc54..bf05795063 100644 --- a/test/reduce/CMakeLists.txt +++ b/test/reduce/CMakeLists.txt @@ -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) diff --git a/test/reduce/reduce_no_index.cpp b/test/reduce/reduce_no_index.cpp index 1ab452442d..9153805889 100644 --- a/test/reduce/reduce_no_index.cpp +++ b/test/reduce/reduce_no_index.cpp @@ -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 #include "ck/library/utility/host_common_util.hpp" #include "profiler/profile_reduce_impl.hpp" - +#include 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 inLengths; - std::vector reduceDims; - std::vector 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(optarg); - break; - case 'R': - if(!optarg) - throw std::runtime_error("Invalid option format!"); - - reduceDims = getTypeValuesFromString(optarg); - break; - case 'S': - if(!optarg) - throw std::runtime_error("Invalid option format!"); - - scales = getTypeValuesFromString(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 inLengths{64, 4, 280, 82}; + std::vector reduceDims{0, 1, 2, 3}; }; -bool test_reduce_no_index(int data_type, - int init_method, - std::vector reduceDims, - std::vector inLengths, - ReduceTensorOp reduceOpId, - bool propagateNan, - float alpha, - float beta) +std::vector> SetGenericReduceDim() { - using ck::profiler::profile_reduce_impl; - - bool result = true; - - if(data_type == 0) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - false, - alpha, - beta); - } - else if(data_type == 1) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - false, - alpha, - beta); - } - else if(data_type == 3) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - false, - alpha, - beta); - } - else if(data_type == 5) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - false, - alpha, - beta); - } - else if(data_type == 6) - { - result = profile_reduce_impl(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 inLengths{64, 4, 280, 80}; - std::vector> 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 +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 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 + void Run() + { + for(auto param : this->params) + { + bool success = ck::profiler::profile_reduce_impl( + 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 +std::vector ReduceWithIndexTest::params = {}; + +using Reduce_float_types = ::testing::Types>; +using Reduce_double_types = ::testing::Types>; +using Reduce_int8t_types = ::testing::Types>; +using Reduce_half_types = ::testing::Types>; +using Reduce_bhalf_float_Types = ::testing::Types>; + +template +class ReduceWithNoIndexFloat : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithNoIndexDouble : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithNoIndexInt8 : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithNoIndexHalf : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithNoIndexBHalfFloat : public ReduceWithIndexTest +{ +}; + +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(); +} + +TYPED_TEST(ReduceWithNoIndexFloat, ReduceWithNoIndexTestFloat_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexFloat, ReduceWithNoIndexTestFloat_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexDouble, ReduceWithNoIndexTestDouble_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexDouble, ReduceWithNoIndexTestDouble_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexDouble, ReduceWithNoIndexTestDouble_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexInt8, ReduceWithNoIndexTestInt8_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexInt8, ReduceWithNoIndexTestInt8_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexInt8, ReduceWithNoIndexTestInt8_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexHalf, ReduceWithNoIndexTestHalf_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexHalf, ReduceWithNoIndexTestHalf_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexHalf, ReduceWithNoIndexTestHalf_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexBHalfFloat, ReduceWithNoIndexTesBtHalfFloat_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexBHalfFloat, ReduceWithNoIndexTestBHalfFloat_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithNoIndexBHalfFloat, ReduceWithNoIndexTestBHalfFloat_MAX) +{ + // trigger Run() -> Generic + this->template Run(); } diff --git a/test/reduce/reduce_with_index.cpp b/test/reduce/reduce_with_index.cpp index 0301669c5a..796211a69a 100644 --- a/test/reduce/reduce_with_index.cpp +++ b/test/reduce/reduce_with_index.cpp @@ -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 #include "ck/library/utility/host_common_util.hpp" #include "profiler/profile_reduce_impl.hpp" - +#include 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 inLengths; - std::vector reduceDims; - std::vector 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(optarg); - break; - case 'R': - if(!optarg) - throw std::runtime_error("Invalid option format!"); - - reduceDims = getTypeValuesFromString(optarg); - break; - case 'S': - if(!optarg) - throw std::runtime_error("Invalid option format!"); - - scales = getTypeValuesFromString(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 inLengths{64, 4, 280, 82}; + std::vector reduceDims{0, 1, 2, 3}; }; -bool test_reduce_with_index(int data_type, - int init_method, - std::vector reduceDims, - std::vector inLengths, - ReduceTensorOp reduceOpId, - bool propagateNan, - float alpha, - float beta) +std::vector> SetGenericReduceDim() { - using ck::profiler::profile_reduce_impl; - - bool result = true; - - if(data_type == 0) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - true, - alpha, - beta); - } - else if(data_type == 1) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - true, - alpha, - beta); - } - else if(data_type == 3) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - true, - alpha, - beta); - } - else if(data_type == 5) - { - result = profile_reduce_impl(true, - init_method, - false, - false, - inLengths, - reduceDims, - reduceOpId, - propagateNan, - true, - alpha, - beta); - } - else if(data_type == 6) - { - result = profile_reduce_impl(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 inLengths{64, 4, 280, 80}; - std::vector> 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 +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 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 + void Run() + { + for(auto param : this->params) + { + bool success = ck::profiler::profile_reduce_impl( + 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 +std::vector ReduceWithIndexTest::params = {}; + +using Reduce_float_types = ::testing::Types>; +using Reduce_double_types = ::testing::Types>; +using Reduce_int8t_types = ::testing::Types>; +using Reduce_half_types = ::testing::Types>; +using Reduce_bhalf_float_Types = ::testing::Types>; + +template +class ReduceWithIndexFloat : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithIndexDouble : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithIndexInt8 : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithIndexHalf : public ReduceWithIndexTest +{ +}; + +template +class ReduceWithIndexBHalfFloat : public ReduceWithIndexTest +{ +}; + +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(); +} + +TYPED_TEST(ReduceWithIndexFloat, ReduceWithIndexTestFloat_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexFloat, ReduceWithIndexTestFloat_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexDouble, ReduceWithIndexTestDouble_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexDouble, ReduceWithIndexTestDouble_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexDouble, ReduceWithIndexTestDouble_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexInt8, ReduceWithIndexTestInt8_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexInt8, ReduceWithIndexTestInt8_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexInt8, ReduceWithIndexTestInt8_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexHalf, ReduceWithIndexTestHalf_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexHalf, ReduceWithIndexTestHalf_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexHalf, ReduceWithIndexTestHalf_MAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexBHalfFloat, ReduceWithIndexTesBtHalfFloat_AMAX) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexBHalfFloat, ReduceWithIndexTestBHalfFloat_MIN) +{ + // trigger Run() -> Generic + this->template Run(); +} + +TYPED_TEST(ReduceWithIndexBHalfFloat, ReduceWithIndexTestBHalfFloat_MAX) +{ + // trigger Run() -> Generic + this->template Run(); }