From 408538549881daa5319ea9a98b379aa0db8dd786 Mon Sep 17 00:00:00 2001 From: Anthony Chang Date: Sat, 14 May 2022 05:54:44 +0800 Subject: [PATCH] Validate examples in CI (#233) * validate examples in ctest runs * format * fix usage of check_err * amend * add example codes to custom target 'check' Co-authored-by: Chao Liu [ROCm/composable_kernel commit: 9f71ff48e28709c8132735d80af57ec90626d4b5] --- CMakeLists.txt | 10 +++--- example/01_gemm/gemm_xdl_bf16.cpp | 2 +- example/01_gemm/gemm_xdl_fp16.cpp | 2 +- example/01_gemm/gemm_xdl_int8.cpp | 2 +- .../gemm_xdl_alpha_beta.cpp | 4 ++- .../03_gemm_bias_relu/gemm_xdl_bias_relu.cpp | 4 ++- .../gemm_xdl_bias_relu_add.cpp | 4 ++- .../conv2d_fwd_xdl_bias_relu.cpp | 5 +-- .../CMakeLists.txt | 3 +- .../conv2d_fwd_xdl_bias_relu_add.cpp | 5 +-- example/09_convnd_fwd/CMakeLists.txt | 6 ++-- example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp | 9 ++--- ...nd_fwd_xdl.cpp => convnd_fwd_xdl_fp32.cpp} | 14 +++++--- example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp | 9 ++--- .../conv2d_bwd_data_xdl.cpp | 6 +++- .../conv2d_bwd_weight_xdl.cpp | 5 ++- example/12_reduce/CMakeLists.txt | 2 +- example/12_reduce/reduce_blockwise.cpp | 7 ++-- example/13_pool2d_fwd/pool2d_fwd.cpp | 8 +++-- .../gemm_xdl_requant_relu_requant_int8.cpp | 2 +- .../15_grouped_gemm/grouped_gemm_xdl_fp16.cpp | 5 +-- .../16_gemm_reduce/gemm_reduce_xdl_fp16.cpp | 19 ++++++++--- .../convnd_bwd_data_xdl.cpp | 6 +++- .../batched_gemm_reduce_xdl_fp16.cpp | 34 ++++++++++++------- example/CMakeLists.txt | 13 +++++-- test/CMakeLists.txt | 1 - 26 files changed, 125 insertions(+), 62 deletions(-) rename example/09_convnd_fwd/{convnd_fwd_xdl.cpp => convnd_fwd_xdl_fp32.cpp} (97%) diff --git a/CMakeLists.txt b/CMakeLists.txt index f18c85c683..a3ec91e3bc 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -245,6 +245,8 @@ if(BUILD_DEV) endif() message("CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}") +add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR}) + add_subdirectory(library) add_subdirectory(example) add_subdirectory(test) @@ -260,14 +262,14 @@ write_basic_package_version_file( COMPATIBILITY AnyNewerVersion ) -configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/Config.cmake.in +configure_package_config_file(${CMAKE_CURRENT_SOURCE_DIR}/Config.cmake.in "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake" - INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel + INSTALL_DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel NO_CHECK_REQUIRED_COMPONENTS_MACRO ) -install(FILES +install(FILES "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfig.cmake" "${CMAKE_CURRENT_BINARY_DIR}/composable_kernelConfigVersion.cmake" - DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel + DESTINATION ${CMAKE_INSTALL_LIBDIR}/cmake/composable_kernel ) diff --git a/example/01_gemm/gemm_xdl_bf16.cpp b/example/01_gemm/gemm_xdl_bf16.cpp index 4077a4f8d8..060750e676 100644 --- a/example/01_gemm/gemm_xdl_bf16.cpp +++ b/example/01_gemm/gemm_xdl_bf16.cpp @@ -232,7 +232,7 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); - ck::utils::check_err(c_m_n_device_f32_result.mData, c_m_n_host_result.mData); + return ck::utils::check_err(c_m_n_device_f32_result.mData, c_m_n_host_result.mData) ? 0 : 1; } return 0; diff --git a/example/01_gemm/gemm_xdl_fp16.cpp b/example/01_gemm/gemm_xdl_fp16.cpp index 4f0228eafe..06523037f9 100644 --- a/example/01_gemm/gemm_xdl_fp16.cpp +++ b/example/01_gemm/gemm_xdl_fp16.cpp @@ -196,7 +196,7 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); - ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData); + return ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData) ? 0 : 1; } return 0; diff --git a/example/01_gemm/gemm_xdl_int8.cpp b/example/01_gemm/gemm_xdl_int8.cpp index d5bf4a8bde..a22c21e40e 100644 --- a/example/01_gemm/gemm_xdl_int8.cpp +++ b/example/01_gemm/gemm_xdl_int8.cpp @@ -219,7 +219,7 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); - ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData); + return ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData) ? 0 : 1; } return 0; diff --git a/example/02_gemm_alpha_beta/gemm_xdl_alpha_beta.cpp b/example/02_gemm_alpha_beta/gemm_xdl_alpha_beta.cpp index 451200e798..1a6e1de4dc 100644 --- a/example/02_gemm_alpha_beta/gemm_xdl_alpha_beta.cpp +++ b/example/02_gemm_alpha_beta/gemm_xdl_alpha_beta.cpp @@ -246,6 +246,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); - ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData); + return ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData) ? 0 : 1; } + + return 0; } diff --git a/example/03_gemm_bias_relu/gemm_xdl_bias_relu.cpp b/example/03_gemm_bias_relu/gemm_xdl_bias_relu.cpp index 308d423ce7..3bf3003c14 100644 --- a/example/03_gemm_bias_relu/gemm_xdl_bias_relu.cpp +++ b/example/03_gemm_bias_relu/gemm_xdl_bias_relu.cpp @@ -232,6 +232,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); - ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData); + return ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData) ? 0 : 1; } + + return 0; } diff --git a/example/04_gemm_bias_relu_add/gemm_xdl_bias_relu_add.cpp b/example/04_gemm_bias_relu_add/gemm_xdl_bias_relu_add.cpp index 012fd21341..73e92f9d11 100644 --- a/example/04_gemm_bias_relu_add/gemm_xdl_bias_relu_add.cpp +++ b/example/04_gemm_bias_relu_add/gemm_xdl_bias_relu_add.cpp @@ -250,6 +250,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); - ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData); + return ck::utils::check_err(c_m_n_device_result.mData, c_m_n_host_result.mData) ? 0 : 1; } + + return 0; } diff --git a/example/06_conv2d_fwd_bias_relu/conv2d_fwd_xdl_bias_relu.cpp b/example/06_conv2d_fwd_bias_relu/conv2d_fwd_xdl_bias_relu.cpp index 342de268e3..d50afb6854 100644 --- a/example/06_conv2d_fwd_bias_relu/conv2d_fwd_xdl_bias_relu.cpp +++ b/example/06_conv2d_fwd_bias_relu/conv2d_fwd_xdl_bias_relu.cpp @@ -305,7 +305,8 @@ int main(int argc, char* argv[]) OutElementOp{}); ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - ck::utils::check_err( - host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f); + return ck::utils::check_err(device_output.mData, host_output.mData) ? 0 : 1; } + + return 0; } diff --git a/example/07_conv2d_fwd_bias_relu_add/CMakeLists.txt b/example/07_conv2d_fwd_bias_relu_add/CMakeLists.txt index 5f6426ff1f..b4dd39d83a 100644 --- a/example/07_conv2d_fwd_bias_relu_add/CMakeLists.txt +++ b/example/07_conv2d_fwd_bias_relu_add/CMakeLists.txt @@ -1,2 +1,3 @@ -add_example_executable(example_conv2d_fwd_xdl_bias_relu_add conv2d_fwd_xdl_bias_relu_add.cpp) +# FIXME: should fix validation failure +add_example_executable_no_testing(example_conv2d_fwd_xdl_bias_relu_add conv2d_fwd_xdl_bias_relu_add.cpp) target_link_libraries(example_conv2d_fwd_xdl_bias_relu_add PRIVATE conv_util) diff --git a/example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp b/example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp index ff4fc66cb8..53d882778a 100644 --- a/example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp +++ b/example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp @@ -320,7 +320,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - ck::utils::check_err( - host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f); + return ck::utils::check_err(device_output.mData, host_output.mData) ? 0 : 1; } + + return 0; } diff --git a/example/09_convnd_fwd/CMakeLists.txt b/example/09_convnd_fwd/CMakeLists.txt index 9ffae06233..ceceb4aedc 100644 --- a/example/09_convnd_fwd/CMakeLists.txt +++ b/example/09_convnd_fwd/CMakeLists.txt @@ -1,6 +1,6 @@ -add_example_executable(example_convnd_fwd_xdl convnd_fwd_xdl.cpp) -target_link_libraries(example_convnd_fwd_xdl PRIVATE conv_util) +add_example_executable(example_convnd_fwd_xdl_fp32 convnd_fwd_xdl_fp32.cpp) add_example_executable(example_convnd_fwd_xdl_int8 convnd_fwd_xdl_int8.cpp) -target_link_libraries(example_convnd_fwd_xdl_int8 PRIVATE conv_util) add_example_executable(example_convnd_fwd_xdl_fp16 convnd_fwd_xdl_fp16.cpp) +target_link_libraries(example_convnd_fwd_xdl_fp32 PRIVATE conv_util) +target_link_libraries(example_convnd_fwd_xdl_int8 PRIVATE conv_util) target_link_libraries(example_convnd_fwd_xdl_fp16 PRIVATE conv_util) diff --git a/example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp b/example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp index 8b658e7790..7ad83d5ad6 100644 --- a/example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp +++ b/example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp @@ -43,10 +43,10 @@ template using DeviceConvNDFwdInstance = ck::tensor_operation::device:: DeviceConvNDFwdXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K< // clang-format off - InDataType, // + InDataType, // WeiDataType, // OutDataType, // - AccDataType, // + AccDataType, // InElementOp, // Input Elementwise Operation WeiElementOp, // Weights Elementwise Operation OutElementOp, // Output Elementwise Operation @@ -312,8 +312,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - ck::utils::check_err( - host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f); + return ck::utils::check_err( + host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f) ? 0 : 1; }; switch(num_dim_spatial) @@ -338,4 +338,5 @@ int main(int argc, char* argv[]) } } } + return 0; } diff --git a/example/09_convnd_fwd/convnd_fwd_xdl.cpp b/example/09_convnd_fwd/convnd_fwd_xdl_fp32.cpp similarity index 97% rename from example/09_convnd_fwd/convnd_fwd_xdl.cpp rename to example/09_convnd_fwd/convnd_fwd_xdl_fp32.cpp index 112d606f56..8a9633d84a 100644 --- a/example/09_convnd_fwd/convnd_fwd_xdl.cpp +++ b/example/09_convnd_fwd/convnd_fwd_xdl_fp32.cpp @@ -39,10 +39,10 @@ template using DeviceConvNDFwdInstance = ck::tensor_operation::device:: DeviceConvNDFwdXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K< // clang-format off - InDataType, // + InDataType, // WeiDataType, // OutDataType, // - AccDataType, // + AccDataType, // InElementOp, // Input Elementwise Operation WeiElementOp, // Weights Elementwise Operation OutElementOp, // Output Elementwise Operation @@ -311,8 +311,13 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - ck::utils::check_err( - host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f); + return ck::utils::check_err(device_output.mData, + host_output.mData, + "Error: incorrect results!", + 1e-5f, + 1e-4f) + ? 0 + : 1; }; switch(num_dim_spatial) @@ -337,4 +342,5 @@ int main(int argc, char* argv[]) } } } + return 0; } diff --git a/example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp b/example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp index e7988d8683..f196d27182 100644 --- a/example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp +++ b/example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp @@ -45,10 +45,10 @@ template using DeviceConvNDFwdInstance = ck::tensor_operation::device:: DeviceConvNDFwdXdl_Input_N_Hi_Wi_C_Weight_K_Y_X_C_Output_N_Ho_Wo_K< // clang-format off - InDataType, // + InDataType, // WeiDataType, // OutDataType, // - AccDataType, // + AccDataType, // InElementOp, // Input Elementwise Operation WeiElementOp, // Weights Elementwise Operation OutElementOp, // Output Elementwise Operation @@ -314,8 +314,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - ck::utils::check_err( - host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f); + return ck::utils::check_err( + host_output.mData, device_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f) ? 0 : 1; }; switch(num_dim_spatial) @@ -340,4 +340,5 @@ int main(int argc, char* argv[]) } } } + return 0; } diff --git a/example/10_conv2d_bwd_data/conv2d_bwd_data_xdl.cpp b/example/10_conv2d_bwd_data/conv2d_bwd_data_xdl.cpp index 73210fa543..2d25f5ac2f 100644 --- a/example/10_conv2d_bwd_data/conv2d_bwd_data_xdl.cpp +++ b/example/10_conv2d_bwd_data/conv2d_bwd_data_xdl.cpp @@ -249,6 +249,10 @@ int main(int argc, char* argv[]) in_device_buf.FromDevice(in_n_c_hi_wi_device_result.mData.data()); - ck::utils::check_err(in_n_c_hi_wi_device_result.mData, in_n_c_hi_wi_host_result.mData); + return ck::utils::check_err(in_n_c_hi_wi_device_result.mData, + in_n_c_hi_wi_host_result.mData) + ? 0 + : 1; } + return 0; } diff --git a/example/11_conv2d_bwd_weight/conv2d_bwd_weight_xdl.cpp b/example/11_conv2d_bwd_weight/conv2d_bwd_weight_xdl.cpp index 0c996dc21b..1578161116 100644 --- a/example/11_conv2d_bwd_weight/conv2d_bwd_weight_xdl.cpp +++ b/example/11_conv2d_bwd_weight/conv2d_bwd_weight_xdl.cpp @@ -291,6 +291,9 @@ int main(int argc, char* argv[]) LogRangeAsType(std::cout << "wei_host : ", wei_k_c_y_x_host_result.mData, ",") << std::endl; } - ck::utils::check_err(wei_k_c_y_x_device_result.mData, wei_k_c_y_x_host_result.mData); + return ck::utils::check_err(wei_k_c_y_x_device_result.mData, wei_k_c_y_x_host_result.mData) + ? 0 + : 1; } + return 0; } diff --git a/example/12_reduce/CMakeLists.txt b/example/12_reduce/CMakeLists.txt index 734c1955d6..d6866abeb8 100644 --- a/example/12_reduce/CMakeLists.txt +++ b/example/12_reduce/CMakeLists.txt @@ -1 +1 @@ -add_example_executable(example_reduce_blockwise reduce_blockwise.cpp) +add_example_executable(example_reduce_blockwise reduce_blockwise.cpp -D 16,64,32,960 -v 1 1 10) diff --git a/example/12_reduce/reduce_blockwise.cpp b/example/12_reduce/reduce_blockwise.cpp index caa93c9df2..b2d312ae8c 100644 --- a/example/12_reduce/reduce_blockwise.cpp +++ b/example/12_reduce/reduce_blockwise.cpp @@ -361,16 +361,17 @@ int main(int argc, char* argv[]) std::cout << "Perf: " << avg_time << " ms, " << gb_per_sec << " GB/s, " << reduce_name << std::endl; + bool pass = true; if(args.do_verification) { out_dev.FromDevice(out.mData.data()); - ck::utils::check_err(out.mData, out_ref.mData); + pass &= ck::utils::check_err(out.mData, out_ref.mData); if(NeedIndices) { out_indices_dev.FromDevice(out_indices.mData.data()); - ck::utils::check_err(out_indices.mData, out_indices_ref.mData); - ; + pass &= ck::utils::check_err(out_indices.mData, out_indices_ref.mData); }; }; + return pass ? 0 : 1; } diff --git a/example/13_pool2d_fwd/pool2d_fwd.cpp b/example/13_pool2d_fwd/pool2d_fwd.cpp index f4eb9d79f6..e6749bf8d7 100644 --- a/example/13_pool2d_fwd/pool2d_fwd.cpp +++ b/example/13_pool2d_fwd/pool2d_fwd.cpp @@ -285,6 +285,7 @@ int main(int argc, char* argv[]) std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s" << std::endl; + bool pass = true; if(do_verification) { pool_host_verify #include #include +#include "check_err.hpp" #include "config.hpp" #include "device.hpp" #include "host_tensor.hpp" @@ -211,6 +212,7 @@ int main(int argc, char* argv[]) std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " << gemm.GetTypeString() << std::endl; + bool pass = true; if(do_verification) { c_device_buf.FromDevice(c_m_n_device_result.mData.data()); @@ -247,10 +249,19 @@ int main(int argc, char* argv[]) d1_m_host_result(m) = ck::type_convert(d1_acc); } - check_error(c_m_n_host_result, c_m_n_device_result); - check_error(d0_m_host_result, d0_m_device_result); - check_error(d1_m_host_result, d1_m_device_result); + pass &= ck::utils::check_err( + c_m_n_device_result.mData, c_m_n_host_result.mData, "Error: Incorrect results c"); + pass &= ck::utils::check_err(d0_m_device_result.mData, + d0_m_host_result.mData, + "Error: Incorrect results d0", + 1e-3, + 1e-3); + pass &= ck::utils::check_err(d1_m_device_result.mData, + d1_m_host_result.mData, + "Error: Incorrect results d1", + 1e-3, + 1e-3); } - return 0; + return pass ? 0 : 1; } diff --git a/example/17_convnd_bwd_data_xdl/convnd_bwd_data_xdl.cpp b/example/17_convnd_bwd_data_xdl/convnd_bwd_data_xdl.cpp index a013f39827..ff2cfac1fa 100644 --- a/example/17_convnd_bwd_data_xdl/convnd_bwd_data_xdl.cpp +++ b/example/17_convnd_bwd_data_xdl/convnd_bwd_data_xdl.cpp @@ -322,7 +322,10 @@ int main(int argc, char* argv[]) in_device_buf.FromDevice(in_n_c_hi_wi_device_result.mData.data()); - check_error(in_n_c_hi_wi_host_result, in_n_c_hi_wi_device_result); + return ck::utils::check_err(in_n_c_hi_wi_device_result.mData, + in_n_c_hi_wi_host_result.mData) + ? 0 + : 1; }; switch(num_dim_spatial) @@ -347,4 +350,5 @@ int main(int argc, char* argv[]) } } } + return 0; } diff --git a/example/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp b/example/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp index f620ee1b20..d993c8e8d1 100644 --- a/example/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp +++ b/example/18_batched_gemm_reduce/batched_gemm_reduce_xdl_fp16.cpp @@ -4,6 +4,7 @@ #include #include #include +#include "check_err.hpp" #include "config.hpp" #include "device.hpp" #include "host_tensor.hpp" @@ -62,13 +63,13 @@ int main(int argc, char* argv[]) bool time_kernel = false; // GEMM shape - ck::index_t M = 3840; - ck::index_t N = 4096; - ck::index_t K = 4096; + ck::index_t M = 2048; + ck::index_t N = 1920; + ck::index_t K = 2048; - ck::index_t StrideA = 4096; - ck::index_t StrideB = 4096; - ck::index_t StrideC = 4096; + ck::index_t StrideA = 2048; + ck::index_t StrideB = 2048; + ck::index_t StrideC = 1920; ck::index_t BatchCount = 4; @@ -96,7 +97,7 @@ int main(int argc, char* argv[]) StrideB = std::stoi(argv[8]); StrideC = std::stoi(argv[9]); - BatchCount = std::stoi(argv[9]); + BatchCount = std::stoi(argv[10]); } else { @@ -224,6 +225,7 @@ int main(int argc, char* argv[]) std::cout << "Perf: " << ave_time << " ms, " << tflops << " TFlops, " << gb_per_sec << " GB/s, " << batched_gemm.GetTypeString() << std::endl; + bool pass = true; if(do_verification) { c_device_buf.FromDevice(c_g_m_n_device_result.mData.data()); @@ -247,7 +249,7 @@ int main(int argc, char* argv[]) for(int n = 0; n < N; ++n) { - float d0_val = ck::type_convert(c_g_m_n_host_result(m, n)); + float d0_val = ck::type_convert(c_g_m_n_host_result(batch, m, n)); float d1_val; d1_element_op(d1_val, d0_val); @@ -260,10 +262,18 @@ int main(int argc, char* argv[]) } } - check_error(c_g_m_n_host_result, c_g_m_n_device_result); - check_error(d0_g_m_host_result, d0_g_m_device_result); - check_error(d1_g_m_host_result, d1_g_m_device_result); + pass &= ck::utils::check_err(c_g_m_n_host_result.mData, c_g_m_n_device_result.mData); + pass &= ck::utils::check_err(d0_g_m_device_result.mData, + d0_g_m_host_result.mData, + "Error: Incorrect results! D0", + 1e-3, + 1e-3); + pass &= ck::utils::check_err(d1_g_m_device_result.mData, + d1_g_m_host_result.mData, + "Error: Incorrect results! D1", + 1e-3, + 1e-3); } - return 0; + return pass ? 0 : 1; } diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 5f04125305..4d81e84c01 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -19,9 +19,18 @@ include_directories(BEFORE add_custom_target(examples) -function(add_example_executable EXAMPLE_NAME) +function(add_example_executable EXAMPLE_NAME FILE_NAME) message("adding example ${EXAMPLE_NAME}") - add_executable(${EXAMPLE_NAME} ${ARGN}) + add_executable(${EXAMPLE_NAME} ${FILE_NAME}) + target_link_libraries(${EXAMPLE_NAME} PRIVATE host_tensor) + add_test(NAME ${EXAMPLE_NAME} COMMAND $ ${ARGN}) + add_dependencies(examples ${EXAMPLE_NAME}) + add_dependencies(check ${EXAMPLE_NAME}) +endfunction(add_example_executable EXAMPLE_NAME) + +function(add_example_executable_no_testing EXAMPLE_NAME FILE_NAME) + message("adding example ${EXAMPLE_NAME}") + add_executable(${EXAMPLE_NAME} ${FILE_NAME}) target_link_libraries(${EXAMPLE_NAME} PRIVATE host_tensor) add_dependencies(examples ${EXAMPLE_NAME}) endfunction(add_example_executable EXAMPLE_NAME) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c696069393..3733563571 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -24,7 +24,6 @@ include_directories(BEFORE include(googletest) -add_custom_target(check COMMAND ${CMAKE_CTEST_COMMAND} --output-on-failure -C ${CMAKE_CFG_INTDIR}) add_custom_target(tests)