Bring back SLN cuda kernel and use provider options to switch to standard implementation (#15660)

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 ''"},

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.
 };

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 <typename T, bool Simplified>
 SkipLayerNorm<T, Simplified>::SkipLayerNorm(const OpKernelInfo& op_kernel_info) : CudaKernel(op_kernel_info) {
   ORT_ENFORCE(op_kernel_info.GetAttr<float>("epsilon", &epsilon_).IsOK());
   ORT_ENFORCE(epsilon_ >= 0);
+
+  const CUDAExecutionProvider* cuda_ep = static_cast<const CUDAExecutionProvider*>(op_kernel_info.GetExecutionProvider());
+
+  strict_ = cuda_ep->IsSkipLayerNormInStrictMode();
 }
 
 template <typename T, bool Simplified>
@@ -109,23 +114,46 @@ Status SkipLayerNorm<T, Simplified>::ComputeInternal(OpKernelContext* ctx) const
     }
   }
 
-  int row_count = gsl::narrow<int>(input->Shape().SizeToDimension(input_dims_size - 1));
+  if (strict_) {
+    int row_count = gsl::narrow<int>(input->Shape().SizeToDimension(input_dims_size - 1));
+    typedef typename ToCudaType<T>::MappedType CudaT;
+    HostApplyLayerNorm<CudaT, float, CudaT, Simplified>(
+        GetDeviceProp(),
+        Stream(ctx),
+        reinterpret_cast<CudaT*>(output->MutableData<T>()),                             // Y_data
+        nullptr,                                                                        // mean_data
+        nullptr,                                                                        // inv_var_data
+        reinterpret_cast<const CudaT*>(input->Data<T>()),                               // X_data
+        row_count,                                                                      // n1
+        hidden_size,                                                                    // n2
+        (double)epsilon_,                                                               // epsilon
+        reinterpret_cast<const CudaT*>(gamma->Data<T>()),                               // gamma
+        (beta != nullptr) ? reinterpret_cast<const CudaT*>(beta->Data<T>()) : nullptr,  // beta
+        reinterpret_cast<const CudaT*>(skip->Data<T>()),                                // skip or residual to add
+        (bias != nullptr) ? reinterpret_cast<const CudaT*>(bias->Data<T>()) : nullptr,  // bias to add
+        skip_input_bias_add_output != nullptr ? reinterpret_cast<CudaT*>(skip_input_bias_add_output->MutableData<T>()) : nullptr);
+
+    CUDA_RETURN_IF_ERROR(cudaGetLastError());
+    return Status::OK();
+  }
+
+  int sequence_length = static_cast<int>(input_dims[1]);
+  int64_t element_count = input_dims[0] * sequence_length * hidden_size;
+  size_t element_size = sizeof(T);
   typedef typename ToCudaType<T>::MappedType CudaT;
-  HostApplyLayerNorm<CudaT, float, CudaT, Simplified>(
-      GetDeviceProp(),
+  return LaunchSkipLayerNormKernel<CudaT, Simplified>(
       Stream(ctx),
-      reinterpret_cast<CudaT*>(output->MutableData<T>()),                             // Y_data
-      nullptr,                                                                        // mean_data
-      nullptr,                                                                        // inv_var_data
-      reinterpret_cast<const CudaT*>(input->Data<T>()),                               // X_data
-      row_count,                                                                      // n1
-      hidden_size,                                                                    // n2
-      (double)epsilon_,                                                               // epsilon
-      reinterpret_cast<const CudaT*>(gamma->Data<T>()),                               // gamma
-      (beta != nullptr) ? reinterpret_cast<const CudaT*>(beta->Data<T>()) : nullptr,  // beta
-      reinterpret_cast<const CudaT*>(skip->Data<T>()),                                // skip or residual to add
-      (bias != nullptr) ? reinterpret_cast<const CudaT*>(bias->Data<T>()) : nullptr,  // bias to add
-      skip_input_bias_add_output != nullptr ? reinterpret_cast<CudaT*>(skip_input_bias_add_output->MutableData<T>()) : nullptr);
+      reinterpret_cast<CudaT*>(output->MutableData<T>()),
+      skip_input_bias_add_output != nullptr ? reinterpret_cast<CudaT*>(skip_input_bias_add_output->MutableData<T>()) : nullptr,
+      reinterpret_cast<const CudaT*>(input->Data<T>()),
+      reinterpret_cast<const CudaT*>(skip->Data<T>()),
+      reinterpret_cast<const CudaT*>(gamma->Data<T>()),
+      (beta != nullptr) ? reinterpret_cast<const CudaT*>(beta->Data<T>()) : nullptr,
+      (bias != nullptr) ? reinterpret_cast<const CudaT*>(bias->Data<T>()) : nullptr,
+      epsilon_,
+      hidden_size,
+      static_cast<int>(element_count),
+      element_size);
 
   CUDA_RETURN_IF_ERROR(cudaGetLastError());
   return Status::OK();

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

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 <cuda_fp16.h>
+
+namespace onnxruntime {
+namespace contrib {
+namespace cuda {
+
+namespace {
+template <typename T>
+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 <typename T, int NumUnroll>
+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<aligned_vector<T, NumUnroll>>::value;
+  return ld % NumUnroll == 0 && reinterpret_cast<uint64_t>(output) % alignment == 0 &&
+         reinterpret_cast<uint64_t>(skip_input_bias_add_output) % alignment == 0 &&
+         reinterpret_cast<uint64_t>(input) % alignment == 0 && reinterpret_cast<uint64_t>(skip) % alignment == 0 &&
+         reinterpret_cast<uint64_t>(gamma) % alignment == 0 && reinterpret_cast<uint64_t>(beta) % alignment == 0 &&
+         reinterpret_cast<uint64_t>(bias) % alignment == 0 && next_size / NumUnroll >= kMinBlockSize &&
+         next_size / NumUnroll <= kMaxBlockSize;
+}
+}  // namespace
+
+template <typename T, unsigned TPB, bool Simplified>
+__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<T, T> 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<T, T>(rldval, rldval * val));
+
+    if (skip_input_bias_add_output != nullptr) {
+      skip_input_bias_add_output[idx] = val;
+    }
+
+    output[idx] = val;
+  }
+  if (Simplified) {
+    SimplifiedLayerNorm<T, TPB>(thread_data.value, ld, offset, gamma, epsilon, output);
+    return;
+  }
+  LayerNorm<T, TPB>(thread_data, ld, offset, beta, gamma, epsilon, output);
+}
+
+// Vectorized kernel
+template <typename T, unsigned TPB, int ILP, bool Simplified>
+__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, ILP>;
+
+  T input_v[ILP], skip_v[ILP], bias_v[ILP], skip_input_bias_add_output_v[ILP];
+
+  VecT* input_val = reinterpret_cast<VecT*>(&input_v);
+  *input_val = *reinterpret_cast<const VecT*>(&input[idx]);
+
+  VecT* skip_val = reinterpret_cast<VecT*>(&skip_v);
+  *skip_val = *reinterpret_cast<const VecT*>(&skip[idx]);
+
+  if (hasBias) {
+    VecT* bias_val = reinterpret_cast<VecT*>(&bias_v);
+    *bias_val = *reinterpret_cast<const VecT*>(&bias[threadIdx.x * ILP]);
+  }
+
+  cub::KeyValuePair<T, T> 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<VecT*>(&skip_input_bias_add_output[idx])) = *reinterpret_cast<VecT*>(&skip_input_bias_add_output_v);
+    }
+
+    thread_data = cub::KeyValuePair<T, T>(rldval_sum, rldvalsq_sum);
+  }
+
+  if (Simplified) {
+    SimplifiedLayerNormSmall<T, TPB, ILP>(input_v, thread_data.value, ld, idx, gamma, epsilon, output);
+    return;
+  }
+  LayerNormSmall<T, TPB, ILP>(input_v, thread_data, ld, idx, beta, gamma, epsilon, output);
+}
+
+template <typename T, bool Simplified>
+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<T, 2>(output, skip_input_bias_add_output, input, skip, gamma, beta, bias, ld, next_size);
+  bool flag_vec4 =
+      CanVectorized<T, 4>(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<T, block_size, num_unroll, Simplified>                                                \
+      <<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias, maybe2half<T>(epsilon), output, \
+                                             skip_input_bias_add_output, hasBias, hasSkipInputBiasAdditionOutput)
+#define LAUNCH_SKIP_LAYER_NORM_KERNEL()                                                       \
+  SkipLayerNormKernel<T, kMaxBlockSize, Simplified><<<grid_size, kMaxBlockSize, 0, stream>>>( \
+      ld, input, skip, beta, gamma, bias, maybe2half<T>(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<T, Simplified>(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

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 <typename T, bool Simplified>
+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

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_);

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;

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);

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<CUDAProviderFactory>(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 {

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);

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 = []

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<OrtCUDAProviderOptionsV2, decltype(api.ReleaseCUDAProviderOptions)>
+            rel_cuda_options(cuda_options, api.ReleaseCUDAProviderOptions);
+        std::vector<const char*> keys{"enable_skip_layer_norm_strict_mode"};
+        std::vector<const char*> 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) {

onnxruntime/test/util/default_providers.cc

@@ -110,6 +110,16 @@ std::unique_ptr<IExecutionProvider> DefaultCudaExecutionProvider() {
   return nullptr;
 }
 
+std::unique_ptr<IExecutionProvider> 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<IExecutionProvider> DefaultDnnlExecutionProvider() {
 #ifdef USE_DNNL
   OrtDnnlProviderOptions dnnl_options;

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<IExecutionProvider> DefaultCpuExecutionProvider(bool enable_arena = true);
 std::unique_ptr<IExecutionProvider> DefaultCudaExecutionProvider();
+std::unique_ptr<IExecutionProvider> CudaExecutionProviderWithOptions(const OrtCUDAProviderOptionsV2* provider_options);
 std::unique_ptr<IExecutionProvider> DefaultDnnlExecutionProvider();
 std::unique_ptr<IExecutionProvider> DnnlExecutionProviderWithOptions(const OrtDnnlProviderOptions* provider_options);
 // std::unique_ptr<IExecutionProvider> DefaultTvmExecutionProvider();