diff --git a/test/verify/test_abs.cpp b/test/verify/test_abs.cpp index 0c7d7ef050c..452a489b55a 100644 --- a/test/verify/test_abs.cpp +++ b/test/verify/test_abs.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,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..8733b9b1ed1 100644 --- a/test/verify/test_acos.cpp +++ b/test/verify/test_acos.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..8b4b924e001 100644 --- a/test/verify/test_acosh.cpp +++ b/test/verify/test_acosh.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..b154a67dbd5 100644 --- a/test/verify/test_add.cpp +++ b/test/verify/test_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..b0998eebf06 --- /dev/null +++ b/test/verify/test_add_bf16.cpp @@ -0,0 +1,42 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..077290b70cb 100644 --- a/test/verify/test_add_dot.cpp +++ b/test/verify/test_add_dot.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..039dfd5d645 --- /dev/null +++ b/test/verify/test_add_gelu_bf16.cpp @@ -0,0 +1,56 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..b8d00637ffe 100644 --- a/test/verify/test_add_mixed_layout.cpp +++ b/test/verify/test_add_mixed_layout.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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_asin.cpp b/test/verify/test_asin.cpp index 34951182059..d9d5f9c7a6d 100644 --- a/test/verify/test_asin.cpp +++ b/test/verify/test_asin.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..d0eed48024d 100644 --- a/test/verify/test_asinh.cpp +++ b/test/verify/test_asinh.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..534cced00f5 100644 --- a/test/verify/test_atan.cpp +++ b/test/verify/test_atan.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..e10441aed9b 100644 --- a/test/verify/test_atanh.cpp +++ b/test/verify/test_atanh.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..0045f94be7b 100644 --- a/test/verify/test_block_reduce_small.cpp +++ b/test/verify/test_block_reduce_small.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ struct test_block_reduce_small : verify_program> 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..09db84b8c6a 100644 --- a/test/verify/test_ceil.cpp +++ b/test/verify/test_ceil.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..5a648fbb70d --- /dev/null +++ b/test/verify/test_ck_gemm_softmax_gemm_bf16.cpp @@ -0,0 +1,57 @@ +// /* +// * The MIT License (MIT) +// * +// * Copyright (c) 2015-2025 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..4e285613d11 100644 --- a/test/verify/test_concat_axis_0.cpp +++ b/test/verify/test_concat_axis_0.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..847aef3ea3f 100644 --- a/test/verify/test_concat_nhwc.cpp +++ b/test/verify/test_concat_nhwc.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..4e3a138c778 100644 --- a/test/verify/test_conv.cpp +++ b/test/verify/test_conv.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..0123b3e1571 100644 --- a/test/verify/test_conv2.cpp +++ b/test/verify/test_conv2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..d7652fc0f32 100644 --- a/test/verify/test_conv_add.cpp +++ b/test/verify/test_conv_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..dcceef5c1ab 100644 --- a/test/verify/test_conv_add_1x1_diff_strides.cpp +++ b/test/verify/test_conv_add_1x1_diff_strides.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..5b5c19859c6 100644 --- a/test/verify/test_conv_add_layernorm_conv.cpp +++ b/test/verify/test_conv_add_layernorm_conv.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..4e580798d65 100644 --- a/test/verify/test_conv_add_relu.cpp +++ b/test/verify/test_conv_add_relu.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..e208b2a872c 100644 --- a/test/verify/test_conv_add_tune.cpp +++ b/test/verify/test_conv_add_tune.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..d1ea4c5f62c 100644 --- a/test/verify/test_conv_bias_clipped_relu.cpp +++ b/test/verify/test_conv_bias_clipped_relu.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..a036aed7dac 100644 --- a/test/verify/test_conv_bn.cpp +++ b/test/verify/test_conv_bn.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..dfe8eab6d42 100644 --- a/test/verify/test_conv_bn_add.cpp +++ b/test/verify/test_conv_bn_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..63816c5d3a9 100644 --- a/test/verify/test_conv_bn_relu_pooling.cpp +++ b/test/verify/test_conv_bn_relu_pooling.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..06f5fa5fde6 100644 --- a/test/verify/test_conv_bn_relu_pooling2.cpp +++ b/test/verify/test_conv_bn_relu_pooling2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..cbfd9b18d77 100644 --- a/test/verify/test_conv_group_add.cpp +++ b/test/verify/test_conv_group_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..85cbc1fbedd 100644 --- a/test/verify/test_conv_pooling.cpp +++ b/test/verify/test_conv_pooling.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..c3dddcf16e4 100644 --- a/test/verify/test_conv_relu.cpp +++ b/test/verify/test_conv_relu.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..132e7bcae46 100644 --- a/test/verify/test_convert.cpp +++ b/test/verify/test_convert.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..a24b343c0d4 100644 --- a/test/verify/test_cos.cpp +++ b/test/verify/test_cos.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..9efe95067cb 100644 --- a/test/verify/test_cosh.cpp +++ b/test/verify/test_cosh.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..34c38b8f453 100644 --- a/test/verify/test_erf.cpp +++ b/test/verify/test_erf.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..e3f310ec70e 100644 --- a/test/verify/test_exp.cpp +++ b/test/verify/test_exp.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..387911fca5f 100644 --- a/test/verify/test_floor.cpp +++ b/test/verify/test_floor.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..ad32b742277 100644 --- a/test/verify/test_fmod_mod.cpp +++ b/test/verify/test_fmod_mod.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..d946d141322 --- /dev/null +++ b/test/verify/test_fp32_bf16_add.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..8b114cd7df9 --- /dev/null +++ b/test/verify/test_fp32_bf16_ladd.cpp @@ -0,0 +1,47 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..a2be6a572fd --- /dev/null +++ b/test/verify/test_fp32_bf16_lall.cpp @@ -0,0 +1,47 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..d8202114225 --- /dev/null +++ b/test/verify/test_fp32_bf16_sub.cpp @@ -0,0 +1,48 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..23a6008e9e9 100644 --- a/test/verify/test_gather.cpp +++ b/test/verify/test_gather.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..83096ab5080 100644 --- a/test/verify/test_gathernd_default.cpp +++ b/test/verify/test_gathernd_default.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,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..85e446028dc 100644 --- a/test/verify/test_gemm.cpp +++ b/test/verify/test_gemm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -44,6 +44,7 @@ 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; diff --git a/test/verify/test_gemm_2args_bmv.cpp b/test/verify/test_gemm_2args_bmv.cpp index 782810fdf79..d3b239c9237 100644 --- a/test/verify/test_gemm_2args_bmv.cpp +++ b/test/verify/test_gemm_2args_bmv.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -51,6 +51,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mm_1.cpp b/test/verify/test_gemm_2args_mm_1.cpp index 7ddb70fc300..c1da92ed351 100644 --- a/test/verify/test_gemm_2args_mm_1.cpp +++ b/test/verify/test_gemm_2args_mm_1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -49,6 +49,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mm_2.cpp b/test/verify/test_gemm_2args_mm_2.cpp index cd2d0df5d39..396693100cd 100644 --- a/test/verify/test_gemm_2args_mm_2.cpp +++ b/test/verify/test_gemm_2args_mm_2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -51,6 +51,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mm_3.cpp b/test/verify/test_gemm_2args_mm_3.cpp index 00be7cee4bc..9b94f3417da 100644 --- a/test/verify/test_gemm_2args_mm_3.cpp +++ b/test/verify/test_gemm_2args_mm_3.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -51,6 +51,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mm_4.cpp b/test/verify/test_gemm_2args_mm_4.cpp index c98bcc1f9c0..fb827b88eb9 100644 --- a/test/verify/test_gemm_2args_mm_4.cpp +++ b/test/verify/test_gemm_2args_mm_4.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..3bfa620aa43 100644 --- a/test/verify/test_gemm_2args_mm_5.cpp +++ b/test/verify/test_gemm_2args_mm_5.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -50,6 +50,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mm_6.cpp b/test/verify/test_gemm_2args_mm_6.cpp index e2d00cb169b..88c37d6b178 100644 --- a/test/verify/test_gemm_2args_mm_6.cpp +++ b/test/verify/test_gemm_2args_mm_6.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -53,6 +53,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mm_7.cpp b/test/verify/test_gemm_2args_mm_7.cpp index 7e03df20abc..65e0bece9a9 100644 --- a/test/verify/test_gemm_2args_mm_7.cpp +++ b/test/verify/test_gemm_2args_mm_7.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -50,6 +50,7 @@ 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; diff --git a/test/verify/test_gemm_2args_mv.cpp b/test/verify/test_gemm_2args_mv.cpp index 5a02d7cd765..72ffdd947ac 100644 --- a/test/verify/test_gemm_2args_mv.cpp +++ b/test/verify/test_gemm_2args_mv.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -49,6 +49,7 @@ 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; diff --git a/test/verify/test_gemm_2args_vbm.cpp b/test/verify/test_gemm_2args_vbm.cpp index b1c560ee331..b8e40457639 100644 --- a/test/verify/test_gemm_2args_vbm.cpp +++ b/test/verify/test_gemm_2args_vbm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -53,6 +53,7 @@ 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; diff --git a/test/verify/test_gemm_2args_vm.cpp b/test/verify/test_gemm_2args_vm.cpp index 5c9b8283a24..255494ef6b4 100644 --- a/test/verify/test_gemm_2args_vm.cpp +++ b/test/verify/test_gemm_2args_vm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -50,6 +50,7 @@ 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; // TODO need hipblaslt support diff --git a/test/verify/test_gemm_2args_vv.cpp b/test/verify/test_gemm_2args_vv.cpp index 47f9d284df9..80f9c93827a 100644 --- a/test/verify/test_gemm_2args_vv.cpp +++ b/test/verify/test_gemm_2args_vv.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -53,6 +53,7 @@ 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; diff --git a/test/verify/test_gemm_add.cpp b/test/verify/test_gemm_add.cpp index cd4231f6871..8905f6ad096 100644 --- a/test/verify/test_gemm_add.cpp +++ b/test/verify/test_gemm_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..c65b1ef9f9e 100644 --- a/test/verify/test_gemm_add_broadcast1.cpp +++ b/test/verify/test_gemm_add_broadcast1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -53,6 +53,7 @@ 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; diff --git a/test/verify/test_gemm_copy.cpp b/test/verify/test_gemm_copy.cpp index c960c6c17da..371e83cbc11 100644 --- a/test/verify/test_gemm_copy.cpp +++ b/test/verify/test_gemm_copy.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -51,6 +51,7 @@ 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; diff --git a/test/verify/test_gemm_ex.cpp b/test/verify/test_gemm_ex.cpp index f269c80f702..37ea0634b9c 100644 --- a/test/verify/test_gemm_ex.cpp +++ b/test/verify/test_gemm_ex.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -43,6 +43,7 @@ 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; diff --git a/test/verify/test_gemm_literal.cpp b/test/verify/test_gemm_literal.cpp index 62d32588166..1e02d03fcec 100644 --- a/test/verify/test_gemm_literal.cpp +++ b/test/verify/test_gemm_literal.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -48,6 +48,7 @@ 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; 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..dee50f1d9b4 --- /dev/null +++ b/test/verify/test_gemm_mul_where_softmax_gemm_bf16.cpp @@ -0,0 +1,57 @@ +// /* +// * The MIT License (MIT) +// * +// * Copyright (c) 2015-2025 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..ef97d914e40 100644 --- a/test/verify/test_gemm_multi_3args.cpp +++ b/test/verify/test_gemm_multi_3args.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ 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; diff --git a/test/verify/test_gemm_multi_3args_alpha0.cpp b/test/verify/test_gemm_multi_3args_alpha0.cpp index a66f55a93fc..cbecdda2cf7 100644 --- a/test/verify/test_gemm_multi_3args_alpha0.cpp +++ b/test/verify/test_gemm_multi_3args_alpha0.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ 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; diff --git a/test/verify/test_gemm_multi_3args_beta0.cpp b/test/verify/test_gemm_multi_3args_beta0.cpp index 6c18d54fdf5..b7d188446d2 100644 --- a/test/verify/test_gemm_multi_3args_beta0.cpp +++ b/test/verify/test_gemm_multi_3args_beta0.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ 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; diff --git a/test/verify/test_gemm_multi_3args_c25.cpp b/test/verify/test_gemm_multi_3args_c25.cpp index b5003bfccfd..1de3bf21b5c 100644 --- a/test/verify/test_gemm_multi_3args_c25.cpp +++ b/test/verify/test_gemm_multi_3args_c25.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ 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; diff --git a/test/verify/test_gemm_multi_dim_2.cpp b/test/verify/test_gemm_multi_dim_2.cpp index 5ef599340e8..3de90106c4b 100644 --- a/test/verify/test_gemm_multi_dim_2.cpp +++ b/test/verify/test_gemm_multi_dim_2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..3d7934b295b 100644 --- a/test/verify/test_gemm_multi_dim_2_3.cpp +++ b/test/verify/test_gemm_multi_dim_2_3.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -48,6 +48,7 @@ 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; diff --git a/test/verify/test_gemm_multi_transpose.cpp b/test/verify/test_gemm_multi_transpose.cpp index c34cf437d3d..e939dede715 100644 --- a/test/verify/test_gemm_multi_transpose.cpp +++ b/test/verify/test_gemm_multi_transpose.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ 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; diff --git a/test/verify/test_gemm_multibroadcast.cpp b/test/verify/test_gemm_multibroadcast.cpp index bd771d30e04..34aece50570 100644 --- a/test/verify/test_gemm_multibroadcast.cpp +++ b/test/verify/test_gemm_multibroadcast.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -46,6 +46,7 @@ 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; diff --git a/test/verify/test_gemm_pointwise.cpp b/test/verify/test_gemm_pointwise.cpp index 572ffbfaf27..2d9b8d066f4 100644 --- a/test/verify/test_gemm_pointwise.cpp +++ b/test/verify/test_gemm_pointwise.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -55,6 +55,7 @@ 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; diff --git a/test/verify/test_gemm_reshapes_add.cpp b/test/verify/test_gemm_reshapes_add.cpp index 7dd8b2dd396..5dc294a9b03 100644 --- a/test/verify/test_gemm_reshapes_add.cpp +++ b/test/verify/test_gemm_reshapes_add.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..c0fbe56dab7 --- /dev/null +++ b/test/verify/test_gemm_softmax_gemm_relu_bf16.cpp @@ -0,0 +1,81 @@ +// /* +// * The MIT License (MIT) +// * +// * Copyright (c) 2015-2025 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..09170951fc1 100644 --- a/test/verify/test_gemm_transpose_add_pooling_sub.cpp +++ b/test/verify/test_gemm_transpose_add_pooling_sub.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -65,6 +65,7 @@ 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..e91af0dfab2 100644 --- a/test/verify/test_gemm_transposea.cpp +++ b/test/verify/test_gemm_transposea.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,7 @@ 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; // TODO need hipblaslt support diff --git a/test/verify/test_gemm_transposea_ex.cpp b/test/verify/test_gemm_transposea_ex.cpp index 7244fc2c544..2a9a57c8cb6 100644 --- a/test/verify/test_gemm_transposea_ex.cpp +++ b/test/verify/test_gemm_transposea_ex.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -46,6 +46,7 @@ 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; diff --git a/test/verify/test_gemm_transposeab.cpp b/test/verify/test_gemm_transposeab.cpp index 63baa4903c8..3f38f08e45a 100644 --- a/test/verify/test_gemm_transposeab.cpp +++ b/test/verify/test_gemm_transposeab.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -46,6 +46,7 @@ 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; diff --git a/test/verify/test_gemm_transposeb.cpp b/test/verify/test_gemm_transposeb.cpp index fef29432136..32fc0d3dfb6 100644 --- a/test/verify/test_gemm_transposeb.cpp +++ b/test/verify/test_gemm_transposeb.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,7 @@ 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; diff --git a/test/verify/test_gemm_transposeb_detect.cpp b/test/verify/test_gemm_transposeb_detect.cpp index 4b83e1c83b7..61b7505c520 100644 --- a/test/verify/test_gemm_transposeb_detect.cpp +++ b/test/verify/test_gemm_transposeb_detect.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -46,6 +46,7 @@ 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; diff --git a/test/verify/test_gemm_transposeb_ex.cpp b/test/verify/test_gemm_transposeb_ex.cpp index ef75d796de1..9dbc5ed6261 100644 --- a/test/verify/test_gemm_transposeb_ex.cpp +++ b/test/verify/test_gemm_transposeb_ex.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -46,6 +46,7 @@ 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; diff --git a/test/verify/test_hsqrt_bf16.cpp b/test/verify/test_hsqrt_bf16.cpp new file mode 100644 index 00000000000..8855e597d8a --- /dev/null +++ b/test/verify/test_hsqrt_bf16.cpp @@ -0,0 +1,42 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..05b093736c7 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-2025 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 @@ -99,3 +99,4 @@ 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..6ba6344525d 100644 --- a/test/verify/test_isinf.cpp +++ b/test/verify/test_isinf.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..89507caef71 100644 --- a/test/verify/test_isnan.cpp +++ b/test/verify/test_isnan.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..8c7f117ce39 100644 --- a/test/verify/test_layernorm.cpp +++ b/test/verify/test_layernorm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -89,6 +89,21 @@ struct test_layernorm_fp16 : verify_program 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; + } + + std::string section() const { return "reduce"; } +}; + struct test_layernorm_fp8_1 : verify_program { migraphx::program create_program() const @@ -232,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_literal_limits.cpp b/test/verify/test_literal_limits.cpp index 9dac1bb7565..d705efcb5e2 100644 --- a/test/verify/test_literal_limits.cpp +++ b/test/verify/test_literal_limits.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..3b92529e51c 100644 --- a/test/verify/test_log.cpp +++ b/test/verify/test_log.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..28e00b9a974 100644 --- a/test/verify/test_log2.cpp +++ b/test/verify/test_log2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..509706f5c45 100644 --- a/test/verify/test_logsoftmax.cpp +++ b/test/verify/test_logsoftmax.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -54,6 +54,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_min_max.cpp b/test/verify/test_min_max.cpp index 12310c2c7a8..88cdd9141ee 100644 --- a/test/verify/test_min_max.cpp +++ b/test/verify/test_min_max.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,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..66cbed65f7d 100644 --- a/test/verify/test_mul_dot_a.cpp +++ b/test/verify/test_mul_dot_a.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -51,6 +51,7 @@ 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; diff --git a/test/verify/test_mul_dot_b.cpp b/test/verify/test_mul_dot_b.cpp index 6941baa9e36..ad8c4bf353c 100644 --- a/test/verify/test_mul_dot_b.cpp +++ b/test/verify/test_mul_dot_b.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,7 @@ 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; diff --git a/test/verify/test_multinomial.cpp b/test/verify/test_multinomial.cpp index 0559d0ec0a5..667e95ffaec 100644 --- a/test/verify/test_multinomial.cpp +++ b/test/verify/test_multinomial.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..165f6be42e7 100644 --- a/test/verify/test_nearbyint.cpp +++ b/test/verify/test_nearbyint.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,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..9d16cdbd326 100644 --- a/test/verify/test_nonzero.cpp +++ b/test/verify/test_nonzero.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,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..2f1c33aab06 100644 --- a/test/verify/test_pad.cpp +++ b/test/verify/test_pad.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..a9a9dd825a7 --- /dev/null +++ b/test/verify/test_pad_highest_bf16.cpp @@ -0,0 +1,46 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..2b867a318db --- /dev/null +++ b/test/verify/test_pad_lowest_bf16.cpp @@ -0,0 +1,46 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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_bf16.cpp b/test/verify/test_pointwise_broadcast_reduce_bf16.cpp new file mode 100644 index 00000000000..4c939e5f2df --- /dev/null +++ b/test/verify/test_pointwise_broadcast_reduce_bf16.cpp @@ -0,0 +1,57 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + }; + + std::string section() const { return "reduce"; } +}; diff --git a/test/verify/test_pow.cpp b/test/verify/test_pow.cpp index 3f1626a08e5..5b768c9fbef 100644 --- a/test/verify/test_pow.cpp +++ b/test/verify/test_pow.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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,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..4c5b5f88839 100644 --- a/test/verify/test_prefix_scan_sum_2d.cpp +++ b/test/verify/test_prefix_scan_sum_2d.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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_bf16.cpp b/test/verify/test_reduce_add_bf16.cpp new file mode 100644 index 00000000000..b10572687d0 --- /dev/null +++ b/test/verify/test_reduce_add_bf16.cpp @@ -0,0 +1,54 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + }; + + 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 new file mode 100644 index 00000000000..4302eefc21f --- /dev/null +++ b/test/verify/test_reduce_mean_bias_bf16.cpp @@ -0,0 +1,50 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + }; + + 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..01c09c27659 --- /dev/null +++ b/test/verify/test_reduce_mean_large_bf16.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + }; + + 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..ef57bc49f11 100644 --- a/test/verify/test_reduce_mean_nhwc.cpp +++ b/test/verify/test_reduce_mean_nhwc.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -49,3 +49,4 @@ struct test_reduce_mean_nhwc : verify_program> 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_variance_bf16.cpp b/test/verify/test_reduce_mean_variance_bf16.cpp new file mode 100644 index 00000000000..4d3bcd1896f --- /dev/null +++ b/test/verify/test_reduce_mean_variance_bf16.cpp @@ -0,0 +1,50 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + }; + + 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..8d3815fa63e --- /dev/null +++ b/test/verify/test_reduce_noop_add_bf16.cpp @@ -0,0 +1,50 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + }; + + std::string section() const { return "reduce"; } +}; diff --git a/test/verify/test_reduce_op_small.cpp b/test/verify/test_reduce_op_small.cpp index b1d824f1c8a..3d27793a21c 100644 --- a/test/verify/test_reduce_op_small.cpp +++ b/test/verify/test_reduce_op_small.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -68,6 +68,15 @@ 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..5b62f6a78ea 100644 --- a/test/verify/test_reverse.cpp +++ b/test/verify/test_reverse.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..8fb4d6c7f9a 100644 --- a/test/verify/test_rnn_sql_1.cpp +++ b/test/verify/test_rnn_sql_1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -84,6 +84,7 @@ 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; diff --git a/test/verify/test_roialign.cpp b/test/verify/test_roialign.cpp index 9f5271f1c5d..57e81df349f 100644 --- a/test/verify/test_roialign.cpp +++ b/test/verify/test_roialign.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..de52a9d8bb5 100644 --- a/test/verify/test_rsqrt.cpp +++ b/test/verify/test_rsqrt.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..2583c97b0e6 100644 --- a/test/verify/test_scatter_elements_none_axis_neg_1.cpp +++ b/test/verify/test_scatter_elements_none_axis_neg_1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..a9771d66ae5 100644 --- a/test/verify/test_scatternd.cpp +++ b/test/verify/test_scatternd.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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_shrink.cpp b/test/verify/test_shrink.cpp index ed3ab4ce86e..e8baf113814 100644 --- a/test/verify/test_shrink.cpp +++ b/test/verify/test_shrink.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..8b54c4fe26e 100644 --- a/test/verify/test_sin.cpp +++ b/test/verify/test_sin.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..4a3b541cf63 --- /dev/null +++ b/test/verify/test_sin_bf16.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..2847f4700e5 100644 --- a/test/verify/test_sinh.cpp +++ b/test/verify/test_sinh.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..95e62e7288f 100644 --- a/test/verify/test_softmax.cpp +++ b/test/verify/test_softmax.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -52,6 +52,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_softmax5.cpp b/test/verify/test_softmax5.cpp new file mode 100644 index 00000000000..c505e2bbe01 --- /dev/null +++ b/test/verify/test_softmax5.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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; + } + + std::string section() const { return "reduce"; } +}; diff --git a/test/verify/test_softmaxcrossentropyloss_2d.cpp b/test/verify/test_softmaxcrossentropyloss_2d.cpp index 684975ea119..fd0451d8e68 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -104,3 +104,11 @@ 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..61b7b346902 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d_mean.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d_mean.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -108,3 +108,11 @@ 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..e03ba58be28 100644 --- a/test/verify/test_softmaxcrossentropyloss_2d_sum.cpp +++ b/test/verify/test_softmaxcrossentropyloss_2d_sum.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -107,3 +107,11 @@ 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..254d4c5b248 100644 --- a/test/verify/test_softmaxcrossentropyloss_kd.cpp +++ b/test/verify/test_softmaxcrossentropyloss_kd.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -130,3 +130,11 @@ 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..3afc709c930 100644 --- a/test/verify/test_softmaxcrossentropyloss_kd_mean.cpp +++ b/test/verify/test_softmaxcrossentropyloss_kd_mean.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -134,3 +134,11 @@ 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..92f0afbcf65 100644 --- a/test/verify/test_softmaxcrossentropyloss_kd_sum.cpp +++ b/test/verify/test_softmaxcrossentropyloss_kd_sum.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -131,3 +131,11 @@ 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..13e348b93d1 100644 --- a/test/verify/test_sqrt.cpp +++ b/test/verify/test_sqrt.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..895a9ee8db6 --- /dev/null +++ b/test/verify/test_sqrt_bf161.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..bccc4cdf5a9 --- /dev/null +++ b/test/verify/test_sqrt_bf162.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..6301a366bd8 --- /dev/null +++ b/test/verify/test_sqrt_bf164.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 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..7faeaee0f50 100644 --- a/test/verify/test_tan.cpp +++ b/test/verify/test_tan.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..3dd595d3545 100644 --- a/test/verify/test_tanh.cpp +++ b/test/verify/test_tanh.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..675d67efbcd 100644 --- a/test/verify/test_topk_0.cpp +++ b/test/verify/test_topk_0.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..fbddb6ece50 100644 --- a/test/verify/test_trans_convert_gemm.cpp +++ b/test/verify/test_trans_convert_gemm.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..1cf01c6cd40 100644 --- a/test/verify/test_transpose_reshape_add_sub_mul.cpp +++ b/test/verify/test_transpose_reshape_add_sub_mul.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..1f66207f439 100644 --- a/test/verify/test_unbatched_gemm_1.cpp +++ b/test/verify/test_unbatched_gemm_1.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..8b010b5532e 100644 --- a/test/verify/test_unbatched_gemm_2.cpp +++ b/test/verify/test_unbatched_gemm_2.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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..0a9057b6f11 100644 --- a/test/verify/test_where.cpp +++ b/test/verify/test_where.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2025 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 @@ -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;