From f4f391955c48d5c44a4e6c18942d4b2032eca1dd Mon Sep 17 00:00:00 2001 From: Adam Osewski <19374865+aosewski@users.noreply.github.com> Date: Sat, 30 Apr 2022 15:50:16 +0200 Subject: [PATCH] Introduce GoogleTest framework. (#204) * Use googletest for tests. Add conv2d_fwd UT. * Add conv1D/3D to gtest UT. * Fix: not duplicate test with CTest. * Convert more tests to googltests. * Fix: GIT_SHALLOW is not allowed for git commit hash. * Clang-format * use integer value for GEMM test Co-authored-by: Adam Osewski Co-authored-by: Chao Liu Co-authored-by: Chao Liu [ROCm/composable_kernel commit: 8eca05a63333d302cbd3bde4a0b83863c08ecc4e] --- CMakeLists.txt | 4 +- cmake/googletest.cmake | 36 +++ test/CMakeLists.txt | 15 ++ test/conv_util/CMakeLists.txt | 2 +- test/conv_util/conv_util.cpp | 217 +++++++++--------- test/convnd_fwd/CMakeLists.txt | 8 +- test/convnd_fwd/conv1d_fwd.cpp | 103 ++++----- test/convnd_fwd/conv2d_fwd.cpp | 115 ++++------ test/convnd_fwd/conv3d_fwd.cpp | 127 ++++------ test/reference_conv_fwd/CMakeLists.txt | 2 +- .../reference_conv_fwd/reference_conv_fwd.cpp | 153 ++++++------ 11 files changed, 378 insertions(+), 404 deletions(-) create mode 100644 cmake/googletest.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index f5da68fa48..2b798e38f3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required(VERSION 3.5) +cmake_minimum_required(VERSION 3.14) # Check support for CUDA/HIP in Cmake project(composable_kernel) @@ -234,6 +234,8 @@ include_directories(BEFORE ${PROJECT_SOURCE_DIR}/library/include ) +include(googletest) + SET(BUILD_DEV ON CACHE BOOL "BUILD_DEV") if(BUILD_DEV) add_compile_options(-Werror) diff --git a/cmake/googletest.cmake b/cmake/googletest.cmake new file mode 100644 index 0000000000..c7e70cc8a9 --- /dev/null +++ b/cmake/googletest.cmake @@ -0,0 +1,36 @@ +include(FetchContent) + +set(GOOGLETEST_DIR "" CACHE STRING "Location of local GoogleTest repo to build against") + +if(GOOGLETEST_DIR) + set(FETCHCONTENT_SOURCE_DIR_GOOGLETEST ${GOOGLETEST_DIR} CACHE STRING "GoogleTest source directory override") +endif() + +message(STATUS "Fetching GoogleTest") + +list(APPEND GTEST_CMAKE_CXX_FLAGS + -Wno-undef + -Wno-reserved-identifier + -Wno-global-constructors + -Wno-missing-noreturn + -Wno-disabled-macro-expansion + -Wno-used-but-marked-unused + -Wno-switch-enum + -Wno-zero-as-null-pointer-constant + -Wno-unused-member-function +) +message(STATUS "Suppressing googltest warnings with flags: ${GTEST_CMAKE_CXX_FLAGS}") + +FetchContent_Declare( + googletest + GIT_REPOSITORY https://github.com/google/googletest.git + GIT_TAG b85864c64758dec007208e56af933fc3f52044ee +) + +# Will be necessary for windows build +# set(gtest_force_shared_crt ON CACHE BOOL "" FORCE) +FetchContent_MakeAvailable(googletest) + +target_compile_options(gtest PRIVATE ${GTEST_CMAKE_CXX_FLAGS}) +target_compile_options(gtest_main PRIVATE ${GTEST_CMAKE_CXX_FLAGS}) + diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index ae9949b8ce..cc0778de4c 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -24,6 +24,7 @@ include_directories(BEFORE add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR}) add_custom_target(tests) + function(add_test_executable TEST_NAME) message("adding test ${TEST_NAME}") add_executable(${TEST_NAME} ${ARGN}) @@ -32,6 +33,20 @@ function(add_test_executable TEST_NAME) add_dependencies(check ${TEST_NAME}) endfunction(add_test_executable TEST_NAME) +include(GoogleTest) + +function(add_gtest_executable TEST_NAME) + message("adding gtest ${TEST_NAME}") + add_executable(${TEST_NAME} ${ARGN}) + add_dependencies(tests ${TEST_NAME}) + add_dependencies(check ${TEST_NAME}) + # suppress gtest warnings + target_compile_options(${TEST_NAME} PRIVATE -Wno-global-constructors) + target_link_libraries(${TEST_NAME} PRIVATE gtest_main) + gtest_discover_tests(${TEST_NAME}) +endfunction(add_gtest_executable TEST_NAME) + + add_subdirectory(magic_number_division) add_subdirectory(space_filling_curve) add_subdirectory(conv_util) diff --git a/test/conv_util/CMakeLists.txt b/test/conv_util/CMakeLists.txt index e3ba9574a2..70b3e851be 100644 --- a/test/conv_util/CMakeLists.txt +++ b/test/conv_util/CMakeLists.txt @@ -1,2 +1,2 @@ -add_test_executable(test_conv_util conv_util.cpp) +add_gtest_executable(test_conv_util conv_util.cpp) target_link_libraries(test_conv_util PRIVATE host_tensor conv_fwd_util) diff --git a/test/conv_util/conv_util.cpp b/test/conv_util/conv_util.cpp index cc487c39e3..453225e800 100644 --- a/test/conv_util/conv_util.cpp +++ b/test/conv_util/conv_util.cpp @@ -1,6 +1,7 @@ #include #include #include +#include "gtest/gtest.h" #include "config.hpp" #include "conv_fwd_util.hpp" @@ -9,196 +10,194 @@ namespace { -bool test_conv_params_get_output_spatial_lengths() +class TestConvUtil : public ::testing::Test { - bool res{true}; - // -------------------------- default 2D ------------------------------------ + public: + void SetNDParams(std::size_t ndims) + { + conv_params.num_dim_spatial = ndims; + conv_params.filter_spatial_lengths = std::vector(ndims, 3); + conv_params.input_spatial_lengths = std::vector(ndims, 71); + conv_params.conv_filter_strides = std::vector(ndims, 2); + conv_params.conv_filter_dilations = std::vector(ndims, 1); + conv_params.input_left_pads = std::vector(ndims, 1); + conv_params.input_right_pads = std::vector(ndims, 1); + } + + protected: + // ------- default 2D ------- // input NCHW {128,192,71,71}, // weights KCYX {256,192,3,3}, // stride {2,2}, // dilations {1,1}, // padding {{1,1}, {1,1}} ck::utils::conv::ConvParams conv_params; +}; + +} // namespace + +TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths2D) +{ + ck::utils::conv::ConvParams conv_params; std::vector out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{36, 36}, - "Error: ConvParams 2D default constructor."); + EXPECT_TRUE(ck::utils::check_err(out_spatial_len, + std::vector{36, 36}, + "Error: ConvParams 2D default constructor.")); conv_params.conv_filter_strides = std::vector{1, 1}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( - out_spatial_len, std::vector{71, 71}, "Error: ConvParams 2D stride {1,1}."); + EXPECT_TRUE(ck::utils::check_err( + out_spatial_len, std::vector{71, 71}, "Error: ConvParams 2D stride {1,1}.")); conv_params.conv_filter_strides = std::vector{2, 2}; conv_params.input_left_pads = std::vector{2, 2}; conv_params.input_right_pads = std::vector{2, 2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{37, 37}, - "Error: ConvParams 2D padding left/right {2,2}."); + EXPECT_TRUE(ck::utils::check_err(out_spatial_len, + std::vector{37, 37}, + "Error: ConvParams 2D padding left/right {2,2}.")); conv_params.conv_filter_dilations = std::vector{2, 2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( - out_spatial_len, std::vector{36, 36}, "Error: ConvParams 2D dilation {2,2}."); + EXPECT_TRUE(ck::utils::check_err( + out_spatial_len, std::vector{36, 36}, "Error: ConvParams 2D dilation {2,2}.")); conv_params.conv_filter_strides = std::vector{3, 3}; conv_params.input_left_pads = std::vector{1, 1}; conv_params.input_right_pads = std::vector{1, 1}; conv_params.conv_filter_dilations = std::vector{2, 2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = + EXPECT_TRUE( ck::utils::check_err(out_spatial_len, std::vector{23, 23}, - "Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}."); + "Error: ConvParams 2D strides{3,3}, padding {1,1}, dilations {2,2}.")); +} - // -------------------------- 1D ------------------------------------ - conv_params.num_dim_spatial = 1; - conv_params.filter_spatial_lengths = std::vector{3}; - conv_params.input_spatial_lengths = std::vector{71}; - conv_params.conv_filter_strides = std::vector{2}; - conv_params.conv_filter_dilations = std::vector{1}; - conv_params.input_left_pads = std::vector{1}; - conv_params.input_right_pads = std::vector{1}; +TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths1D) +{ + SetNDParams(1); - out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( - out_spatial_len, std::vector{36}, "Error: ConvParams 1D."); + std::vector out_spatial_len = conv_params.GetOutputSpatialLengths(); + EXPECT_TRUE(ck::utils::check_err( + out_spatial_len, std::vector{36}, "Error: ConvParams 1D.")); conv_params.conv_filter_strides = std::vector{1}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( - out_spatial_len, std::vector{71}, "Error: ConvParams 1D stride {1}."); + EXPECT_TRUE(ck::utils::check_err( + out_spatial_len, std::vector{71}, "Error: ConvParams 1D stride {1}.")); conv_params.conv_filter_strides = std::vector{2}; conv_params.input_left_pads = std::vector{2}; conv_params.input_right_pads = std::vector{2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{37}, - "Error: ConvParams 1D padding left/right {2}."); + EXPECT_TRUE(ck::utils::check_err(out_spatial_len, + std::vector{37}, + "Error: ConvParams 1D padding left/right {2}.")); conv_params.conv_filter_dilations = std::vector{2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( - out_spatial_len, std::vector{36}, "Error: ConvParams 1D dilation {2}."); + EXPECT_TRUE(ck::utils::check_err( + out_spatial_len, std::vector{36}, "Error: ConvParams 1D dilation {2}.")); conv_params.conv_filter_strides = std::vector{3}; conv_params.input_left_pads = std::vector{1}; conv_params.input_right_pads = std::vector{1}; conv_params.conv_filter_dilations = std::vector{2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{23}, - "Error: ConvParams 1D strides{3}, padding {1}, dilations {2}."); + EXPECT_TRUE( + ck::utils::check_err(out_spatial_len, + std::vector{23}, + "Error: ConvParams 1D strides{3}, padding {1}, dilations {2}.")); +} - // -------------------------- 3D ------------------------------------ - conv_params.num_dim_spatial = 3; - conv_params.filter_spatial_lengths = std::vector{3, 3, 3}; - conv_params.input_spatial_lengths = std::vector{71, 71, 71}; - conv_params.conv_filter_strides = std::vector{2, 2, 2}; - conv_params.conv_filter_dilations = std::vector{1, 1, 1}; - conv_params.input_left_pads = std::vector{1, 1, 1}; - conv_params.input_right_pads = std::vector{1, 1, 1}; +TEST_F(TestConvUtil, ConvParamsGetOutputSpatialLengths3D) +{ + SetNDParams(3); - out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( - out_spatial_len, std::vector{36, 36, 36}, "Error: ConvParams 3D."); + std::vector out_spatial_len = conv_params.GetOutputSpatialLengths(); + EXPECT_TRUE(ck::utils::check_err( + out_spatial_len, std::vector{36, 36, 36}, "Error: ConvParams 3D.")); conv_params.conv_filter_strides = std::vector{1, 1, 1}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{71, 71, 71}, - "Error: ConvParams 3D stride {1, 1, 1}."); + EXPECT_TRUE(ck::utils::check_err(out_spatial_len, + std::vector{71, 71, 71}, + "Error: ConvParams 3D stride {1, 1, 1}.")); conv_params.conv_filter_strides = std::vector{2, 2, 2}; conv_params.input_left_pads = std::vector{2, 2, 2}; conv_params.input_right_pads = std::vector{2, 2, 2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{37, 37, 37}, - "Error: ConvParams 3D padding left/right {2, 2, 2}."); + EXPECT_TRUE(ck::utils::check_err(out_spatial_len, + std::vector{37, 37, 37}, + "Error: ConvParams 3D padding left/right {2, 2, 2}.")); conv_params.conv_filter_dilations = std::vector{2, 2, 2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err(out_spatial_len, - std::vector{36, 36, 36}, - "Error: ConvParams 3D dilation {2, 2, 2}."); + EXPECT_TRUE(ck::utils::check_err(out_spatial_len, + std::vector{36, 36, 36}, + "Error: ConvParams 3D dilation {2, 2, 2}.")); conv_params.conv_filter_strides = std::vector{3, 3, 3}; conv_params.input_left_pads = std::vector{1, 1, 1}; conv_params.input_right_pads = std::vector{1, 1, 1}; conv_params.conv_filter_dilations = std::vector{2, 2, 2}; out_spatial_len = conv_params.GetOutputSpatialLengths(); - res = ck::utils::check_err( + EXPECT_TRUE(ck::utils::check_err( out_spatial_len, std::vector{23, 23, 23}, - "Error: ConvParams 3D strides{3, 3, 3}, padding {1, 1, 1}, dilations {2, 2, 2}."); - - return res; + "Error: ConvParams 3D strides{3, 3, 3}, padding {1, 1, 1}, dilations {2, 2, 2}.")); } -bool test_get_host_tensor_descriptor() +TEST(ConvUtil, GetHostTensorDescriptor) { - bool res{true}; namespace tl = ck::tensor_layout::convolution; std::vector dims{2, 3, 4, 5}; HostTensorDescriptor h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NHWC{}); - res = - ck::utils::check_err(h.GetLengths(), {2, 3, 4, 5}, "Error: wrong NHWC dimensions lengths!"); - res = ck::utils::check_err( - h.GetStrides(), {3 * 4 * 5, 1, 3 * 5, 3}, "Error: wrong NHWC dimensions strides!"); + EXPECT_TRUE(ck::utils::check_err( + h.GetLengths(), {2, 3, 4, 5}, "Error: wrong NHWC dimensions lengths!")); + EXPECT_TRUE(ck::utils::check_err( + h.GetStrides(), {3 * 4 * 5, 1, 3 * 5, 3}, "Error: wrong NHWC dimensions strides!")); h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCHW{}); - res = - ck::utils::check_err(h.GetLengths(), {2, 3, 4, 5}, "Error: wrong NCHW dimensions lengths!"); - res = ck::utils::check_err( - h.GetStrides(), {3 * 4 * 5, 4 * 5, 5, 1}, "Error: wrong NCHW dimensions strides!"); + EXPECT_TRUE(ck::utils::check_err( + h.GetLengths(), {2, 3, 4, 5}, "Error: wrong NCHW dimensions lengths!")); + EXPECT_TRUE(ck::utils::check_err( + h.GetStrides(), {3 * 4 * 5, 4 * 5, 5, 1}, "Error: wrong NCHW dimensions strides!")); dims = std::vector{2, 3, 4}; h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NWC{}); - res = ck::utils::check_err(h.GetLengths(), {2, 3, 4}, "Error: wrong NWC dimensions lengths!"); - res = - ck::utils::check_err(h.GetStrides(), {3 * 4, 1, 3}, "Error: wrong NWC dimensions strides!"); + EXPECT_TRUE( + ck::utils::check_err(h.GetLengths(), {2, 3, 4}, "Error: wrong NWC dimensions lengths!")); + EXPECT_TRUE(ck::utils::check_err( + h.GetStrides(), {3 * 4, 1, 3}, "Error: wrong NWC dimensions strides!")); - h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCW{}); - res = ck::utils::check_err(h.GetLengths(), {2, 3, 4}, "Error: wrong NCW dimensions lengths!"); - res = - ck::utils::check_err(h.GetStrides(), {3 * 4, 4, 1}, "Error: wrong NCW dimensions strides!"); + h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCW{}); + EXPECT_TRUE( + ck::utils::check_err(h.GetLengths(), {2, 3, 4}, "Error: wrong NCW dimensions lengths!")); + EXPECT_TRUE(ck::utils::check_err( + h.GetStrides(), {3 * 4, 4, 1}, "Error: wrong NCW dimensions strides!")); dims = std::vector{2, 3, 4, 5, 6}; h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NDHWC{}); - res = ck::utils::check_err(h.GetLengths(), dims, "Error: wrong NDHWC dimensions lengths!"); - res = ck::utils::check_err(h.GetStrides(), - {3 * 4 * 5 * 6, // N - 1, // C - 3 * 5 * 6, // D - 3 * 6, // H - 3}, // W - "Error: wrong NDHWC dimensions strides!"); + EXPECT_TRUE( + ck::utils::check_err(h.GetLengths(), dims, "Error: wrong NDHWC dimensions lengths!")); + EXPECT_TRUE(ck::utils::check_err(h.GetStrides(), + {3 * 4 * 5 * 6, // N + 1, // C + 3 * 5 * 6, // D + 3 * 6, // H + 3}, // W + "Error: wrong NDHWC dimensions strides!")); - h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCDHW{}); - res = ck::utils::check_err(h.GetLengths(), dims, "Error: wrong NCDHW dimensions lengths!"); - res = ck::utils::check_err(h.GetStrides(), - {3 * 4 * 5 * 6, // N - 4 * 5 * 6, // C - 5 * 6, // D - 6, // H - 1}, // W - "Error: wrong NCDHW dimensions strides!"); - - return res; -} - -} // namespace - -int main(void) -{ - bool res = test_conv_params_get_output_spatial_lengths(); - std::cout << "test_conv_params_get_output_spatial_lengths ..... " - << (res ? "SUCCESS" : "FAILURE") << std::endl; - res = test_get_host_tensor_descriptor(); - std::cout << "test_get_host_tensor_descriptor ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - return res ? 0 : 1; + h = ck::utils::conv::get_host_tensor_descriptor(dims, tl::NCDHW{}); + EXPECT_TRUE( + ck::utils::check_err(h.GetLengths(), dims, "Error: wrong NCDHW dimensions lengths!")); + EXPECT_TRUE(ck::utils::check_err(h.GetStrides(), + {3 * 4 * 5 * 6, // N + 4 * 5 * 6, // C + 5 * 6, // D + 6, // H + 1}, // W + "Error: wrong NCDHW dimensions strides!")); } diff --git a/test/convnd_fwd/CMakeLists.txt b/test/convnd_fwd/CMakeLists.txt index 442c45dc8c..1d2ae3e4e3 100644 --- a/test/convnd_fwd/CMakeLists.txt +++ b/test/convnd_fwd/CMakeLists.txt @@ -1,15 +1,13 @@ add_custom_target(test_convnd_fwd) -add_test_executable(test_conv1d_fwd conv1d_fwd.cpp) +add_gtest_executable(test_conv1d_fwd conv1d_fwd.cpp) target_link_libraries(test_conv1d_fwd PRIVATE host_tensor device_conv1d_fwd_instance conv_fwd_util) -target_link_libraries(test_conv1d_fwd PRIVATE ) add_dependencies(test_convnd_fwd test_conv1d_fwd) -add_test_executable(test_conv2d_fwd conv2d_fwd.cpp) +add_gtest_executable(test_conv2d_fwd conv2d_fwd.cpp) target_link_libraries(test_conv2d_fwd PRIVATE host_tensor device_conv2d_fwd_instance conv_fwd_util) add_dependencies(test_convnd_fwd test_conv2d_fwd) -add_test_executable(test_conv3d_fwd conv3d_fwd.cpp) +add_gtest_executable(test_conv3d_fwd conv3d_fwd.cpp) target_link_libraries(test_conv3d_fwd PRIVATE host_tensor device_conv3d_fwd_instance conv_fwd_util) add_dependencies(test_convnd_fwd test_conv3d_fwd) - diff --git a/test/convnd_fwd/conv1d_fwd.cpp b/test/convnd_fwd/conv1d_fwd.cpp index df3b3a2945..c161b2795e 100644 --- a/test/convnd_fwd/conv1d_fwd.cpp +++ b/test/convnd_fwd/conv1d_fwd.cpp @@ -2,6 +2,7 @@ #include #include #include +#include "gtest/gtest.h" #include "data_type.hpp" #include "element_wise_operation.hpp" @@ -10,7 +11,33 @@ namespace { -bool test_conv1D_nwc() +template +bool test_conv1d_nwc_instances(const std::vector& conv_ptrs) +{ + using namespace std::placeholders; + using namespace ck::utils; + namespace ctl = ck::tensor_layout::convolution; + + ck::utils::conv::ConvParams params; + params.num_dim_spatial = 1; + params.filter_spatial_lengths = std::vector{3}; + params.input_spatial_lengths = std::vector{71}; + params.conv_filter_strides = std::vector{2}; + params.conv_filter_dilations = std::vector{1}; + params.input_left_pads = std::vector{1}; + params.input_right_pads = std::vector{1}; + + conv::ConvFwdOpInstance conv_instance(params); + + auto reference_conv_fwd_fun = + std::bind(conv::run_reference_convolution_forward<1, T, T, T>, params, _1, _2, _3); + OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); + return run_engine.Test(conv_ptrs); +} + +} // anonymous namespace + +TEST(Conv1DFwdNWC, TestConv1D) { using namespace std::placeholders; using namespace ck::utils; @@ -38,77 +65,29 @@ bool test_conv1D_nwc() OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); run_engine.SetAtol(1e-5); run_engine.SetRtol(1e-4); - return run_engine.Test(conv_ptrs); + EXPECT_TRUE(run_engine.Test(conv_ptrs)); } -template -bool test_conv1d_nwc_instances(const std::vector& conv_ptrs) +TEST(Conv1DFwdNWC, Bf16Iinstances) { - using namespace std::placeholders; - using namespace ck::utils; - namespace ctl = ck::tensor_layout::convolution; - - ck::utils::conv::ConvParams params; - params.num_dim_spatial = 1; - params.filter_spatial_lengths = std::vector{3}; - params.input_spatial_lengths = std::vector{71}; - params.conv_filter_strides = std::vector{2}; - params.conv_filter_dilations = std::vector{1}; - params.input_left_pads = std::vector{1}; - params.input_right_pads = std::vector{1}; - - conv::ConvFwdOpInstance conv_instance(params); - - auto reference_conv_fwd_fun = - std::bind(conv::run_reference_convolution_forward<1, T, T, T>, params, _1, _2, _3); - OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); - return run_engine.Test(conv_ptrs); + EXPECT_TRUE(test_conv1d_nwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<1>())); } -bool test_conv1d_nwc_bf16_instances() +TEST(Conv1DFwdNWC, F16Instances) { - return test_conv1d_nwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<1>()); + EXPECT_TRUE(test_conv1d_nwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<1>())); } -bool test_conv1d_nwc_f16_instances() +TEST(Conv1DFwdNWC, F32Instances) { - return test_conv1d_nwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<1>()); + EXPECT_TRUE(test_conv1d_nwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<1>())); } -bool test_conv1d_nwc_f32_instances() +TEST(Conv1DFwdNWC, Int8Instances) { - return test_conv1d_nwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<1>()); -} - -bool test_conv1d_nwc_int8_instances() -{ - return test_conv1d_nwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<1>()); -} - -} // anonymous namespace - -int main() -{ - bool res{true}; - res = test_conv1D_nwc(); - std::cout << "test_conv1D_nwc ..... " << (res ? "SUCCESS" : "FAILURE") << std::endl; - - res = test_conv1d_nwc_bf16_instances(); - std::cout << "\nTestConv1DNWCBF16Instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv1d_nwc_f16_instances(); - std::cout << "\ntest_conv1d_nwc_f16_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv1d_nwc_f32_instances(); - std::cout << "\ntest_conv1d_nwc_f32_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv1d_nwc_int8_instances(); - std::cout << "\ntest_conv1d_nwc_int8_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - - return res ? 0 : 1; + EXPECT_TRUE(test_conv1d_nwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<1>())); } diff --git a/test/convnd_fwd/conv2d_fwd.cpp b/test/convnd_fwd/conv2d_fwd.cpp index f35c69bbd0..e3815f778a 100644 --- a/test/convnd_fwd/conv2d_fwd.cpp +++ b/test/convnd_fwd/conv2d_fwd.cpp @@ -2,6 +2,7 @@ #include #include #include +#include "gtest/gtest.h" #include "data_type.hpp" #include "element_wise_operation.hpp" @@ -10,30 +11,6 @@ namespace { -bool test_conv2d_nhwc() -{ - using namespace std::placeholders; - using namespace ck::utils; - - ck::utils::conv::ConvParams params; - params.N = 2; - params.K = 16; - params.C = 4; - params.input_spatial_lengths = std::vector{16, 16}; - params.conv_filter_strides = std::vector{1, 1}; - - std::vector conv_ptrs; - test::conv::get_test_convolution_fwd_instance<2>(conv_ptrs); - conv::ConvFwdOpInstance conv_instance(params); - - auto reference_conv_fwd_fun = std::bind( - conv::run_reference_convolution_forward<2, float, float, float>, params, _1, _2, _3); - OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); - run_engine.SetAtol(1e-5); - run_engine.SetRtol(1e-4); - return run_engine.Test(conv_ptrs); -} - template bool test_conv2d_nhwc_instances(const std::vector& conv_ptrs) { @@ -57,50 +34,58 @@ bool test_conv2d_nhwc_instances(const std::vector( - ck::utils::conv::ConvolutionFwdInstances::Get<2>()); -} - -bool test_conv2d_nhwc_f16_instances() -{ - return test_conv2d_nhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<2>()); -} - -bool test_conv2d_nhwc_f32_instances() -{ - return test_conv2d_nhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<2>()); -} - -bool test_conv2d_nhwc_int8_instances() -{ - return test_conv2d_nhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<2>()); -} - } // anonymous namespace -int main() +TEST(Conv2DFwdNHWC, TestConv2D) { - bool res{true}; - res = test_conv2d_nhwc(); - std::cout << "test_conv2d_nhwc ..... " << (res ? "SUCCESS" : "FAILURE") << std::endl; + using namespace std::placeholders; + using namespace ck::utils; - res = test_conv2d_nhwc_bf16_instances(); - std::cout << "\ntest_conv2d_nhwc_bf16_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv2d_nhwc_f16_instances(); - std::cout << "\ntest_conv2d_nhwc_f16_instances ....." << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv2d_nhwc_f32_instances(); - std::cout << "\ntest_conv2d_nhwc_f32_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv2d_nhwc_int8_instances(); - std::cout << "\ntest_conv2d_nhwc_int8_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; + ck::utils::conv::ConvParams params; + params.N = 2; + params.K = 16; + params.C = 4; + params.input_spatial_lengths = std::vector{16, 16}; + params.conv_filter_strides = std::vector{1, 1}; - return res ? 0 : 1; + std::vector conv_ptrs; + test::conv::get_test_convolution_fwd_instance<2>(conv_ptrs); + conv::ConvFwdOpInstance conv_instance(params); + + auto reference_conv_fwd_fun = std::bind( + conv::run_reference_convolution_forward<2, float, float, float>, params, _1, _2, _3); + OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); + run_engine.SetAtol(1e-5); + run_engine.SetRtol(1e-4); + EXPECT_TRUE(run_engine.Test(conv_ptrs)); +} + +TEST(Conv2DFwdNHWC, Bf16Instances) +{ + EXPECT_TRUE(test_conv2d_nhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<2>())); +} + +TEST(Conv2DFwdNHWC, F16Instances) +{ + EXPECT_TRUE(test_conv2d_nhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<2>())); +} + +TEST(Conv2DFwdNHWC, BF32Instances) +{ + EXPECT_TRUE(test_conv2d_nhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<2>())); +} + +TEST(Conv2DFwdNHWC, F32Instances) +{ + EXPECT_TRUE(test_conv2d_nhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<2>())); +} + +TEST(Conv2DFwdNHWC, Int8Instances) +{ + EXPECT_TRUE(test_conv2d_nhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<2>())); } diff --git a/test/convnd_fwd/conv3d_fwd.cpp b/test/convnd_fwd/conv3d_fwd.cpp index 2375148753..fc3da3e9c7 100644 --- a/test/convnd_fwd/conv3d_fwd.cpp +++ b/test/convnd_fwd/conv3d_fwd.cpp @@ -3,6 +3,7 @@ #include #include #include +#include "gtest/gtest.h" #include "data_type.hpp" #include "element_wise_operation.hpp" @@ -11,7 +12,34 @@ namespace { -bool test_conv3d_ndhwc() +template +bool test_conv3d_ndhwc_instances(const std::vector& conv_ptrs) +{ + using namespace std::placeholders; + using namespace ck::utils; + namespace ctl = ck::tensor_layout::convolution; + + conv::ConvParams params; + params.N = 64; + params.num_dim_spatial = 3; + params.filter_spatial_lengths = std::vector{3, 3, 2}; + params.input_spatial_lengths = std::vector{32, 32, 2}; + params.conv_filter_strides = std::vector{2, 2, 2}; + params.conv_filter_dilations = std::vector{1, 1, 1}; + params.input_left_pads = std::vector{1, 1, 1}; + params.input_right_pads = std::vector{1, 1, 1}; + + conv::ConvFwdOpInstance conv_instance(params); + + auto reference_conv_fwd_fun = + std::bind(conv::run_reference_convolution_forward<3, T, T, T>, params, _1, _2, _3); + OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); + return run_engine.Test(conv_ptrs); +} + +} // anonymous namespace + +TEST(Conv3DFwdNDHWC, TestConv3D) { using namespace std::placeholders; using namespace ck::utils; @@ -39,10 +67,10 @@ bool test_conv3d_ndhwc() OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); run_engine.SetAtol(1e-5); run_engine.SetRtol(1e-4); - return run_engine.Test(conv_ptrs); + EXPECT_TRUE(run_engine.Test(conv_ptrs)); } -bool test_conv3d_ndhwc_2gb_input() +TEST(Conv3DFwdNDHWC, InputOver2GB) { using PassThrough = ck::tensor_operation::element_wise::PassThrough; using namespace ck::utils; @@ -79,10 +107,10 @@ bool test_conv3d_ndhwc_2gb_input() PassThrough{}, PassThrough{}, PassThrough{}); - return !(conv_ptrs.back()->IsSupportedArgument(arg.get())); + EXPECT_FALSE(conv_ptrs.back()->IsSupportedArgument(arg.get())); } -bool test_conv3d_ndhwc_2gb_filters() +TEST(Conv3DFwdNDHWC, FiltersOver2GB) { using PassThrough = ck::tensor_operation::element_wise::PassThrough; using namespace ck::utils; @@ -119,10 +147,10 @@ bool test_conv3d_ndhwc_2gb_filters() PassThrough{}, PassThrough{}, PassThrough{}); - return !(conv_ptrs.back()->IsSupportedArgument(arg.get())); + EXPECT_FALSE(conv_ptrs.back()->IsSupportedArgument(arg.get())); } -bool test_conv3d_ndhwc_2gb_output() +TEST(Conv3DFwdNDHWC, OutputOver2GB) { using PassThrough = ck::tensor_operation::element_wise::PassThrough; using namespace ck::utils; @@ -158,88 +186,29 @@ bool test_conv3d_ndhwc_2gb_output() PassThrough{}, PassThrough{}, PassThrough{}); - return !(conv_ptrs.back()->IsSupportedArgument(arg.get())); + EXPECT_FALSE(conv_ptrs.back()->IsSupportedArgument(arg.get())); } -template -bool test_conv3d_ndhwc_instances(const std::vector& conv_ptrs) +TEST(Conv3DFwdNDHWC, Bf16Instances) { - using namespace std::placeholders; - using namespace ck::utils; - namespace ctl = ck::tensor_layout::convolution; - - conv::ConvParams params; - params.N = 64; - params.num_dim_spatial = 3; - params.filter_spatial_lengths = std::vector{3, 3, 2}; - params.input_spatial_lengths = std::vector{32, 32, 2}; - params.conv_filter_strides = std::vector{2, 2, 2}; - params.conv_filter_dilations = std::vector{1, 1, 1}; - params.input_left_pads = std::vector{1, 1, 1}; - params.input_right_pads = std::vector{1, 1, 1}; - - conv::ConvFwdOpInstance conv_instance(params); - - auto reference_conv_fwd_fun = - std::bind(conv::run_reference_convolution_forward<3, T, T, T>, params, _1, _2, _3); - OpInstanceRunEngine run_engine(conv_instance, reference_conv_fwd_fun); - return run_engine.Test(conv_ptrs); + EXPECT_TRUE(test_conv3d_ndhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<3>())); } -bool test_conv3d_ndhwc_bf16_instances() +TEST(Conv3DFwdNDHWC, F16Instances) { - return test_conv3d_ndhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<3>()); + EXPECT_TRUE(test_conv3d_ndhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<3>())); } -bool test_conv3d_ndhwc_f16_instances() +TEST(Conv3DFwdNDHWC, F32Instances) { - return test_conv3d_ndhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<3>()); + EXPECT_TRUE(test_conv3d_ndhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<3>())); } -bool test_conv3d_ndhwc_f32_instances() +TEST(Conv3DFwdNDHWC, Int8Instances) { - return test_conv3d_ndhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<3>()); -} - -bool test_conv3d_ndhwc_int8_instances() -{ - return test_conv3d_ndhwc_instances( - ck::utils::conv::ConvolutionFwdInstances::Get<3>()); -} - -} // anonymous namespace - -int main() -{ - bool res{true}; - res = test_conv3d_ndhwc(); - std::cout << "test_conv3d_ndhwc ..... " << (res ? "SUCCESS" : "FAILURE") << std::endl; - - res = test_conv3d_ndhwc_2gb_input(); - std::cout << "\ntest_conv3d_ndhwc_2gb_input ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv3d_ndhwc_2gb_filters(); - std::cout << "\ntest_conv3d_ndhwc_2gb_filters ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv3d_ndhwc_2gb_output(); - std::cout << "\ntest_conv3d_ndhwc_2gb_output ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - - res = test_conv3d_ndhwc_bf16_instances(); - std::cout << "\ntest_conv3d_ndhwc_bf16_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv3d_ndhwc_f16_instances(); - std::cout << "\ntest_conv3d_ndhwc_f16_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv3d_ndhwc_f32_instances(); - std::cout << "\ntest_conv3d_ndhwc_f32_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - res = test_conv3d_ndhwc_int8_instances(); - std::cout << "\ntest_conv3d_ndhwc_int8_instances ..... " << (res ? "SUCCESS" : "FAILURE") - << std::endl; - - return res ? 0 : 1; + EXPECT_TRUE(test_conv3d_ndhwc_instances( + ck::utils::conv::ConvolutionFwdInstances::Get<3>())); } diff --git a/test/reference_conv_fwd/CMakeLists.txt b/test/reference_conv_fwd/CMakeLists.txt index 9d0bf45ef5..e5a7b31aff 100644 --- a/test/reference_conv_fwd/CMakeLists.txt +++ b/test/reference_conv_fwd/CMakeLists.txt @@ -1,2 +1,2 @@ -add_test_executable(test_reference_conv_fwd reference_conv_fwd.cpp) +add_gtest_executable(test_reference_conv_fwd reference_conv_fwd.cpp) target_link_libraries(test_reference_conv_fwd PRIVATE host_tensor conv_fwd_util) diff --git a/test/reference_conv_fwd/reference_conv_fwd.cpp b/test/reference_conv_fwd/reference_conv_fwd.cpp index e163298041..f660559e62 100644 --- a/test/reference_conv_fwd/reference_conv_fwd.cpp +++ b/test/reference_conv_fwd/reference_conv_fwd.cpp @@ -4,6 +4,7 @@ #include #include #include +#include "gtest/gtest.h" #include "check_err.hpp" #include "config.hpp" @@ -82,13 +83,13 @@ run_reference_convolution_forward(const ck::utils::conv::ConvParams& params, OutElementOp{}); ref_invoker.Run(ref_argument); - // std::cout <<"output: " << host_output.mDesc << std::endl << host_output.mData << std::endl; return host_output; } -bool test_conv2d_nhwc() +} // anonymous namespace + +TEST(ReferenceConvolutionFWD, Conv2DNHWC) { - bool res{true}; ck::utils::conv::ConvParams params; params.N = 1; params.K = 1; @@ -118,11 +119,14 @@ bool test_conv2d_nhwc() 472.5, 490.5, 508.5}; - res = res && ck::utils::check_err(out_tensor.mDesc.GetLengths(), - ref_dims, - "Error: wrong output tensor dimensions!"); - res = res && ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!"); + EXPECT_TRUE(ck::utils::check_err( + out_tensor.mDesc.GetLengths(), ref_dims, "Error: wrong output tensor dimensions!")); + EXPECT_TRUE(ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!")); +} +TEST(ReferenceConvolutionFWD, Conv2DNHWCStridesDilationsPadding) +{ + ck::utils::conv::ConvParams params; params.N = 1; params.K = 2; params.C = 2; @@ -133,25 +137,21 @@ bool test_conv2d_nhwc() params.input_left_pads = std::vector{1, 1}; params.input_right_pads = std::vector{1, 1}; - out_tensor = run_reference_convolution_forward<2>(params); - ref_dims = std::vector{1, 2, 5, 5}; - ref_data = std::vector{ + auto out_tensor = run_reference_convolution_forward<2>(params); + std::vector ref_dims = std::vector{1, 2, 5, 5}; + std::vector ref_data{ 210., 210., 327., 327., 351., 351., 375., 375., 399., 399., 459., 459., 706.5, 706.5, 742.5, 742.5, 778.5, 778.5, 814.5, 814.5, 747., 747., 1138.5, 1138.5, 1174.5, 1174.5, 1210.5, 1210.5, 1246.5, 1246.5, 1035., 1035., 1570.5, 1570.5, 1606.5, 1606.5, 1642.5, 1642.5, 1678.5, 1678.5, 1323., 1323., 2002.5, 2002.5, 2038.5, 2038.5, 2074.5, 2074.5, 2110.5, 2110.5}; - res = res && ck::utils::check_err(out_tensor.mDesc.GetLengths(), - ref_dims, - "Error: wrong output tensor dimensions!"); - res = res && ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!"); - - return res; + EXPECT_TRUE(ck::utils::check_err( + out_tensor.mDesc.GetLengths(), ref_dims, "Error: wrong output tensor dimensions!")); + EXPECT_TRUE(ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!")); } -bool test_conv1d_nwc() +TEST(ReferenceConvolutionFWD, Conv1DNWC) { - bool res{true}; ck::utils::conv::ConvParams params; params.num_dim_spatial = 1; params.N = 1; @@ -174,11 +174,14 @@ bool test_conv1d_nwc() ck::tensor_layout::convolution::NWK>(params); std::vector ref_dims{1, 1, 4}; std::vector ref_data{7.5, 13.5, 19.5, 25.5}; - res = res && ck::utils::check_err(out_tensor.mDesc.GetLengths(), - ref_dims, - "Error: wrong output tensor dimensions!"); - res = res && ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!"); + EXPECT_TRUE(ck::utils::check_err( + out_tensor.mDesc.GetLengths(), ref_dims, "Error: wrong output tensor dimensions!")); + EXPECT_TRUE(ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!")); +} +TEST(ReferenceConvolutionFWD, Conv1DNWCStridesDilationsPadding) +{ + ck::utils::conv::ConvParams params; params.num_dim_spatial = 1; params.N = 1; params.K = 2; @@ -190,20 +193,24 @@ bool test_conv1d_nwc() params.input_left_pads = std::vector{1}; params.input_right_pads = std::vector{1}; - out_tensor = run_reference_convolution_forward<1, - float, - float, - float, - ck::tensor_layout::convolution::NWC, - ck::tensor_layout::convolution::KXC, - ck::tensor_layout::convolution::NWK>(params); - ref_dims = std::vector{1, 2, 5}; - ref_data = std::vector{9., 9., 19.5, 19.5, 31.5, 31.5, 43.5, 43.5, 55.5, 55.5}; - res = res && ck::utils::check_err(out_tensor.mDesc.GetLengths(), - ref_dims, - "Error: wrong output tensor dimensions!"); - res = res && ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!"); + auto out_tensor = + run_reference_convolution_forward<1, + float, + float, + float, + ck::tensor_layout::convolution::NWC, + ck::tensor_layout::convolution::KXC, + ck::tensor_layout::convolution::NWK>(params); + std::vector ref_dims{1, 2, 5}; + std::vector ref_data{9., 9., 19.5, 19.5, 31.5, 31.5, 43.5, 43.5, 55.5, 55.5}; + EXPECT_TRUE(ck::utils::check_err( + out_tensor.mDesc.GetLengths(), ref_dims, "Error: wrong output tensor dimensions!")); + EXPECT_TRUE(ck::utils::check_err(out_tensor.mData, ref_data, "Error: incorrect results!")); +} +TEST(ReferenceConvolutionFWD, Conv1DNWCSameOutputSize) +{ + ck::utils::conv::ConvParams params; params.num_dim_spatial = 1; params.N = 2; params.K = 16; @@ -224,8 +231,8 @@ bool test_conv1d_nwc() ck::tensor_layout::convolution::NWK>( params, ck::utils::FillMonotonicSeq{0.f, 0.1f}); - ref_dims = std::vector{2, 16, 16}; - ref_data = std::vector{ + std::vector ref_dims{2, 16, 16}; + std::vector ref_data{ 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 1.4, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3, 3.3, @@ -290,17 +297,13 @@ bool test_conv1d_nwc() 72.9, 72.9, 72.9, 72.9, 72.9, 72.9, 72.9, 72.9, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4, 49.4}; - res = res && ck::utils::check_err(out_tensor2.mDesc.GetLengths(), - ref_dims, - "Error: wrong output tensor dimensions!"); - res = res && ck::utils::check_err(out_tensor2.mData, ref_data, "Error: incorrect results!"); - - return res; + EXPECT_TRUE(ck::utils::check_err( + out_tensor2.mDesc.GetLengths(), ref_dims, "Error: wrong output tensor dimensions!")); + EXPECT_TRUE(ck::utils::check_err(out_tensor2.mData, ref_data, "Error: incorrect results!")); } -bool test_conv3d_ncdhw() +TEST(ReferenceConvolutionFWD, Conv3DNCDHW) { - bool res{true}; ck::utils::conv::ConvParams params; params.num_dim_spatial = 3; params.N = 1; @@ -331,12 +334,17 @@ bool test_conv3d_ncdhw() 634.5, 637.2, 639.9, 642.60004, 650.7, 653.4, 656.10004, 658.8, 699.3, 702., 704.7, 707.4, 715.5, 718.2, 720.9, 723.60004, 731.7, 734.4001, 737.10004, 739.8, 747.9001, 750.60004, 753.3, 756.}; - res = res && ck::utils::check_err(out_tensor.mDesc.GetLengths(), - ref_dims, - "Error [case 1]: wrong output tensor dimensions!"); - res = res && - ck::utils::check_err(out_tensor.mData, ref_data, "Error [case 1]: incorrect results!"); + EXPECT_TRUE(ck::utils::check_err(out_tensor.mDesc.GetLengths(), + ref_dims, + "Error [case 1]: wrong output tensor dimensions!")); + EXPECT_TRUE( + ck::utils::check_err(out_tensor.mData, ref_data, "Error [case 1]: incorrect results!")); +} +TEST(ReferenceConvolutionFWD, Conv3DNCDHWStridesDilations) +{ + ck::utils::conv::ConvParams params; + params.num_dim_spatial = 3; params.N = 1; params.K = 2; params.C = 2; @@ -347,16 +355,16 @@ bool test_conv3d_ncdhw() params.input_left_pads = std::vector{0, 0, 0}; params.input_right_pads = std::vector{0, 0, 0}; - out_tensor = run_reference_convolution_forward<3, - float, - float, - float, - ck::tensor_layout::convolution::NCDHW, - ck::tensor_layout::convolution::KCZYX, - ck::tensor_layout::convolution::NKDHW>( + auto out_tensor = run_reference_convolution_forward<3, + float, + float, + float, + ck::tensor_layout::convolution::NCDHW, + ck::tensor_layout::convolution::KCZYX, + ck::tensor_layout::convolution::NKDHW>( params, ck::utils::FillMonotonicSeq{0.f, 0.1f}); - ref_dims = std::vector{1, 2, 4, 4, 4}; - ref_data = std::vector{ + std::vector ref_dims{1, 2, 4, 4, 4}; + std::vector ref_data{ 2756.7002, 2764.7998, 2772.9001, 2781., 2853.9001, 2862., 2870.1, 2878.2002, 2951.1, 2959.2002, 2967.2998, 2975.4001, 3048.2998, 3056.4001, 3064.5, 3072.6, 3923.1, 3931.2, 3939.2998, 3947.4, 4020.2998, 4028.4001, 4036.5002, 4044.5999, @@ -373,26 +381,9 @@ bool test_conv3d_ncdhw() 5283.9004, 5292., 5300.0996, 5308.2, 5381.0996, 5389.2, 5397.3, 5405.4004, 6255.9004, 6264.0005, 6272.1, 6280.2, 6353.1, 6361.2, 6369.301, 6377.4, 6450.301, 6458.4, 6466.5, 6474.6, 6547.5, 6555.6, 6563.699, 6571.801}; - res = res && ck::utils::check_err(out_tensor.mDesc.GetLengths(), - ref_dims, - "Error [case 2]: wrong output tensor dimensions!"); - res = - res && ck::utils::check_err( - out_tensor.mData, ref_data, "Error [case 2]: incorrect results!", 1e-4f, 1e-6f); - - return res; -} - -} // anonymous namespace - -int main(void) -{ - bool res{true}; - res = test_conv2d_nhwc(); - std::cout << "test_conv2d_nhwc ..... " << (res ? "SUCCESS" : "FAILURE") << std::endl; - res = test_conv1d_nwc(); - std::cout << "TestConv1DNHWC ..... " << (res ? "SUCCESS" : "FAILURE") << std::endl; - res = test_conv3d_ncdhw(); - std::cout << "test_conv3d_ncdhw ..... " << (res ? "SUCCESS" : "FAILURE") << std::endl; - return res ? 0 : 1; + EXPECT_TRUE(ck::utils::check_err(out_tensor.mDesc.GetLengths(), + ref_dims, + "Error [case 2]: wrong output tensor dimensions!")); + EXPECT_TRUE(ck::utils::check_err( + out_tensor.mData, ref_data, "Error [case 2]: incorrect results!", 1e-4f, 1e-6f)); }