From f14b258a5c9ffb586aa27d3837b6f7778dfaff29 Mon Sep 17 00:00:00 2001 From: Hector Li Date: Wed, 6 Feb 2019 14:17:36 -0800 Subject: [PATCH] Fix float 16 type support for some CUDA kernels (#436) * Correct the Consts::Zero & Consts::One for half type * 1. Fix the CreateConstantOnes for float16 type 2. Add cuda kernel code in the BatchNorm for float 16 type, there's issue to run cudnnBatchNormalizationForwardInference with float 16 type 3. Add float 16 test case for Gemm & BatchNorm CUDA kernel only * Fix build * fix Linux build * fix build * Update the fix for BatchNorm, still use cuddn API cudnnBatchNormalizationForwardInference. The root case is, for half type, should use alpha, beta, scale, B, mean, var with float type. * fix build * enable 2 fp16 models for GPU test * enable fp16 test for MaxPool * Need to adjust per_sample_tolerance configuration in the model test --- onnxruntime/core/providers/cuda/cuda_utils.cu | 3 +- .../core/providers/cuda/nn/batch_norm.cc | 45 ++++++++- onnxruntime/test/onnx/main.cc | 6 +- .../test/providers/cpu/math/gemm_test.cc | 50 ++++++++-- .../providers/cpu/nn/batch_norm_op_test.cc | 93 +++++++++++++++++++ .../test/providers/cpu/nn/pool_op_test.cc | 54 +++++++++++ .../test/providers/provider_test_utils.cc | 23 +++++ .../test/providers/provider_test_utils.h | 16 ++++ 8 files changed, 275 insertions(+), 15 deletions(-) diff --git a/onnxruntime/core/providers/cuda/cuda_utils.cu b/onnxruntime/core/providers/cuda/cuda_utils.cu index 3790d4e745..01bb6a73df 100644 --- a/onnxruntime/core/providers/cuda/cuda_utils.cu +++ b/onnxruntime/core/providers/cuda/cuda_utils.cu @@ -9,6 +9,7 @@ #include #include "core/providers/cuda/shared_inc/cuda_utils.h" #include "core/providers/cuda/cu_inc/common.cuh" +#include "cudnn_common.h" namespace onnxruntime { namespace cuda { @@ -31,7 +32,7 @@ class ConstantBufferImpl : public IConstantBuffer { template std::unique_ptr> CreateConstantOnes() { - return std::make_unique>((T)1); + return std::make_unique>(Consts::One); } template std::unique_ptr> CreateConstantOnes(); diff --git a/onnxruntime/core/providers/cuda/nn/batch_norm.cc b/onnxruntime/core/providers/cuda/nn/batch_norm.cc index 4b81bd5f6e..1c40d850e5 100644 --- a/onnxruntime/core/providers/cuda/nn/batch_norm.cc +++ b/onnxruntime/core/providers/cuda/nn/batch_norm.cc @@ -5,6 +5,8 @@ #include "core/providers/common.h" #include "core/providers/cuda/cudnn_common.h" #include "core/providers/cpu/nn/batch_norm_helper.h" +#include "core/providers/cuda/math/unary_elementwise_ops_impl.h" + using namespace std; namespace onnxruntime { namespace cuda { @@ -46,17 +48,54 @@ Status BatchNorm::ComputeInternal(OpKernelContext* p_op_kernel_context) const auto mean_data = reinterpret_cast(mean->template Data()); auto var_data = reinterpret_cast(var->template Data()); + const auto alpha = Consts::One; + const auto beta = Consts::Zero; + CudnnTensor data_desc; vector new_dims; BatchNormHelper::NormalizeDims(x_shape, new_dims); ORT_RETURN_IF_ERROR(data_desc.Set(new_dims, CudnnTensor::GetDataType())); + // For half data type, the alpha, beta, scale, B, mean, var need to be float type + if (X->DataType() == DataTypeImpl::GetType()) { + CudnnTensor scale_desc; + ORT_RETURN_IF_ERROR(scale_desc.Set(new_dims, CudnnTensor::GetDataType())); + CudnnTensor bn_tensor_desc; + ORT_RETURN_IF_ERROR(bn_tensor_desc.Set(data_desc, cudnn_batch_norm_mode_)); + + // Convert the scale, B, mean, var to float + const int64_t C = x_shape.GetDims()[1]; + auto f_scale = GetScratchBuffer(C); + auto f_B = GetScratchBuffer(C); + auto f_mean = GetScratchBuffer(C); + auto f_var = GetScratchBuffer(C); + Impl_Cast(scale_data, f_scale.get(), C); + Impl_Cast(b_data, f_B.get(), C); + Impl_Cast(mean_data, f_mean.get(), C); + Impl_Cast(var_data, f_var.get(), C); + + CUDNN_RETURN_IF_ERROR(cudnnBatchNormalizationForwardInference( + CudnnHandle(), + cudnn_batch_norm_mode_, + &alpha, + &beta, + data_desc, + x_data, + data_desc, + y_data, + bn_tensor_desc, + f_scale.get(), + f_B.get(), + f_mean.get(), + f_var.get(), + epsilon_)); + + return Status::OK(); + } + CudnnTensor bn_tensor_desc; ORT_RETURN_IF_ERROR(bn_tensor_desc.Set(data_desc, cudnn_batch_norm_mode_)); - const auto alpha = Consts::One; - const auto beta = Consts::Zero; - CUDNN_RETURN_IF_ERROR(cudnnBatchNormalizationForwardInference( CudnnHandle(), cudnn_batch_norm_mode_, diff --git a/onnxruntime/test/onnx/main.cc b/onnxruntime/test/onnx/main.cc index 58000ee52f..be3f1fb79b 100644 --- a/onnxruntime/test/onnx/main.cc +++ b/onnxruntime/test/onnx/main.cc @@ -341,9 +341,9 @@ int real_main(int argc, char* argv[]) { broken_tests["maxpool_3d_default"] = "cudnn pooling only support input dimension >= 3"; broken_tests["maxpool_1d_default"] = "cudnn pooling only support input dimension >= 3"; - broken_tests["fp16_tiny_yolov2"] = "unknown failure on CUDA"; - broken_tests["fp16_shufflenet"] = "unknown failure on CUDA"; - broken_tests["fp16_inception_v1"] = "unknown failure on CUDA"; + broken_tests["fp16_tiny_yolov2"] = "Need to adjust the per_sample_tolerance: 0.2"; + broken_tests["fp16_shufflenet"] = "still have issue on Linux"; + broken_tests["fp16_inception_v1"] = "need to adjust the per_sample_tolerance: 0.002"; #endif int result = 0; diff --git a/onnxruntime/test/providers/cpu/math/gemm_test.cc b/onnxruntime/test/providers/cpu/math/gemm_test.cc index 8c63564a86..2fb17717fa 100644 --- a/onnxruntime/test/providers/cpu/math/gemm_test.cc +++ b/onnxruntime/test/providers/cpu/math/gemm_test.cc @@ -7,7 +7,7 @@ namespace onnxruntime { namespace test { -TEST(MathOpTest, GemmNoTrans) { +TEST(GemmOpTest, GemmNoTrans) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)0); @@ -26,7 +26,41 @@ TEST(MathOpTest, GemmNoTrans) { test.Run(); } -TEST(MathOpTest, GemmBroadcast) { +// Only CUDA kernel has float 16 support +#ifdef USE_CUDA +TEST(GemmOpTest, GemmNoTrans_f16) { + OpTester test("Gemm"); + + test.AddAttribute("transA", (int64_t)0); + test.AddAttribute("transB", (int64_t)0); + test.AddAttribute("alpha", 1.0f); + test.AddAttribute("beta", 1.0f); + + std::vector A{1.0f, 2.0f, 3.0f, 4.0f, + -1.0f, -2.0f, -3.0f, -4.0f}; + std::vector B(12, 1.0f); + std::vector C(6, 1.0f); + std::vector Y{11.0f, 11.0f, 11.0f, + -9.0f, -9.0f, -9.0f}; + + std::vector f_A(8); + std::vector f_B(12); + std::vector f_C(6); + std::vector f_Y(6); + ConvertFloatToMLFloat16(A.data(), f_A.data(), 8); + ConvertFloatToMLFloat16(B.data(), f_B.data(), 12); + ConvertFloatToMLFloat16(C.data(), f_C.data(), 6); + ConvertFloatToMLFloat16(Y.data(), f_Y.data(), 6); + + test.AddInput("A", {2, 4}, f_A); + test.AddInput("B", {4, 3}, f_B); + test.AddInput("C", {2, 3}, f_C); + test.AddOutput("Y", {2, 3}, f_Y); + test.Run(); +} +#endif + +TEST(GemmOpTest, GemmBroadcast) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)0); @@ -45,7 +79,7 @@ TEST(MathOpTest, GemmBroadcast) { test.Run(); } -TEST(MathOpTest, GemmTrans) { +TEST(GemmOpTest, GemmTrans) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)1); @@ -66,7 +100,7 @@ TEST(MathOpTest, GemmTrans) { test.Run(); } -TEST(MathOpTest, GemmAlphaBeta) { +TEST(GemmOpTest, GemmAlphaBeta) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)0); @@ -85,7 +119,7 @@ TEST(MathOpTest, GemmAlphaBeta) { test.Run(); } -TEST(MathOpTest, GemmNaN) { +TEST(GemmOpTest, GemmNaN) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)0); @@ -104,7 +138,7 @@ TEST(MathOpTest, GemmNaN) { test.Run(); } -TEST(MathOpTest, GemmScalarBroadcast) { +TEST(GemmOpTest, GemmScalarBroadcast) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)0); @@ -142,7 +176,7 @@ TEST(MathOpTest, Gemm2DBroadcast) { test.Run(); } -TEST(MathOpTest, GemmFalseBroadcast) { +TEST(GemmOpTest, GemmFalseBroadcast) { OpTester test("Gemm"); test.AddAttribute("transA", (int64_t)0); @@ -161,7 +195,7 @@ TEST(MathOpTest, GemmFalseBroadcast) { test.Run(); } -TEST(MathOpTest, GemmEmptyTensor) { +TEST(GemmOpTest, GemmEmptyTensor) { OpTester test("Gemm"); test.AddAttribute("transA", static_cast(0)); diff --git a/onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc b/onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc index bf69252542..5b4e78ab75 100644 --- a/onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc +++ b/onnxruntime/test/providers/cpu/nn/batch_norm_op_test.cc @@ -512,5 +512,98 @@ TEST(BatchNormTest, InvalidVarDim) { OpTester::ExpectResult::kExpectFailure, "Invalid input var"); } + +// Only CUDA kernel has float 16 support +#ifdef USE_CUDA +TEST(BatchNormTest, BatchNorm2d_fp16) { + vector X{-0.91221f, -0.283559f, 0.937637f, 2.09818f, -0.100199f, -0.608113f, 0.444562f, -1.07505f, 0.940591f, + -0.922262f, 0.0931303f, 0.69611f, 1.55187f, 0.159808f, 0.914874f, -1.24856f, -1.98928f, -0.331621f, + 2.33131f, 0.260409f, 0.0944811f, 0.442397f, 0.76461f, -0.203334f, -0.244228f, -0.387267f, -1.65039f, + -0.815409f, 0.931696f, -1.15328f, 0.773952f, -1.28195f, -0.437349f, 0.0644882f, -0.087637f, 1.74999f, + 0.640154f, -0.505641f, -1.84014f, -0.00135415f, 0.782006f, -1.21172f, -0.621273f, -0.0977471f, + -0.941333f, -0.170302f, 0.18923f, 0.436322f, 0.870412f, -0.582312f, 0.679017f, 0.510252f, 0.0786005f, + 0.160138f, -2.61889f, 0.402828f, 0.551144f, -1.39366f, -1.15191f, 0.160008f, -0.57735f, 0.210758f, + 1.0541f, -2.12569f, 0.101656f, 1.10223f, 0.725811f, -1.5019f, -0.0892582f, 0.063546f, 0.822734f, + 1.67707f, 0.478121f, -1.07438f, -0.0487855f, 0.0972885f, -1.54122f, 2.47422f, 0.596108f, 0.0026957f, + -0.967677f, -2.08882f, 0.469692f, 0.630784f, 0.196915f, -1.91331f, 1.26255f, 0.0491993f, -0.358415f, + 0.720588f, 0.976776f, -0.418116f, 1.70979f, 2.49971f, 1.30942f, -1.18304f, -1.64901f, -1.11048f, + 1.41467f, -0.275486f, -1.20602f, -0.545566f, -0.918059f, 1.48513f, 2.04224f, -0.96909f, -1.92804f, + 0.634147f, -1.02079f, -0.000786079f, 0.72428f, 0.893569f, 1.14604f, -1.3423f, -1.05061f, -0.617524f, + -0.12619f, -0.203127f, -0.941956f, 2.06916f, 2.03025f, 0.37269f, -0.340471f, -1.27962f, 0.159472f, + 0.643999f, 0.881773f, -0.50873f, 1.04599f, -0.287968f, 1.84344f, -0.728637f, 0.668021f, -2.00452f, + -0.585523f, -0.24982f, -0.379091f, 0.213692f, 1.21336f, -0.499157f, -1.50841f, -1.01256f, 0.745338f, + 0.591107f, 1.33781f, -0.258927f, -1.87304f, 0.884799f, 1.63174f, 0.500887f, 1.80608f, -1.25441f, + -0.655316f, 1.22439f, 0.384174f, 0.401395f, 1.43172f, 1.85338f, -0.644909f, -1.46975f, -1.06138f, + 1.09724f, 0.013438f, 0.589742f, -0.695768f, 0.758401f, 0.924533f, -0.0988563f, -0.197066f, 1.01118f, + 0.195163f, 0.975466f, 1.7682f, 0.977977f, -0.88963f, -0.251431f, 0.115828f, -0.230065f, -1.08882f, + 1.62318f, 0.502684f, 0.789724f, 1.13057f, -0.890021f, -0.614755f, 1.11055f, -1.21681f, 0.133085f, + 0.564458f, 0.723117f, 1.67088f, -0.111012f, 1.39732f, -0.846095f, -0.194408f, -0.381931f, -0.735943f, + -0.788814f, -0.910318f, 1.16345f, -1.98542f, 0.742905f, -0.749476f, -0.110805f, 0.307949f, -1.66811f, + 0.294031f, -0.522837f, -0.774399f, -0.264072f, -0.426894f, 0.965971f, 0.173348f, -0.991018f, 1.9406f, + 0.0853744f}; + vector scale{0.736494f, 0.580251f, 0.374834f}; + vector B{0.0f, 0.0f, 0.0f}; + vector mean{0.0f, 0.0f, 0.0f}; + vector var{1.0f, 1.0f, 1.0f}; + + vector expected_output{-0.671834f, -0.208838f, 0.69056f, 1.54529f, -0.0737958f, -0.447869f, 0.327415f, -0.791764f, + 0.692736f, -0.679237f, 0.0685895f, 0.512678f, 1.14294f, 0.117697f, 0.673796f, -0.91955f, + -1.46508f, -0.244235f, 1.71699f, 0.191789f, 0.0695843f, 0.325821f, 0.563128f, -0.149753f, + -0.179871f, -0.285218f, -1.2155f, -0.60054f, 0.686185f, -0.84938f, 0.570008f, -0.944146f, + -0.322103f, 0.0474949f, -0.0645437f, 1.28885f, 0.371448f, -0.293397f, -1.06774f, -0.00078574f, + 0.453757f, -0.703098f, -0.360492f, -0.0567175f, -0.546206f, -0.0988174f, 0.1098f, 0.253175f, + 0.505055f, -0.337885f, 0.393998f, 0.296073f, 0.0456077f, 0.0929198f, -1.5196f, 0.23374f, 0.3198f, + -0.808671f, -0.668391f, 0.0928441f, -0.335006f, 0.122292f, 0.611642f, -1.23343f, 0.0589855f, + 0.639564f, 0.42115f, -0.871472f, -0.0517919f, 0.0368724f, 0.477389f, 0.973116f, 0.179215f, + -0.402711f, -0.0182864f, 0.0364668f, -0.577699f, 0.927416f, 0.22344f, 0.00101043f, -0.362716f, + -0.782957f, 0.176056f, 0.236438f, 0.0738101f, -0.717172f, 0.473243f, 0.0184415f, -0.134346f, + 0.2701f, 0.366127f, -0.156723f, 0.640883f, 0.936974f, 0.490814f, -0.443443f, -0.618102f, + -0.416243f, 0.530263f, -0.103261f, -0.452055f, -0.204496f, -0.344118f, 0.556675f, 0.765497f, + -0.363246f, -0.722693f, 0.237699f, -0.751804f, -0.00057894f, 0.533425f, 0.658105f, 0.844047f, + -0.988591f, -0.773767f, -0.4548f, -0.0929374f, -0.149601f, -0.693741f, 1.52392f, 1.49526f, + 0.274482f, -0.250754f, -0.942429f, 0.11745f, 0.474299f, 0.649417f, -0.374675f, 0.770361f, + -0.212086f, 1.35768f, -0.536634f, 0.491991f, -1.47631f, -0.431231f, -0.18399f, -0.279197f, + 0.157382f, 0.89363f, -0.367624f, -1.11093f, -0.745738f, 0.548934f, 0.435344f, 0.776261f, + -0.150242f, -1.08683f, 0.513403f, 0.946813f, 0.290638f, 1.04797f, -0.72787f, -0.380246f, + 0.710451f, 0.222916f, 0.232908f, 0.830753f, 1.07542f, -0.374207f, -0.852818f, -0.615864f, + 0.63667f, 0.00779735f, 0.342196f, -0.403718f, 0.44006f, 0.536458f, -0.0573611f, -0.114347f, + 0.586737f, 0.113243f, 0.566012f, 1.02599f, 0.567469f, -0.516206f, -0.145892f, 0.0672092f, + -0.133495f, -0.631785f, 0.941846f, 0.188422f, 0.296014f, 0.423775f, -0.333609f, -0.23043f, + 0.416269f, -0.456101f, 0.0498845f, 0.211577f, 0.271048f, 0.626298f, -0.0416111f, 0.523762f, + -0.317144f, -0.0728705f, -0.14316f, -0.275855f, -0.295673f, -0.341217f, 0.436097f, -0.7442f, + 0.278465f, -0.280928f, -0.0415335f, 0.115429f, -0.625263f, 0.110212f, -0.195976f, -0.29027f, + -0.0989828f, -0.160014f, 0.362077f, 0.0649763f, -0.371465f, 0.727401f, 0.0320011f}; + float epsilon = 1e-05f; + + OpTester test("BatchNormalization"); + test.AddAttribute("epsilon", epsilon); + + vector input_shape{2, 3, 6, 6}; + int input_size = 2 * 3 * 6 * 6; + + vector f_X(input_size); + vector f_output(input_size); + vector f_scale(3); + vector f_B(3); + vector f_mean(3); + vector f_var(3); + + ConvertFloatToMLFloat16(X.data(), f_X.data(), input_size); + ConvertFloatToMLFloat16(scale.data(), f_scale.data(), 3); + ConvertFloatToMLFloat16(B.data(), f_B.data(), 3); + ConvertFloatToMLFloat16(mean.data(), f_mean.data(), 3); + ConvertFloatToMLFloat16(var.data(), f_var.data(), 3); + ConvertFloatToMLFloat16(expected_output.data(), f_output.data(), input_size); + + test.AddInput("X", input_shape, f_X); + test.AddInput("scale", {3}, f_scale); + test.AddInput("B", {3}, f_B); + test.AddInput("mean", {3}, f_mean); + test.AddInput("var", {3}, f_var); + test.AddOutput("output", input_shape, f_output); + test.Run(OpTester::ExpectResult::kExpectSuccess, ""); +} +#endif + } // namespace test } // namespace onnxruntime diff --git a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc index f2b08680d6..3d92157c2b 100644 --- a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc +++ b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc @@ -52,6 +52,60 @@ TEST(PoolTest, MaxPool) { test.Run(); } +// Only CUDA kernel has float 16 support +// Disable for now, still investigating the issue with cudnn lib +#ifdef USE_CUDA +TEST(PoolTest, MaxPool_F16) { + OpTester test("MaxPool"); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{1, 1}); + test.AddAttribute("pads", vector{0, 0, 0, 0}); + test.AddAttribute("kernel_shape", vector{8, 8}); + + std::vector x_vals = { + 0.19151945412158966, 0.6221087574958801, 0.43772774934768677, 0.7853586077690125, 0.7799758315086365, 0.27259260416030884, 0.2764642536640167, 0.801872193813324, + 0.9581393599510193, 0.8759326338768005, 0.35781726241111755, 0.5009950995445251, 0.683462917804718, 0.7127020359039307, 0.37025076150894165, 0.5611962080001831, + 0.5030831694602966, 0.013768449425697327, 0.772826611995697, 0.8826411962509155, 0.36488598585128784, 0.6153962016105652, 0.07538124173879623, 0.3688240051269531, + 0.9331400990486145, 0.6513781547546387, 0.39720258116722107, 0.7887301445007324, 0.3168361186981201, 0.5680986642837524, 0.8691273927688599, 0.4361734092235565, + 0.802147626876831, 0.14376682043075562, 0.7042609453201294, 0.7045813202857971, 0.2187921106815338, 0.9248676300048828, 0.44214075803756714, 0.9093159437179565, + 0.05980922281742096, 0.18428708612918854, 0.047355279326438904, 0.6748809218406677, 0.5946247577667236, 0.5333101749420166, 0.043324064463377, 0.5614330768585205, + 0.32966843247413635, 0.5029668211936951, 0.11189431697130203, 0.6071937084197998, 0.5659446716308594, 0.006764062214642763, 0.617441713809967, 0.912122905254364, + 0.7905241250991821, 0.9920814633369446, 0.9588017463684082, 0.7919641137123108, 0.2852509617805481, 0.6249167323112488, 0.47809380292892456, 0.19567517936229706, + + 0.382317453622818, 0.053873684257268906, 0.45164841413497925, 0.9820047616958618, 0.12394270300865173, 0.1193808987736702, 0.7385230660438538, 0.587303638458252, + 0.47163254022598267, 0.10712681710720062, 0.22921857237815857, 0.8999651670455933, 0.41675353050231934, 0.5358516573905945, 0.0062085166573524475, 0.3006417155265808, + 0.43689316511154175, 0.6121490001678467, 0.9181980490684509, 0.625736653804779, 0.7059975862503052, 0.14983370900154114, 0.7460634112358093, 0.8310070037841797, + 0.6337257623672485, 0.4383098781108856, 0.15257278084754944, 0.5684096217155457, 0.5282242894172668, 0.9514287710189819, 0.48035916686058044, 0.5025595426559448, + 0.5368781685829163, 0.8192020654678345, 0.05711563676595688, 0.6694217324256897, 0.7671166062355042, 0.7081153392791748, 0.7968671917915344, 0.5577608346939087, + 0.9658365249633789, 0.14715689420700073, 0.02964700013399124, 0.5938934683799744, 0.11406569927930832, 0.9508098363876343, 0.32570740580558777, 0.19361868500709534, + 0.4578116536140442, 0.9204025864601135, 0.8790691494941711, 0.252615749835968, 0.34800878167152405, 0.18258872628211975, 0.9017960429191589, 0.7065281867980957, + 0.7266584634780884, 0.900087833404541, 0.7791637778282166, 0.5991547703742981, 0.29112523794174194, 0.1513952612876892, 0.33517464995384216, 0.6575517654418945, + + 0.07334254682064056, 0.055006396025419235, 0.32319480180740356, 0.5904818177223206, 0.8538985848426819, 0.2870624363422394, 0.17306722700595856, 0.13402120769023895, + 0.9946538209915161, 0.1794978678226471, 0.3175468146800995, 0.568291425704956, 0.009348574094474316, 0.9006485939025879, 0.9772414565086365, 0.5568946599960327, + 0.08477384597063065, 0.3330024778842926, 0.7284286618232727, 0.14243537187576294, 0.5524689555168152, 0.2730432450771332, 0.9744951128959656, 0.6677868962287903, + 0.2556532919406891, 0.1083114966750145, 0.7761807441711426, 0.7824779748916626, 0.7616038918495178, 0.9144031405448914, 0.6586228013038635, 0.568367600440979, + 0.20175568759441376, 0.6982963681221008, 0.952195405960083, 0.8899632692337036, 0.9935673475265503, 0.8187035322189331, 0.5451221466064453, 0.45125406980514526, + 0.8905571699142456, 0.9732648134231567, 0.5934113264083862, 0.36607450246810913, 0.3230946958065033, 0.8714232444763184, 0.2156340628862381, 0.7349451780319214, + 0.36561909317970276, 0.8016026020050049, 0.7827355861663818, 0.7013553977012634, 0.6227765679359436, 0.4936826527118683, 0.8405377268791199, 0.7120969891548157, + 0.4439089894294739, 0.031034860759973526, 0.36323976516723633, 0.7307217717170715, 0.475566565990448, 0.3444169759750366, 0.6408804059028625, 0.12620532512664795}; + std::vector x_dims = {1, 3, 8, 8}; + int x_size = 1 * 3 * 8 * 8; + std::vector expected_dims = {1, 3, 1, 1}; + std::vector expected_vals = {0.9920814633369446, 0.9820047616958618, 0.9946538209915161}; + + std::vector f_X(x_size); + std::vector f_Y(3); + ConvertFloatToMLFloat16(x_vals.data(), f_X.data(), x_size); + ConvertFloatToMLFloat16(expected_vals.data(), f_Y.data(), 3); + + test.AddInput("X", x_dims, f_X); + test.AddOutput("Y", expected_dims, f_Y); + test.Run(); +} +#endif + static void MaxPool_8_WithIndexTest(bool has_index, int64_t storage_order=0) { OpTester test("MaxPool", 8); diff --git a/onnxruntime/test/providers/provider_test_utils.cc b/onnxruntime/test/providers/provider_test_utils.cc index 4cc35d1ef9..62f4f3b62d 100644 --- a/onnxruntime/test/providers/provider_test_utils.cc +++ b/onnxruntime/test/providers/provider_test_utils.cc @@ -68,6 +68,29 @@ void Check(const OpTester::Data& expected_data, const Tensor& output_tens } } +template <> +void Check(const OpTester::Data& expected_data, const Tensor& output_tensor, const std::string& provider_type) { + auto& expected_tensor = expected_data.data_.Get(); + auto* expected = expected_tensor.template Data(); + auto* output = output_tensor.template Data(); + auto size = output_tensor.Shape().Size(); + + std::vector f_expected(size); + std::vector f_output(size); + ConvertMLFloat16ToFloat(expected, f_expected.data(), static_cast(size)); + ConvertMLFloat16ToFloat(output, f_output.data(), static_cast(size)); + + float threshold = 0.001f; + for (int i = 0; i < size; ++i) { + if (std::isinf(f_expected[i])) // Test infinity for equality + EXPECT_EQ(f_expected[i], f_output[i]); + else { + // the default for existing tests + EXPECT_NEAR(f_expected[i], f_output[i], threshold) << "provider_type: " << provider_type; + } + } +} + template void CheckDispatch(MLDataType type, const OpTester::Data& expected_data, const Tensor& output_tensor, const std::string& provider_type) { if (type == DataTypeImpl::GetType()) diff --git a/onnxruntime/test/providers/provider_test_utils.h b/onnxruntime/test/providers/provider_test_utils.h index fe356b29fe..11caaf4e2b 100644 --- a/onnxruntime/test/providers/provider_test_utils.h +++ b/onnxruntime/test/providers/provider_test_utils.h @@ -20,6 +20,7 @@ #include "gmock/gmock.h" #include "gtest/gtest.h" #include +#include "core/util/math_cpuonly.h" namespace onnxruntime { class InferenceSession; @@ -347,5 +348,20 @@ void DebugTrap(); void Check(const OpTester::Data& expected_data, const Tensor& output_tensor, const std::string& provider_type); +// Only used for CUDA test since no toher kernel has float 16 support +#ifdef USE_CUDA +inline void ConvertFloatToMLFloat16(const float* f_datat, MLFloat16* h_data, int input_size) { + auto in_vector = ConstEigenVectorMap(f_datat, input_size); + auto output_vector = EigenVectorMap(static_cast(static_cast(h_data)), input_size); + output_vector = in_vector.template cast(); +} +#endif + +inline void ConvertMLFloat16ToFloat(const MLFloat16* h_data, float* f_data, int input_size) { + auto in_vector = ConstEigenVectorMap(static_cast(static_cast(h_data)), input_size); + auto output_vector = EigenVectorMap(f_data, input_size); + output_vector = in_vector.template cast(); +} + } // namespace test } // namespace onnxruntime