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