From 0a7edb031f01358e5c2ec1d7816dd4c7c00be8ce Mon Sep 17 00:00:00 2001 From: Aleksandr Eremin Date: Tue, 19 Sep 2023 16:21:18 +0300 Subject: [PATCH 1/5] remove direct std::random_device usage --- test/gtest/bn_test_data.hpp | 25 ++++++++++--------------- test/gtest/group_conv3d_bwd.hpp | 9 +++------ test/gtest/group_conv3d_fwd.hpp | 9 +++------ test/gtest/group_conv3d_wrw.hpp | 10 ++++------ test/gtest/group_solver.hpp | 8 +++----- test/gtest/log.cpp | 16 ++++++++-------- test/gtest/na.hpp | 16 +++++++--------- 7 files changed, 38 insertions(+), 55 deletions(-) diff --git a/test/gtest/bn_test_data.hpp b/test/gtest/bn_test_data.hpp index 36d7813fd2..755b1d4be6 100644 --- a/test/gtest/bn_test_data.hpp +++ b/test/gtest/bn_test_data.hpp @@ -25,8 +25,6 @@ *******************************************************************************/ #pragma once -#include - #include #include #include @@ -35,6 +33,7 @@ #include "tensor_util.hpp" #include "get_handle.hpp" +#include "random.hpp" struct BNTestCase { @@ -130,13 +129,9 @@ struct BNTestData void InitTensorsWithRandValue() { - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_int_distribution<> d{0, 100}; - auto gen_value = [&](auto...) { - return 1e-2 * static_cast(d(gen)) * ((d(gen) % 2 == 1) ? -1 : 1); - }; - input.generate(gen_value); + input.generate([](auto...) { + return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); + }); } void SetDirection() { direction = bn_config.Direction; } @@ -204,17 +199,17 @@ struct BNInferTestData : public BNTestData void InitTensorsWithRandValue() { - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_int_distribution<> d{0, 100}; - auto gen_value = [&](auto...) { - return 1e-2 * static_cast(d(gen)) * ((d(gen) % 2 == 1) ? -1 : 1); + auto gen_value = [](auto...) { + return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); }; scale.generate(gen_value); shift.generate(gen_value); estMean.generate(gen_value); - auto gen_var = [&](auto...) { return 1e-2 * (static_cast(d(gen)) + 1); }; + auto gen_var = [](auto...) { + return static_cast(1e-2) * + static_cast(prng::gen_0_to_B(100) + 1); + }; estVariance.generate(gen_var); } void WriteToGPU() diff --git a/test/gtest/group_conv3d_bwd.hpp b/test/gtest/group_conv3d_bwd.hpp index 410d71e6d0..9110fa34a0 100644 --- a/test/gtest/group_conv3d_bwd.hpp +++ b/test/gtest/group_conv3d_bwd.hpp @@ -25,8 +25,6 @@ *******************************************************************************/ #pragma once -#include - #include "get_handle.hpp" #include @@ -145,10 +143,9 @@ struct ConvBwdSolverTest weights = tensor{miopen_type{}, tensor_layout, conv_config.GetWeights()}; SetTensorLayout(input.desc); SetTensorLayout(weights.desc); - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_real_distribution<> d{-3, 3}; - auto gen_value = [&](auto...) { return d(gen); }; + auto gen_value = [](auto...) { + return prng::gen_A_to_B(static_cast(-3.0), static_cast(3.0)); + }; std::fill(input.begin(), input.end(), std::numeric_limits::quiet_NaN()); weights.generate(gen_value); conv_desc = conv_config.GetConv(); diff --git a/test/gtest/group_conv3d_fwd.hpp b/test/gtest/group_conv3d_fwd.hpp index 983f897d78..8783135560 100644 --- a/test/gtest/group_conv3d_fwd.hpp +++ b/test/gtest/group_conv3d_fwd.hpp @@ -25,8 +25,6 @@ *******************************************************************************/ #pragma once -#include - #include "get_handle.hpp" #include @@ -151,10 +149,9 @@ struct ConvFwdSolverTest weights = tensor{miopen_type{}, tensor_layout, conv_config.GetWeights()}; SetTensorLayout(input.desc); SetTensorLayout(weights.desc); - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_real_distribution<> d{-3, 3}; - auto gen_value = [&](auto...) { return d(gen); }; + auto gen_value = [](auto...) { + return prng::gen_A_to_B(static_cast(-3.0), static_cast(3.0)); + }; input.generate(gen_value); weights.generate(gen_value); conv_desc = conv_config.GetConv(); diff --git a/test/gtest/group_conv3d_wrw.hpp b/test/gtest/group_conv3d_wrw.hpp index 76d8ae5d90..92ada458ad 100644 --- a/test/gtest/group_conv3d_wrw.hpp +++ b/test/gtest/group_conv3d_wrw.hpp @@ -25,8 +25,6 @@ *******************************************************************************/ #pragma once -#include - #include "get_handle.hpp" #include @@ -145,10 +143,10 @@ struct ConvWrwSolverTest std::tie(algo, conv_config, tensor_layout) = GetParam(); input = tensor{miopen_type{}, tensor_layout, conv_config.GetInput()}; weights = tensor{miopen_type{}, tensor_layout, conv_config.GetWeights()}; - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_real_distribution<> d{-3, 3}; - auto gen_value = [&](auto...) { return d(gen); }; + + auto gen_value = [](auto...) { + return prng::gen_A_to_B(static_cast(-3.0), static_cast(3.0)); + }; input.generate(gen_value); std::fill(weights.begin(), weights.end(), 0); diff --git a/test/gtest/group_solver.hpp b/test/gtest/group_solver.hpp index 6fe02e00da..a473262750 100644 --- a/test/gtest/group_solver.hpp +++ b/test/gtest/group_solver.hpp @@ -134,11 +134,9 @@ struct ConvFwdSolverTest weights = tensor{miopen_type{}, tensor_layout, conv_config.GetWeights()}; SetTensorLayout(input.desc); SetTensorLayout(weights.desc); - - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_real_distribution<> d{-3, 3}; - auto gen_value = [&](auto...) { return d(gen); }; + auto gen_value = [](auto...) { + return prng::gen_A_to_B(static_cast(-3.0), static_cast(3.0)); + }; input.generate(gen_value); weights.generate(gen_value); diff --git a/test/gtest/log.cpp b/test/gtest/log.cpp index fcb5457d27..01a7253c25 100644 --- a/test/gtest/log.cpp +++ b/test/gtest/log.cpp @@ -29,8 +29,8 @@ #include #include -#include -#include +#include +#include "random.hpp" #if MIOPEN_BACKEND_OPENCL #define BKEND "OpenCL" @@ -226,17 +226,17 @@ struct CreateBNormFusionPlan shift = tensor{input_lens}; estMean = tensor{input_lens}; estVariance = tensor{input_lens}; - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_int_distribution<> d{0, 100}; - auto gen_value = [&](auto...) { - return 1e-2 * static_cast(d(gen)) * ((d(gen) % 2 == 1) ? -1 : 1); + + auto gen_value = [](auto...) { + return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); }; input.generate(gen_value); scale.generate(gen_value); shift.generate(gen_value); estMean.generate(gen_value); - auto gen_var = [&](auto...) { return 1e-2 * (static_cast(d(gen)) + 1); }; + auto gen_var = [](auto...) { + return static_cast(1e-2) * static_cast(prng::gen_0_to_B(100) + 1); + }; estVariance.generate(gen_var); activ_desc = {activ_mode, activ_alpha, activ_beta, activ_gamma}; output = tensor{input_lens}; diff --git a/test/gtest/na.hpp b/test/gtest/na.hpp index 02bc98d0bc..c518da592a 100644 --- a/test/gtest/na.hpp +++ b/test/gtest/na.hpp @@ -25,8 +25,6 @@ *******************************************************************************/ #pragma once -#include - #include #include #include @@ -35,7 +33,7 @@ #include "tensor_util.hpp" #include "get_handle.hpp" - +#include "random.hpp" struct BNTestCase { size_t N; @@ -107,17 +105,17 @@ struct BNActivInferTest shift = tensor{derivedBnDesc.GetLengths()}; estMean = tensor{derivedBnDesc.GetLengths()}; estVariance = tensor{derivedBnDesc.GetLengths()}; - std::random_device rd{}; - std::mt19937 gen{rd()}; - std::uniform_int_distribution<> d{0, 100}; - auto gen_value = [&](auto...) { - return 1e-2 * static_cast(d(gen)) * ((d(gen) % 2 == 1) ? -1 : 1); + + auto gen_value = [](auto...) { + return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); }; input.generate(gen_value); scale.generate(gen_value); shift.generate(gen_value); estMean.generate(gen_value); - auto gen_var = [&](auto...) { return 1e-2 * (static_cast(d(gen)) + 1); }; + auto gen_var = [](auto...) { + return static_cast(1e-2) * static_cast(prng::gen_0_to_B(100) + 1); + }; estVariance.generate(gen_var); activ_desc = {activ_mode, activ_alpha, activ_beta, activ_gamma}; output = tensor{bn_config.GetInput()}; From 2ff5f089f6e77eb1c6bd9471e9b39d2fd2f91e88 Mon Sep 17 00:00:00 2001 From: Aleksandr Eremin Date: Sat, 14 Oct 2023 00:51:57 +0300 Subject: [PATCH 2/5] fix fp16 precision lost due to fp32 conversion --- test/bn_3d_peract_test.cpp | 20 ++++++++++---------- test/bn_3d_spatial_test.cpp | 20 ++++++++++---------- test/bn_peract_test.cpp | 20 ++++++++++---------- test/bn_spatial_test.cpp | 20 ++++++++++---------- test/cba_inference.cpp | 2 +- test/gru_common.hpp | 12 ++++++------ test/gtest/bn_test_data.hpp | 20 ++++++++------------ test/gtest/log.cpp | 8 ++------ test/gtest/na.hpp | 8 ++------ test/lstm_common.hpp | 16 ++++++++-------- test/na_inference.cpp | 12 ++++++------ test/random.hpp | 9 +++++---- test/rnn_vanilla_common.hpp | 12 ++++++------ 13 files changed, 84 insertions(+), 95 deletions(-) diff --git a/test/bn_3d_peract_test.cpp b/test/bn_3d_peract_test.cpp index b82ec2153b..19fd15e7ce 100644 --- a/test/bn_3d_peract_test.cpp +++ b/test/bn_3d_peract_test.cpp @@ -108,11 +108,11 @@ struct verify_forward_train_3d_bn_per_activation runMean = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -243,11 +243,11 @@ struct verify_forward_train_3d_bn_per_activation runMean = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1034,15 +1034,15 @@ struct batch_norm_3d_per_activation_driver : test_driver scale = tensor{ssn, ssc, ssd, ssh, ssw}; shift = tensor{ssn, ssc, ssd, ssh, ssw}; - const PREC_TYPE Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < scale.desc.GetElementSize(); i++) { - scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } for(std::size_t i = 0; i < input.desc.GetElementSize(); i++) { - input[i] = prng::gen_descreet_uniform_sign(static_cast(1e-4), 100); + input[i] = prng::gen_descreet_uniform_sign(1e-4, 100); } } diff --git a/test/bn_3d_spatial_test.cpp b/test/bn_3d_spatial_test.cpp index cac0cb5c7a..8d428fca2b 100644 --- a/test/bn_3d_spatial_test.cpp +++ b/test/bn_3d_spatial_test.cpp @@ -102,11 +102,11 @@ struct verify_forward_train_3d_bn_spatial runMean = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } auto saveMean = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; @@ -289,11 +289,11 @@ struct verify_forward_train_3d_bn_spatial runMean = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_depth, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1236,15 +1236,15 @@ struct batch_norm_3d_spatial_driver : test_driver scale = tensor{ssn, ssc, ssd, ssh, ssw}; shift = tensor{ssn, ssc, ssd, ssh, ssw}; - const PREC_TYPE Data_scale = static_cast(1e-4); + const double Data_scale = 1e-4; for(std::size_t i = 0; i < scale.desc.GetElementSize(); i++) { - scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } for(std::size_t i = 0; i < input.desc.GetElementSize(); i++) { - input[i] = prng::gen_descreet_uniform_sign(static_cast(1e-5), 100); + input[i] = prng::gen_descreet_uniform_sign(1e-5, 100); } } diff --git a/test/bn_peract_test.cpp b/test/bn_peract_test.cpp index 7dce9aabf5..6622230666 100644 --- a/test/bn_peract_test.cpp +++ b/test/bn_peract_test.cpp @@ -107,11 +107,11 @@ struct verify_forward_train_bn_per_activation runMean = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -233,11 +233,11 @@ struct verify_forward_train_bn_per_activation runMean = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -992,15 +992,15 @@ struct batch_norm_per_activation_driver : test_driver scale = tensor{ssn, ssc, ssh, ssw}; shift = tensor{ssn, ssc, ssh, ssw}; - const PREC_TYPE Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < scale.desc.GetElementSize(); i++) { - scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } for(std::size_t i = 0; i < input.desc.GetElementSize(); i++) { - input[i] = prng::gen_descreet_uniform_sign(static_cast(1e-4), 100); + input[i] = prng::gen_descreet_uniform_sign(1e-4, 100); } } diff --git a/test/bn_spatial_test.cpp b/test/bn_spatial_test.cpp index 8c8d7bc2eb..82d1cc271b 100644 --- a/test/bn_spatial_test.cpp +++ b/test/bn_spatial_test.cpp @@ -101,11 +101,11 @@ struct verify_forward_train_bn_spatial runMean = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } auto saveMean = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; @@ -270,11 +270,11 @@ struct verify_forward_train_bn_spatial runMean = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; runVar = tensor{rs_n_batch, rs_channels, rs_height, rs_width}; - const U Data_scale = static_cast(0.001); + const double Data_scale = 0.001; for(std::size_t i = 0; i < runMean.desc.GetElementSize(); i++) { - runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); + runMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + runVar[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1150,15 +1150,15 @@ struct batch_norm_spatial_driver : test_driver scale = tensor{ssn, ssc, ssh, ssw}; shift = tensor{ssn, ssc, ssh, ssw}; - const PREC_TYPE Data_scale = static_cast(1e-4); + const double Data_scale = 1e-4; for(std::size_t i = 0; i < scale.desc.GetElementSize(); i++) { - scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } for(std::size_t i = 0; i < input.desc.GetElementSize(); i++) { - input[i] = prng::gen_descreet_uniform_sign(static_cast(1e-5), 100); + input[i] = prng::gen_descreet_uniform_sign(1e-5, 100); } } diff --git a/test/cba_inference.cpp b/test/cba_inference.cpp index fb8879997d..502d9bbc11 100644 --- a/test/cba_inference.cpp +++ b/test/cba_inference.cpp @@ -386,7 +386,7 @@ struct cba_fusion_driver : test_driver bias = tensor{1, output.desc.GetLengths()[1], 1, 1}; for(std::size_t i = 0; i < bias.desc.GetElementSize(); i++) { - bias[i] = prng::gen_descreet_uniform_sign(static_cast(0.1), 100); + bias[i] = prng::gen_descreet_uniform_sign(0.1, 100); } } diff --git a/test/gru_common.hpp b/test/gru_common.hpp index d6c3c569f7..51e3c09012 100644 --- a/test/gru_common.hpp +++ b/test/gru_common.hpp @@ -2945,7 +2945,7 @@ struct gru_basic_driver : test_driver void run() { - const T Data_scale = static_cast(0.001); + const double Data_scale = 0.001; #if(MIOPEN_BACKEND_OPENCL == 1) if(type == miopenHalf) exit(EXIT_SUCCESS); // NOLINT (concurrency-mt-unsafe) @@ -3058,7 +3058,7 @@ struct gru_basic_driver : test_driver std::vector input(in_sz); for(std::size_t i = 0; i < in_sz; i++) { - input[i] = prng::gen_descreet_unsigned(Data_scale, 100); + input[i] = prng::gen_descreet_unsigned(Data_scale, 100); } std::size_t hx_sz = ((dirMode != 0) ? 2ULL : 1ULL) * hiddenSize * batchSize * numLayers; @@ -3076,7 +3076,7 @@ struct gru_basic_driver : test_driver std::vector weights(wei_sz); for(std::size_t i = 0; i < wei_sz; i++) { - weights[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + weights[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } #if(MIO_GRU_TEST_DEBUG > 0) @@ -3093,7 +3093,7 @@ struct gru_basic_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - hx[i] = prng::gen_descreet_unsigned(Data_scale, 100); + hx[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -3101,7 +3101,7 @@ struct gru_basic_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - dhyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dhyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -3165,7 +3165,7 @@ struct gru_basic_driver : test_driver std::vector dyin(yin.size()); for(std::size_t i = 0; i < yin.size(); i++) { - dyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } #if(MIO_GRU_TEST_DEBUG > 0) diff --git a/test/gtest/bn_test_data.hpp b/test/gtest/bn_test_data.hpp index 755412146a..799f6ae4b8 100644 --- a/test/gtest/bn_test_data.hpp +++ b/test/gtest/bn_test_data.hpp @@ -130,9 +130,8 @@ struct BNTestData void InitTensorsWithRandValue() { - input.generate([](auto...) { - return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); - }); + input.generate( + [](auto...) { return prng::gen_descreet_uniform_sign(1e-2, 100); }); } void SetDirection() { direction = bn_config.Direction; } @@ -201,15 +200,14 @@ struct BNInferTestData : public BNTestData void InitTensorsWithRandValue() { auto gen_value = [](auto...) { - return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); + return prng::gen_descreet_uniform_sign(1e-2, 100); }; scale.generate(gen_value); shift.generate(gen_value); estMean.generate(gen_value); auto gen_var = [](auto...) { - return static_cast(1e-2) * - static_cast(prng::gen_0_to_B(100) + 1); + return static_cast(1e-2 * (prng::gen_0_to_B(100) + 1)); }; estVariance.generate(gen_var); } @@ -303,15 +301,14 @@ struct BNBwdTestData : public BNTestData void InitTensorsWithRandValue() { auto gen_value = [](auto...) { - return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); + return prng::gen_descreet_uniform_sign(1e-2, 100); }; dy.generate(gen_value); bnScale.generate(gen_value); savedMean.generate(gen_value); auto gen_var = [](auto...) { - return static_cast(1e-2) * - static_cast(prng::gen_0_to_B(100) + 1); + return static_cast(1e-2 * (prng::gen_0_to_B(100) + 1)); }; savedInvVar.generate(gen_var); @@ -409,14 +406,13 @@ struct BNFwdTrainTestData : public BNTestData void InitTensorsWithRandValue() { auto gen_value = [](auto...) { - return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); + return prng::gen_descreet_uniform_sign(1e-2, 100); }; scale.generate(gen_value); shift.generate(gen_value); auto gen_var = [](auto...) { - return static_cast(1e-2) * - static_cast(prng::gen_0_to_B(100) + 1); + return static_cast(1e-2 * (prng::gen_0_to_B(100) + 1)); }; runMean.generate(gen_var); runVariance.generate(gen_var); diff --git a/test/gtest/log.cpp b/test/gtest/log.cpp index 01a7253c25..3f07e12366 100644 --- a/test/gtest/log.cpp +++ b/test/gtest/log.cpp @@ -227,16 +227,12 @@ struct CreateBNormFusionPlan estMean = tensor{input_lens}; estVariance = tensor{input_lens}; - auto gen_value = [](auto...) { - return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); - }; + auto gen_value = [](auto...) { return prng::gen_descreet_uniform_sign(1e-2, 100); }; input.generate(gen_value); scale.generate(gen_value); shift.generate(gen_value); estMean.generate(gen_value); - auto gen_var = [](auto...) { - return static_cast(1e-2) * static_cast(prng::gen_0_to_B(100) + 1); - }; + auto gen_var = [](auto...) { return static_cast(1e-2 * (prng::gen_0_to_B(100) + 1)); }; estVariance.generate(gen_var); activ_desc = {activ_mode, activ_alpha, activ_beta, activ_gamma}; output = tensor{input_lens}; diff --git a/test/gtest/na.hpp b/test/gtest/na.hpp index c518da592a..387f88e12c 100644 --- a/test/gtest/na.hpp +++ b/test/gtest/na.hpp @@ -106,16 +106,12 @@ struct BNActivInferTest estMean = tensor{derivedBnDesc.GetLengths()}; estVariance = tensor{derivedBnDesc.GetLengths()}; - auto gen_value = [](auto...) { - return prng::gen_descreet_uniform_sign(static_cast(1e-2), 100); - }; + auto gen_value = [](auto...) { return prng::gen_descreet_uniform_sign(1e-2, 100); }; input.generate(gen_value); scale.generate(gen_value); shift.generate(gen_value); estMean.generate(gen_value); - auto gen_var = [](auto...) { - return static_cast(1e-2) * static_cast(prng::gen_0_to_B(100) + 1); - }; + auto gen_var = [](auto...) { return static_cast(1e-2 * (prng::gen_0_to_B(100) + 1)); }; estVariance.generate(gen_var); activ_desc = {activ_mode, activ_alpha, activ_beta, activ_gamma}; output = tensor{bn_config.GetInput()}; diff --git a/test/lstm_common.hpp b/test/lstm_common.hpp index 950c49f4cc..1f0d5f98cb 100644 --- a/test/lstm_common.hpp +++ b/test/lstm_common.hpp @@ -1503,7 +1503,7 @@ struct lstm_basic_driver : test_driver void run() { - const T Data_scale = static_cast(0.001); + const double Data_scale = 0.001; #if(MIOPEN_BACKEND_OPENCL == 1) #if WORKAROUND_ISSUE_692 == 1 std::cout << "Skip test for Issue #692: " << std::endl; @@ -1634,7 +1634,7 @@ struct lstm_basic_driver : test_driver std::vector input(in_sz); for(std::size_t i = 0; i < in_sz; i++) { - input[i] = prng::gen_descreet_unsigned(Data_scale, 100); + input[i] = prng::gen_descreet_unsigned(Data_scale, 100); } std::size_t hx_sz = ((dirMode != 0) ? 2ULL : 1ULL) * hiddenSize * batchSize * numLayers; @@ -1654,7 +1654,7 @@ struct lstm_basic_driver : test_driver std::vector weights(wei_sz); for(std::size_t i = 0; i < wei_sz; i++) { - weights[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + weights[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } #if(MIO_LSTM_TEST_DEBUG > 0) @@ -1679,7 +1679,7 @@ struct lstm_basic_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - hx[i] = prng::gen_descreet_unsigned(Data_scale, 100); + hx[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1687,7 +1687,7 @@ struct lstm_basic_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - dhyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dhyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1695,7 +1695,7 @@ struct lstm_basic_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - cx[i] = prng::gen_descreet_unsigned(Data_scale, 100); + cx[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1703,7 +1703,7 @@ struct lstm_basic_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - dcyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dcyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -1791,7 +1791,7 @@ struct lstm_basic_driver : test_driver std::vector dyin(out_sz); for(std::size_t i = 0; i < out_sz; i++) { - dyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } #if(MIO_LSTM_TEST_DEBUG > 0) diff --git a/test/na_inference.cpp b/test/na_inference.cpp index 129c970a3f..5ea5daf15b 100644 --- a/test/na_inference.cpp +++ b/test/na_inference.cpp @@ -263,17 +263,17 @@ struct na_fusion_driver : test_driver estVariance = tensor{ ssn, ssc, ssh, ssw}; //.generate( tensor_elem_gen_integer{max_value});; - PREC_TYPE Data_scale{static_cast(0.01)}; + const double Data_scale = 0.01; for(std::size_t i = 0; i < scale.desc.GetElementSize(); i++) { - scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - estMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); - estVariance[i] = Data_scale * static_cast(prng::gen_off_range(1, 100)); + scale[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + shift[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + estMean[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); + estVariance[i] = static_cast(Data_scale * prng::gen_off_range(1, 100)); } for(std::size_t i = 0; i < input.desc.GetElementSize(); i++) { - input[i] = prng::gen_descreet_uniform_sign(static_cast(Data_scale), 100); + input[i] = prng::gen_descreet_uniform_sign(Data_scale, 100); } auto&& handle = get_handle(); diff --git a/test/random.hpp b/test/random.hpp index a6988e0bf4..f8fad740f8 100644 --- a/test/random.hpp +++ b/test/random.hpp @@ -30,15 +30,16 @@ namespace prng { template -inline T gen_descreet_uniform_sign(T scale, int32_t range) +inline T gen_descreet_uniform_sign(double scale, int32_t range) { - return (gen_canonical() ? -scale : scale) * static_cast(gen_0_to_B(range)); + return static_cast((gen_canonical() ? -scale : scale) * + static_cast(gen_0_to_B(range))); } template -inline T gen_descreet_unsigned(T scale, int32_t range) +inline T gen_descreet_unsigned(double scale, int32_t range) { - return scale * static_cast(gen_0_to_B(range)); + return static_cast(scale * static_cast(gen_0_to_B(range))); } } // namespace prng #endif // GUARD_MIOPEN_TEST_RANDOM_HPP diff --git a/test/rnn_vanilla_common.hpp b/test/rnn_vanilla_common.hpp index 3287b13e7e..ef5c1088c3 100644 --- a/test/rnn_vanilla_common.hpp +++ b/test/rnn_vanilla_common.hpp @@ -2354,7 +2354,7 @@ struct rnn_basic_vanilla_driver : test_driver void run() { - const T Data_scale = static_cast(0.001); + const double Data_scale = 0.001; #if(MIOPEN_BACKEND_OPENCL == 1) if(type == miopenHalf) exit(EXIT_SUCCESS); // NOLINT (concurrency-mt-unsafe) @@ -2468,7 +2468,7 @@ struct rnn_basic_vanilla_driver : test_driver std::vector input(in_sz); for(std::size_t i = 0; i < in_sz; i++) { - input[i] = prng::gen_descreet_unsigned(Data_scale, 100); + input[i] = prng::gen_descreet_unsigned(Data_scale, 100); } std::size_t hx_sz = ((dirMode != 0) ? 2ULL : 1ULL) * hiddenSize * batchSize * numLayers; @@ -2491,7 +2491,7 @@ struct rnn_basic_vanilla_driver : test_driver std::vector weights(wei_sz); for(std::size_t i = 0; i < wei_sz; i++) { - weights[i] = prng::gen_descreet_uniform_sign(Data_scale / 10, 100); + weights[i] = prng::gen_descreet_uniform_sign(Data_scale / 10, 100); } #if(MIO_RNN_TEST_DEBUG > 0) @@ -2514,7 +2514,7 @@ struct rnn_basic_vanilla_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - hx[i] = prng::gen_descreet_unsigned(Data_scale, 100); + hx[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -2522,7 +2522,7 @@ struct rnn_basic_vanilla_driver : test_driver { for(std::size_t i = 0; i < hx_sz; i++) { - dhyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dhyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } } @@ -2587,7 +2587,7 @@ struct rnn_basic_vanilla_driver : test_driver std::vector dyin(yin.size()); for(std::size_t i = 0; i < yin.size(); i++) { - dyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); + dyin[i] = prng::gen_descreet_unsigned(Data_scale, 100); } #if(MIO_RNN_TEST_DEBUG > 0) printf("Running backward data RNN.\n"); From 3a95ba3d43e463329e666e60bec70cb5d3d095bc Mon Sep 17 00:00:00 2001 From: Aleksandr Eremin Date: Tue, 17 Oct 2023 13:03:00 +0300 Subject: [PATCH 3/5] make distribution less biased to 0 --- test/random.hpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/test/random.hpp b/test/random.hpp index f8fad740f8..9b4815bc1d 100644 --- a/test/random.hpp +++ b/test/random.hpp @@ -32,8 +32,7 @@ namespace prng { template inline T gen_descreet_uniform_sign(double scale, int32_t range) { - return static_cast((gen_canonical() ? -scale : scale) * - static_cast(gen_0_to_B(range))); + return static_cast(scale * prng::gen_A_to_B(-range + 1, range)); } template From d26a84c6ff57f0823e9da7beed43c6206b3768aa Mon Sep 17 00:00:00 2001 From: Aleksandr Eremin Date: Tue, 17 Oct 2023 13:05:54 +0300 Subject: [PATCH 4/5] increase batchnorm test threshold since it's very sensitive to input data --- test/gtest/bn.hpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/test/gtest/bn.hpp b/test/gtest/bn.hpp index 22f8391fe6..49a81a359b 100644 --- a/test/gtest/bn.hpp +++ b/test/gtest/bn.hpp @@ -146,10 +146,9 @@ struct BNBwdTest : public ::testing::TestWithParam(bn_bwd_test_data); - // using tolerance = 1e-4 since this the tolerance CK uses - test::CompareTensor(bn_bwd_test_data.output, bn_bwd_test_data.ref_out, 1e-4); - test::CompareTensor(bn_bwd_test_data.dScale, bn_bwd_test_data.dScale_ref, 1e-4); - test::CompareTensor(bn_bwd_test_data.dBias, bn_bwd_test_data.dBias_ref, 1e-4); + test::CompareTensor(bn_bwd_test_data.output, bn_bwd_test_data.ref_out, 5e-4); + test::CompareTensor(bn_bwd_test_data.dScale, bn_bwd_test_data.dScale_ref, 5e-4); + test::CompareTensor(bn_bwd_test_data.dBias, bn_bwd_test_data.dBias_ref, 5e-4); } BNTestCase bn_config; From ced07958c80153a2aad06c2447eb70336a196d87 Mon Sep 17 00:00:00 2001 From: Alex Eremin Date: Tue, 24 Oct 2023 13:07:50 +0200 Subject: [PATCH 5/5] Fix includes to comply CG SF.12 --- test/gtest/log.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/gtest/log.cpp b/test/gtest/log.cpp index 3f07e12366..80116c55df 100644 --- a/test/gtest/log.cpp +++ b/test/gtest/log.cpp @@ -30,7 +30,7 @@ #include #include #include -#include "random.hpp" +#include "../random.hpp" #if MIOPEN_BACKEND_OPENCL #define BKEND "OpenCL"