From 296ed2b1c3d4f8584d2439ab613c6e1e8818092c Mon Sep 17 00:00:00 2001 From: Anthony Chang Date: Thu, 12 May 2022 10:19:53 +0800 Subject: [PATCH 1/5] validate examples in ctest runs --- example/01_gemm/gemm_xdl_bf16.cpp | 2 +- example/01_gemm/gemm_xdl_fp16.cpp | 2 +- example/01_gemm/gemm_xdl_int8.cpp | 2 +- example/02_gemm_alpha_beta/gemm_xdl_alpha_beta.cpp | 4 +++- example/03_gemm_bias_relu/gemm_xdl_bias_relu.cpp | 4 +++- .../gemm_xdl_bias_relu_add.cpp | 4 +++- .../conv2d_fwd_xdl_bias_relu.cpp | 8 ++++++-- example/07_conv2d_fwd_bias_relu_add/CMakeLists.txt | 3 ++- .../conv2d_fwd_xdl_bias_relu_add.cpp | 9 +++++++-- example/09_convnd_fwd/convnd_fwd_xdl.cpp | 14 ++++++++++---- example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp | 9 +++++---- example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp | 9 +++++---- example/10_conv2d_bwd_data/conv2d_bwd_data_xdl.cpp | 6 +++++- .../11_conv2d_bwd_weight/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 +- example/15_grouped_gemm/grouped_gemm_xdl_fp16.cpp | 5 +++-- example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp | 10 ++++++---- .../17_convnd_bwd_data_xdl/convnd_bwd_data_xdl.cpp | 6 +++++- .../batched_gemm_reduce_xdl_fp16.cpp | 14 ++++++++------ example/CMakeLists.txt | 12 ++++++++++-- 23 files changed, 99 insertions(+), 48 deletions(-) diff --git a/example/01_gemm/gemm_xdl_bf16.cpp b/example/01_gemm/gemm_xdl_bf16.cpp index 4077a4f8d85..060750e6768 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 4f0228eafe3..06523037f96 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 d5bf4a8bde4..a22c21e40e2 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 451200e798b..1a6e1de4dcf 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 308d423ce7c..3bf3003c147 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 012fd21341b..73e92f9d116 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 342de268e35..f845989c55a 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,11 @@ 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(host_output.mData, + device_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 5f6426ff1f2..b4dd39d83a7 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 ff4fc66cb85..5ad82e2dd46 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,12 @@ 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!") + ? 0 + : 1; } + + return 0; } diff --git a/example/09_convnd_fwd/convnd_fwd_xdl.cpp b/example/09_convnd_fwd/convnd_fwd_xdl.cpp index 112d606f56b..70e5b261049 100644 --- a/example/09_convnd_fwd/convnd_fwd_xdl.cpp +++ b/example/09_convnd_fwd/convnd_fwd_xdl.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(host_output.mData, + device_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_fp16.cpp b/example/09_convnd_fwd/convnd_fwd_xdl_fp16.cpp index 8b658e77908..7ad83d5ad63 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_int8.cpp b/example/09_convnd_fwd/convnd_fwd_xdl_int8.cpp index e7988d8683e..f196d271828 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 73210fa543e..2d25f5ac2f1 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 0c996dc21b5..1578161116c 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 734c1955d6f..d6866abeb85 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 caa93c9df26..b2d312ae8cd 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 f4eb9d79f69..e6749bf8d7c 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,10 @@ 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_host_result.mData, c_m_n_device_result.mData, "Error: Incorrect results c"); + pass &= ck::utils::check_err(d0_m_host_result.mData, d0_m_device_result.mData, "Error: Incorrect results d0", 1e-3, 1e-3); + pass &= ck::utils::check_err(d1_m_host_result.mData, d1_m_device_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 a013f39827d..568888c2063 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_host_result.mData, + in_n_c_hi_wi_device_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 f620ee1b200..5810ca81dc1 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" @@ -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,10 @@ 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_host_result.mData, d0_g_m_device_result.mData, "Error: Incorrect results!", 1e-3, 1e-3); + pass &= ck::utils::check_err(d1_g_m_host_result.mData, d1_g_m_device_result.mData, "Error: Incorrect results!", 1e-3, 1e-3); } - return 0; + return pass ? 0 : 1; } diff --git a/example/CMakeLists.txt b/example/CMakeLists.txt index 5f041253056..ff07c55aec2 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -19,9 +19,17 @@ 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}) +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) From d40b7e5160dade301ce44c0a975036082568dc6b Mon Sep 17 00:00:00 2001 From: Anthony Chang Date: Thu, 12 May 2022 17:15:07 +0800 Subject: [PATCH 2/5] format --- .../conv2d_fwd_xdl_bias_relu.cpp | 5 +---- .../conv2d_fwd_xdl_bias_relu_add.cpp | 5 ++--- example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp | 15 ++++++++++++--- .../batched_gemm_reduce_xdl_fp16.cpp | 12 ++++++++++-- 4 files changed, 25 insertions(+), 12 deletions(-) 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 f845989c55a..50ca01a9458 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,10 +305,7 @@ int main(int argc, char* argv[]) OutElementOp{}); ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - return ck::utils::check_err(host_output.mData, - device_output.mData) - ? 0 - : 1; + return ck::utils::check_err(host_output.mData, device_output.mData) ? 0 : 1; } return 0; 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 5ad82e2dd46..75fdf2ae28b 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,9 +320,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - return ck::utils::check_err(host_output.mData, - device_output.mData, - "Error: incorrect results!") + return ck::utils::check_err( + host_output.mData, device_output.mData, "Error: incorrect results!") ? 0 : 1; } diff --git a/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp b/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp index 6b7c4b2fe19..51e5aba576e 100644 --- a/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp +++ b/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp @@ -249,9 +249,18 @@ int main(int argc, char* argv[]) d1_m_host_result(m) = ck::type_convert(d1_acc); } - pass &= ck::utils::check_err(c_m_n_host_result.mData, c_m_n_device_result.mData, "Error: Incorrect results c"); - pass &= ck::utils::check_err(d0_m_host_result.mData, d0_m_device_result.mData, "Error: Incorrect results d0", 1e-3, 1e-3); - pass &= ck::utils::check_err(d1_m_host_result.mData, d1_m_device_result.mData, "Error: Incorrect results d1", 1e-3, 1e-3); + pass &= ck::utils::check_err( + c_m_n_host_result.mData, c_m_n_device_result.mData, "Error: Incorrect results c"); + pass &= ck::utils::check_err(d0_m_host_result.mData, + d0_m_device_result.mData, + "Error: Incorrect results d0", + 1e-3, + 1e-3); + pass &= ck::utils::check_err(d1_m_host_result.mData, + d1_m_device_result.mData, + "Error: Incorrect results d1", + 1e-3, + 1e-3); } return pass ? 0 : 1; 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 5810ca81dc1..82a258b364f 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 @@ -263,8 +263,16 @@ int main(int argc, char* argv[]) } 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_host_result.mData, d0_g_m_device_result.mData, "Error: Incorrect results!", 1e-3, 1e-3); - pass &= ck::utils::check_err(d1_g_m_host_result.mData, d1_g_m_device_result.mData, "Error: Incorrect results!", 1e-3, 1e-3); + pass &= ck::utils::check_err(d0_g_m_host_result.mData, + d0_g_m_device_result.mData, + "Error: Incorrect results!", + 1e-3, + 1e-3); + pass &= ck::utils::check_err(d1_g_m_host_result.mData, + d1_g_m_device_result.mData, + "Error: Incorrect results!", + 1e-3, + 1e-3); } return pass ? 0 : 1; From 353bdc4cd62394e72d2b6f4b66564215f574001c Mon Sep 17 00:00:00 2001 From: Chao Liu Date: Thu, 12 May 2022 17:28:18 +0000 Subject: [PATCH 3/5] fix usage of check_err --- .../conv2d_fwd_xdl_bias_relu.cpp | 2 +- .../conv2d_fwd_xdl_bias_relu_add.cpp | 5 +--- example/09_convnd_fwd/CMakeLists.txt | 6 ++--- ...nd_fwd_xdl.cpp => convnd_fwd_xdl_fp32.cpp} | 4 ++-- .../16_gemm_reduce/gemm_reduce_xdl_fp16.cpp | 10 ++++---- .../convnd_bwd_data_xdl.cpp | 4 ++-- .../batched_gemm_reduce_xdl_fp16.cpp | 24 +++++++++---------- 7 files changed, 26 insertions(+), 29 deletions(-) rename example/09_convnd_fwd/{convnd_fwd_xdl.cpp => convnd_fwd_xdl_fp32.cpp} (99%) 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 50ca01a9458..d50afb6854c 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,7 @@ int main(int argc, char* argv[]) OutElementOp{}); ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - return ck::utils::check_err(host_output.mData, device_output.mData) ? 0 : 1; + 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/conv2d_fwd_xdl_bias_relu_add.cpp b/example/07_conv2d_fwd_bias_relu_add/conv2d_fwd_xdl_bias_relu_add.cpp index 75fdf2ae28b..e88c114c904 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,10 +320,7 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - return ck::utils::check_err( - host_output.mData, device_output.mData, "Error: incorrect results!") - ? 0 - : 1; + return ck::utils::check_err(device_output.mData, host_output.mData); } return 0; diff --git a/example/09_convnd_fwd/CMakeLists.txt b/example/09_convnd_fwd/CMakeLists.txt index 9ffae06233e..ceceb4aedc9 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.cpp b/example/09_convnd_fwd/convnd_fwd_xdl_fp32.cpp similarity index 99% rename from example/09_convnd_fwd/convnd_fwd_xdl.cpp rename to example/09_convnd_fwd/convnd_fwd_xdl_fp32.cpp index 70e5b261049..8a9633d84a9 100644 --- a/example/09_convnd_fwd/convnd_fwd_xdl.cpp +++ b/example/09_convnd_fwd/convnd_fwd_xdl_fp32.cpp @@ -311,8 +311,8 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - return ck::utils::check_err(host_output.mData, - device_output.mData, + return ck::utils::check_err(device_output.mData, + host_output.mData, "Error: incorrect results!", 1e-5f, 1e-4f) diff --git a/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp b/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp index 51e5aba576e..860d9eea2ac 100644 --- a/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp +++ b/example/16_gemm_reduce/gemm_reduce_xdl_fp16.cpp @@ -250,14 +250,14 @@ int main(int argc, char* argv[]) } pass &= ck::utils::check_err( - c_m_n_host_result.mData, c_m_n_device_result.mData, "Error: Incorrect results c"); - pass &= ck::utils::check_err(d0_m_host_result.mData, - d0_m_device_result.mData, + 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_host_result.mData, - d1_m_device_result.mData, + pass &= ck::utils::check_err(d1_m_device_result.mData, + d1_m_host_result.mData, "Error: Incorrect results d1", 1e-3, 1e-3); 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 568888c2063..ff2cfac1fa7 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,8 +322,8 @@ int main(int argc, char* argv[]) in_device_buf.FromDevice(in_n_c_hi_wi_device_result.mData.data()); - return ck::utils::check_err(in_n_c_hi_wi_host_result.mData, - in_n_c_hi_wi_device_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; }; 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 82a258b364f..d993c8e8d1b 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 @@ -63,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; @@ -263,14 +263,14 @@ int main(int argc, char* argv[]) } 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_host_result.mData, - d0_g_m_device_result.mData, - "Error: Incorrect results!", + 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_host_result.mData, - d1_g_m_device_result.mData, - "Error: Incorrect results!", + 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); } From 190899be77829852d3250b8a40ef89b1b9ee831b Mon Sep 17 00:00:00 2001 From: Anthony Chang Date: Fri, 13 May 2022 08:26:31 +0800 Subject: [PATCH 4/5] amend --- .../conv2d_fwd_xdl_bias_relu_add.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 e88c114c904..53d882778a2 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,7 @@ int main(int argc, char* argv[]) ref_invoker.Run(ref_argument); out_device_buf.FromDevice(device_output.mData.data()); - return ck::utils::check_err(device_output.mData, host_output.mData); + return ck::utils::check_err(device_output.mData, host_output.mData) ? 0 : 1; } return 0; From daed5eff7d19a653dfe5a354b87f031ca155e170 Mon Sep 17 00:00:00 2001 From: Anthony Chang Date: Fri, 13 May 2022 14:44:20 +0800 Subject: [PATCH 5/5] add example codes to custom target 'check' --- CMakeLists.txt | 10 ++++++---- example/CMakeLists.txt | 1 + test/CMakeLists.txt | 1 - 3 files changed, 7 insertions(+), 5 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index f18c85c6839..a3ec91e3bcb 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/CMakeLists.txt b/example/CMakeLists.txt index ff07c55aec2..4d81e84c01c 100644 --- a/example/CMakeLists.txt +++ b/example/CMakeLists.txt @@ -25,6 +25,7 @@ function(add_example_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) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c696069393b..37335635712 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)