From 63d58bf9e99a1f8f191ec0d7d455f068fa6456a2 Mon Sep 17 00:00:00 2001 From: richagadgil Date: Wed, 11 Dec 2024 19:25:40 -0600 Subject: [PATCH 1/3] bf16 verify tests --- test/verify/CMakeLists.txt | 2 +- test/verify/test_abs.cpp | 1 + test/verify/test_acos.cpp | 1 + test/verify/test_acosh.cpp | 1 + test/verify/test_add.cpp | 1 + test/verify/test_add_bf16.cpp | 42 ++++++++++ test/verify/test_add_dot.cpp | 1 + test/verify/test_add_gelu_bf16.cpp | 56 +++++++++++++ test/verify/test_add_mixed_layout.cpp | 1 + test/verify/test_arg_ops.cpp | 2 - test/verify/test_asin.cpp | 1 + test/verify/test_asinh.cpp | 1 + test/verify/test_atan.cpp | 1 + test/verify/test_atanh.cpp | 1 + test/verify/test_block_reduce_small.cpp | 3 +- test/verify/test_ceil.cpp | 1 + .../verify/test_ck_gemm_softmax_gemm_bf16.cpp | 57 +++++++++++++ test/verify/test_concat_axis_0.cpp | 1 + test/verify/test_concat_nhwc.cpp | 1 + test/verify/test_conv.cpp | 1 + test/verify/test_conv2.cpp | 1 + test/verify/test_conv_add.cpp | 1 + .../verify/test_conv_add_1x1_diff_strides.cpp | 1 + test/verify/test_conv_add_layernorm_conv.cpp | 1 + test/verify/test_conv_add_relu.cpp | 1 + test/verify/test_conv_add_tune.cpp | 1 + test/verify/test_conv_bias_clipped_relu.cpp | 1 + test/verify/test_conv_bn.cpp | 1 + test/verify/test_conv_bn_add.cpp | 1 + test/verify/test_conv_bn_relu_pooling.cpp | 1 + test/verify/test_conv_bn_relu_pooling2.cpp | 1 + test/verify/test_conv_group_add.cpp | 1 + test/verify/test_conv_pooling.cpp | 1 + test/verify/test_conv_relu.cpp | 1 + test/verify/test_convert.cpp | 2 + test/verify/test_cos.cpp | 1 + test/verify/test_cosh.cpp | 1 + test/verify/test_erf.cpp | 1 + test/verify/test_exp.cpp | 1 + test/verify/test_floor.cpp | 1 + test/verify/test_fmod_mod.cpp | 2 + test/verify/test_fp32_bf16_add.cpp | 48 +++++++++++ test/verify/test_fp32_bf16_ladd.cpp | 47 +++++++++++ test/verify/test_fp32_bf16_lall.cpp | 47 +++++++++++ test/verify/test_fp32_bf16_sub.cpp | 48 +++++++++++ test/verify/test_gather.cpp | 2 + test/verify/test_gathernd_default.cpp | 1 + test/verify/test_gemm.cpp | 3 +- test/verify/test_gemm_2args_bmv.cpp | 3 +- test/verify/test_gemm_2args_mm_1.cpp | 3 +- test/verify/test_gemm_2args_mm_2.cpp | 3 +- test/verify/test_gemm_2args_mm_3.cpp | 3 +- test/verify/test_gemm_2args_mm_4.cpp | 1 + test/verify/test_gemm_2args_mm_5.cpp | 3 +- test/verify/test_gemm_2args_mm_6.cpp | 3 +- test/verify/test_gemm_2args_mm_7.cpp | 3 +- test/verify/test_gemm_2args_mm_8.cpp | 2 +- test/verify/test_gemm_2args_mv.cpp | 3 +- test/verify/test_gemm_2args_vbm.cpp | 3 +- test/verify/test_gemm_2args_vm.cpp | 3 +- test/verify/test_gemm_2args_vv.cpp | 3 +- test/verify/test_gemm_add.cpp | 1 + test/verify/test_gemm_add_broadcast1.cpp | 3 +- test/verify/test_gemm_add_broadcast2.cpp | 2 +- test/verify/test_gemm_copy.cpp | 3 +- test/verify/test_gemm_ex.cpp | 3 +- test/verify/test_gemm_literal.cpp | 3 +- .../test_gemm_mul_where_softmax_gemm_bf16.cpp | 56 +++++++++++++ test/verify/test_gemm_multi_3args.cpp | 3 +- test/verify/test_gemm_multi_3args_alpha0.cpp | 3 +- test/verify/test_gemm_multi_3args_beta0.cpp | 3 +- test/verify/test_gemm_multi_3args_c25.cpp | 3 +- test/verify/test_gemm_multi_dim_2.cpp | 1 + test/verify/test_gemm_multi_dim_2_3.cpp | 3 +- test/verify/test_gemm_multi_transpose.cpp | 3 +- test/verify/test_gemm_multibroadcast.cpp | 3 +- test/verify/test_gemm_pointwise.cpp | 3 +- test/verify/test_gemm_reshapes_add.cpp | 1 + .../test_gemm_softmax_gemm_relu_bf16.cpp | 80 +++++++++++++++++++ .../test_gemm_transpose_add_pooling_sub.cpp | 3 +- test/verify/test_gemm_transposea.cpp | 3 +- test/verify/test_gemm_transposea_ex.cpp | 3 +- test/verify/test_gemm_transposeab.cpp | 3 +- test/verify/test_gemm_transposeb.cpp | 3 +- test/verify/test_gemm_transposeb_detect.cpp | 3 +- test/verify/test_gemm_transposeb_ex.cpp | 3 +- test/verify/test_hsqrt_bf16.cpp | 42 ++++++++++ test/verify/test_instancenorm.cpp | 7 +- test/verify/test_isinf.cpp | 1 + test/verify/test_isnan.cpp | 1 + test/verify/test_layernorm.cpp | 35 +++----- test/verify/test_literal_limits.cpp | 1 + test/verify/test_log.cpp | 1 + test/verify/test_log2.cpp | 1 + test/verify/test_logsoftmax.cpp | 7 +- test/verify/test_logsoftmax1.cpp | 4 +- test/verify/test_min_max.cpp | 2 + test/verify/test_mul_dot_a.cpp | 3 +- test/verify/test_mul_dot_b.cpp | 3 +- test/verify/test_multinomial.cpp | 1 + test/verify/test_nearbyint.cpp | 1 + test/verify/test_nonzero.cpp | 1 + test/verify/test_pad.cpp | 1 + test/verify/test_pad_highest_bf16.cpp | 46 +++++++++++ test/verify/test_pad_lowest_bf16.cpp | 46 +++++++++++ .../test_pointwise_broadcast_reduce.cpp | 2 - .../test_pointwise_broadcast_reduce_bf16.cpp | 55 +++++++++++++ test/verify/test_pow.cpp | 1 + test/verify/test_prefix_scan_sum_2d.cpp | 2 + test/verify/test_reduce_add.cpp | 2 - test/verify/test_reduce_add_bf16.cpp | 52 ++++++++++++ test/verify/test_reduce_mean_bias_bf16.cpp | 48 +++++++++++ test/verify/test_reduce_mean_bias_half.cpp | 4 +- test/verify/test_reduce_mean_large_bf16.cpp | 41 ++++++++++ test/verify/test_reduce_mean_large_half.cpp | 4 +- test/verify/test_reduce_mean_nhwc.cpp | 3 +- test/verify/test_reduce_mean_reduce_sum.cpp | 2 - test/verify/test_reduce_mean_variance.cpp | 2 - .../verify/test_reduce_mean_variance_bf16.cpp | 48 +++++++++++ test/verify/test_reduce_noop_add.cpp | 4 +- test/verify/test_reduce_noop_add_bf16.cpp | 48 +++++++++++ test/verify/test_reduce_op_large.cpp | 2 - test/verify/test_reduce_op_small.cpp | 12 ++- test/verify/test_reverse.cpp | 1 + test/verify/test_rnn_sql_1.cpp | 3 +- test/verify/test_roialign.cpp | 1 + test/verify/test_rsqrt.cpp | 1 + .../test_scatter_elements_none_axis_neg_1.cpp | 1 + test/verify/test_scatternd.cpp | 1 + test/verify/test_select_module_reduce.cpp | 4 +- test/verify/test_shrink.cpp | 1 + test/verify/test_sin.cpp | 1 + test/verify/test_sin_bf16.cpp | 41 ++++++++++ test/verify/test_sinh.cpp | 1 + test/verify/test_softmax.cpp | 7 +- test/verify/test_softmax1.cpp | 4 +- test/verify/test_softmax2.cpp | 4 +- test/verify/test_softmax3.cpp | 4 +- test/verify/test_softmax4.cpp | 4 +- test/verify/test_softmax5.cpp | 41 ++++++++++ test/verify/test_softmax_large1.cpp | 4 +- test/verify/test_softmax_large2.cpp | 4 +- test/verify/test_softmax_large3.cpp | 4 +- .../test_softmaxcrossentropyloss_2d.cpp | 11 ++- .../test_softmaxcrossentropyloss_2d_mean.cpp | 10 ++- .../test_softmaxcrossentropyloss_2d_sum.cpp | 10 ++- .../test_softmaxcrossentropyloss_kd.cpp | 10 ++- .../test_softmaxcrossentropyloss_kd_mean.cpp | 11 ++- .../test_softmaxcrossentropyloss_kd_sum.cpp | 10 ++- test/verify/test_sqrt.cpp | 1 + test/verify/test_sqrt_bf161.cpp | 43 ++++++++++ test/verify/test_sqrt_bf162.cpp | 44 ++++++++++ test/verify/test_sqrt_bf164.cpp | 43 ++++++++++ test/verify/test_tan.cpp | 1 + test/verify/test_tanh.cpp | 1 + test/verify/test_topk_0.cpp | 1 + test/verify/test_trans_convert_gemm.cpp | 1 + .../test_transpose_reshape_add_sub_mul.cpp | 1 + test/verify/test_unbatched_gemm_1.cpp | 1 + test/verify/test_unbatched_gemm_2.cpp | 1 + test/verify/test_where.cpp | 1 + 161 files changed, 1363 insertions(+), 135 deletions(-) create mode 100644 test/verify/test_add_bf16.cpp create mode 100644 test/verify/test_add_gelu_bf16.cpp create mode 100644 test/verify/test_ck_gemm_softmax_gemm_bf16.cpp create mode 100644 test/verify/test_fp32_bf16_add.cpp create mode 100644 test/verify/test_fp32_bf16_ladd.cpp create mode 100644 test/verify/test_fp32_bf16_lall.cpp create mode 100644 test/verify/test_fp32_bf16_sub.cpp create mode 100644 test/verify/test_gemm_mul_where_softmax_gemm_bf16.cpp create mode 100644 test/verify/test_gemm_softmax_gemm_relu_bf16.cpp create mode 100644 test/verify/test_hsqrt_bf16.cpp create mode 100644 test/verify/test_pad_highest_bf16.cpp create mode 100644 test/verify/test_pad_lowest_bf16.cpp create mode 100644 test/verify/test_pointwise_broadcast_reduce_bf16.cpp create mode 100644 test/verify/test_reduce_add_bf16.cpp create mode 100644 test/verify/test_reduce_mean_bias_bf16.cpp create mode 100644 test/verify/test_reduce_mean_large_bf16.cpp create mode 100644 test/verify/test_reduce_mean_variance_bf16.cpp create mode 100644 test/verify/test_reduce_noop_add_bf16.cpp create mode 100644 test/verify/test_sin_bf16.cpp create mode 100644 test/verify/test_softmax5.cpp create mode 100644 test/verify/test_sqrt_bf161.cpp create mode 100644 test/verify/test_sqrt_bf162.cpp create mode 100644 test/verify/test_sqrt_bf164.cpp diff --git a/test/verify/CMakeLists.txt b/test/verify/CMakeLists.txt index bd57abea883..3e45e9bc3a9 100644 --- a/test/verify/CMakeLists.txt +++ b/test/verify/CMakeLists.txt @@ -31,7 +31,7 @@ target_link_libraries(test_verify migraphx migraphx_all_targets) target_include_directories(test_verify PUBLIC ../include) rocm_clang_tidy_check(test_verify) -foreach(SECTION general reduce rnn conv gemm) +foreach(SECTION general rnn conv gemm) rocm_add_test(NAME test_verify_${SECTION} COMMAND test_verify ${SECTION}) set_tests_properties(test_verify_${SECTION} PROPERTIES COST 100 diff --git a/test/verify/test_abs.cpp b/test/verify/test_abs.cpp index 0c7d7ef050c..a02c84967bb 100644 --- a/test/verify/test_abs.cpp +++ b/test/verify/test_abs.cpp @@ -41,6 +41,7 @@ struct test_abs : verify_program> }; template struct test_abs; +template struct test_abs; template struct test_abs; template struct test_abs; template struct test_abs; diff --git a/test/verify/test_acos.cpp b/test/verify/test_acos.cpp index c22357afa77..34b3391fafa 100644 --- a/test/verify/test_acos.cpp +++ b/test/verify/test_acos.cpp @@ -42,6 +42,7 @@ struct test_acos : verify_program> }; template struct test_acos; +template struct test_acos; template struct test_acos; template struct test_acos; template struct test_acos; diff --git a/test/verify/test_acosh.cpp b/test/verify/test_acosh.cpp index 51f10fc8556..4fa6ce6edab 100644 --- a/test/verify/test_acosh.cpp +++ b/test/verify/test_acosh.cpp @@ -52,6 +52,7 @@ struct test_acosh : verify_program> template struct test_acosh; template struct test_acosh; +template struct test_acosh; template struct test_acosh; template struct test_acosh; template struct test_acosh; diff --git a/test/verify/test_add.cpp b/test/verify/test_add.cpp index 54dbedf23ad..9ed820863d4 100644 --- a/test/verify/test_add.cpp +++ b/test/verify/test_add.cpp @@ -43,6 +43,7 @@ struct test_add : verify_program> }; template struct test_add; +template struct test_add; template struct test_add; template struct test_add; template struct test_add; diff --git a/test/verify/test_add_bf16.cpp b/test/verify/test_add_bf16.cpp new file mode 100644 index 00000000000..2a1722e650d --- /dev/null +++ b/test/verify/test_add_bf16.cpp @@ -0,0 +1,42 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_add_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {3}}; + auto x = mm->add_parameter("x", s); + auto y = mm->add_parameter("y", s); + mm->add_instruction(migraphx::make_op("add"), x, y); + return p; + } +}; diff --git a/test/verify/test_add_dot.cpp b/test/verify/test_add_dot.cpp index ad23cc5acf6..cebddb8c380 100644 --- a/test/verify/test_add_dot.cpp +++ b/test/verify/test_add_dot.cpp @@ -46,4 +46,5 @@ struct test_add_dot : verify_program> }; template struct test_add_dot; +template struct test_add_dot; template struct test_add_dot; diff --git a/test/verify/test_add_gelu_bf16.cpp b/test/verify/test_add_gelu_bf16.cpp new file mode 100644 index 00000000000..8e625568565 --- /dev/null +++ b/test/verify/test_add_gelu_bf16.cpp @@ -0,0 +1,56 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_add_gelu_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + std::vector input_lens{1, 1, 5}; + auto x = mm->add_parameter("x", {migraphx::shape::bf16_type, input_lens}); + auto y = mm->add_parameter("y", {migraphx::shape::bf16_type, input_lens}); + auto bf16 = mm->add_literal(migraphx::literal{{migraphx::shape::bf16_type}, {0.5f}}); + auto one = mm->add_literal(migraphx::literal{{migraphx::shape::bf16_type}, {1.0f}}); + auto sqrt2 = mm->add_literal(migraphx::literal{{migraphx::shape::bf16_type}, {M_SQRT2}}); + auto add = mm->add_instruction(migraphx::make_op("add"), x, y); + auto bf16_mbcast = mm->add_instruction( + migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), bf16); + auto mul_bf16 = mm->add_instruction(migraphx::make_op("mul"), add, bf16_mbcast); + auto sqrt2_mbcast = mm->add_instruction( + migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), sqrt2); + auto div = mm->add_instruction(migraphx::make_op("div"), add, sqrt2_mbcast); + auto erf = mm->add_instruction(migraphx::make_op("erf"), div); + auto one_mbcast = mm->add_instruction( + migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), one); + auto add_one = mm->add_instruction(migraphx::make_op("add"), erf, one_mbcast); + mm->add_instruction(migraphx::make_op("mul"), mul_bf16, add_one); + return p; + } +}; diff --git a/test/verify/test_add_mixed_layout.cpp b/test/verify/test_add_mixed_layout.cpp index c4e94feda45..c6e88ddfb03 100644 --- a/test/verify/test_add_mixed_layout.cpp +++ b/test/verify/test_add_mixed_layout.cpp @@ -44,4 +44,5 @@ struct test_add_mixed_layout : verify_program> }; template struct test_add_mixed_layout; +template struct test_add_mixed_layout; template struct test_add_mixed_layout; diff --git a/test/verify/test_arg_ops.cpp b/test/verify/test_arg_ops.cpp index 674e8cb8969..ac52da56a23 100644 --- a/test/verify/test_arg_ops.cpp +++ b/test/verify/test_arg_ops.cpp @@ -57,8 +57,6 @@ struct test_arg_ops : verify_programadd_instruction(T{Axis, LastIndex}, param); return p; } - - std::string section() const { return "reduce"; } }; // transpose argmax tests template struct test_arg_ops; diff --git a/test/verify/test_asin.cpp b/test/verify/test_asin.cpp index 34951182059..f41cd8d9bc1 100644 --- a/test/verify/test_asin.cpp +++ b/test/verify/test_asin.cpp @@ -42,6 +42,7 @@ struct test_asin : verify_program> }; template struct test_asin; +template struct test_asin; template struct test_asin; template struct test_asin; template struct test_asin; diff --git a/test/verify/test_asinh.cpp b/test/verify/test_asinh.cpp index 4bc6680dd2d..642a90f2186 100644 --- a/test/verify/test_asinh.cpp +++ b/test/verify/test_asinh.cpp @@ -42,6 +42,7 @@ struct test_asinh : verify_program> }; template struct test_asinh; +template struct test_asinh; template struct test_asinh; template struct test_asinh; template struct test_asinh; diff --git a/test/verify/test_atan.cpp b/test/verify/test_atan.cpp index 63f8167d000..271d3bb2335 100644 --- a/test/verify/test_atan.cpp +++ b/test/verify/test_atan.cpp @@ -43,6 +43,7 @@ struct test_atan : verify_program> template struct test_atan; template struct test_atan; +template struct test_atan; template struct test_atan; template struct test_atan; template struct test_atan; diff --git a/test/verify/test_atanh.cpp b/test/verify/test_atanh.cpp index 8175f4f001a..dcc09ffdd37 100644 --- a/test/verify/test_atanh.cpp +++ b/test/verify/test_atanh.cpp @@ -54,6 +54,7 @@ struct test_atanh : verify_program> template struct test_atanh; template struct test_atanh; +template struct test_atanh; template struct test_atanh; template struct test_atanh; template struct test_atanh; diff --git a/test/verify/test_block_reduce_small.cpp b/test/verify/test_block_reduce_small.cpp index c97e5b7002d..2345591b682 100644 --- a/test/verify/test_block_reduce_small.cpp +++ b/test/verify/test_block_reduce_small.cpp @@ -46,12 +46,11 @@ struct test_block_reduce_small : verify_program> mm->add_return({add}); return p; }; - - std::string section() const { return "reduce"; } }; template struct test_block_reduce_small_for_size : test_block_reduce_small, + test_block_reduce_small, test_block_reduce_small, test_block_reduce_small { diff --git a/test/verify/test_ceil.cpp b/test/verify/test_ceil.cpp index d6640165542..32847b674ac 100644 --- a/test/verify/test_ceil.cpp +++ b/test/verify/test_ceil.cpp @@ -44,6 +44,7 @@ struct test_ceil : verify_program> template struct test_ceil; template struct test_ceil; +template struct test_ceil; template struct test_ceil; template struct test_ceil; template struct test_ceil; diff --git a/test/verify/test_ck_gemm_softmax_gemm_bf16.cpp b/test/verify/test_ck_gemm_softmax_gemm_bf16.cpp new file mode 100644 index 00000000000..e439443f676 --- /dev/null +++ b/test/verify/test_ck_gemm_softmax_gemm_bf16.cpp @@ -0,0 +1,57 @@ +// /* +// * The MIT License (MIT) +// * +// * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. +// * +// * Permission is hereby granted, free of charge, to any person obtaining a copy +// * of this software and associated documentation files (the "Software"), to deal +// * in the Software without restriction, including without limitation the rights +// * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// * copies of the Software, and to permit persons to whom the Software is +// * furnished to do so, subject to the following conditions: +// * +// * The above copyright notice and this permission notice shall be included in +// * all copies or substantial portions of the Software. +// * +// * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// * THE SOFTWARE. +// */ + +// #include "verify_program.hpp" +// #include +// #include +// #include + +// struct test_ck_gemm_softmax_gemm_bf16 : verify_program +// { +// migraphx::program create_program() const +// { +// migraphx::program p; +// auto* mm = p.get_main_module(); +// migraphx::shape m1_shape{migraphx::shape::bf16_type, {1, 12, 256, 256}}; +// migraphx::shape m2_shape{migraphx::shape::bf16_type, {1, 12, 256, 256}}; +// auto m2_elements = m2_shape.elements(); +// auto a = mm->add_parameter("1", m1_shape); +// auto b = mm->add_parameter("2", m1_shape); +// auto b1 = mm->add_parameter("3", m1_shape); +// std::vector eights(m2_elements, 0.125); +// auto eight = mm->add_literal(migraphx::literal{m2_shape, eights}); +// std::vector zeros(m2_elements, 0); +// auto zero = mm->add_literal(migraphx::literal{m2_shape, zeros}); + +// b = mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), b); +// auto gemm1 = mm->add_instruction(migraphx::make_op("dot"), a, b); +// auto scale = mm->add_instruction(migraphx::make_op("mul"), gemm1, eight); +// auto bias = mm->add_instruction(migraphx::make_op("add"), scale, zero); +// auto softmax = mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), bias); +// mm->add_instruction(migraphx::make_op("dot"), softmax, b1); + +// return p; +// } +// std::string section() const { return "gemm"; } +// }; diff --git a/test/verify/test_concat_axis_0.cpp b/test/verify/test_concat_axis_0.cpp index 4d4801bf038..aac455518b4 100644 --- a/test/verify/test_concat_axis_0.cpp +++ b/test/verify/test_concat_axis_0.cpp @@ -47,6 +47,7 @@ struct test_concat_axis_0 : verify_program> }; template struct test_concat_axis_0; +template struct test_concat_axis_0; template struct test_concat_axis_0; template struct test_concat_axis_0; template struct test_concat_axis_0; diff --git a/test/verify/test_concat_nhwc.cpp b/test/verify/test_concat_nhwc.cpp index 8006386fdba..21b1eb69b01 100644 --- a/test/verify/test_concat_nhwc.cpp +++ b/test/verify/test_concat_nhwc.cpp @@ -44,5 +44,6 @@ struct test_concat_nhwc : verify_program> template struct test_concat_nhwc; template struct test_concat_nhwc; +template struct test_concat_nhwc; template struct test_concat_nhwc; template struct test_concat_nhwc; diff --git a/test/verify/test_conv.cpp b/test/verify/test_conv.cpp index d75edf4e830..c76db1e194f 100644 --- a/test/verify/test_conv.cpp +++ b/test/verify/test_conv.cpp @@ -43,6 +43,7 @@ struct test_conv : verify_program> }; template struct test_conv; +template struct test_conv; template struct test_conv; template struct test_conv; template struct test_conv; diff --git a/test/verify/test_conv2.cpp b/test/verify/test_conv2.cpp index eb14a550fab..928a4c7da12 100644 --- a/test/verify/test_conv2.cpp +++ b/test/verify/test_conv2.cpp @@ -46,6 +46,7 @@ struct test_conv2 : verify_program> std::string section() const { return "conv"; } }; template struct test_conv2; +template struct test_conv2; template struct test_conv2; template struct test_conv2; template struct test_conv2; diff --git a/test/verify/test_conv_add.cpp b/test/verify/test_conv_add.cpp index 2c8739cdc5c..dc3f344055e 100644 --- a/test/verify/test_conv_add.cpp +++ b/test/verify/test_conv_add.cpp @@ -48,6 +48,7 @@ struct test_conv_add : verify_program> }; template struct test_conv_add; +template struct test_conv_add; template struct test_conv_add; template struct test_conv_add; template struct test_conv_add; diff --git a/test/verify/test_conv_add_1x1_diff_strides.cpp b/test/verify/test_conv_add_1x1_diff_strides.cpp index 7faf9840451..cb429c13e85 100644 --- a/test/verify/test_conv_add_1x1_diff_strides.cpp +++ b/test/verify/test_conv_add_1x1_diff_strides.cpp @@ -54,6 +54,7 @@ struct test_conv_add_1x1_diff_strides : verify_program; +template struct test_conv_add_1x1_diff_strides; template struct test_conv_add_1x1_diff_strides; template struct test_conv_add_1x1_diff_strides; template struct test_conv_add_1x1_diff_strides; diff --git a/test/verify/test_conv_add_layernorm_conv.cpp b/test/verify/test_conv_add_layernorm_conv.cpp index 989073aacb0..70b4f7c6296 100644 --- a/test/verify/test_conv_add_layernorm_conv.cpp +++ b/test/verify/test_conv_add_layernorm_conv.cpp @@ -70,4 +70,5 @@ struct test_conv_add_layernorm_conv : verify_program; template struct test_conv_add_layernorm_conv; +template struct test_conv_add_layernorm_conv; template struct test_conv_add_layernorm_conv; diff --git a/test/verify/test_conv_add_relu.cpp b/test/verify/test_conv_add_relu.cpp index 9839ac509cc..048956793c8 100644 --- a/test/verify/test_conv_add_relu.cpp +++ b/test/verify/test_conv_add_relu.cpp @@ -52,6 +52,7 @@ struct test_conv_add_relu : verify_program> }; template struct test_conv_add_relu; +template struct test_conv_add_relu; template struct test_conv_add_relu; template struct test_conv_add_relu; template struct test_conv_add_relu; diff --git a/test/verify/test_conv_add_tune.cpp b/test/verify/test_conv_add_tune.cpp index 1a2a0efa326..732fdf8a6df 100644 --- a/test/verify/test_conv_add_tune.cpp +++ b/test/verify/test_conv_add_tune.cpp @@ -73,6 +73,7 @@ struct test_conv_add_tune : verify_program> template struct test_conv_add_tune; template struct test_conv_add_tune; +template struct test_conv_add_tune; template struct test_conv_add_tune; template struct test_conv_add_tune; template struct test_conv_add_tune; diff --git a/test/verify/test_conv_bias_clipped_relu.cpp b/test/verify/test_conv_bias_clipped_relu.cpp index fa42b428fef..a8891b276f9 100644 --- a/test/verify/test_conv_bias_clipped_relu.cpp +++ b/test/verify/test_conv_bias_clipped_relu.cpp @@ -58,6 +58,7 @@ struct test_conv_bias_clipped_relu : verify_program; +template struct test_conv_bias_clipped_relu; template struct test_conv_bias_clipped_relu; template struct test_conv_bias_clipped_relu; template struct test_conv_bias_clipped_relu; diff --git a/test/verify/test_conv_bn.cpp b/test/verify/test_conv_bn.cpp index 99c1b6f1c0c..c3b82aa07ff 100644 --- a/test/verify/test_conv_bn.cpp +++ b/test/verify/test_conv_bn.cpp @@ -86,6 +86,7 @@ struct test_conv_bn : verify_program> }; template struct test_conv_bn; +template struct test_conv_bn; template struct test_conv_bn; template struct test_conv_bn; template struct test_conv_bn; diff --git a/test/verify/test_conv_bn_add.cpp b/test/verify/test_conv_bn_add.cpp index e8fd3a77de2..52ceeee69ba 100644 --- a/test/verify/test_conv_bn_add.cpp +++ b/test/verify/test_conv_bn_add.cpp @@ -93,6 +93,7 @@ struct test_conv_bn_add : verify_program> }; template struct test_conv_bn_add; +template struct test_conv_bn_add; template struct test_conv_bn_add; template struct test_conv_bn_add; template struct test_conv_bn_add; diff --git a/test/verify/test_conv_bn_relu_pooling.cpp b/test/verify/test_conv_bn_relu_pooling.cpp index 5fe0d8fa4a1..22e70be69d0 100644 --- a/test/verify/test_conv_bn_relu_pooling.cpp +++ b/test/verify/test_conv_bn_relu_pooling.cpp @@ -92,6 +92,7 @@ struct test_conv_bn_relu_pooling : verify_program; +template struct test_conv_bn_relu_pooling; template struct test_conv_bn_relu_pooling; template struct test_conv_bn_relu_pooling; template struct test_conv_bn_relu_pooling; diff --git a/test/verify/test_conv_bn_relu_pooling2.cpp b/test/verify/test_conv_bn_relu_pooling2.cpp index d003e146cd0..25b4b20773e 100644 --- a/test/verify/test_conv_bn_relu_pooling2.cpp +++ b/test/verify/test_conv_bn_relu_pooling2.cpp @@ -108,6 +108,7 @@ struct test_conv_bn_relu_pooling2 : verify_program; +template struct test_conv_bn_relu_pooling2; template struct test_conv_bn_relu_pooling2; template struct test_conv_bn_relu_pooling2; template struct test_conv_bn_relu_pooling2; diff --git a/test/verify/test_conv_group_add.cpp b/test/verify/test_conv_group_add.cpp index e9171b68b47..d17e383512a 100644 --- a/test/verify/test_conv_group_add.cpp +++ b/test/verify/test_conv_group_add.cpp @@ -47,5 +47,6 @@ struct test_conv_group_add : verify_program> std::string section() const { return "conv"; } }; template struct test_conv_group_add; +template struct test_conv_group_add; // TODO grouped convolutions are not supported with MLIR therefore disable it // template struct test_conv_group_add; diff --git a/test/verify/test_conv_pooling.cpp b/test/verify/test_conv_pooling.cpp index 5194fb4977e..138208448d9 100644 --- a/test/verify/test_conv_pooling.cpp +++ b/test/verify/test_conv_pooling.cpp @@ -47,6 +47,7 @@ struct test_conv_pooling : verify_program> }; template struct test_conv_pooling; +template struct test_conv_pooling; template struct test_conv_pooling; template struct test_conv_pooling; template struct test_conv_pooling; diff --git a/test/verify/test_conv_relu.cpp b/test/verify/test_conv_relu.cpp index 10973c0f7e0..81f115d0f26 100644 --- a/test/verify/test_conv_relu.cpp +++ b/test/verify/test_conv_relu.cpp @@ -44,6 +44,7 @@ struct test_conv_relu : verify_program> }; template struct test_conv_relu; template struct test_conv_relu; +template struct test_conv_relu; template struct test_conv_relu; template struct test_conv_relu; template struct test_conv_relu; diff --git a/test/verify/test_convert.cpp b/test/verify/test_convert.cpp index 92bcb58dba6..af5f6980325 100644 --- a/test/verify/test_convert.cpp +++ b/test/verify/test_convert.cpp @@ -56,3 +56,5 @@ template struct test_convert; template struct test_convert; template struct test_convert; +template struct test_convert; +template struct test_convert; diff --git a/test/verify/test_cos.cpp b/test/verify/test_cos.cpp index 2482f0dad14..6874e8b0d25 100644 --- a/test/verify/test_cos.cpp +++ b/test/verify/test_cos.cpp @@ -43,6 +43,7 @@ struct test_cos : verify_program> template struct test_cos; template struct test_cos; +template struct test_cos; template struct test_cos; template struct test_cos; template struct test_cos; diff --git a/test/verify/test_cosh.cpp b/test/verify/test_cosh.cpp index 1e8dd150a59..a07f5fad6b6 100644 --- a/test/verify/test_cosh.cpp +++ b/test/verify/test_cosh.cpp @@ -43,6 +43,7 @@ struct test_cosh : verify_program> template struct test_cosh; template struct test_cosh; +template struct test_cosh; template struct test_cosh; template struct test_cosh; template struct test_cosh; diff --git a/test/verify/test_erf.cpp b/test/verify/test_erf.cpp index 1edcc199b9b..e8b69aa63ba 100644 --- a/test/verify/test_erf.cpp +++ b/test/verify/test_erf.cpp @@ -43,6 +43,7 @@ struct test_erf : verify_program> template struct test_erf; template struct test_erf; +template struct test_erf; template struct test_erf; template struct test_erf; template struct test_erf; diff --git a/test/verify/test_exp.cpp b/test/verify/test_exp.cpp index c6437603e17..1333db5fda4 100644 --- a/test/verify/test_exp.cpp +++ b/test/verify/test_exp.cpp @@ -43,6 +43,7 @@ struct test_exp : verify_program> template struct test_exp; template struct test_exp; +template struct test_exp; template struct test_exp; template struct test_exp; template struct test_exp; diff --git a/test/verify/test_floor.cpp b/test/verify/test_floor.cpp index 1f5c4baad97..1e7562c9c4c 100644 --- a/test/verify/test_floor.cpp +++ b/test/verify/test_floor.cpp @@ -44,6 +44,7 @@ struct test_floor : verify_program> template struct test_floor; template struct test_floor; +template struct test_floor; template struct test_floor; template struct test_floor; template struct test_floor; diff --git a/test/verify/test_fmod_mod.cpp b/test/verify/test_fmod_mod.cpp index e92cb15406c..281131f5df1 100644 --- a/test/verify/test_fmod_mod.cpp +++ b/test/verify/test_fmod_mod.cpp @@ -61,6 +61,7 @@ struct test_fmod : verify_program> }; template struct test_fmod; template struct test_fmod; +template struct test_fmod; template struct test_fmod; template struct test_fmod; template struct test_fmod; @@ -83,6 +84,7 @@ struct test_mod : verify_program> template struct test_mod; template struct test_mod; +template struct test_mod; template struct test_mod; template struct test_mod; template struct test_mod; diff --git a/test/verify/test_fp32_bf16_add.cpp b/test/verify/test_fp32_bf16_add.cpp new file mode 100644 index 00000000000..da0ba65599f --- /dev/null +++ b/test/verify/test_fp32_bf16_add.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +#include + +struct test_fp32_bf16_add : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::float_type, {2, 3}}; + auto p1 = mm->add_parameter("x", s); + auto p2 = mm->add_parameter("y", s); + auto sum = mm->add_instruction(migraphx::make_op("add"), p1, p2); + auto diff = mm->add_instruction(migraphx::make_op("sub"), sum, p2); + mm->add_instruction(migraphx::make_op("add"), diff, p1); + migraphx::quantize_bf16(p, {"add"}); + + return p; + }; +}; diff --git a/test/verify/test_fp32_bf16_ladd.cpp b/test/verify/test_fp32_bf16_ladd.cpp new file mode 100644 index 00000000000..b4881931979 --- /dev/null +++ b/test/verify/test_fp32_bf16_ladd.cpp @@ -0,0 +1,47 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +#include + +struct test_fp32_bf16_ladd : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::float_type, {2, 3}}; + std::vector data(2 * 3); + std::iota(data.begin(), data.end(), 1.0f); + auto l1 = mm->add_literal(migraphx::literal(s, data)); + auto l2 = mm->add_parameter("p2", s); + mm->add_instruction(migraphx::make_op("add"), l1, l2); + migraphx::quantize_bf16(p, {"add"}); + return p; + }; +}; diff --git a/test/verify/test_fp32_bf16_lall.cpp b/test/verify/test_fp32_bf16_lall.cpp new file mode 100644 index 00000000000..0bb01207ab4 --- /dev/null +++ b/test/verify/test_fp32_bf16_lall.cpp @@ -0,0 +1,47 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +#include + +struct test_fp32_bf16_lall : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::float_type, {2, 3}}; + std::vector data(2 * 3); + std::iota(data.begin(), data.end(), 1.0f); + auto l1 = mm->add_literal(migraphx::literal(s, data)); + auto l2 = mm->add_parameter("p2", s); + mm->add_instruction(migraphx::make_op("add"), l1, l2); + migraphx::quantize_bf16(p, {"all"}); + return p; + }; +}; diff --git a/test/verify/test_fp32_bf16_sub.cpp b/test/verify/test_fp32_bf16_sub.cpp new file mode 100644 index 00000000000..f771fe13cba --- /dev/null +++ b/test/verify/test_fp32_bf16_sub.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +#include + +struct test_fp32_bf16_sub : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::float_type, {2, 3}}; + auto p1 = mm->add_parameter("x", s); + auto p2 = mm->add_parameter("y", s); + auto sum = mm->add_instruction(migraphx::make_op("add"), p1, p2); + auto diff = mm->add_instruction(migraphx::make_op("sub"), sum, p2); + mm->add_instruction(migraphx::make_op("add"), diff, p1); + migraphx::quantize_bf16(p, {"sub"}); + + return p; + }; +}; diff --git a/test/verify/test_gather.cpp b/test/verify/test_gather.cpp index 19d8e6ae77c..2c4dd7d220a 100644 --- a/test/verify/test_gather.cpp +++ b/test/verify/test_gather.cpp @@ -48,6 +48,7 @@ struct test_gather : verify_program> // Standard gather test template struct test_gather<0, migraphx::shape::float_type>; template struct test_gather<0, migraphx::shape::half_type>; +template struct test_gather<0, migraphx::shape::bf16_type>; template struct test_gather<0, migraphx::shape::fp8e4m3fnuz_type>; template struct test_gather<0, migraphx::shape::fp8e5m2fnuz_type>; template struct test_gather<0, migraphx::shape::fp8e4m3fn_type>; @@ -55,6 +56,7 @@ template struct test_gather<0, migraphx::shape::fp8e5m2_type>; // Test Negative axis template struct test_gather<-2, migraphx::shape::float_type>; template struct test_gather<-2, migraphx::shape::half_type>; +template struct test_gather<-2, migraphx::shape::bf16_type>; template struct test_gather<-2, migraphx::shape::fp8e4m3fnuz_type>; template struct test_gather<-2, migraphx::shape::fp8e5m2fnuz_type>; template struct test_gather<-2, migraphx::shape::fp8e4m3fn_type>; diff --git a/test/verify/test_gathernd_default.cpp b/test/verify/test_gathernd_default.cpp index 5ff2ec40358..9c9e1fadd62 100644 --- a/test/verify/test_gathernd_default.cpp +++ b/test/verify/test_gathernd_default.cpp @@ -45,6 +45,7 @@ struct test_gathernd_default : verify_program> template struct test_gathernd_default; template struct test_gathernd_default; +template struct test_gathernd_default; template struct test_gathernd_default; template struct test_gathernd_default; template struct test_gathernd_default; diff --git a/test/verify/test_gemm.cpp b/test/verify/test_gemm.cpp index c6d93f29ac4..5952a82ea23 100644 --- a/test/verify/test_gemm.cpp +++ b/test/verify/test_gemm.cpp @@ -44,7 +44,8 @@ struct test_gemm : verify_program> template struct test_gemm; template struct test_gemm; +template struct test_gemm; template struct test_gemm; -template struct test_gemm; +// template struct test_gemm; template struct test_gemm; template struct test_gemm; diff --git a/test/verify/test_gemm_2args_bmv.cpp b/test/verify/test_gemm_2args_bmv.cpp index 782810fdf79..94c65c318d4 100644 --- a/test/verify/test_gemm_2args_bmv.cpp +++ b/test/verify/test_gemm_2args_bmv.cpp @@ -51,7 +51,8 @@ struct test_gemm_2args_bmv : verify_program> template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; +template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; -template struct test_gemm_2args_bmv; +// template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; diff --git a/test/verify/test_gemm_2args_mm_1.cpp b/test/verify/test_gemm_2args_mm_1.cpp index 7ddb70fc300..f029a21a4d3 100644 --- a/test/verify/test_gemm_2args_mm_1.cpp +++ b/test/verify/test_gemm_2args_mm_1.cpp @@ -49,8 +49,9 @@ struct test_gemm_2args_mm_1 : verify_program> }; template struct test_gemm_2args_mm_1; +template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; -template struct test_gemm_2args_mm_1; +// template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; diff --git a/test/verify/test_gemm_2args_mm_2.cpp b/test/verify/test_gemm_2args_mm_2.cpp index cd2d0df5d39..9efc07fca3d 100644 --- a/test/verify/test_gemm_2args_mm_2.cpp +++ b/test/verify/test_gemm_2args_mm_2.cpp @@ -51,7 +51,8 @@ struct test_gemm_2args_mm_2 : verify_program> template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; +template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; -template struct test_gemm_2args_mm_2; +// template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; diff --git a/test/verify/test_gemm_2args_mm_3.cpp b/test/verify/test_gemm_2args_mm_3.cpp index 00be7cee4bc..d63a17b8e40 100644 --- a/test/verify/test_gemm_2args_mm_3.cpp +++ b/test/verify/test_gemm_2args_mm_3.cpp @@ -51,7 +51,8 @@ struct test_gemm_2args_mm_3 : verify_program> template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; +template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; -template struct test_gemm_2args_mm_3; +// template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; diff --git a/test/verify/test_gemm_2args_mm_4.cpp b/test/verify/test_gemm_2args_mm_4.cpp index c98bcc1f9c0..476ba8a2daa 100644 --- a/test/verify/test_gemm_2args_mm_4.cpp +++ b/test/verify/test_gemm_2args_mm_4.cpp @@ -51,6 +51,7 @@ struct test_gemm_2args_mm_4 : verify_program> template struct test_gemm_2args_mm_4; template struct test_gemm_2args_mm_4; +template struct test_gemm_2args_mm_4; template struct test_gemm_2args_mm_4; template struct test_gemm_2args_mm_4; template struct test_gemm_2args_mm_4; diff --git a/test/verify/test_gemm_2args_mm_5.cpp b/test/verify/test_gemm_2args_mm_5.cpp index 45336a1c2f2..eda7518e239 100644 --- a/test/verify/test_gemm_2args_mm_5.cpp +++ b/test/verify/test_gemm_2args_mm_5.cpp @@ -50,7 +50,8 @@ struct test_gemm_2args_mm_5 : verify_program> template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; +template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; -template struct test_gemm_2args_mm_5; +// template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; diff --git a/test/verify/test_gemm_2args_mm_6.cpp b/test/verify/test_gemm_2args_mm_6.cpp index e2d00cb169b..100235feea9 100644 --- a/test/verify/test_gemm_2args_mm_6.cpp +++ b/test/verify/test_gemm_2args_mm_6.cpp @@ -53,7 +53,8 @@ struct test_gemm_2args_mm_6 : verify_program> template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; +template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; -template struct test_gemm_2args_mm_6; +// template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; diff --git a/test/verify/test_gemm_2args_mm_7.cpp b/test/verify/test_gemm_2args_mm_7.cpp index 7e03df20abc..7c78e7be1d5 100644 --- a/test/verify/test_gemm_2args_mm_7.cpp +++ b/test/verify/test_gemm_2args_mm_7.cpp @@ -50,7 +50,8 @@ struct test_gemm_2args_mm_7 : verify_program> template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; +template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; -template struct test_gemm_2args_mm_7; +// template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; diff --git a/test/verify/test_gemm_2args_mm_8.cpp b/test/verify/test_gemm_2args_mm_8.cpp index 175105d1003..bb0256e85ef 100644 --- a/test/verify/test_gemm_2args_mm_8.cpp +++ b/test/verify/test_gemm_2args_mm_8.cpp @@ -51,6 +51,6 @@ struct test_gemm_2args_mm_8 : verify_program> template struct test_gemm_2args_mm_8; // template struct test_gemm_2args_mm_8; // fails with CK, issue#2514 template struct test_gemm_2args_mm_8; -template struct test_gemm_2args_mm_8; +// template struct test_gemm_2args_mm_8; template struct test_gemm_2args_mm_8; template struct test_gemm_2args_mm_8; diff --git a/test/verify/test_gemm_2args_mv.cpp b/test/verify/test_gemm_2args_mv.cpp index 5a02d7cd765..070f9157870 100644 --- a/test/verify/test_gemm_2args_mv.cpp +++ b/test/verify/test_gemm_2args_mv.cpp @@ -49,7 +49,8 @@ struct test_gemm_2args_mv : verify_program> template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; +template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; -template struct test_gemm_2args_mv; +// template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; diff --git a/test/verify/test_gemm_2args_vbm.cpp b/test/verify/test_gemm_2args_vbm.cpp index b1c560ee331..20a49ef149d 100644 --- a/test/verify/test_gemm_2args_vbm.cpp +++ b/test/verify/test_gemm_2args_vbm.cpp @@ -53,7 +53,8 @@ struct test_gemm_2args_vbm : verify_program> template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; +template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; -template struct test_gemm_2args_vbm; +// template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; diff --git a/test/verify/test_gemm_2args_vm.cpp b/test/verify/test_gemm_2args_vm.cpp index 5c9b8283a24..76712f85403 100644 --- a/test/verify/test_gemm_2args_vm.cpp +++ b/test/verify/test_gemm_2args_vm.cpp @@ -50,8 +50,9 @@ struct test_gemm_2args_vm : verify_program> template struct test_gemm_2args_vm; template struct test_gemm_2args_vm; +template struct test_gemm_2args_vm; template struct test_gemm_2args_vm; -template struct test_gemm_2args_vm; +// template struct test_gemm_2args_vm; // TODO need hipblaslt support // template struct test_gemm_2args_vm; // template struct test_gemm_2args_vm; diff --git a/test/verify/test_gemm_2args_vv.cpp b/test/verify/test_gemm_2args_vv.cpp index 47f9d284df9..4550f801527 100644 --- a/test/verify/test_gemm_2args_vv.cpp +++ b/test/verify/test_gemm_2args_vv.cpp @@ -53,7 +53,8 @@ struct test_gemm_2args_vv : verify_program> template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; +template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; -template struct test_gemm_2args_vv; +// template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; diff --git a/test/verify/test_gemm_add.cpp b/test/verify/test_gemm_add.cpp index cd4231f6871..f52264c9e02 100644 --- a/test/verify/test_gemm_add.cpp +++ b/test/verify/test_gemm_add.cpp @@ -57,6 +57,7 @@ struct test_gemm_add : verify_program> template struct test_gemm_add; template struct test_gemm_add; +template struct test_gemm_add; // TODO template struct test_gemm_add; // TODO template struct test_gemm_add; // TODO template struct test_gemm_add; diff --git a/test/verify/test_gemm_add_broadcast1.cpp b/test/verify/test_gemm_add_broadcast1.cpp index 1106442aba7..a1ce9ac9e6f 100644 --- a/test/verify/test_gemm_add_broadcast1.cpp +++ b/test/verify/test_gemm_add_broadcast1.cpp @@ -53,7 +53,8 @@ struct test_gemm_add_broadcast1 : verify_program template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; +template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; -template struct test_gemm_add_broadcast1; +// template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; diff --git a/test/verify/test_gemm_add_broadcast2.cpp b/test/verify/test_gemm_add_broadcast2.cpp index 88674901d0f..7e086f423fd 100644 --- a/test/verify/test_gemm_add_broadcast2.cpp +++ b/test/verify/test_gemm_add_broadcast2.cpp @@ -55,6 +55,6 @@ template struct test_gemm_add_broadcast2; // template struct test_gemm_add_broadcast2; // fails with CK, // issue#2514 template struct test_gemm_add_broadcast2; -template struct test_gemm_add_broadcast2; +// template struct test_gemm_add_broadcast2; template struct test_gemm_add_broadcast2; template struct test_gemm_add_broadcast2; diff --git a/test/verify/test_gemm_copy.cpp b/test/verify/test_gemm_copy.cpp index c960c6c17da..f868cb2f1cb 100644 --- a/test/verify/test_gemm_copy.cpp +++ b/test/verify/test_gemm_copy.cpp @@ -51,7 +51,8 @@ struct test_gemm_copy : verify_program> template struct test_gemm_copy; template struct test_gemm_copy; +template struct test_gemm_copy; template struct test_gemm_copy; -template struct test_gemm_copy; +// template struct test_gemm_copy; template struct test_gemm_copy; template struct test_gemm_copy; diff --git a/test/verify/test_gemm_ex.cpp b/test/verify/test_gemm_ex.cpp index f269c80f702..86834de54a6 100644 --- a/test/verify/test_gemm_ex.cpp +++ b/test/verify/test_gemm_ex.cpp @@ -43,7 +43,8 @@ struct test_gemm_ex : verify_program> }; template struct test_gemm_ex; template struct test_gemm_ex; +template struct test_gemm_ex; template struct test_gemm_ex; -template struct test_gemm_ex; +// template struct test_gemm_ex; template struct test_gemm_ex; template struct test_gemm_ex; diff --git a/test/verify/test_gemm_literal.cpp b/test/verify/test_gemm_literal.cpp index 62d32588166..02532ce01df 100644 --- a/test/verify/test_gemm_literal.cpp +++ b/test/verify/test_gemm_literal.cpp @@ -48,7 +48,8 @@ struct test_gemm_literal : verify_program> template struct test_gemm_literal; template struct test_gemm_literal; +template struct test_gemm_literal; template struct test_gemm_literal; -template struct test_gemm_literal; +// template struct test_gemm_literal; template struct test_gemm_literal; template struct test_gemm_literal; diff --git a/test/verify/test_gemm_mul_where_softmax_gemm_bf16.cpp b/test/verify/test_gemm_mul_where_softmax_gemm_bf16.cpp new file mode 100644 index 00000000000..2f11b8465a6 --- /dev/null +++ b/test/verify/test_gemm_mul_where_softmax_gemm_bf16.cpp @@ -0,0 +1,56 @@ +// /* +// * The MIT License (MIT) +// * +// * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. +// * +// * Permission is hereby granted, free of charge, to any person obtaining a copy +// * of this software and associated documentation files (the "Software"), to deal +// * in the Software without restriction, including without limitation the rights +// * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// * copies of the Software, and to permit persons to whom the Software is +// * furnished to do so, subject to the following conditions: +// * +// * The above copyright notice and this permission notice shall be included in +// * all copies or substantial portions of the Software. +// * +// * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// * THE SOFTWARE. +// */ + +// #include "verify_program.hpp" +// #include +// #include +// #include + +// struct test_gemm_mul_where_softmax_gemm_bf16 : verify_program +// { +// migraphx::program create_program() const +// { +// migraphx::program p; +// auto* mm = p.get_main_module(); +// migraphx::shape m1_shape{migraphx::shape::bf16_type, {1, 12, 256, 256}}; +// migraphx::shape m2_shape{migraphx::shape::bool_type, {1, 12, 256, 256}}; +// auto m1_elements = m1_shape.elements(); +// auto a = mm->add_parameter("1", m1_shape); +// auto b = mm->add_parameter("2", m1_shape); +// auto b1 = mm->add_parameter("3", m1_shape); +// auto select = mm->add_parameter("4", m2_shape); +// std::vector eights(m1_elements, 0.125); +// std::vector tens(m1_elements, 10); +// auto eight = mm->add_literal(migraphx::literal{m1_shape, eights}); +// auto ten = mm->add_literal(migraphx::literal{m1_shape, tens}); +// b = mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), b); +// auto gemm1 = mm->add_instruction(migraphx::make_op("dot"), a, b); +// auto scale = mm->add_instruction(migraphx::make_op("mul"), gemm1, eight); +// auto where = mm->add_instruction(migraphx::make_op("where"), select, scale, ten); +// auto softmax = mm->add_instruction(migraphx::make_op("softmax", {{"axis", 3}}), where); +// mm->add_instruction(migraphx::make_op("dot"), softmax, b1); +// return p; +// } +// std::string section() const { return "gemm"; } +// }; diff --git a/test/verify/test_gemm_multi_3args.cpp b/test/verify/test_gemm_multi_3args.cpp index f61eee5ed03..99828d1f3f4 100644 --- a/test/verify/test_gemm_multi_3args.cpp +++ b/test/verify/test_gemm_multi_3args.cpp @@ -52,7 +52,8 @@ struct test_gemm_multi_3args : verify_program> template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; +template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; -template struct test_gemm_multi_3args; +// template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; diff --git a/test/verify/test_gemm_multi_3args_alpha0.cpp b/test/verify/test_gemm_multi_3args_alpha0.cpp index a66f55a93fc..69bb6799601 100644 --- a/test/verify/test_gemm_multi_3args_alpha0.cpp +++ b/test/verify/test_gemm_multi_3args_alpha0.cpp @@ -52,7 +52,8 @@ struct test_gemm_multi_3args_alpha0 : verify_program; template struct test_gemm_multi_3args_alpha0; +template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; -template struct test_gemm_multi_3args_alpha0; +// template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; diff --git a/test/verify/test_gemm_multi_3args_beta0.cpp b/test/verify/test_gemm_multi_3args_beta0.cpp index 6c18d54fdf5..3db08c23c76 100644 --- a/test/verify/test_gemm_multi_3args_beta0.cpp +++ b/test/verify/test_gemm_multi_3args_beta0.cpp @@ -52,7 +52,8 @@ struct test_gemm_multi_3args_beta0 : verify_program; template struct test_gemm_multi_3args_beta0; +template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; -template struct test_gemm_multi_3args_beta0; +// template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; diff --git a/test/verify/test_gemm_multi_3args_c25.cpp b/test/verify/test_gemm_multi_3args_c25.cpp index b5003bfccfd..8d2b2a28d72 100644 --- a/test/verify/test_gemm_multi_3args_c25.cpp +++ b/test/verify/test_gemm_multi_3args_c25.cpp @@ -52,7 +52,8 @@ struct test_gemm_multi_3args_c25 : verify_program; template struct test_gemm_multi_3args_c25; +template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; -template struct test_gemm_multi_3args_c25; +// template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; diff --git a/test/verify/test_gemm_multi_dim_2.cpp b/test/verify/test_gemm_multi_dim_2.cpp index 5ef599340e8..cc802e4dcbd 100644 --- a/test/verify/test_gemm_multi_dim_2.cpp +++ b/test/verify/test_gemm_multi_dim_2.cpp @@ -48,4 +48,5 @@ struct test_gemm_multi_dim_2 : verify_program> template struct test_gemm_multi_dim_2; template struct test_gemm_multi_dim_2; +template struct test_gemm_multi_dim_2; template struct test_gemm_multi_dim_2; diff --git a/test/verify/test_gemm_multi_dim_2_3.cpp b/test/verify/test_gemm_multi_dim_2_3.cpp index 1c9f936d332..e59cdcef648 100644 --- a/test/verify/test_gemm_multi_dim_2_3.cpp +++ b/test/verify/test_gemm_multi_dim_2_3.cpp @@ -48,7 +48,8 @@ struct test_gemm_multi_dim_2_3 : verify_program> template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; +template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; -template struct test_gemm_multi_dim_2_3; +// template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; diff --git a/test/verify/test_gemm_multi_transpose.cpp b/test/verify/test_gemm_multi_transpose.cpp index c34cf437d3d..fb3cf22e477 100644 --- a/test/verify/test_gemm_multi_transpose.cpp +++ b/test/verify/test_gemm_multi_transpose.cpp @@ -52,7 +52,8 @@ struct test_gemm_multi_transpose : verify_program; template struct test_gemm_multi_transpose; +template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; -template struct test_gemm_multi_transpose; +// template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; diff --git a/test/verify/test_gemm_multibroadcast.cpp b/test/verify/test_gemm_multibroadcast.cpp index bd771d30e04..db9899156c2 100644 --- a/test/verify/test_gemm_multibroadcast.cpp +++ b/test/verify/test_gemm_multibroadcast.cpp @@ -46,7 +46,8 @@ struct test_gemm_multibroadcast : verify_program template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; +template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; -template struct test_gemm_multibroadcast; +// template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; diff --git a/test/verify/test_gemm_pointwise.cpp b/test/verify/test_gemm_pointwise.cpp index 572ffbfaf27..c35cbc2e8fe 100644 --- a/test/verify/test_gemm_pointwise.cpp +++ b/test/verify/test_gemm_pointwise.cpp @@ -55,7 +55,8 @@ struct test_gemm_pointwise : verify_program> template struct test_gemm_pointwise; template struct test_gemm_pointwise; +template struct test_gemm_pointwise; template struct test_gemm_pointwise; -template struct test_gemm_pointwise; +// template struct test_gemm_pointwise; template struct test_gemm_pointwise; template struct test_gemm_pointwise; diff --git a/test/verify/test_gemm_reshapes_add.cpp b/test/verify/test_gemm_reshapes_add.cpp index 7dd8b2dd396..0139211b184 100644 --- a/test/verify/test_gemm_reshapes_add.cpp +++ b/test/verify/test_gemm_reshapes_add.cpp @@ -60,3 +60,4 @@ struct test_gemm_reshapes_add : verify_program> template struct test_gemm_reshapes_add; template struct test_gemm_reshapes_add; +template struct test_gemm_reshapes_add; diff --git a/test/verify/test_gemm_softmax_gemm_relu_bf16.cpp b/test/verify/test_gemm_softmax_gemm_relu_bf16.cpp new file mode 100644 index 00000000000..1bb182624fa --- /dev/null +++ b/test/verify/test_gemm_softmax_gemm_relu_bf16.cpp @@ -0,0 +1,80 @@ +// /* +// * The MIT License (MIT) +// * +// * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. +// * +// * Permission is hereby granted, free of charge, to any person obtaining a copy +// * of this software and associated documentation files (the "Software"), to deal +// * in the Software without restriction, including without limitation the rights +// * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// * copies of the Software, and to permit persons to whom the Software is +// * furnished to do so, subject to the following conditions: +// * +// * The above copyright notice and this permission notice shall be included in +// * all copies or substantial portions of the Software. +// * +// * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// * THE SOFTWARE. +// */ + +// #include "verify_program.hpp" +// #include +// #include +// #include + +// enum class bias +// { +// without, +// with, +// with_standard_shape +// }; + +// template +// struct test_gemm_softmax_gemm_relu_bf16 : verify_program> +// { +// migraphx::program create_program() const +// { +// migraphx::program p; +// auto* mm = p.get_main_module(); +// migraphx::shape m1_shape{migraphx::shape::bf16_type, {1, 12, 256, 256}}; +// auto m2_elements = m1_shape.elements(); +// auto a = mm->add_parameter("1", m1_shape); +// auto b = mm->add_parameter("2", m1_shape); +// auto b1 = mm->add_parameter("3", m1_shape); +// std::vector eights(m2_elements, 0.125); +// auto eight = mm->add_literal(migraphx::literal{m1_shape, eights}); + +// b = mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), b); +// auto gemm1 = mm->add_instruction(migraphx::make_op("dot"), a, b); +// auto scale = mm->add_instruction(migraphx::make_op("mul"), gemm1, eight); + +// std::optional add_bias{std::nullopt}; +// if constexpr(Config == bias::with or Config == bias::with_standard_shape) +// { +// auto bias_shape = m1_shape; +// if(Config != bias::with_standard_shape) +// { +// bias_shape = migraphx::shape::from_permutation( +// bias_shape.type(), bias_shape.lens(), {0, 1, 3, 2}); +// } +// auto bias_term = mm->add_parameter("4", bias_shape); +// add_bias = mm->add_instruction(migraphx::make_op("add"), scale, bias_term); +// } + +// auto softmax = mm->add_instruction(migraphx::make_op("softmax", {{"axis", 3}}), +// Config == bias::without ? scale : add_bias.value()); +// auto gemm2 = mm->add_instruction(migraphx::make_op("dot"), softmax, b1); +// mm->add_instruction(migraphx::make_op("relu"), gemm2); +// return p; +// } +// std::string section() const { return "gemm"; } +// }; + +// template struct test_gemm_softmax_gemm_relu_bf16; +// template struct test_gemm_softmax_gemm_relu_bf16; +// template struct test_gemm_softmax_gemm_relu_bf16; diff --git a/test/verify/test_gemm_transpose_add_pooling_sub.cpp b/test/verify/test_gemm_transpose_add_pooling_sub.cpp index 7ca7ba931b9..5c234df7b3f 100644 --- a/test/verify/test_gemm_transpose_add_pooling_sub.cpp +++ b/test/verify/test_gemm_transpose_add_pooling_sub.cpp @@ -65,7 +65,8 @@ struct test_gemm_transpose_add_pooling_sub template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; +template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; -template struct test_gemm_transpose_add_pooling_sub; +// template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; diff --git a/test/verify/test_gemm_transposea.cpp b/test/verify/test_gemm_transposea.cpp index 527a537ea0c..86686408c1f 100644 --- a/test/verify/test_gemm_transposea.cpp +++ b/test/verify/test_gemm_transposea.cpp @@ -45,8 +45,9 @@ struct test_gemm_transposea : verify_program> template struct test_gemm_transposea; template struct test_gemm_transposea; +template struct test_gemm_transposea; template struct test_gemm_transposea; -template struct test_gemm_transposea; +// template struct test_gemm_transposea; // TODO need hipblaslt support // template struct test_gemm_transposea; // template struct test_gemm_transposea; diff --git a/test/verify/test_gemm_transposea_ex.cpp b/test/verify/test_gemm_transposea_ex.cpp index 7244fc2c544..c5004183a12 100644 --- a/test/verify/test_gemm_transposea_ex.cpp +++ b/test/verify/test_gemm_transposea_ex.cpp @@ -46,7 +46,8 @@ struct test_gemm_transposea_ex : verify_program> template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; +template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; -template struct test_gemm_transposea_ex; +// template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; diff --git a/test/verify/test_gemm_transposeab.cpp b/test/verify/test_gemm_transposeab.cpp index 63baa4903c8..0ce2c4db0fc 100644 --- a/test/verify/test_gemm_transposeab.cpp +++ b/test/verify/test_gemm_transposeab.cpp @@ -46,7 +46,8 @@ struct test_gemm_transposeab : verify_program> template struct test_gemm_transposeab; template struct test_gemm_transposeab; +template struct test_gemm_transposeab; template struct test_gemm_transposeab; -template struct test_gemm_transposeab; +// template struct test_gemm_transposeab; template struct test_gemm_transposeab; template struct test_gemm_transposeab; diff --git a/test/verify/test_gemm_transposeb.cpp b/test/verify/test_gemm_transposeb.cpp index fef29432136..1e58deca006 100644 --- a/test/verify/test_gemm_transposeb.cpp +++ b/test/verify/test_gemm_transposeb.cpp @@ -45,7 +45,8 @@ struct test_gemm_transposeb : verify_program> template struct test_gemm_transposeb; template struct test_gemm_transposeb; +template struct test_gemm_transposeb; template struct test_gemm_transposeb; -template struct test_gemm_transposeb; +// template struct test_gemm_transposeb; template struct test_gemm_transposeb; template struct test_gemm_transposeb; diff --git a/test/verify/test_gemm_transposeb_detect.cpp b/test/verify/test_gemm_transposeb_detect.cpp index 4b83e1c83b7..bd8a987c35e 100644 --- a/test/verify/test_gemm_transposeb_detect.cpp +++ b/test/verify/test_gemm_transposeb_detect.cpp @@ -46,7 +46,8 @@ struct test_gemm_transposeb_detect : verify_program; template struct test_gemm_transposeb_detect; +template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; -template struct test_gemm_transposeb_detect; +// template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; diff --git a/test/verify/test_gemm_transposeb_ex.cpp b/test/verify/test_gemm_transposeb_ex.cpp index ef75d796de1..1408c56c29b 100644 --- a/test/verify/test_gemm_transposeb_ex.cpp +++ b/test/verify/test_gemm_transposeb_ex.cpp @@ -46,7 +46,8 @@ struct test_gemm_transposeb_ex : verify_program> template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; +template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; -template struct test_gemm_transposeb_ex; +// template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; diff --git a/test/verify/test_hsqrt_bf16.cpp b/test/verify/test_hsqrt_bf16.cpp new file mode 100644 index 00000000000..5f8fbb12755 --- /dev/null +++ b/test/verify/test_hsqrt_bf16.cpp @@ -0,0 +1,42 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_hsqrt_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {2, 3, 4, 6}}; + auto param = mm->add_parameter("x", s); + auto param_abs = mm->add_instruction(migraphx::make_op("abs"), param); + mm->add_instruction(migraphx::make_op("sqrt"), param_abs); + return p; + } +}; diff --git a/test/verify/test_instancenorm.cpp b/test/verify/test_instancenorm.cpp index 77a5541fbae..47d951aea5a 100644 --- a/test/verify/test_instancenorm.cpp +++ b/test/verify/test_instancenorm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -75,8 +75,6 @@ struct test_instancenorm : verify_program> add_instancenorm(*mm, x, {1, 2, 1, 1}); return p; } - - std::string section() const { return "reduce"; } }; template struct test_instancenorm; template struct test_instancenorm; @@ -93,9 +91,8 @@ struct test_instancenorm_large_3d : verify_program; template struct test_instancenorm_large_3d; +template struct test_instancenorm_large_3d; diff --git a/test/verify/test_isinf.cpp b/test/verify/test_isinf.cpp index dcdf700379e..c6e4c798fb9 100644 --- a/test/verify/test_isinf.cpp +++ b/test/verify/test_isinf.cpp @@ -49,4 +49,5 @@ struct test_isinf : verify_program> }; template struct test_isinf; +template struct test_isinf; template struct test_isinf; diff --git a/test/verify/test_isnan.cpp b/test/verify/test_isnan.cpp index 786fd87f798..702d7f3c954 100644 --- a/test/verify/test_isnan.cpp +++ b/test/verify/test_isnan.cpp @@ -44,6 +44,7 @@ struct test_isnan : verify_program> template struct test_isnan; template struct test_isnan; +template struct test_isnan; template struct test_isnan; template struct test_isnan; template struct test_isnan; diff --git a/test/verify/test_layernorm.cpp b/test/verify/test_layernorm.cpp index 725834bad1e..013970d2c89 100644 --- a/test/verify/test_layernorm.cpp +++ b/test/verify/test_layernorm.cpp @@ -40,8 +40,6 @@ struct test_layernorm : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm2 : verify_program @@ -55,8 +53,6 @@ struct test_layernorm2 : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_large : verify_program @@ -70,8 +66,6 @@ struct test_layernorm_large : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_fp16 : verify_program @@ -85,8 +79,19 @@ struct test_layernorm_fp16 : verify_program add_layernorm(*mm, x, dims); return p; } +}; - std::string section() const { return "reduce"; } +struct test_layernorm_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + std::vector dims = {1, 24, 64}; + auto x = mm->add_parameter("x", migraphx::shape{migraphx::shape::bf16_type, dims}); + add_layernorm(*mm, x, dims); + return p; + } }; struct test_layernorm_fp8_1 : verify_program @@ -100,8 +105,6 @@ struct test_layernorm_fp8_1 : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_2 : verify_program @@ -115,8 +118,6 @@ struct test_layernorm_fp8_2 : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_3 : verify_program @@ -130,8 +131,6 @@ struct test_layernorm_fp8_3 : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_4 : verify_program @@ -145,8 +144,6 @@ struct test_layernorm_fp8_4 : verify_program add_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_eps : verify_program @@ -160,8 +157,6 @@ struct test_layernorm_eps : verify_program add_layernorm(*mm, x, dims, 1e-5f); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_triadd : verify_program @@ -179,8 +174,6 @@ struct test_layernorm_triadd : verify_program add_layernorm(*mm, add2, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_layernorm_triadd_large : verify_program @@ -198,8 +191,6 @@ struct test_layernorm_triadd_large : verify_program add_layernorm(*mm, add2, dims); return p; } - - std::string section() const { return "reduce"; } }; struct test_add_layernorm_add_gemm_nonstd : verify_program @@ -232,6 +223,4 @@ struct test_pw_layernorm : verify_program add_pointwise_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_literal_limits.cpp b/test/verify/test_literal_limits.cpp index 9dac1bb7565..609ccbb6300 100644 --- a/test/verify/test_literal_limits.cpp +++ b/test/verify/test_literal_limits.cpp @@ -55,6 +55,7 @@ struct test_literal_limits : verify_program> template struct test_literal_limits; template struct test_literal_limits; template struct test_literal_limits; +template struct test_literal_limits; template struct test_literal_limits; template struct test_literal_limits; template struct test_literal_limits; diff --git a/test/verify/test_log.cpp b/test/verify/test_log.cpp index 366e5c7988a..2cea7e777a1 100644 --- a/test/verify/test_log.cpp +++ b/test/verify/test_log.cpp @@ -43,6 +43,7 @@ struct test_log : verify_program> template struct test_log; template struct test_log; +template struct test_log; template struct test_log; template struct test_log; template struct test_log; diff --git a/test/verify/test_log2.cpp b/test/verify/test_log2.cpp index bcbf149e4e3..fda642bff91 100644 --- a/test/verify/test_log2.cpp +++ b/test/verify/test_log2.cpp @@ -43,6 +43,7 @@ struct test_log2 : verify_program> template struct test_log2; template struct test_log2; +template struct test_log2; template struct test_log2; template struct test_log2; template struct test_log2; diff --git a/test/verify/test_logsoftmax.cpp b/test/verify/test_logsoftmax.cpp index afa1ec43aa1..ff9a7bc456b 100644 --- a/test/verify/test_logsoftmax.cpp +++ b/test/verify/test_logsoftmax.cpp @@ -40,8 +40,6 @@ struct test_logsoftmax : verify_program> return p; } - - std::string section() const { return "reduce"; } }; template struct test_logsoftmax<0, migraphx::shape::float_type>; @@ -54,6 +52,11 @@ template struct test_logsoftmax<0, migraphx::shape::half_type>; template struct test_logsoftmax<2, migraphx::shape::half_type>; template struct test_logsoftmax<3, migraphx::shape::half_type>; +template struct test_logsoftmax<1, migraphx::shape::bf16_type>; +template struct test_logsoftmax<0, migraphx::shape::bf16_type>; +template struct test_logsoftmax<2, migraphx::shape::bf16_type>; +template struct test_logsoftmax<3, migraphx::shape::bf16_type>; + template struct test_logsoftmax<1, migraphx::shape::fp8e4m3fnuz_type>; template struct test_logsoftmax<3, migraphx::shape::fp8e4m3fnuz_type>; diff --git a/test/verify/test_logsoftmax1.cpp b/test/verify/test_logsoftmax1.cpp index d6a9daf83b5..a51a83e832b 100644 --- a/test/verify/test_logsoftmax1.cpp +++ b/test/verify/test_logsoftmax1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -40,6 +40,4 @@ struct test_logsoftmax1 : verify_program mm->add_return({r}); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_min_max.cpp b/test/verify/test_min_max.cpp index 12310c2c7a8..b1985347049 100644 --- a/test/verify/test_min_max.cpp +++ b/test/verify/test_min_max.cpp @@ -45,6 +45,7 @@ struct test_min_max : verify_program> template struct test_min_max; template struct test_min_max; +template struct test_min_max; template struct test_min_max; template struct test_min_max; template struct test_min_max; @@ -53,6 +54,7 @@ template struct test_min_max; template struct test_min_max; template struct test_min_max; +template struct test_min_max; template struct test_min_max; template struct test_min_max; template struct test_min_max; diff --git a/test/verify/test_mul_dot_a.cpp b/test/verify/test_mul_dot_a.cpp index d46a613d646..e5854937734 100644 --- a/test/verify/test_mul_dot_a.cpp +++ b/test/verify/test_mul_dot_a.cpp @@ -51,7 +51,8 @@ struct test_mul_dot_a : verify_program> template struct test_mul_dot_a; template struct test_mul_dot_a; +template struct test_mul_dot_a; template struct test_mul_dot_a; -template struct test_mul_dot_a; +// template struct test_mul_dot_a; template struct test_mul_dot_a; template struct test_mul_dot_a; diff --git a/test/verify/test_mul_dot_b.cpp b/test/verify/test_mul_dot_b.cpp index 6941baa9e36..07d0781b755 100644 --- a/test/verify/test_mul_dot_b.cpp +++ b/test/verify/test_mul_dot_b.cpp @@ -52,7 +52,8 @@ struct test_mul_dot_b : verify_program> template struct test_mul_dot_b; template struct test_mul_dot_b; +template struct test_mul_dot_b; template struct test_mul_dot_b; -template struct test_mul_dot_b; +// template struct test_mul_dot_b; template struct test_mul_dot_b; template struct test_mul_dot_b; diff --git a/test/verify/test_multinomial.cpp b/test/verify/test_multinomial.cpp index 0559d0ec0a5..53bc2d1f4a2 100644 --- a/test/verify/test_multinomial.cpp +++ b/test/verify/test_multinomial.cpp @@ -62,6 +62,7 @@ struct test_multinomial : verify_program> template struct test_multinomial; template struct test_multinomial; +template struct test_multinomial; // TODO This fails, need to figure out why // template struct test_multinomial; // template struct test_multinomial; diff --git a/test/verify/test_nearbyint.cpp b/test/verify/test_nearbyint.cpp index 2f6b3163c4a..f9bdb7a65a3 100644 --- a/test/verify/test_nearbyint.cpp +++ b/test/verify/test_nearbyint.cpp @@ -45,6 +45,7 @@ struct test_nearbyint : verify_program> }; template struct test_nearbyint; +template struct test_nearbyint; template struct test_nearbyint; template struct test_nearbyint; template struct test_nearbyint; diff --git a/test/verify/test_nonzero.cpp b/test/verify/test_nonzero.cpp index c689c1e7485..ffd45b230a0 100644 --- a/test/verify/test_nonzero.cpp +++ b/test/verify/test_nonzero.cpp @@ -45,6 +45,7 @@ struct test_nonzero : verify_program> template struct test_nonzero; template struct test_nonzero; +template struct test_nonzero; template struct test_nonzero; template struct test_nonzero; template struct test_nonzero; diff --git a/test/verify/test_pad.cpp b/test/verify/test_pad.cpp index d4351d32fd4..e88d48ce1f3 100644 --- a/test/verify/test_pad.cpp +++ b/test/verify/test_pad.cpp @@ -51,3 +51,4 @@ struct test_pad : verify_program> template struct test_pad; template struct test_pad; template struct test_pad; +template struct test_pad; diff --git a/test/verify/test_pad_highest_bf16.cpp b/test/verify/test_pad_highest_bf16.cpp new file mode 100644 index 00000000000..e17189f4cec --- /dev/null +++ b/test/verify/test_pad_highest_bf16.cpp @@ -0,0 +1,46 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_pad_highest_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + std::vector data0(4); + std::iota(data0.begin(), data0.end(), 0); + migraphx::shape s0{migraphx::shape::bf16_type, {2, 2}}; + auto l0 = mm->add_literal(migraphx::literal{s0, data0}); + migraphx::op::pad op{}; + op.value = std::numeric_limits::max(); + op.pads = {0, 0, 1, 1}; + mm->add_instruction(op, l0); + return p; + } +}; diff --git a/test/verify/test_pad_lowest_bf16.cpp b/test/verify/test_pad_lowest_bf16.cpp new file mode 100644 index 00000000000..ada6705f9c0 --- /dev/null +++ b/test/verify/test_pad_lowest_bf16.cpp @@ -0,0 +1,46 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_pad_lowest_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + std::vector data0(4); + std::iota(data0.begin(), data0.end(), 0); + migraphx::shape s0{migraphx::shape::bf16_type, {2, 2}}; + auto l0 = mm->add_literal(migraphx::literal{s0, data0}); + migraphx::op::pad op{}; + op.value = std::numeric_limits::lowest(); + op.pads = {0, 0, 1, 1}; + mm->add_instruction(op, l0); + return p; + } +}; diff --git a/test/verify/test_pointwise_broadcast_reduce.cpp b/test/verify/test_pointwise_broadcast_reduce.cpp index c0269625457..ffcc658838b 100644 --- a/test/verify/test_pointwise_broadcast_reduce.cpp +++ b/test/verify/test_pointwise_broadcast_reduce.cpp @@ -52,6 +52,4 @@ struct test_pointwise_broadcast_reduce : verify_programadd_return({reshape}); return p; }; - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_pointwise_broadcast_reduce_bf16.cpp b/test/verify/test_pointwise_broadcast_reduce_bf16.cpp new file mode 100644 index 00000000000..65300bd6189 --- /dev/null +++ b/test/verify/test_pointwise_broadcast_reduce_bf16.cpp @@ -0,0 +1,55 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include +#include + +struct test_pointwise_broadcast_reduce_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::shape s{migraphx::shape::bf16_type, {2, 32, 96}}; + migraphx::shape rs{migraphx::shape::bf16_type, {2, 1, 1}}; + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_parameter("x", rs); + auto y = mm->add_parameter("y", s); + auto abs = mm->add_instruction(migraphx::make_op("abs"), x); + auto sqrt = mm->add_instruction(migraphx::make_op("sqrt"), abs); + auto sqrtb = mm->add_instruction( + migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), sqrt); + auto add = mm->add_instruction(migraphx::make_op("add"), y, sqrtb); + auto rsum = mm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", {1, 2}}}), add); + auto rsumb = mm->add_instruction( + migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), rsum); + auto sub = mm->add_instruction(migraphx::make_op("sub"), rsumb, add); + auto reshape = + mm->add_instruction(migraphx::make_op("reshape", {{"dims", {s.elements()}}}), sub); + mm->add_return({reshape}); + return p; + }; +}; diff --git a/test/verify/test_pow.cpp b/test/verify/test_pow.cpp index 3f1626a08e5..71765b7d7d9 100644 --- a/test/verify/test_pow.cpp +++ b/test/verify/test_pow.cpp @@ -45,6 +45,7 @@ struct test_pow : verify_program> }; template struct test_pow; template struct test_pow; +template struct test_pow; template struct test_pow; template struct test_pow; template struct test_pow; diff --git a/test/verify/test_prefix_scan_sum_2d.cpp b/test/verify/test_prefix_scan_sum_2d.cpp index a7530e08009..ae9245c5329 100644 --- a/test/verify/test_prefix_scan_sum_2d.cpp +++ b/test/verify/test_prefix_scan_sum_2d.cpp @@ -46,6 +46,7 @@ struct test_prefix_scan_sum_2d_small : verify_program; template struct test_prefix_scan_sum_2d_small; +template struct test_prefix_scan_sum_2d_small; template struct test_prefix_scan_sum_2d_small; template struct test_prefix_scan_sum_2d_small; template struct test_prefix_scan_sum_2d_small; @@ -68,3 +69,4 @@ struct test_prefix_scan_sum_2d_large : verify_program; template struct test_prefix_scan_sum_2d_large; +template struct test_prefix_scan_sum_2d_large; diff --git a/test/verify/test_reduce_add.cpp b/test/verify/test_reduce_add.cpp index f1d8d6f3fd3..c0e2b5ddf93 100644 --- a/test/verify/test_reduce_add.cpp +++ b/test/verify/test_reduce_add.cpp @@ -47,8 +47,6 @@ struct test_reduce_add : verify_program> mm->add_return({add}); return p; }; - - std::string section() const { return "reduce"; } }; template struct test_reduce_add; diff --git a/test/verify/test_reduce_add_bf16.cpp b/test/verify/test_reduce_add_bf16.cpp new file mode 100644 index 00000000000..aaa91be5bcc --- /dev/null +++ b/test/verify/test_reduce_add_bf16.cpp @@ -0,0 +1,52 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include +#include +#include + +template +struct test_reduce_add_bf16 : verify_program> +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{DType, {4, 1000, 2, 2}}; + migraphx::shape bs{migraphx::shape::bf16_type, {1, 32, 128}}; + auto x = mm->add_parameter("x", s); + auto reduce_mean = + mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2, 3}}}), x); + auto reduce_max = + mm->add_instruction(migraphx::make_op("reduce_max", {{"axes", {2, 3}}}), x); + auto add = mm->add_instruction(migraphx::make_op("add"), reduce_mean, reduce_max); + mm->add_return({add}); + return p; + }; +}; + +template struct test_reduce_add_bf16; diff --git a/test/verify/test_reduce_mean_bias_bf16.cpp b/test/verify/test_reduce_mean_bias_bf16.cpp new file mode 100644 index 00000000000..df5839966ec --- /dev/null +++ b/test/verify/test_reduce_mean_bias_bf16.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include +#include + +struct test_reduce_mean_bias_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {1, 32, 128}}; + migraphx::shape bs{migraphx::shape::bf16_type, {1, 32, 128}}; + auto x = mm->add_parameter("x", s); + auto reduce = mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x); + auto bias = mm->add_parameter("bias", reduce->get_shape()); + auto add = mm->add_instruction(migraphx::make_op("add"), reduce, bias); + auto abs = mm->add_instruction(migraphx::make_op("abs"), add); + auto sqrt = mm->add_instruction(migraphx::make_op("sqrt"), abs); + mm->add_return({sqrt}); + return p; + }; +}; diff --git a/test/verify/test_reduce_mean_bias_half.cpp b/test/verify/test_reduce_mean_bias_half.cpp index e70d9dc1dc4..c83d7565de7 100644 --- a/test/verify/test_reduce_mean_bias_half.cpp +++ b/test/verify/test_reduce_mean_bias_half.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -45,6 +45,4 @@ struct test_reduce_mean_bias_half : verify_program mm->add_return({sqrt}); return p; }; - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_large_bf16.cpp b/test/verify/test_reduce_mean_large_bf16.cpp new file mode 100644 index 00000000000..927b1827838 --- /dev/null +++ b/test/verify/test_reduce_mean_large_bf16.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_reduce_mean_large_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {1, 32, 65536}}; + auto x = mm->add_parameter("x", s); + mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x); + return p; + }; +}; diff --git a/test/verify/test_reduce_mean_large_half.cpp b/test/verify/test_reduce_mean_large_half.cpp index 5e40f9e9229..f8925dbe577 100644 --- a/test/verify/test_reduce_mean_large_half.cpp +++ b/test/verify/test_reduce_mean_large_half.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -38,6 +38,4 @@ struct test_reduce_mean_large_half : verify_program mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x); return p; }; - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_nhwc.cpp b/test/verify/test_reduce_mean_nhwc.cpp index 275c7d218d3..0307af20988 100644 --- a/test/verify/test_reduce_mean_nhwc.cpp +++ b/test/verify/test_reduce_mean_nhwc.cpp @@ -43,9 +43,8 @@ struct test_reduce_mean_nhwc : verify_program> mm->add_return({sqrt}); return p; }; - - std::string section() const { return "reduce"; } }; template struct test_reduce_mean_nhwc; template struct test_reduce_mean_nhwc; +template struct test_reduce_mean_nhwc; diff --git a/test/verify/test_reduce_mean_reduce_sum.cpp b/test/verify/test_reduce_mean_reduce_sum.cpp index b5df0f0147a..9dd29f8cffe 100644 --- a/test/verify/test_reduce_mean_reduce_sum.cpp +++ b/test/verify/test_reduce_mean_reduce_sum.cpp @@ -51,6 +51,4 @@ struct test_reduce_mean_reduce_sum : verify_program mm->add_return({mean_div2}); return p; }; - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_variance.cpp b/test/verify/test_reduce_mean_variance.cpp index 6843ae7e459..dccd74a84c3 100644 --- a/test/verify/test_reduce_mean_variance.cpp +++ b/test/verify/test_reduce_mean_variance.cpp @@ -45,6 +45,4 @@ struct test_reduce_mean_variance : verify_program mm->add_return({add}); return p; }; - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_variance_bf16.cpp b/test/verify/test_reduce_mean_variance_bf16.cpp new file mode 100644 index 00000000000..a0694a4b5a0 --- /dev/null +++ b/test/verify/test_reduce_mean_variance_bf16.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include +#include + +struct test_reduce_mean_variance_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {1, 32, 128}}; + auto x = mm->add_parameter("x", s); + auto mean = mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x); + auto x2 = mm->add_instruction(migraphx::make_op("mul"), x, x); + auto mean_x2 = mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x2); + auto mean_2 = mm->add_instruction(migraphx::make_op("mul"), mean, mean); + auto variance = mm->add_instruction(migraphx::make_op("sub"), mean_x2, mean_2); + auto add = mm->add_instruction(migraphx::make_op("add"), mean, variance); + mm->add_return({add}); + return p; + }; +}; diff --git a/test/verify/test_reduce_noop_add.cpp b/test/verify/test_reduce_noop_add.cpp index c72d7f967f0..9e642da2f71 100644 --- a/test/verify/test_reduce_noop_add.cpp +++ b/test/verify/test_reduce_noop_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -45,6 +45,4 @@ struct test_reduce_noop_add : verify_program mm->add_return({add}); return p; }; - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_noop_add_bf16.cpp b/test/verify/test_reduce_noop_add_bf16.cpp new file mode 100644 index 00000000000..4ee3692e89b --- /dev/null +++ b/test/verify/test_reduce_noop_add_bf16.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include +#include + +struct test_reduce_noop_add_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::float_type, {4, 1000, 1, 1}}; + migraphx::shape bs{migraphx::shape::bf16_type, {1, 32, 128}}; + auto x = mm->add_parameter("x", s); + auto reduce_mean = + mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2, 3}}}), x); + auto reduce_max = + mm->add_instruction(migraphx::make_op("reduce_max", {{"axes", {2, 3}}}), x); + auto add = mm->add_instruction(migraphx::make_op("add"), reduce_mean, reduce_max); + mm->add_return({add}); + return p; + }; +}; diff --git a/test/verify/test_reduce_op_large.cpp b/test/verify/test_reduce_op_large.cpp index edcd550ef66..473c762d23a 100644 --- a/test/verify/test_reduce_op_large.cpp +++ b/test/verify/test_reduce_op_large.cpp @@ -45,8 +45,6 @@ struct test_reduce_op_large : verify_program> mm->add_instruction(Op{{Axis}}, x); return p; }; - - std::string section() const { return "reduce"; } }; template struct test_reduce_op_large; diff --git a/test/verify/test_reduce_op_small.cpp b/test/verify/test_reduce_op_small.cpp index b1d824f1c8a..81298039f77 100644 --- a/test/verify/test_reduce_op_small.cpp +++ b/test/verify/test_reduce_op_small.cpp @@ -45,8 +45,6 @@ struct test_reduce_op_small : verify_program> mm->add_instruction(Op{{Axis}}, x); return p; }; - - std::string section() const { return "reduce"; } }; template struct test_reduce_op_small; @@ -68,6 +66,16 @@ template struct test_reduce_op_small; template struct test_reduce_op_small; + +template struct test_reduce_op_small; +template struct test_reduce_op_small; +template struct test_reduce_op_small; +template struct test_reduce_op_small; +template struct test_reduce_op_small; +template struct test_reduce_op_small; +template struct test_reduce_op_small; +template struct test_reduce_op_small; + template struct test_reduce_op_small; diff --git a/test/verify/test_reverse.cpp b/test/verify/test_reverse.cpp index c2c2f037c55..2f5b6769df2 100644 --- a/test/verify/test_reverse.cpp +++ b/test/verify/test_reverse.cpp @@ -43,6 +43,7 @@ struct test_reverse : verify_program> template struct test_reverse; template struct test_reverse; +template struct test_reverse; template struct test_reverse; template struct test_reverse; template struct test_reverse; diff --git a/test/verify/test_rnn_sql_1.cpp b/test/verify/test_rnn_sql_1.cpp index 87218a81b74..c9d72fd99a8 100644 --- a/test/verify/test_rnn_sql_1.cpp +++ b/test/verify/test_rnn_sql_1.cpp @@ -84,7 +84,8 @@ struct test_rnn_sql_1 : verify_program> template struct test_rnn_sql_1; template struct test_rnn_sql_1; +template struct test_rnn_sql_1; template struct test_rnn_sql_1; -template struct test_rnn_sql_1; +// template struct test_rnn_sql_1; template struct test_rnn_sql_1; template struct test_rnn_sql_1; diff --git a/test/verify/test_roialign.cpp b/test/verify/test_roialign.cpp index 9f5271f1c5d..a3c59aa513c 100644 --- a/test/verify/test_roialign.cpp +++ b/test/verify/test_roialign.cpp @@ -60,3 +60,4 @@ struct test_roialign : verify_program> template struct test_roialign; template struct test_roialign; +template struct test_roialign; diff --git a/test/verify/test_rsqrt.cpp b/test/verify/test_rsqrt.cpp index cc10878788a..86abf4088c3 100644 --- a/test/verify/test_rsqrt.cpp +++ b/test/verify/test_rsqrt.cpp @@ -54,6 +54,7 @@ struct test_rsqrt : verify_program> template struct test_rsqrt; template struct test_rsqrt; +template struct test_rsqrt; template struct test_rsqrt; template struct test_rsqrt; template struct test_rsqrt; diff --git a/test/verify/test_scatter_elements_none_axis_neg_1.cpp b/test/verify/test_scatter_elements_none_axis_neg_1.cpp index 354e9f85332..d16d7950b35 100644 --- a/test/verify/test_scatter_elements_none_axis_neg_1.cpp +++ b/test/verify/test_scatter_elements_none_axis_neg_1.cpp @@ -52,6 +52,7 @@ struct test_scatter_elements_none_axis_neg_1 template struct test_scatter_elements_none_axis_neg_1; template struct test_scatter_elements_none_axis_neg_1; +template struct test_scatter_elements_none_axis_neg_1; template struct test_scatter_elements_none_axis_neg_1; template struct test_scatter_elements_none_axis_neg_1; template struct test_scatter_elements_none_axis_neg_1; diff --git a/test/verify/test_scatternd.cpp b/test/verify/test_scatternd.cpp index f33bd9ab05d..04a34d45476 100644 --- a/test/verify/test_scatternd.cpp +++ b/test/verify/test_scatternd.cpp @@ -55,4 +55,5 @@ struct test_scatternd : verify_program> template struct test_scatternd; template struct test_scatternd; +template struct test_scatternd; template struct test_scatternd; diff --git a/test/verify/test_select_module_reduce.cpp b/test/verify/test_select_module_reduce.cpp index 5e3d8ad1d63..aeffd90417e 100644 --- a/test/verify/test_select_module_reduce.cpp +++ b/test/verify/test_select_module_reduce.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -66,6 +66,4 @@ struct test_select_module_reduce : verify_program return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_shrink.cpp b/test/verify/test_shrink.cpp index ed3ab4ce86e..14c0709b056 100644 --- a/test/verify/test_shrink.cpp +++ b/test/verify/test_shrink.cpp @@ -76,6 +76,7 @@ struct test_shrink : verify_program> template struct test_shrink; template struct test_shrink; template struct test_shrink; +template struct test_shrink; template struct test_shrink; template struct test_shrink; template struct test_shrink; diff --git a/test/verify/test_sin.cpp b/test/verify/test_sin.cpp index 90d0e6da26b..626aa225be6 100644 --- a/test/verify/test_sin.cpp +++ b/test/verify/test_sin.cpp @@ -43,6 +43,7 @@ struct test_sin : verify_program> template struct test_sin; template struct test_sin; +template struct test_sin; template struct test_sin; template struct test_sin; template struct test_sin; diff --git a/test/verify/test_sin_bf16.cpp b/test/verify/test_sin_bf16.cpp new file mode 100644 index 00000000000..64e6589ad44 --- /dev/null +++ b/test/verify/test_sin_bf16.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_sin_bf16 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {10}}; + auto x = mm->add_parameter("x", s); + mm->add_instruction(migraphx::make_op("sin"), x); + return p; + } +}; diff --git a/test/verify/test_sinh.cpp b/test/verify/test_sinh.cpp index 91bceea24ca..541001c4d46 100644 --- a/test/verify/test_sinh.cpp +++ b/test/verify/test_sinh.cpp @@ -43,6 +43,7 @@ struct test_sinh : verify_program> template struct test_sinh; template struct test_sinh; +template struct test_sinh; template struct test_sinh; template struct test_sinh; template struct test_sinh; diff --git a/test/verify/test_softmax.cpp b/test/verify/test_softmax.cpp index 3d666a240ad..b6e27771756 100644 --- a/test/verify/test_softmax.cpp +++ b/test/verify/test_softmax.cpp @@ -40,8 +40,6 @@ struct test_softmax : verify_program> return p; } - - std::string section() const { return "reduce"; } }; template struct test_softmax<0, migraphx::shape::float_type>; @@ -52,6 +50,11 @@ template struct test_softmax<1, migraphx::shape::half_type>; template struct test_softmax<2, migraphx::shape::half_type>; template struct test_softmax<3, migraphx::shape::half_type>; +template struct test_softmax<0, migraphx::shape::bf16_type>; +template struct test_softmax<1, migraphx::shape::bf16_type>; +template struct test_softmax<2, migraphx::shape::bf16_type>; +template struct test_softmax<3, migraphx::shape::bf16_type>; + template struct test_softmax<1, migraphx::shape::fp8e4m3fnuz_type>; template struct test_softmax<3, migraphx::shape::fp8e4m3fnuz_type>; diff --git a/test/verify/test_softmax1.cpp b/test/verify/test_softmax1.cpp index ef8c71871bb..4742e03607a 100644 --- a/test/verify/test_softmax1.cpp +++ b/test/verify/test_softmax1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -37,6 +37,4 @@ struct test_softmax1 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", 0}}), x); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax2.cpp b/test/verify/test_softmax2.cpp index c39dcec6c35..804f8f04f0b 100644 --- a/test/verify/test_softmax2.cpp +++ b/test/verify/test_softmax2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -38,6 +38,4 @@ struct test_softmax2 : verify_program mm->add_instruction(migraphx::make_op("softmax"), x); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax3.cpp b/test/verify/test_softmax3.cpp index c8ca6bc95ef..027c4d64ad3 100644 --- a/test/verify/test_softmax3.cpp +++ b/test/verify/test_softmax3.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -41,6 +41,4 @@ struct test_softmax3 : verify_program mm->add_return({r}); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax4.cpp b/test/verify/test_softmax4.cpp index dae139227a8..73117df57c5 100644 --- a/test/verify/test_softmax4.cpp +++ b/test/verify/test_softmax4.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -38,6 +38,4 @@ struct test_softmax4 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", 3}}), x); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax5.cpp b/test/verify/test_softmax5.cpp new file mode 100644 index 00000000000..499cc497931 --- /dev/null +++ b/test/verify/test_softmax5.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +struct test_softmax5 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = + mm->add_parameter("x", migraphx::shape{migraphx::shape::bf16_type, {1, 12, 384, 384}}); + mm->add_instruction(migraphx::make_op("softmax", {{"axis", 3}}), x); + return p; + } +}; diff --git a/test/verify/test_softmax_large1.cpp b/test/verify/test_softmax_large1.cpp index 9669b095807..e8d98194da8 100644 --- a/test/verify/test_softmax_large1.cpp +++ b/test/verify/test_softmax_large1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -40,6 +40,4 @@ struct test_softmax_large1 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), add); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax_large2.cpp b/test/verify/test_softmax_large2.cpp index 12874f647b9..c4311f6f819 100644 --- a/test/verify/test_softmax_large2.cpp +++ b/test/verify/test_softmax_large2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -40,6 +40,4 @@ struct test_softmax_large2 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), add); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax_large3.cpp b/test/verify/test_softmax_large3.cpp index 900d52ca90d..d91bece8048 100644 --- a/test/verify/test_softmax_large3.cpp +++ b/test/verify/test_softmax_large3.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -40,6 +40,4 @@ struct test_softmax_large3 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), add); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmaxcrossentropyloss_2d.cpp b/test/verify/test_softmaxcrossentropyloss_2d.cpp index 684975ea119..dd92f9cf73c 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d.cpp @@ -81,8 +81,6 @@ struct test_softmaxcrossentropyloss_2d return p; } - - std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_2d; +template struct test_softmaxcrossentropyloss_2d; +template struct test_softmaxcrossentropyloss_2d; + diff --git a/test/verify/test_softmaxcrossentropyloss_2d_mean.cpp b/test/verify/test_softmaxcrossentropyloss_2d_mean.cpp index e4baa424743..3d1b379234c 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d_mean.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d_mean.cpp @@ -84,8 +84,6 @@ struct test_softmaxcrossentropyloss_2d_mean return p; } - - std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_2d_mean; +template struct test_softmaxcrossentropyloss_2d_mean; +template struct test_softmaxcrossentropyloss_2d_mean; diff --git a/test/verify/test_softmaxcrossentropyloss_2d_sum.cpp b/test/verify/test_softmaxcrossentropyloss_2d_sum.cpp index cfc7cfef893..722fe3a8a61 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d_sum.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d_sum.cpp @@ -83,8 +83,6 @@ struct test_softmaxcrossentropyloss_2d_sum return p; } - - std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_2d_sum; +template struct test_softmaxcrossentropyloss_2d_sum; +template struct test_softmaxcrossentropyloss_2d_sum; diff --git a/test/verify/test_softmaxcrossentropyloss_kd.cpp b/test/verify/test_softmaxcrossentropyloss_kd.cpp index 638f7053021..7cae01e0c53 100644 --- a/test/verify/test_softmaxcrossentropyloss_kd.cpp +++ b/test/verify/test_softmaxcrossentropyloss_kd.cpp @@ -107,8 +107,6 @@ struct test_softmaxcrossentropyloss_kd return p; } - - std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_kd; +template struct test_softmaxcrossentropyloss_kd; +template struct test_softmaxcrossentropyloss_kd; diff --git a/test/verify/test_softmaxcrossentropyloss_kd_mean.cpp b/test/verify/test_softmaxcrossentropyloss_kd_mean.cpp index e35d69a8ffd..4a6fa922974 100644 --- a/test/verify/test_softmaxcrossentropyloss_kd_mean.cpp +++ b/test/verify/test_softmaxcrossentropyloss_kd_mean.cpp @@ -115,8 +115,6 @@ struct test_softmaxcrossentropyloss_kd_mean return p; } - - std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_kd_mean; +template struct test_softmaxcrossentropyloss_kd_mean; +template struct test_softmaxcrossentropyloss_kd_mean; + diff --git a/test/verify/test_softmaxcrossentropyloss_kd_sum.cpp b/test/verify/test_softmaxcrossentropyloss_kd_sum.cpp index 140a10fa84c..e41dbe5453c 100644 --- a/test/verify/test_softmaxcrossentropyloss_kd_sum.cpp +++ b/test/verify/test_softmaxcrossentropyloss_kd_sum.cpp @@ -111,8 +111,6 @@ struct test_softmaxcrossentropyloss_kd_sum_reduction return p; } - - std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_kd_sum_reduction; +template struct test_softmaxcrossentropyloss_kd_sum_reduction; +template struct test_softmaxcrossentropyloss_kd_sum_reduction; diff --git a/test/verify/test_sqrt.cpp b/test/verify/test_sqrt.cpp index 9916e7a82bd..83ae758eee8 100644 --- a/test/verify/test_sqrt.cpp +++ b/test/verify/test_sqrt.cpp @@ -44,6 +44,7 @@ struct test_sqrt : verify_program> template struct test_sqrt; template struct test_sqrt; +template struct test_sqrt; template struct test_sqrt; template struct test_sqrt; template struct test_sqrt; diff --git a/test/verify/test_sqrt_bf161.cpp b/test/verify/test_sqrt_bf161.cpp new file mode 100644 index 00000000000..e2079fd1216 --- /dev/null +++ b/test/verify/test_sqrt_bf161.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +// math op on bf16-precision float with odd size tensor can't fit bf162 packing +struct test_sqrt_bf161 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {5}}; + auto param = mm->add_parameter("x", s); + auto param_abs = mm->add_instruction(migraphx::make_op("abs"), param); + mm->add_instruction(migraphx::make_op("sqrt"), param_abs); + return p; + } +}; diff --git a/test/verify/test_sqrt_bf162.cpp b/test/verify/test_sqrt_bf162.cpp new file mode 100644 index 00000000000..dbb26d76ebf --- /dev/null +++ b/test/verify/test_sqrt_bf162.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +// math op on bf16-precision float with tensor size that's divisible by 2, +// but not divisible by 4 +struct test_sqrt_bf162 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {6}}; + auto param = mm->add_parameter("x", s); + auto param_abs = mm->add_instruction(migraphx::make_op("abs"), param); + mm->add_instruction(migraphx::make_op("sqrt"), param_abs); + return p; + } +}; diff --git a/test/verify/test_sqrt_bf164.cpp b/test/verify/test_sqrt_bf164.cpp new file mode 100644 index 00000000000..077afe2a63e --- /dev/null +++ b/test/verify/test_sqrt_bf164.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#include "verify_program.hpp" +#include +#include +#include + +// math op on bf16-precision float with tensor size that fits into bf164 packing +struct test_sqrt_bf164 : verify_program +{ + migraphx::program create_program() const + { + migraphx::program p; + auto* mm = p.get_main_module(); + migraphx::shape s{migraphx::shape::bf16_type, {8}}; + auto param = mm->add_parameter("x", s); + auto param_abs = mm->add_instruction(migraphx::make_op("abs"), param); + mm->add_instruction(migraphx::make_op("sqrt"), param_abs); + return p; + } +}; diff --git a/test/verify/test_tan.cpp b/test/verify/test_tan.cpp index b870c1c43d3..d2a7cb2a47c 100644 --- a/test/verify/test_tan.cpp +++ b/test/verify/test_tan.cpp @@ -43,6 +43,7 @@ struct test_tan : verify_program> template struct test_tan; template struct test_tan; +template struct test_tan; template struct test_tan; template struct test_tan; template struct test_tan; diff --git a/test/verify/test_tanh.cpp b/test/verify/test_tanh.cpp index 8458ea80a38..af1bd468775 100644 --- a/test/verify/test_tanh.cpp +++ b/test/verify/test_tanh.cpp @@ -42,6 +42,7 @@ struct test_tanh : verify_program> template struct test_tanh; template struct test_tanh; +template struct test_tanh; template struct test_tanh; template struct test_tanh; template struct test_tanh; diff --git a/test/verify/test_topk_0.cpp b/test/verify/test_topk_0.cpp index 625e19b6cac..f42a8b88d15 100644 --- a/test/verify/test_topk_0.cpp +++ b/test/verify/test_topk_0.cpp @@ -47,3 +47,4 @@ struct test_topk_0 : verify_program> template struct test_topk_0; template struct test_topk_0; +template struct test_topk_0; diff --git a/test/verify/test_trans_convert_gemm.cpp b/test/verify/test_trans_convert_gemm.cpp index 45701988598..bb150ed34a6 100644 --- a/test/verify/test_trans_convert_gemm.cpp +++ b/test/verify/test_trans_convert_gemm.cpp @@ -47,6 +47,7 @@ struct test_trans_convert_gemm : verify_program> template struct test_trans_convert_gemm; template struct test_trans_convert_gemm; +template struct test_trans_convert_gemm; template struct test_trans_convert_gemm; template struct test_trans_convert_gemm; template struct test_trans_convert_gemm; diff --git a/test/verify/test_transpose_reshape_add_sub_mul.cpp b/test/verify/test_transpose_reshape_add_sub_mul.cpp index 203d8916fd6..e3c242d9d22 100644 --- a/test/verify/test_transpose_reshape_add_sub_mul.cpp +++ b/test/verify/test_transpose_reshape_add_sub_mul.cpp @@ -56,6 +56,7 @@ struct test_transpose_reshape_add_sub_mul template struct test_transpose_reshape_add_sub_mul; template struct test_transpose_reshape_add_sub_mul; +template struct test_transpose_reshape_add_sub_mul; template struct test_transpose_reshape_add_sub_mul; template struct test_transpose_reshape_add_sub_mul; template struct test_transpose_reshape_add_sub_mul; diff --git a/test/verify/test_unbatched_gemm_1.cpp b/test/verify/test_unbatched_gemm_1.cpp index 05aa138efcf..fd8f8a68a5f 100644 --- a/test/verify/test_unbatched_gemm_1.cpp +++ b/test/verify/test_unbatched_gemm_1.cpp @@ -62,6 +62,7 @@ struct test_unbatched_gemm_1 : verify_program> template struct test_unbatched_gemm_1; template struct test_unbatched_gemm_1; +template struct test_unbatched_gemm_1; template struct test_unbatched_gemm_1; template struct test_unbatched_gemm_1; template struct test_unbatched_gemm_1; diff --git a/test/verify/test_unbatched_gemm_2.cpp b/test/verify/test_unbatched_gemm_2.cpp index 1b9640a694b..eaf188f48f7 100644 --- a/test/verify/test_unbatched_gemm_2.cpp +++ b/test/verify/test_unbatched_gemm_2.cpp @@ -50,6 +50,7 @@ struct test_unbatched_gemm_2 : verify_program> template struct test_unbatched_gemm_2; template struct test_unbatched_gemm_2; +template struct test_unbatched_gemm_2; template struct test_unbatched_gemm_2; template struct test_unbatched_gemm_2; template struct test_unbatched_gemm_2; diff --git a/test/verify/test_where.cpp b/test/verify/test_where.cpp index 86c8c2cc942..feac464efd7 100644 --- a/test/verify/test_where.cpp +++ b/test/verify/test_where.cpp @@ -48,6 +48,7 @@ struct test_where : verify_program> template struct test_where; template struct test_where; +template struct test_where; template struct test_where; template struct test_where; template struct test_where; From e2024fa65be118f90d63b785994e5dd243ea0ed6 Mon Sep 17 00:00:00 2001 From: richagadgil Date: Wed, 11 Dec 2024 19:44:32 -0600 Subject: [PATCH 2/3] align tests with dev branch --- test/verify/CMakeLists.txt | 2 +- test/verify/test_arg_ops.cpp | 2 ++ test/verify/test_block_reduce_small.cpp | 2 ++ test/verify/test_gemm.cpp | 2 +- test/verify/test_gemm_2args_bmv.cpp | 2 +- test/verify/test_gemm_2args_mm_1.cpp | 2 +- test/verify/test_gemm_2args_mm_2.cpp | 2 +- test/verify/test_gemm_2args_mm_3.cpp | 2 +- test/verify/test_gemm_2args_mm_5.cpp | 2 +- test/verify/test_gemm_2args_mm_6.cpp | 2 +- test/verify/test_gemm_2args_mm_7.cpp | 2 +- test/verify/test_gemm_2args_mm_8.cpp | 2 +- test/verify/test_gemm_2args_mv.cpp | 2 +- test/verify/test_gemm_2args_vbm.cpp | 2 +- test/verify/test_gemm_2args_vm.cpp | 2 +- test/verify/test_gemm_2args_vv.cpp | 2 +- test/verify/test_gemm_add_broadcast1.cpp | 2 +- test/verify/test_gemm_add_broadcast2.cpp | 2 +- test/verify/test_gemm_copy.cpp | 2 +- test/verify/test_gemm_ex.cpp | 2 +- test/verify/test_gemm_literal.cpp | 2 +- test/verify/test_gemm_multi_3args.cpp | 2 +- test/verify/test_gemm_multi_3args_alpha0.cpp | 2 +- test/verify/test_gemm_multi_3args_beta0.cpp | 2 +- test/verify/test_gemm_multi_3args_c25.cpp | 2 +- test/verify/test_gemm_multi_dim_2_3.cpp | 2 +- test/verify/test_gemm_multi_transpose.cpp | 2 +- test/verify/test_gemm_multibroadcast.cpp | 2 +- test/verify/test_gemm_pointwise.cpp | 2 +- .../test_gemm_transpose_add_pooling_sub.cpp | 2 +- test/verify/test_gemm_transposea.cpp | 2 +- test/verify/test_gemm_transposea_ex.cpp | 2 +- test/verify/test_gemm_transposeab.cpp | 2 +- test/verify/test_gemm_transposeb.cpp | 2 +- test/verify/test_gemm_transposeb_detect.cpp | 2 +- test/verify/test_gemm_transposeb_ex.cpp | 2 +- test/verify/test_layernorm.cpp | 22 +++++++++++++++++++ test/verify/test_logsoftmax.cpp | 2 ++ test/verify/test_mul_dot_a.cpp | 2 +- test/verify/test_mul_dot_b.cpp | 2 +- .../test_pointwise_broadcast_reduce.cpp | 2 ++ .../test_pointwise_broadcast_reduce_bf16.cpp | 2 ++ test/verify/test_reduce_add.cpp | 2 ++ test/verify/test_reduce_add_bf16.cpp | 2 ++ test/verify/test_reduce_mean_bias_bf16.cpp | 2 ++ test/verify/test_reduce_mean_bias_half.cpp | 2 ++ test/verify/test_reduce_mean_large_bf16.cpp | 2 ++ test/verify/test_reduce_mean_large_half.cpp | 2 ++ test/verify/test_reduce_mean_nhwc.cpp | 2 ++ test/verify/test_reduce_mean_reduce_sum.cpp | 2 ++ test/verify/test_reduce_mean_variance.cpp | 2 ++ .../verify/test_reduce_mean_variance_bf16.cpp | 2 ++ test/verify/test_reduce_noop_add.cpp | 2 ++ test/verify/test_reduce_noop_add_bf16.cpp | 2 ++ test/verify/test_reduce_op_large.cpp | 2 ++ test/verify/test_reduce_op_small.cpp | 2 ++ test/verify/test_rnn_sql_1.cpp | 2 +- test/verify/test_select_module_reduce.cpp | 2 ++ test/verify/test_softmax.cpp | 2 ++ test/verify/test_softmax1.cpp | 2 ++ test/verify/test_softmax2.cpp | 2 ++ test/verify/test_softmax3.cpp | 2 ++ test/verify/test_softmax4.cpp | 2 ++ test/verify/test_softmax5.cpp | 2 ++ test/verify/test_softmax_large1.cpp | 2 ++ test/verify/test_softmax_large2.cpp | 2 ++ test/verify/test_softmax_large3.cpp | 2 ++ .../test_softmaxcrossentropyloss_2d.cpp | 2 ++ .../test_softmaxcrossentropyloss_2d_mean.cpp | 2 ++ .../test_softmaxcrossentropyloss_2d_sum.cpp | 2 ++ .../test_softmaxcrossentropyloss_kd.cpp | 2 ++ .../test_softmaxcrossentropyloss_kd_mean.cpp | 2 ++ .../test_softmaxcrossentropyloss_kd_sum.cpp | 2 ++ 73 files changed, 129 insertions(+), 37 deletions(-) diff --git a/test/verify/CMakeLists.txt b/test/verify/CMakeLists.txt index 3e45e9bc3a9..bd57abea883 100644 --- a/test/verify/CMakeLists.txt +++ b/test/verify/CMakeLists.txt @@ -31,7 +31,7 @@ target_link_libraries(test_verify migraphx migraphx_all_targets) target_include_directories(test_verify PUBLIC ../include) rocm_clang_tidy_check(test_verify) -foreach(SECTION general rnn conv gemm) +foreach(SECTION general reduce rnn conv gemm) rocm_add_test(NAME test_verify_${SECTION} COMMAND test_verify ${SECTION}) set_tests_properties(test_verify_${SECTION} PROPERTIES COST 100 diff --git a/test/verify/test_arg_ops.cpp b/test/verify/test_arg_ops.cpp index ac52da56a23..674e8cb8969 100644 --- a/test/verify/test_arg_ops.cpp +++ b/test/verify/test_arg_ops.cpp @@ -57,6 +57,8 @@ struct test_arg_ops : verify_programadd_instruction(T{Axis, LastIndex}, param); return p; } + + std::string section() const { return "reduce"; } }; // transpose argmax tests template struct test_arg_ops; diff --git a/test/verify/test_block_reduce_small.cpp b/test/verify/test_block_reduce_small.cpp index 2345591b682..847e59bc1c3 100644 --- a/test/verify/test_block_reduce_small.cpp +++ b/test/verify/test_block_reduce_small.cpp @@ -46,6 +46,8 @@ struct test_block_reduce_small : verify_program> mm->add_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; template diff --git a/test/verify/test_gemm.cpp b/test/verify/test_gemm.cpp index 5952a82ea23..6a1b1bb1a58 100644 --- a/test/verify/test_gemm.cpp +++ b/test/verify/test_gemm.cpp @@ -46,6 +46,6 @@ template struct test_gemm; template struct test_gemm; template struct test_gemm; template struct test_gemm; -// template struct test_gemm; +template struct test_gemm; template struct test_gemm; template struct test_gemm; diff --git a/test/verify/test_gemm_2args_bmv.cpp b/test/verify/test_gemm_2args_bmv.cpp index 94c65c318d4..333c116adc0 100644 --- a/test/verify/test_gemm_2args_bmv.cpp +++ b/test/verify/test_gemm_2args_bmv.cpp @@ -53,6 +53,6 @@ template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; -// template struct test_gemm_2args_bmv; +template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; template struct test_gemm_2args_bmv; diff --git a/test/verify/test_gemm_2args_mm_1.cpp b/test/verify/test_gemm_2args_mm_1.cpp index f029a21a4d3..fdc21b40872 100644 --- a/test/verify/test_gemm_2args_mm_1.cpp +++ b/test/verify/test_gemm_2args_mm_1.cpp @@ -52,6 +52,6 @@ template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; -// template struct test_gemm_2args_mm_1; +template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; template struct test_gemm_2args_mm_1; diff --git a/test/verify/test_gemm_2args_mm_2.cpp b/test/verify/test_gemm_2args_mm_2.cpp index 9efc07fca3d..8094acdb1bc 100644 --- a/test/verify/test_gemm_2args_mm_2.cpp +++ b/test/verify/test_gemm_2args_mm_2.cpp @@ -53,6 +53,6 @@ template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; -// template struct test_gemm_2args_mm_2; +template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; template struct test_gemm_2args_mm_2; diff --git a/test/verify/test_gemm_2args_mm_3.cpp b/test/verify/test_gemm_2args_mm_3.cpp index d63a17b8e40..646ac56cf6a 100644 --- a/test/verify/test_gemm_2args_mm_3.cpp +++ b/test/verify/test_gemm_2args_mm_3.cpp @@ -53,6 +53,6 @@ template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; -// template struct test_gemm_2args_mm_3; +template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; template struct test_gemm_2args_mm_3; diff --git a/test/verify/test_gemm_2args_mm_5.cpp b/test/verify/test_gemm_2args_mm_5.cpp index eda7518e239..d45ccb5d1be 100644 --- a/test/verify/test_gemm_2args_mm_5.cpp +++ b/test/verify/test_gemm_2args_mm_5.cpp @@ -52,6 +52,6 @@ template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; -// template struct test_gemm_2args_mm_5; +template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; template struct test_gemm_2args_mm_5; diff --git a/test/verify/test_gemm_2args_mm_6.cpp b/test/verify/test_gemm_2args_mm_6.cpp index 100235feea9..f6a562eaccf 100644 --- a/test/verify/test_gemm_2args_mm_6.cpp +++ b/test/verify/test_gemm_2args_mm_6.cpp @@ -55,6 +55,6 @@ template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; -// template struct test_gemm_2args_mm_6; +template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; template struct test_gemm_2args_mm_6; diff --git a/test/verify/test_gemm_2args_mm_7.cpp b/test/verify/test_gemm_2args_mm_7.cpp index 7c78e7be1d5..7889e62e42d 100644 --- a/test/verify/test_gemm_2args_mm_7.cpp +++ b/test/verify/test_gemm_2args_mm_7.cpp @@ -52,6 +52,6 @@ template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; -// template struct test_gemm_2args_mm_7; +template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; template struct test_gemm_2args_mm_7; diff --git a/test/verify/test_gemm_2args_mm_8.cpp b/test/verify/test_gemm_2args_mm_8.cpp index bb0256e85ef..175105d1003 100644 --- a/test/verify/test_gemm_2args_mm_8.cpp +++ b/test/verify/test_gemm_2args_mm_8.cpp @@ -51,6 +51,6 @@ struct test_gemm_2args_mm_8 : verify_program> template struct test_gemm_2args_mm_8; // template struct test_gemm_2args_mm_8; // fails with CK, issue#2514 template struct test_gemm_2args_mm_8; -// template struct test_gemm_2args_mm_8; +template struct test_gemm_2args_mm_8; template struct test_gemm_2args_mm_8; template struct test_gemm_2args_mm_8; diff --git a/test/verify/test_gemm_2args_mv.cpp b/test/verify/test_gemm_2args_mv.cpp index 070f9157870..00282367e0e 100644 --- a/test/verify/test_gemm_2args_mv.cpp +++ b/test/verify/test_gemm_2args_mv.cpp @@ -51,6 +51,6 @@ template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; -// template struct test_gemm_2args_mv; +template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; template struct test_gemm_2args_mv; diff --git a/test/verify/test_gemm_2args_vbm.cpp b/test/verify/test_gemm_2args_vbm.cpp index 20a49ef149d..3063311811f 100644 --- a/test/verify/test_gemm_2args_vbm.cpp +++ b/test/verify/test_gemm_2args_vbm.cpp @@ -55,6 +55,6 @@ template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; -// template struct test_gemm_2args_vbm; +template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; template struct test_gemm_2args_vbm; diff --git a/test/verify/test_gemm_2args_vm.cpp b/test/verify/test_gemm_2args_vm.cpp index 76712f85403..c3314faa615 100644 --- a/test/verify/test_gemm_2args_vm.cpp +++ b/test/verify/test_gemm_2args_vm.cpp @@ -52,7 +52,7 @@ template struct test_gemm_2args_vm; template struct test_gemm_2args_vm; template struct test_gemm_2args_vm; template struct test_gemm_2args_vm; -// template struct test_gemm_2args_vm; +template struct test_gemm_2args_vm; // TODO need hipblaslt support // template struct test_gemm_2args_vm; // template struct test_gemm_2args_vm; diff --git a/test/verify/test_gemm_2args_vv.cpp b/test/verify/test_gemm_2args_vv.cpp index 4550f801527..d09b6d7695a 100644 --- a/test/verify/test_gemm_2args_vv.cpp +++ b/test/verify/test_gemm_2args_vv.cpp @@ -55,6 +55,6 @@ template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; -// template struct test_gemm_2args_vv; +template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; template struct test_gemm_2args_vv; diff --git a/test/verify/test_gemm_add_broadcast1.cpp b/test/verify/test_gemm_add_broadcast1.cpp index a1ce9ac9e6f..33f22780528 100644 --- a/test/verify/test_gemm_add_broadcast1.cpp +++ b/test/verify/test_gemm_add_broadcast1.cpp @@ -55,6 +55,6 @@ template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; -// template struct test_gemm_add_broadcast1; +template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; template struct test_gemm_add_broadcast1; diff --git a/test/verify/test_gemm_add_broadcast2.cpp b/test/verify/test_gemm_add_broadcast2.cpp index 7e086f423fd..88674901d0f 100644 --- a/test/verify/test_gemm_add_broadcast2.cpp +++ b/test/verify/test_gemm_add_broadcast2.cpp @@ -55,6 +55,6 @@ template struct test_gemm_add_broadcast2; // template struct test_gemm_add_broadcast2; // fails with CK, // issue#2514 template struct test_gemm_add_broadcast2; -// template struct test_gemm_add_broadcast2; +template struct test_gemm_add_broadcast2; template struct test_gemm_add_broadcast2; template struct test_gemm_add_broadcast2; diff --git a/test/verify/test_gemm_copy.cpp b/test/verify/test_gemm_copy.cpp index f868cb2f1cb..144fd5238ae 100644 --- a/test/verify/test_gemm_copy.cpp +++ b/test/verify/test_gemm_copy.cpp @@ -53,6 +53,6 @@ template struct test_gemm_copy; template struct test_gemm_copy; template struct test_gemm_copy; template struct test_gemm_copy; -// template struct test_gemm_copy; +template struct test_gemm_copy; template struct test_gemm_copy; template struct test_gemm_copy; diff --git a/test/verify/test_gemm_ex.cpp b/test/verify/test_gemm_ex.cpp index 86834de54a6..2a4546d0b40 100644 --- a/test/verify/test_gemm_ex.cpp +++ b/test/verify/test_gemm_ex.cpp @@ -45,6 +45,6 @@ template struct test_gemm_ex; template struct test_gemm_ex; template struct test_gemm_ex; template struct test_gemm_ex; -// template struct test_gemm_ex; +template struct test_gemm_ex; template struct test_gemm_ex; template struct test_gemm_ex; diff --git a/test/verify/test_gemm_literal.cpp b/test/verify/test_gemm_literal.cpp index 02532ce01df..e57085d868a 100644 --- a/test/verify/test_gemm_literal.cpp +++ b/test/verify/test_gemm_literal.cpp @@ -50,6 +50,6 @@ template struct test_gemm_literal; template struct test_gemm_literal; template struct test_gemm_literal; template struct test_gemm_literal; -// template struct test_gemm_literal; +template struct test_gemm_literal; template struct test_gemm_literal; template struct test_gemm_literal; diff --git a/test/verify/test_gemm_multi_3args.cpp b/test/verify/test_gemm_multi_3args.cpp index 99828d1f3f4..36b610ead21 100644 --- a/test/verify/test_gemm_multi_3args.cpp +++ b/test/verify/test_gemm_multi_3args.cpp @@ -54,6 +54,6 @@ template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; -// template struct test_gemm_multi_3args; +template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; template struct test_gemm_multi_3args; diff --git a/test/verify/test_gemm_multi_3args_alpha0.cpp b/test/verify/test_gemm_multi_3args_alpha0.cpp index 69bb6799601..b28c6191d94 100644 --- a/test/verify/test_gemm_multi_3args_alpha0.cpp +++ b/test/verify/test_gemm_multi_3args_alpha0.cpp @@ -54,6 +54,6 @@ template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; -// template struct test_gemm_multi_3args_alpha0; +template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; template struct test_gemm_multi_3args_alpha0; diff --git a/test/verify/test_gemm_multi_3args_beta0.cpp b/test/verify/test_gemm_multi_3args_beta0.cpp index 3db08c23c76..5d31154c0b7 100644 --- a/test/verify/test_gemm_multi_3args_beta0.cpp +++ b/test/verify/test_gemm_multi_3args_beta0.cpp @@ -54,6 +54,6 @@ template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; -// template struct test_gemm_multi_3args_beta0; +template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; template struct test_gemm_multi_3args_beta0; diff --git a/test/verify/test_gemm_multi_3args_c25.cpp b/test/verify/test_gemm_multi_3args_c25.cpp index 8d2b2a28d72..6fe280b3b2a 100644 --- a/test/verify/test_gemm_multi_3args_c25.cpp +++ b/test/verify/test_gemm_multi_3args_c25.cpp @@ -54,6 +54,6 @@ template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; -// template struct test_gemm_multi_3args_c25; +template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; template struct test_gemm_multi_3args_c25; diff --git a/test/verify/test_gemm_multi_dim_2_3.cpp b/test/verify/test_gemm_multi_dim_2_3.cpp index e59cdcef648..43f3ce2753c 100644 --- a/test/verify/test_gemm_multi_dim_2_3.cpp +++ b/test/verify/test_gemm_multi_dim_2_3.cpp @@ -50,6 +50,6 @@ template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; -// template struct test_gemm_multi_dim_2_3; +template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; template struct test_gemm_multi_dim_2_3; diff --git a/test/verify/test_gemm_multi_transpose.cpp b/test/verify/test_gemm_multi_transpose.cpp index fb3cf22e477..d974e13e8c7 100644 --- a/test/verify/test_gemm_multi_transpose.cpp +++ b/test/verify/test_gemm_multi_transpose.cpp @@ -54,6 +54,6 @@ template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; -// template struct test_gemm_multi_transpose; +template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; template struct test_gemm_multi_transpose; diff --git a/test/verify/test_gemm_multibroadcast.cpp b/test/verify/test_gemm_multibroadcast.cpp index db9899156c2..46074adcfd7 100644 --- a/test/verify/test_gemm_multibroadcast.cpp +++ b/test/verify/test_gemm_multibroadcast.cpp @@ -48,6 +48,6 @@ template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; -// template struct test_gemm_multibroadcast; +template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; template struct test_gemm_multibroadcast; diff --git a/test/verify/test_gemm_pointwise.cpp b/test/verify/test_gemm_pointwise.cpp index c35cbc2e8fe..857ea949f1b 100644 --- a/test/verify/test_gemm_pointwise.cpp +++ b/test/verify/test_gemm_pointwise.cpp @@ -57,6 +57,6 @@ template struct test_gemm_pointwise; template struct test_gemm_pointwise; template struct test_gemm_pointwise; template struct test_gemm_pointwise; -// template struct test_gemm_pointwise; +template struct test_gemm_pointwise; template struct test_gemm_pointwise; template struct test_gemm_pointwise; diff --git a/test/verify/test_gemm_transpose_add_pooling_sub.cpp b/test/verify/test_gemm_transpose_add_pooling_sub.cpp index 5c234df7b3f..9f06598723b 100644 --- a/test/verify/test_gemm_transpose_add_pooling_sub.cpp +++ b/test/verify/test_gemm_transpose_add_pooling_sub.cpp @@ -67,6 +67,6 @@ template struct test_gemm_transpose_add_pooling_sub template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; -// template struct test_gemm_transpose_add_pooling_sub; +template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; template struct test_gemm_transpose_add_pooling_sub; diff --git a/test/verify/test_gemm_transposea.cpp b/test/verify/test_gemm_transposea.cpp index 86686408c1f..cc4306ba681 100644 --- a/test/verify/test_gemm_transposea.cpp +++ b/test/verify/test_gemm_transposea.cpp @@ -47,7 +47,7 @@ template struct test_gemm_transposea; template struct test_gemm_transposea; template struct test_gemm_transposea; template struct test_gemm_transposea; -// template struct test_gemm_transposea; +template struct test_gemm_transposea; // TODO need hipblaslt support // template struct test_gemm_transposea; // template struct test_gemm_transposea; diff --git a/test/verify/test_gemm_transposea_ex.cpp b/test/verify/test_gemm_transposea_ex.cpp index c5004183a12..a515e499e93 100644 --- a/test/verify/test_gemm_transposea_ex.cpp +++ b/test/verify/test_gemm_transposea_ex.cpp @@ -48,6 +48,6 @@ template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; -// template struct test_gemm_transposea_ex; +template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; template struct test_gemm_transposea_ex; diff --git a/test/verify/test_gemm_transposeab.cpp b/test/verify/test_gemm_transposeab.cpp index 0ce2c4db0fc..54cf150b934 100644 --- a/test/verify/test_gemm_transposeab.cpp +++ b/test/verify/test_gemm_transposeab.cpp @@ -48,6 +48,6 @@ template struct test_gemm_transposeab; template struct test_gemm_transposeab; template struct test_gemm_transposeab; template struct test_gemm_transposeab; -// template struct test_gemm_transposeab; +template struct test_gemm_transposeab; template struct test_gemm_transposeab; template struct test_gemm_transposeab; diff --git a/test/verify/test_gemm_transposeb.cpp b/test/verify/test_gemm_transposeb.cpp index 1e58deca006..eb214ce0129 100644 --- a/test/verify/test_gemm_transposeb.cpp +++ b/test/verify/test_gemm_transposeb.cpp @@ -47,6 +47,6 @@ template struct test_gemm_transposeb; template struct test_gemm_transposeb; template struct test_gemm_transposeb; template struct test_gemm_transposeb; -// template struct test_gemm_transposeb; +template struct test_gemm_transposeb; template struct test_gemm_transposeb; template struct test_gemm_transposeb; diff --git a/test/verify/test_gemm_transposeb_detect.cpp b/test/verify/test_gemm_transposeb_detect.cpp index bd8a987c35e..f1e34d25645 100644 --- a/test/verify/test_gemm_transposeb_detect.cpp +++ b/test/verify/test_gemm_transposeb_detect.cpp @@ -48,6 +48,6 @@ template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; -// template struct test_gemm_transposeb_detect; +template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; template struct test_gemm_transposeb_detect; diff --git a/test/verify/test_gemm_transposeb_ex.cpp b/test/verify/test_gemm_transposeb_ex.cpp index 1408c56c29b..89cffaee6a3 100644 --- a/test/verify/test_gemm_transposeb_ex.cpp +++ b/test/verify/test_gemm_transposeb_ex.cpp @@ -48,6 +48,6 @@ template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; -// template struct test_gemm_transposeb_ex; +template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; template struct test_gemm_transposeb_ex; diff --git a/test/verify/test_layernorm.cpp b/test/verify/test_layernorm.cpp index 013970d2c89..653a113f2ee 100644 --- a/test/verify/test_layernorm.cpp +++ b/test/verify/test_layernorm.cpp @@ -40,6 +40,8 @@ struct test_layernorm : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm2 : verify_program @@ -53,6 +55,8 @@ struct test_layernorm2 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_large : verify_program @@ -66,6 +70,8 @@ struct test_layernorm_large : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_fp16 : verify_program @@ -79,6 +85,8 @@ struct test_layernorm_fp16 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_bf16 : verify_program @@ -92,6 +100,8 @@ struct test_layernorm_bf16 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_1 : verify_program @@ -105,6 +115,8 @@ struct test_layernorm_fp8_1 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_2 : verify_program @@ -131,6 +143,8 @@ struct test_layernorm_fp8_3 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_4 : verify_program @@ -157,6 +171,8 @@ struct test_layernorm_eps : verify_program add_layernorm(*mm, x, dims, 1e-5f); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_triadd : verify_program @@ -174,6 +190,8 @@ struct test_layernorm_triadd : verify_program add_layernorm(*mm, add2, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_triadd_large : verify_program @@ -191,6 +209,8 @@ struct test_layernorm_triadd_large : verify_program add_layernorm(*mm, add2, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_add_layernorm_add_gemm_nonstd : verify_program @@ -223,4 +243,6 @@ struct test_pw_layernorm : verify_program add_pointwise_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_logsoftmax.cpp b/test/verify/test_logsoftmax.cpp index ff9a7bc456b..35ab9e88b8b 100644 --- a/test/verify/test_logsoftmax.cpp +++ b/test/verify/test_logsoftmax.cpp @@ -40,6 +40,8 @@ struct test_logsoftmax : verify_program> return p; } + + std::string section() const { return "reduce"; } }; template struct test_logsoftmax<0, migraphx::shape::float_type>; diff --git a/test/verify/test_mul_dot_a.cpp b/test/verify/test_mul_dot_a.cpp index e5854937734..402ca1611ba 100644 --- a/test/verify/test_mul_dot_a.cpp +++ b/test/verify/test_mul_dot_a.cpp @@ -53,6 +53,6 @@ template struct test_mul_dot_a; template struct test_mul_dot_a; template struct test_mul_dot_a; template struct test_mul_dot_a; -// template struct test_mul_dot_a; +template struct test_mul_dot_a; template struct test_mul_dot_a; template struct test_mul_dot_a; diff --git a/test/verify/test_mul_dot_b.cpp b/test/verify/test_mul_dot_b.cpp index 07d0781b755..442004671ec 100644 --- a/test/verify/test_mul_dot_b.cpp +++ b/test/verify/test_mul_dot_b.cpp @@ -54,6 +54,6 @@ template struct test_mul_dot_b; template struct test_mul_dot_b; template struct test_mul_dot_b; template struct test_mul_dot_b; -// template struct test_mul_dot_b; +template struct test_mul_dot_b; template struct test_mul_dot_b; template struct test_mul_dot_b; diff --git a/test/verify/test_pointwise_broadcast_reduce.cpp b/test/verify/test_pointwise_broadcast_reduce.cpp index ffcc658838b..c0269625457 100644 --- a/test/verify/test_pointwise_broadcast_reduce.cpp +++ b/test/verify/test_pointwise_broadcast_reduce.cpp @@ -52,4 +52,6 @@ struct test_pointwise_broadcast_reduce : verify_programadd_return({reshape}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_pointwise_broadcast_reduce_bf16.cpp b/test/verify/test_pointwise_broadcast_reduce_bf16.cpp index 65300bd6189..049d3efedfa 100644 --- a/test/verify/test_pointwise_broadcast_reduce_bf16.cpp +++ b/test/verify/test_pointwise_broadcast_reduce_bf16.cpp @@ -52,4 +52,6 @@ struct test_pointwise_broadcast_reduce_bf16 : verify_programadd_return({reshape}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_add.cpp b/test/verify/test_reduce_add.cpp index c0e2b5ddf93..f1d8d6f3fd3 100644 --- a/test/verify/test_reduce_add.cpp +++ b/test/verify/test_reduce_add.cpp @@ -47,6 +47,8 @@ struct test_reduce_add : verify_program> mm->add_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; template struct test_reduce_add; diff --git a/test/verify/test_reduce_add_bf16.cpp b/test/verify/test_reduce_add_bf16.cpp index aaa91be5bcc..2a3b3ebdbab 100644 --- a/test/verify/test_reduce_add_bf16.cpp +++ b/test/verify/test_reduce_add_bf16.cpp @@ -47,6 +47,8 @@ struct test_reduce_add_bf16 : verify_program> mm->add_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; template struct test_reduce_add_bf16; diff --git a/test/verify/test_reduce_mean_bias_bf16.cpp b/test/verify/test_reduce_mean_bias_bf16.cpp index df5839966ec..3c9b5f40a35 100644 --- a/test/verify/test_reduce_mean_bias_bf16.cpp +++ b/test/verify/test_reduce_mean_bias_bf16.cpp @@ -45,4 +45,6 @@ struct test_reduce_mean_bias_bf16 : verify_program mm->add_return({sqrt}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_bias_half.cpp b/test/verify/test_reduce_mean_bias_half.cpp index c83d7565de7..d74757873df 100644 --- a/test/verify/test_reduce_mean_bias_half.cpp +++ b/test/verify/test_reduce_mean_bias_half.cpp @@ -45,4 +45,6 @@ struct test_reduce_mean_bias_half : verify_program mm->add_return({sqrt}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_large_bf16.cpp b/test/verify/test_reduce_mean_large_bf16.cpp index 927b1827838..42ce24b6357 100644 --- a/test/verify/test_reduce_mean_large_bf16.cpp +++ b/test/verify/test_reduce_mean_large_bf16.cpp @@ -38,4 +38,6 @@ struct test_reduce_mean_large_bf16 : verify_program mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_large_half.cpp b/test/verify/test_reduce_mean_large_half.cpp index f8925dbe577..1a88bc57abb 100644 --- a/test/verify/test_reduce_mean_large_half.cpp +++ b/test/verify/test_reduce_mean_large_half.cpp @@ -38,4 +38,6 @@ struct test_reduce_mean_large_half : verify_program mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {2}}}), x); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_nhwc.cpp b/test/verify/test_reduce_mean_nhwc.cpp index 0307af20988..9327df71d92 100644 --- a/test/verify/test_reduce_mean_nhwc.cpp +++ b/test/verify/test_reduce_mean_nhwc.cpp @@ -43,6 +43,8 @@ struct test_reduce_mean_nhwc : verify_program> mm->add_return({sqrt}); return p; }; + + std::string section() const { return "reduce"; } }; template struct test_reduce_mean_nhwc; diff --git a/test/verify/test_reduce_mean_reduce_sum.cpp b/test/verify/test_reduce_mean_reduce_sum.cpp index 9dd29f8cffe..b5df0f0147a 100644 --- a/test/verify/test_reduce_mean_reduce_sum.cpp +++ b/test/verify/test_reduce_mean_reduce_sum.cpp @@ -51,4 +51,6 @@ struct test_reduce_mean_reduce_sum : verify_program mm->add_return({mean_div2}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_variance.cpp b/test/verify/test_reduce_mean_variance.cpp index dccd74a84c3..6843ae7e459 100644 --- a/test/verify/test_reduce_mean_variance.cpp +++ b/test/verify/test_reduce_mean_variance.cpp @@ -45,4 +45,6 @@ struct test_reduce_mean_variance : verify_program mm->add_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_mean_variance_bf16.cpp b/test/verify/test_reduce_mean_variance_bf16.cpp index a0694a4b5a0..1e65fdc077e 100644 --- a/test/verify/test_reduce_mean_variance_bf16.cpp +++ b/test/verify/test_reduce_mean_variance_bf16.cpp @@ -45,4 +45,6 @@ struct test_reduce_mean_variance_bf16 : verify_programadd_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_noop_add.cpp b/test/verify/test_reduce_noop_add.cpp index 9e642da2f71..48ce85c535b 100644 --- a/test/verify/test_reduce_noop_add.cpp +++ b/test/verify/test_reduce_noop_add.cpp @@ -45,4 +45,6 @@ struct test_reduce_noop_add : verify_program mm->add_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_noop_add_bf16.cpp b/test/verify/test_reduce_noop_add_bf16.cpp index 4ee3692e89b..f435fc6f62f 100644 --- a/test/verify/test_reduce_noop_add_bf16.cpp +++ b/test/verify/test_reduce_noop_add_bf16.cpp @@ -45,4 +45,6 @@ struct test_reduce_noop_add_bf16 : verify_program mm->add_return({add}); return p; }; + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_reduce_op_large.cpp b/test/verify/test_reduce_op_large.cpp index 473c762d23a..edcd550ef66 100644 --- a/test/verify/test_reduce_op_large.cpp +++ b/test/verify/test_reduce_op_large.cpp @@ -45,6 +45,8 @@ struct test_reduce_op_large : verify_program> mm->add_instruction(Op{{Axis}}, x); return p; }; + + std::string section() const { return "reduce"; } }; template struct test_reduce_op_large; diff --git a/test/verify/test_reduce_op_small.cpp b/test/verify/test_reduce_op_small.cpp index 81298039f77..b89dcd0f080 100644 --- a/test/verify/test_reduce_op_small.cpp +++ b/test/verify/test_reduce_op_small.cpp @@ -45,6 +45,8 @@ struct test_reduce_op_small : verify_program> mm->add_instruction(Op{{Axis}}, x); return p; }; + + std::string section() const { return "reduce"; } }; template struct test_reduce_op_small; diff --git a/test/verify/test_rnn_sql_1.cpp b/test/verify/test_rnn_sql_1.cpp index c9d72fd99a8..9e4c7f519c8 100644 --- a/test/verify/test_rnn_sql_1.cpp +++ b/test/verify/test_rnn_sql_1.cpp @@ -86,6 +86,6 @@ template struct test_rnn_sql_1; template struct test_rnn_sql_1; template struct test_rnn_sql_1; template struct test_rnn_sql_1; -// template struct test_rnn_sql_1; +template struct test_rnn_sql_1; template struct test_rnn_sql_1; template struct test_rnn_sql_1; diff --git a/test/verify/test_select_module_reduce.cpp b/test/verify/test_select_module_reduce.cpp index aeffd90417e..33c26b56a9d 100644 --- a/test/verify/test_select_module_reduce.cpp +++ b/test/verify/test_select_module_reduce.cpp @@ -66,4 +66,6 @@ struct test_select_module_reduce : verify_program return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax.cpp b/test/verify/test_softmax.cpp index b6e27771756..7dc6447744c 100644 --- a/test/verify/test_softmax.cpp +++ b/test/verify/test_softmax.cpp @@ -40,6 +40,8 @@ struct test_softmax : verify_program> return p; } + + std::string section() const { return "reduce"; } }; template struct test_softmax<0, migraphx::shape::float_type>; diff --git a/test/verify/test_softmax1.cpp b/test/verify/test_softmax1.cpp index 4742e03607a..2f95fb5392e 100644 --- a/test/verify/test_softmax1.cpp +++ b/test/verify/test_softmax1.cpp @@ -37,4 +37,6 @@ struct test_softmax1 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", 0}}), x); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax2.cpp b/test/verify/test_softmax2.cpp index 804f8f04f0b..222433cd759 100644 --- a/test/verify/test_softmax2.cpp +++ b/test/verify/test_softmax2.cpp @@ -38,4 +38,6 @@ struct test_softmax2 : verify_program mm->add_instruction(migraphx::make_op("softmax"), x); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax3.cpp b/test/verify/test_softmax3.cpp index 027c4d64ad3..061847a2f5b 100644 --- a/test/verify/test_softmax3.cpp +++ b/test/verify/test_softmax3.cpp @@ -41,4 +41,6 @@ struct test_softmax3 : verify_program mm->add_return({r}); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax4.cpp b/test/verify/test_softmax4.cpp index 73117df57c5..6ceeb5009f4 100644 --- a/test/verify/test_softmax4.cpp +++ b/test/verify/test_softmax4.cpp @@ -38,4 +38,6 @@ struct test_softmax4 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", 3}}), x); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax5.cpp b/test/verify/test_softmax5.cpp index 499cc497931..18c767d60f3 100644 --- a/test/verify/test_softmax5.cpp +++ b/test/verify/test_softmax5.cpp @@ -38,4 +38,6 @@ struct test_softmax5 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", 3}}), x); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax_large1.cpp b/test/verify/test_softmax_large1.cpp index e8d98194da8..349d3685219 100644 --- a/test/verify/test_softmax_large1.cpp +++ b/test/verify/test_softmax_large1.cpp @@ -40,4 +40,6 @@ struct test_softmax_large1 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), add); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax_large2.cpp b/test/verify/test_softmax_large2.cpp index c4311f6f819..5a8c88ea1da 100644 --- a/test/verify/test_softmax_large2.cpp +++ b/test/verify/test_softmax_large2.cpp @@ -40,4 +40,6 @@ struct test_softmax_large2 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), add); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmax_large3.cpp b/test/verify/test_softmax_large3.cpp index d91bece8048..ed6f9b200a2 100644 --- a/test/verify/test_softmax_large3.cpp +++ b/test/verify/test_softmax_large3.cpp @@ -40,4 +40,6 @@ struct test_softmax_large3 : verify_program mm->add_instruction(migraphx::make_op("softmax", {{"axis", -1}}), add); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_softmaxcrossentropyloss_2d.cpp b/test/verify/test_softmaxcrossentropyloss_2d.cpp index dd92f9cf73c..ae8cd76519c 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d.cpp @@ -81,6 +81,8 @@ struct test_softmaxcrossentropyloss_2d return p; } + + std::string section() const { return "reduce"; } }; // template struct test_softmaxcrossentropyloss_2d Date: Wed, 11 Dec 2024 20:00:34 -0600 Subject: [PATCH 3/3] cleanup --- test/verify/test_add_bf16.cpp | 2 +- test/verify/test_add_gelu_bf16.cpp | 2 +- test/verify/test_block_reduce_small.cpp | 2 +- test/verify/test_fp32_bf16_add.cpp | 2 +- test/verify/test_fp32_bf16_ladd.cpp | 2 +- test/verify/test_fp32_bf16_lall.cpp | 2 +- test/verify/test_fp32_bf16_sub.cpp | 2 +- test/verify/test_hsqrt_bf16.cpp | 2 +- test/verify/test_instancenorm.cpp | 6 +++++- test/verify/test_layernorm.cpp | 6 ++++-- test/verify/test_logsoftmax1.cpp | 4 +++- test/verify/test_pad_highest_bf16.cpp | 2 +- test/verify/test_pad_lowest_bf16.cpp | 2 +- test/verify/test_reduce_mean_bias_bf16.cpp | 2 +- test/verify/test_reduce_mean_bias_half.cpp | 2 +- test/verify/test_reduce_mean_large_bf16.cpp | 2 +- test/verify/test_reduce_mean_large_half.cpp | 2 +- test/verify/test_reduce_noop_add.cpp | 2 +- test/verify/test_reduce_noop_add_bf16.cpp | 2 +- test/verify/test_select_module_reduce.cpp | 2 +- test/verify/test_sin_bf16.cpp | 2 +- test/verify/test_softmax1.cpp | 2 +- test/verify/test_softmax2.cpp | 2 +- test/verify/test_softmax3.cpp | 2 +- test/verify/test_softmax4.cpp | 2 +- test/verify/test_softmax5.cpp | 2 +- test/verify/test_softmax_large1.cpp | 2 +- test/verify/test_softmax_large2.cpp | 2 +- test/verify/test_softmax_large3.cpp | 2 +- test/verify/test_sqrt_bf161.cpp | 2 +- test/verify/test_sqrt_bf162.cpp | 2 +- test/verify/test_sqrt_bf164.cpp | 2 +- 32 files changed, 41 insertions(+), 33 deletions(-) diff --git a/test/verify/test_add_bf16.cpp b/test/verify/test_add_bf16.cpp index 2a1722e650d..99aa3f60bf3 100644 --- a/test/verify/test_add_bf16.cpp +++ b/test/verify/test_add_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_add_gelu_bf16.cpp b/test/verify/test_add_gelu_bf16.cpp index 8e625568565..11339f0c7e0 100644 --- a/test/verify/test_add_gelu_bf16.cpp +++ b/test/verify/test_add_gelu_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_block_reduce_small.cpp b/test/verify/test_block_reduce_small.cpp index 847e59bc1c3..2776b9a759f 100644 --- a/test/verify/test_block_reduce_small.cpp +++ b/test/verify/test_block_reduce_small.cpp @@ -46,7 +46,7 @@ struct test_block_reduce_small : verify_program> mm->add_return({add}); return p; }; - + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_fp32_bf16_add.cpp b/test/verify/test_fp32_bf16_add.cpp index da0ba65599f..21ed909bec3 100644 --- a/test/verify/test_fp32_bf16_add.cpp +++ b/test/verify/test_fp32_bf16_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_fp32_bf16_ladd.cpp b/test/verify/test_fp32_bf16_ladd.cpp index b4881931979..8cace09a460 100644 --- a/test/verify/test_fp32_bf16_ladd.cpp +++ b/test/verify/test_fp32_bf16_ladd.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_fp32_bf16_lall.cpp b/test/verify/test_fp32_bf16_lall.cpp index 0bb01207ab4..e767a2cb47d 100644 --- a/test/verify/test_fp32_bf16_lall.cpp +++ b/test/verify/test_fp32_bf16_lall.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_fp32_bf16_sub.cpp b/test/verify/test_fp32_bf16_sub.cpp index f771fe13cba..0cfd67029ea 100644 --- a/test/verify/test_fp32_bf16_sub.cpp +++ b/test/verify/test_fp32_bf16_sub.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_hsqrt_bf16.cpp b/test/verify/test_hsqrt_bf16.cpp index 5f8fbb12755..f79754b7fbe 100644 --- a/test/verify/test_hsqrt_bf16.cpp +++ b/test/verify/test_hsqrt_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_instancenorm.cpp b/test/verify/test_instancenorm.cpp index 47d951aea5a..e5b87b1b52c 100644 --- a/test/verify/test_instancenorm.cpp +++ b/test/verify/test_instancenorm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -75,6 +75,8 @@ struct test_instancenorm : verify_program> add_instancenorm(*mm, x, {1, 2, 1, 1}); return p; } + + std::string section() const { return "reduce"; } }; template struct test_instancenorm; template struct test_instancenorm; @@ -91,6 +93,8 @@ struct test_instancenorm_large_3d : verify_program; diff --git a/test/verify/test_layernorm.cpp b/test/verify/test_layernorm.cpp index 653a113f2ee..cea8f9e0f2f 100644 --- a/test/verify/test_layernorm.cpp +++ b/test/verify/test_layernorm.cpp @@ -130,6 +130,8 @@ struct test_layernorm_fp8_2 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_fp8_3 : verify_program @@ -158,6 +160,8 @@ struct test_layernorm_fp8_4 : verify_program add_layernorm(*mm, x, dims); return p; } + + std::string section() const { return "reduce"; } }; struct test_layernorm_eps : verify_program @@ -243,6 +247,4 @@ struct test_pw_layernorm : verify_program add_pointwise_layernorm(*mm, x, dims); return p; } - - std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_logsoftmax1.cpp b/test/verify/test_logsoftmax1.cpp index a51a83e832b..d6a9daf83b5 100644 --- a/test/verify/test_logsoftmax1.cpp +++ b/test/verify/test_logsoftmax1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -40,4 +40,6 @@ struct test_logsoftmax1 : verify_program mm->add_return({r}); return p; } + + std::string section() const { return "reduce"; } }; diff --git a/test/verify/test_pad_highest_bf16.cpp b/test/verify/test_pad_highest_bf16.cpp index e17189f4cec..f5b1d3adab3 100644 --- a/test/verify/test_pad_highest_bf16.cpp +++ b/test/verify/test_pad_highest_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_pad_lowest_bf16.cpp b/test/verify/test_pad_lowest_bf16.cpp index ada6705f9c0..b3191c63d79 100644 --- a/test/verify/test_pad_lowest_bf16.cpp +++ b/test/verify/test_pad_lowest_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_reduce_mean_bias_bf16.cpp b/test/verify/test_reduce_mean_bias_bf16.cpp index 3c9b5f40a35..8e1cdc75f68 100644 --- a/test/verify/test_reduce_mean_bias_bf16.cpp +++ b/test/verify/test_reduce_mean_bias_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_reduce_mean_bias_half.cpp b/test/verify/test_reduce_mean_bias_half.cpp index d74757873df..e70d9dc1dc4 100644 --- a/test/verify/test_reduce_mean_bias_half.cpp +++ b/test/verify/test_reduce_mean_bias_half.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_reduce_mean_large_bf16.cpp b/test/verify/test_reduce_mean_large_bf16.cpp index 42ce24b6357..48d4ff389fa 100644 --- a/test/verify/test_reduce_mean_large_bf16.cpp +++ b/test/verify/test_reduce_mean_large_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_reduce_mean_large_half.cpp b/test/verify/test_reduce_mean_large_half.cpp index 1a88bc57abb..5e40f9e9229 100644 --- a/test/verify/test_reduce_mean_large_half.cpp +++ b/test/verify/test_reduce_mean_large_half.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_reduce_noop_add.cpp b/test/verify/test_reduce_noop_add.cpp index 48ce85c535b..c72d7f967f0 100644 --- a/test/verify/test_reduce_noop_add.cpp +++ b/test/verify/test_reduce_noop_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_reduce_noop_add_bf16.cpp b/test/verify/test_reduce_noop_add_bf16.cpp index f435fc6f62f..c7528ccbe1b 100644 --- a/test/verify/test_reduce_noop_add_bf16.cpp +++ b/test/verify/test_reduce_noop_add_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_select_module_reduce.cpp b/test/verify/test_select_module_reduce.cpp index 33c26b56a9d..5e3d8ad1d63 100644 --- a/test/verify/test_select_module_reduce.cpp +++ b/test/verify/test_select_module_reduce.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_sin_bf16.cpp b/test/verify/test_sin_bf16.cpp index 64e6589ad44..26943d099bb 100644 --- a/test/verify/test_sin_bf16.cpp +++ b/test/verify/test_sin_bf16.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax1.cpp b/test/verify/test_softmax1.cpp index 2f95fb5392e..ef8c71871bb 100644 --- a/test/verify/test_softmax1.cpp +++ b/test/verify/test_softmax1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax2.cpp b/test/verify/test_softmax2.cpp index 222433cd759..c39dcec6c35 100644 --- a/test/verify/test_softmax2.cpp +++ b/test/verify/test_softmax2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax3.cpp b/test/verify/test_softmax3.cpp index 061847a2f5b..c8ca6bc95ef 100644 --- a/test/verify/test_softmax3.cpp +++ b/test/verify/test_softmax3.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax4.cpp b/test/verify/test_softmax4.cpp index 6ceeb5009f4..dae139227a8 100644 --- a/test/verify/test_softmax4.cpp +++ b/test/verify/test_softmax4.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax5.cpp b/test/verify/test_softmax5.cpp index 18c767d60f3..11bc904afde 100644 --- a/test/verify/test_softmax5.cpp +++ b/test/verify/test_softmax5.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax_large1.cpp b/test/verify/test_softmax_large1.cpp index 349d3685219..9669b095807 100644 --- a/test/verify/test_softmax_large1.cpp +++ b/test/verify/test_softmax_large1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax_large2.cpp b/test/verify/test_softmax_large2.cpp index 5a8c88ea1da..12874f647b9 100644 --- a/test/verify/test_softmax_large2.cpp +++ b/test/verify/test_softmax_large2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_softmax_large3.cpp b/test/verify/test_softmax_large3.cpp index ed6f9b200a2..900d52ca90d 100644 --- a/test/verify/test_softmax_large3.cpp +++ b/test/verify/test_softmax_large3.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_sqrt_bf161.cpp b/test/verify/test_sqrt_bf161.cpp index e2079fd1216..4ebb36db1a8 100644 --- a/test/verify/test_sqrt_bf161.cpp +++ b/test/verify/test_sqrt_bf161.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_sqrt_bf162.cpp b/test/verify/test_sqrt_bf162.cpp index dbb26d76ebf..237c2997df4 100644 --- a/test/verify/test_sqrt_bf162.cpp +++ b/test/verify/test_sqrt_bf162.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/verify/test_sqrt_bf164.cpp b/test/verify/test_sqrt_bf164.cpp index 077afe2a63e..a6a9439242e 100644 --- a/test/verify/test_sqrt_bf164.cpp +++ b/test/verify/test_sqrt_bf164.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal