diff --git a/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs b/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs index 5132e267b1..0e1e5ea5ea 100644 --- a/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs +++ b/csharp/test/Microsoft.ML.OnnxRuntime.Tests.NetCoreApp/InferenceTest.netcore.cs @@ -308,7 +308,7 @@ namespace Microsoft.ML.OnnxRuntime.Tests // the expansion of Softplus uses Exp(1). ORT has a Softplus kernel, so testing the expansion is // unnecessary and fails as ORT support for Exp started at opset 6 (as ORT didn't exist until opset 7). - + { "test_clip_default_int8_max_expanded", "Could not find an implementation for Less(13) nodeMeta with name ''" }, { "test_softplus_expanded", "Could not find an implementation for Exp(1) node with name ''"}, { "test_softplus_example_expanded", "Could not find an implementation for Exp(1) node with name ''"}, diff --git a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h index 52d522b312..18dcfeea68 100644 --- a/include/onnxruntime/core/providers/cuda/cuda_provider_options.h +++ b/include/onnxruntime/core/providers/cuda/cuda_provider_options.h @@ -29,4 +29,6 @@ struct OrtCUDAProviderOptionsV2 { int cudnn_conv1d_pad_to_nc1d; // flag specifying if pad Conv1D's input [N,C,D] to [N,C,1,D] or [N,C,D,1]. int tunable_op_enable; // flag specifying if TunableOp is enabled. int tunable_op_tuning_enable; // flag specifying if TunableOp is enabled for tuning, this relies on TunableOp is enabled. + int enable_skip_layer_norm_strict_mode; // flag specifying if SkipLayerNorm is in strict mode. If true, use LayerNormalization kernel. + // The strict mode has better accuracy but lower performance. }; diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc index 4cf45b5d33..62f93b5e5b 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc @@ -4,6 +4,7 @@ #include "core/providers/cuda/cuda_common.h" #include "core/providers/cuda/nn/layer_norm_impl.h" #include "skip_layer_norm.h" +#include "skip_layer_norm_impl.h" namespace onnxruntime { namespace contrib { @@ -38,6 +39,10 @@ template SkipLayerNorm::SkipLayerNorm(const OpKernelInfo& op_kernel_info) : CudaKernel(op_kernel_info) { ORT_ENFORCE(op_kernel_info.GetAttr("epsilon", &epsilon_).IsOK()); ORT_ENFORCE(epsilon_ >= 0); + + const CUDAExecutionProvider* cuda_ep = static_cast(op_kernel_info.GetExecutionProvider()); + + strict_ = cuda_ep->IsSkipLayerNormInStrictMode(); } template @@ -109,23 +114,46 @@ Status SkipLayerNorm::ComputeInternal(OpKernelContext* ctx) const } } - int row_count = gsl::narrow(input->Shape().SizeToDimension(input_dims_size - 1)); + if (strict_) { + int row_count = gsl::narrow(input->Shape().SizeToDimension(input_dims_size - 1)); + typedef typename ToCudaType::MappedType CudaT; + HostApplyLayerNorm( + GetDeviceProp(), + Stream(ctx), + reinterpret_cast(output->MutableData()), // Y_data + nullptr, // mean_data + nullptr, // inv_var_data + reinterpret_cast(input->Data()), // X_data + row_count, // n1 + hidden_size, // n2 + (double)epsilon_, // epsilon + reinterpret_cast(gamma->Data()), // gamma + (beta != nullptr) ? reinterpret_cast(beta->Data()) : nullptr, // beta + reinterpret_cast(skip->Data()), // skip or residual to add + (bias != nullptr) ? reinterpret_cast(bias->Data()) : nullptr, // bias to add + skip_input_bias_add_output != nullptr ? reinterpret_cast(skip_input_bias_add_output->MutableData()) : nullptr); + + CUDA_RETURN_IF_ERROR(cudaGetLastError()); + return Status::OK(); + } + + int sequence_length = static_cast(input_dims[1]); + int64_t element_count = input_dims[0] * sequence_length * hidden_size; + size_t element_size = sizeof(T); typedef typename ToCudaType::MappedType CudaT; - HostApplyLayerNorm( - GetDeviceProp(), + return LaunchSkipLayerNormKernel( Stream(ctx), - reinterpret_cast(output->MutableData()), // Y_data - nullptr, // mean_data - nullptr, // inv_var_data - reinterpret_cast(input->Data()), // X_data - row_count, // n1 - hidden_size, // n2 - (double)epsilon_, // epsilon - reinterpret_cast(gamma->Data()), // gamma - (beta != nullptr) ? reinterpret_cast(beta->Data()) : nullptr, // beta - reinterpret_cast(skip->Data()), // skip or residual to add - (bias != nullptr) ? reinterpret_cast(bias->Data()) : nullptr, // bias to add - skip_input_bias_add_output != nullptr ? reinterpret_cast(skip_input_bias_add_output->MutableData()) : nullptr); + reinterpret_cast(output->MutableData()), + skip_input_bias_add_output != nullptr ? reinterpret_cast(skip_input_bias_add_output->MutableData()) : nullptr, + reinterpret_cast(input->Data()), + reinterpret_cast(skip->Data()), + reinterpret_cast(gamma->Data()), + (beta != nullptr) ? reinterpret_cast(beta->Data()) : nullptr, + (bias != nullptr) ? reinterpret_cast(bias->Data()) : nullptr, + epsilon_, + hidden_size, + static_cast(element_count), + element_size); CUDA_RETURN_IF_ERROR(cudaGetLastError()); return Status::OK(); diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.h b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.h index 9d47373fba..06b3945427 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.h +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.h @@ -19,6 +19,7 @@ class SkipLayerNorm final : public CudaKernel { private: float epsilon_; + bool strict_; }; } // namespace cuda diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu new file mode 100644 index 0000000000..0a815f7f49 --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu @@ -0,0 +1,241 @@ +/* + The implementation of this file is based on skipLayerNorm plugin in TensorRT demo: + https://github.com/NVIDIA/TensorRT/tree/release/5.1/demo/BERT/ + +Copyright 2019 NVIDIA Corporation + +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + + http://www.apache.org/licenses/LICENSE-2.0 + +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +*/ + +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +// Modifications: Add SkipLayerNormKernelVec to +// leverage vectorized load/write. +// and templatize ComputeSkipLayerNorm for different +// data types. +// Copyright (c) Advanced Micro Devices, Inc. All rights reserved. +// Licensed under the MIT License. + +#include "contrib_ops/cuda/bert/layer_norm.cuh" +#include "contrib_ops/cuda/bert/skip_layer_norm_impl.h" +#include + +namespace onnxruntime { +namespace contrib { +namespace cuda { + +namespace { +template +T maybe2half(float x); + +template <> +float maybe2half(float x) { + return x; +} + +template <> +half maybe2half(float x) { + return __float2half_rn(x); +} + +// Using only power of 2 numbers will lead to waste of compute for same size such as 768, which is a very common case +// in BERT. Ideally we can step by wrap_size * num_unroll, but listing too many steps will cause long compile time. +constexpr int kSizes[] = {32, 64, 128, 384, 768, 1024, 2048}; +constexpr int kMinBlockSize = 32; +constexpr int kMaxBlockSize = 256; + +int NextSize(int x) { + size_t len = sizeof(kSizes) / sizeof(kSizes[0]); + for (size_t i = 0; i < len; ++i) { + if (x <= kSizes[i]) { + return kSizes[i]; + } + } + return kSizes[len - 1]; +} + +template +bool CanVectorized(T* output, T* skip_input_bias_add_output, const T* input, const T* skip, const T* gamma, + const T* beta, const T* bias, const int ld, const int next_size) { + constexpr int alignment = std::alignment_of>::value; + return ld % NumUnroll == 0 && reinterpret_cast(output) % alignment == 0 && + reinterpret_cast(skip_input_bias_add_output) % alignment == 0 && + reinterpret_cast(input) % alignment == 0 && reinterpret_cast(skip) % alignment == 0 && + reinterpret_cast(gamma) % alignment == 0 && reinterpret_cast(beta) % alignment == 0 && + reinterpret_cast(bias) % alignment == 0 && next_size / NumUnroll >= kMinBlockSize && + next_size / NumUnroll <= kMaxBlockSize; +} +} // namespace + +template +__global__ void SkipLayerNormKernel( + const int ld, const T* input, const T* skip, + const T* beta, const T* gamma, const T* bias, + const T epsilon, T* output, T* skip_input_bias_add_output) { + const T reverse_ld = T(1.f / ld); + const int offset = blockIdx.x * ld; + + KeyValuePairSum pair_sum; + // reduce x and x^2 + cub::KeyValuePair thread_data(0, 0); + + for (int i = threadIdx.x; i < ld; i += TPB) { + const int idx = offset + i; + + const T val = (bias == nullptr) ? input[idx] + skip[idx] : input[idx] + skip[idx] + bias[i]; + const T rldval = reverse_ld * val; + thread_data = pair_sum(thread_data, cub::KeyValuePair(rldval, rldval * val)); + + if (skip_input_bias_add_output != nullptr) { + skip_input_bias_add_output[idx] = val; + } + + output[idx] = val; + } + if (Simplified) { + SimplifiedLayerNorm(thread_data.value, ld, offset, gamma, epsilon, output); + return; + } + LayerNorm(thread_data, ld, offset, beta, gamma, epsilon, output); +} + +// Vectorized kernel +template +__global__ void SkipLayerNormKernelSmall( + const int ld, const T* input, const T* skip, const T* beta, const T* gamma, + const T* bias, const T epsilon, T* output, T* skip_input_bias_add_output, + bool hasBias, bool hasSkipInputBiasAdditionOutput) { + const T rld = T(1.f / ld); + const int idx = blockIdx.x * ld + threadIdx.x * ILP; // grid_size = n / ld + + using VecT = aligned_vector; + + T input_v[ILP], skip_v[ILP], bias_v[ILP], skip_input_bias_add_output_v[ILP]; + + VecT* input_val = reinterpret_cast(&input_v); + *input_val = *reinterpret_cast(&input[idx]); + + VecT* skip_val = reinterpret_cast(&skip_v); + *skip_val = *reinterpret_cast(&skip[idx]); + + if (hasBias) { + VecT* bias_val = reinterpret_cast(&bias_v); + *bias_val = *reinterpret_cast(&bias[threadIdx.x * ILP]); + } + + cub::KeyValuePair thread_data(T(0.f), T(0.f)); + + if (ILP * threadIdx.x < ld) { + T rldval_sum = T(0.f); + T rldvalsq_sum = T(0.f); +#pragma unroll + for (int i = 0; i < ILP; i++) { + input_v[i] += hasBias ? skip_v[i] + bias_v[i] : skip_v[i]; + + if (hasSkipInputBiasAdditionOutput) { + skip_input_bias_add_output_v[i] = input_v[i]; + } + + const T rldval = rld * input_v[i]; + rldval_sum += rldval; + rldvalsq_sum += rldval * input_v[i]; + } + + if (hasSkipInputBiasAdditionOutput) { + *(reinterpret_cast(&skip_input_bias_add_output[idx])) = *reinterpret_cast(&skip_input_bias_add_output_v); + } + + thread_data = cub::KeyValuePair(rldval_sum, rldvalsq_sum); + } + + if (Simplified) { + SimplifiedLayerNormSmall(input_v, thread_data.value, ld, idx, gamma, epsilon, output); + return; + } + LayerNormSmall(input_v, thread_data, ld, idx, beta, gamma, epsilon, output); +} + +template +Status LaunchSkipLayerNormKernel( + cudaStream_t stream, T* output, T* skip_input_bias_add_output, const T* input, const T* skip, const T* gamma, + const T* beta, const T* bias, float epsilon, const int ld, const int element_count, + size_t element_size) { + // this must be true because n is the total size of the tensor + assert(element_count % ld == 0); + bool hasBias = (bias == nullptr) ? false : true; + bool hasSkipInputBiasAdditionOutput = (skip_input_bias_add_output == nullptr) ? false : true; + + const int next_size = NextSize(ld); + const int grid_size = element_count / ld; + bool flag_vec2 = + CanVectorized(output, skip_input_bias_add_output, input, skip, gamma, beta, bias, ld, next_size); + bool flag_vec4 = + CanVectorized(output, skip_input_bias_add_output, input, skip, gamma, beta, bias, ld, next_size); + + switch (next_size) { +#define LAUNCH_SKIP_LAYER_NORM_KERNEL_SMALL(num_unroll) \ + SkipLayerNormKernelSmall \ + <<>>(ld, input, skip, beta, gamma, bias, maybe2half(epsilon), output, \ + skip_input_bias_add_output, hasBias, hasSkipInputBiasAdditionOutput) +#define LAUNCH_SKIP_LAYER_NORM_KERNEL() \ + SkipLayerNormKernel<<>>( \ + ld, input, skip, beta, gamma, bias, maybe2half(epsilon), output, skip_input_bias_add_output) +#define CASE_NEXT_SIZE(next_size_value) \ + case next_size_value: { \ + if (flag_vec4) { \ + constexpr int block_size = next_size_value / 4; \ + LAUNCH_SKIP_LAYER_NORM_KERNEL_SMALL(4); \ + } else if (flag_vec2) { \ + constexpr int block_size = next_size_value / 2; \ + LAUNCH_SKIP_LAYER_NORM_KERNEL_SMALL(2); \ + } else { \ + if (next_size_value <= kMaxBlockSize) { \ + constexpr int block_size = next_size_value; \ + LAUNCH_SKIP_LAYER_NORM_KERNEL_SMALL(1); \ + } else { \ + LAUNCH_SKIP_LAYER_NORM_KERNEL(); \ + } \ + } \ + } break + CASE_NEXT_SIZE(kSizes[0]); + CASE_NEXT_SIZE(kSizes[1]); + CASE_NEXT_SIZE(kSizes[2]); + CASE_NEXT_SIZE(kSizes[3]); + CASE_NEXT_SIZE(kSizes[4]); + CASE_NEXT_SIZE(kSizes[5]); + CASE_NEXT_SIZE(kSizes[6]); +#undef CASE_NEXT_SIZE +#undef LAUNCH_SKIP_LAYER_NORM_KERNEL +#undef LAUNCH_SKIP_LAYER_NORM_KERNEL_SMALL + } + + return CUDA_CALL(cudaGetLastError()); +} + +#define SKIPLAYERNORM_IMPL(T, Simplified) \ + template Status LaunchSkipLayerNormKernel(cudaStream_t stream, T * output, \ + T * skip_input_bias_add_output, \ + const T* input, const T* skip, const T* gamma, \ + const T* beta, const T* bias, float epsilon, \ + const int ld, const int element_count, \ + size_t element_size); + +SKIPLAYERNORM_IMPL(float, true); +SKIPLAYERNORM_IMPL(float, false); +SKIPLAYERNORM_IMPL(half, true); +SKIPLAYERNORM_IMPL(half, false); + +} // namespace cuda +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.h b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.h new file mode 100644 index 0000000000..da2894928c --- /dev/null +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.h @@ -0,0 +1,28 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once +#include "core/common/common.h" + +namespace onnxruntime { +namespace contrib { +namespace cuda { + +template +Status LaunchSkipLayerNormKernel( + cudaStream_t stream, + T* output, // normalized output tensor + T* skip_input_bias_add_output, // sum of the input and skip (and bias if it exists) tensors output + const T* input, // input tensor + const T* skip, // skip tensor + const T* gamma, // Layer normalization gamma tensor + const T* beta, // Layer normalization beta tensor + const T* bias, // Layer normalization beta tensor + float epsilon, // Layer normalization epsilon + int hidden_size, // hidden size, it is the leading dimension (ld) + int element_count, // number of elements in input tensor + size_t element_size); + +} // namespace cuda +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.h b/onnxruntime/core/providers/cuda/cuda_execution_provider.h index 50f5e65110..63bf2a712e 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.h @@ -96,6 +96,7 @@ class CUDAExecutionProvider : public IExecutionProvider { bool DoCopyOnDefaultStream() const { return info_.do_copy_in_default_stream; } bool GetCudnnConvUseMaxWorkspace() const { return info_.cudnn_conv_use_max_workspace; } bool GetCudnnConv1dPadToNc1d() const { return info_.cudnn_conv1d_pad_to_nc1d; } + bool IsSkipLayerNormInStrictMode() const { return info_.enable_skip_layer_norm_strict_mode; } ProviderOptions GetProviderOptions() const override { return CUDAExecutionProviderInfo::ToProviderOptions(info_); diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc index eb257a652f..bb840e9de2 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.cc @@ -26,6 +26,7 @@ constexpr const char* kEnableCudaGraph = "enable_cuda_graph"; constexpr const char* kCudnnConv1dPadToNc1d = "cudnn_conv1d_pad_to_nc1d"; constexpr const char* kTunableOpEnable = "tunable_op_enable"; constexpr const char* kTunableOpTuningEnable = "tunable_op_tuning_enable"; +constexpr const char* kEnableSkipLayerNormStrictMode = "enable_skip_layer_norm_strict_mode"; } // namespace provider_option_names } // namespace cuda @@ -94,6 +95,7 @@ CUDAExecutionProviderInfo CUDAExecutionProviderInfo::FromProviderOptions(const P .AddAssignmentToReference(cuda::provider_option_names::kCudnnConvUseMaxWorkspace, info.cudnn_conv_use_max_workspace) .AddAssignmentToReference(cuda::provider_option_names::kEnableCudaGraph, info.enable_cuda_graph) .AddAssignmentToReference(cuda::provider_option_names::kCudnnConv1dPadToNc1d, info.cudnn_conv1d_pad_to_nc1d) + .AddAssignmentToReference(cuda::provider_option_names::kEnableSkipLayerNormStrictMode, info.enable_skip_layer_norm_strict_mode) .AddValueParser( cuda::provider_option_names::kTunableOpEnable, [&info](const std::string& value_str) -> Status { @@ -130,6 +132,7 @@ ProviderOptions CUDAExecutionProviderInfo::ToProviderOptions(const CUDAExecution {cuda::provider_option_names::kCudnnConv1dPadToNc1d, MakeStringWithClassicLocale(info.cudnn_conv1d_pad_to_nc1d)}, {cuda::provider_option_names::kTunableOpEnable, MakeStringWithClassicLocale(info.tunable_op.enable)}, {cuda::provider_option_names::kTunableOpTuningEnable, MakeStringWithClassicLocale(info.tunable_op.tuning_enable)}, + {cuda::provider_option_names::kEnableSkipLayerNormStrictMode, MakeStringWithClassicLocale(info.enable_skip_layer_norm_strict_mode)}, }; return options; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h index 6440967345..c268518974 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider_info.h @@ -69,6 +69,8 @@ struct CUDAExecutionProviderInfo { cuda::TunableOpInfo tunable_op{}; + bool enable_skip_layer_norm_strict_mode{false}; + static CUDAExecutionProviderInfo FromProviderOptions(const ProviderOptions& options); static ProviderOptions ToProviderOptions(const CUDAExecutionProviderInfo& info); static ProviderOptions ToProviderOptions(const OrtCUDAProviderOptionsV2& info); diff --git a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc index 5fb4f24014..6e8c4d0ff0 100644 --- a/onnxruntime/core/providers/cuda/cuda_provider_factory.cc +++ b/onnxruntime/core/providers/cuda/cuda_provider_factory.cc @@ -252,6 +252,7 @@ struct CUDA_Provider : Provider { info.cudnn_conv1d_pad_to_nc1d = params->cudnn_conv1d_pad_to_nc1d != 0; info.tunable_op.enable = params->tunable_op_enable; info.tunable_op.tuning_enable = params->tunable_op_tuning_enable; + info.enable_skip_layer_norm_strict_mode = params->enable_skip_layer_norm_strict_mode != 0; return std::make_shared(info); } @@ -271,6 +272,7 @@ struct CUDA_Provider : Provider { cuda_options.cudnn_conv_use_max_workspace = internal_options.cudnn_conv_use_max_workspace; cuda_options.enable_cuda_graph = internal_options.enable_cuda_graph; cuda_options.cudnn_conv1d_pad_to_nc1d = internal_options.cudnn_conv1d_pad_to_nc1d; + cuda_options.enable_skip_layer_norm_strict_mode = internal_options.enable_skip_layer_norm_strict_mode; } ProviderOptions GetProviderOptions(const void* provider_options) override { diff --git a/onnxruntime/core/session/provider_bridge_ort.cc b/onnxruntime/core/session/provider_bridge_ort.cc index 6f0fda7758..9a4fe6900c 100644 --- a/onnxruntime/core/session/provider_bridge_ort.cc +++ b/onnxruntime/core/session/provider_bridge_ort.cc @@ -1239,6 +1239,7 @@ OrtCUDAProviderOptionsV2 OrtCUDAProviderOptionsToOrtCUDAProviderOptionsV2(const cuda_options_converted.cudnn_conv_use_max_workspace = 1; cuda_options_converted.enable_cuda_graph = 0; cuda_options_converted.cudnn_conv1d_pad_to_nc1d = 0; + cuda_options_converted.enable_skip_layer_norm_strict_mode = 0; return cuda_options_converted; } @@ -1796,6 +1797,7 @@ ORT_API_STATUS_IMPL(OrtApis::CreateCUDAProviderOptions, _Outptr_ OrtCUDAProvider (*out)->cudnn_conv_use_max_workspace = 1; (*out)->enable_cuda_graph = 0; (*out)->cudnn_conv1d_pad_to_nc1d = 0; + (*out)->enable_skip_layer_norm_strict_mode = 0; return nullptr; #else ORT_UNUSED_PARAMETER(out); diff --git a/onnxruntime/python/tools/transformers/convert_generation.py b/onnxruntime/python/tools/transformers/convert_generation.py index 3de7968e94..5862806a65 100644 --- a/onnxruntime/python/tools/transformers/convert_generation.py +++ b/onnxruntime/python/tools/transformers/convert_generation.py @@ -8,25 +8,36 @@ This converts GPT2 or T5 model to onnx with beam search operator. Example 1: convert gpt2 model with beam search: python convert_generation.py -m gpt2 --output gpt2_beam_search.onnx -Example 2: convert T5 model with beam search in two steps: +Example 2: convert gpt2 model with beam search containing specific cuda optimizations: + python convert_generation.py -m gpt2 --output gpt2_beam_search.onnx --use_gpu \ + --past_present_share_buffer --use_decoder_masked_attention + +Example 3: convert gpt2 model with beam search with mixed precision and enable SkipLayerNorm strict mode: + python convert_generation.py -m gpt2 --output gpt2_beam_search.onnx --use_gpu -p fp16 --use_sln_strict_mode + +Example 4: convert T5 model with beam search in two steps: cd ./models/t5 python convert_to_onnx.py -m t5-small cd ../.. - python convert_generation.py -m t5-small --model_type t5 \ + python convert_generation.py -m t5-small --model_type t5 \ --decoder_onnx ./models/t5/onnx_models/t5-small_decoder.onnx \ --encoder_decoder_init_onnx ./models/t5/onnx_models/t5-small_encoder_decoder_init.onnx \ --output ./models/t5/onnx_models/t5_small_beam_search.onnx -Example 3: convert T5 model with beam search. All in one step: +Example 5: convert T5 model with beam search. All in one step: python convert_generation.py -m t5-small --model_type t5 --output ./models/t5/onnx_models/t5_small_beam_search.onnx -Example 4: convert MT5 model with external data file like mt5-base-beamsearch.onnx.data in below example. +Example 6: convert T5 model with beam search containing specific cuda optimizations. All in one step: + python convert_generation.py -m t5-small --model_type t5 --output ./models/t5/onnx_models/t5_small_beam_search.onnx \ + --use_gpu --past_present_share_buffer --use_decoder_masked_attention + +Example 7: convert MT5 model with external data file like mt5-base-beamsearch.onnx.data in below example. python convert_generation.py -m google/mt5-base --model_type mt5 --output mt5-base-beamsearch.onnx -e -Example 5: convert gpt2 model with greedy search: +Example 8: convert gpt2 model with greedy search: python convert_generation.py -m gpt2 --output gpt2_greedy_search.onnx --num_beams 1 --num_return_sequences 1 -Example 6: convert gpt2 model with sampling: +Example 9: convert gpt2 model with sampling: python convert_generation.py -m gpt2 --output gpt2_sampling.onnx --num_beams 1 --num_return_sequences 1 --top_p 0.6 """ @@ -412,6 +423,14 @@ def parse_arguments(argv: Optional[List[str]] = None) -> argparse.Namespace: test_group = parser.add_argument_group("Other options for testing parity and performance") + test_group.add_argument( + "--use_sln_strict_mode", + required=False, + action="store_true", + help="Enable strict mode for SLN in CUDA provider. This ensures a better accuracy but will be slower.", + ) + test_group.set_defaults(use_sln_strict_mode=False) + test_group.add_argument( "--use_gpu", required=False, action="store_true", help="use GPU for inference. Required for fp16." ) @@ -632,12 +651,13 @@ def pad_weights_of_logits_matmul(onnx_path: str, use_external_data_format: bool return True -def create_ort_session(model_path: str, use_gpu: bool) -> InferenceSession: +def create_ort_session(model_path: str, use_gpu: bool, use_sln_strict_mode: bool) -> InferenceSession: """Create OnnxRuntime session. Args: model_path (str): onnx model path use_gpu (bool): use GPU or not + use_sln_strict_mode (bool): use strict mode for skip layer normalization or not Raises: RuntimeError: CUDAExecutionProvider is not available when --use_gpu is specified. @@ -653,6 +673,12 @@ def create_ort_session(model_path: str, use_gpu: bool) -> InferenceSession: raise RuntimeError("CUDAExecutionProvider is not available for --use_gpu!") else: logger.info("use CUDAExecutionProvider") + if use_sln_strict_mode: + cuda_provider_options = {"enable_skip_layer_norm_strict_mode": True} + provider_options = {"CUDAExecutionProvider": cuda_provider_options} + execution_providers = [ + (name, provider_options[name]) if name in provider_options else name for name in execution_providers + ] ort_session = InferenceSession(model_path, sess_options, providers=execution_providers) return ort_session @@ -2346,7 +2372,7 @@ def test_gpt_model(args: argparse.Namespace, sentences: Optional[List[str]] = No return logger.debug("Creating ort session......") - ort_session = create_ort_session(args.output, args.use_gpu) + ort_session = create_ort_session(args.output, args.use_gpu, args.use_sln_strict_mode) logger.debug("Run ort session......") result = ort_session.run(None, inputs) @@ -2548,7 +2574,7 @@ def test_t5_model(args: argparse.Namespace, sentences: Optional[List[str]] = Non logger.debug("ORT inputs", inputs) # noqa: PLE1205 - ort_session = create_ort_session(args.output, args.use_gpu) + ort_session = create_ort_session(args.output, args.use_gpu, args.use_sln_strict_mode) # Test performance latency = [] diff --git a/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc b/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc index c3078bcbaa..67e279dc08 100644 --- a/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc +++ b/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "gtest/gtest.h" +#include "core/session/onnxruntime_cxx_api.h" #include "test/common/tensor_op_test_utils.h" #include "test/common/cuda_op_test_utils.h" #include "test/providers/provider_test_utils.h" @@ -25,7 +26,8 @@ static void RunTest( bool use_float16 = false, bool no_beta = false, bool simplified = false, - bool use_token_count = false) { + bool use_token_count = false, + bool strict = false) { // Input and output shapes // Input 0 - input: (batch_size, sequence_length, hidden_size) or (batch_size * sequence_length, hidden_size) // Input 1 - skip : (batch_size, sequence_length, hidden_size) or (batch_size * sequence_length, hidden_size) @@ -113,7 +115,19 @@ static void RunTest( } else if (rocm_ep != nullptr) { execution_providers.push_back(DefaultRocmExecutionProvider()); } else { - execution_providers.push_back(DefaultCudaExecutionProvider()); + if (strict) { + const auto& api = Ort::GetApi(); + OrtCUDAProviderOptionsV2* cuda_options = nullptr; + ASSERT_TRUE(api.CreateCUDAProviderOptions(&cuda_options) == nullptr); + std::unique_ptr + rel_cuda_options(cuda_options, api.ReleaseCUDAProviderOptions); + std::vector keys{"enable_skip_layer_norm_strict_mode"}; + std::vector values{"1"}; + ASSERT_TRUE(api.UpdateCUDAProviderOptions(rel_cuda_options.get(), keys.data(), values.data(), 1) == nullptr); + execution_providers.push_back(CudaExecutionProviderWithOptions(std::move(rel_cuda_options.get()))); + } else { + execution_providers.push_back(DefaultCudaExecutionProvider()); + } } test.Run(OpTester::ExpectResult::kExpectSuccess, "", {}, nullptr, &execution_providers); @@ -325,7 +339,11 @@ TEST(SkipLayerNormTest, SkipLayerNormBatch1_Float16_vec) { batch_size, sequence_length, hidden_size, - true); + true /*use_float16*/, + false /*no_beta*/, + false /*simplified*/, + false /*use_token_count*/, + true /*strict*/); } TEST(SkipLayerNormTest, SkipLayerNormBatch1_NoBeta) { @@ -610,10 +628,11 @@ TEST(SkipLayerNormTest, SkipLayerNormBatch1_Float16_vec_token_count) { batch_size, sequence_length, hidden_size, - true, - false, - false, - true); + true /*use_float16*/, + false /*no_beta*/, + false /*simplified*/, + true /*use_token_count*/, + true /*strict*/); } TEST(SkipLayerNormTest, SkipLayerNormBatch2_TokenCount) { diff --git a/onnxruntime/test/util/default_providers.cc b/onnxruntime/test/util/default_providers.cc index 5cb80f9edb..9f76c888c6 100644 --- a/onnxruntime/test/util/default_providers.cc +++ b/onnxruntime/test/util/default_providers.cc @@ -110,6 +110,16 @@ std::unique_ptr DefaultCudaExecutionProvider() { return nullptr; } +std::unique_ptr CudaExecutionProviderWithOptions(const OrtCUDAProviderOptionsV2* provider_options) { +#ifdef USE_CUDA + if (auto factory = CudaProviderFactoryCreator::Create(provider_options)) + return factory->CreateProvider(); +#else + ORT_UNUSED_PARAMETER(provider_options); +#endif + return nullptr; +} + std::unique_ptr DefaultDnnlExecutionProvider() { #ifdef USE_DNNL OrtDnnlProviderOptions dnnl_options; diff --git a/onnxruntime/test/util/include/default_providers.h b/onnxruntime/test/util/include/default_providers.h index e6ddf642b4..d6c9339af0 100644 --- a/onnxruntime/test/util/include/default_providers.h +++ b/onnxruntime/test/util/include/default_providers.h @@ -36,6 +36,7 @@ namespace test { // unique_ptr providers with default values for session registration std::unique_ptr DefaultCpuExecutionProvider(bool enable_arena = true); std::unique_ptr DefaultCudaExecutionProvider(); +std::unique_ptr CudaExecutionProviderWithOptions(const OrtCUDAProviderOptionsV2* provider_options); std::unique_ptr DefaultDnnlExecutionProvider(); std::unique_ptr DnnlExecutionProviderWithOptions(const OrtDnnlProviderOptions* provider_options); // std::unique_ptr DefaultTvmExecutionProvider();