From 835ecb264d3027b70a8d0cab9cb04ba190e15f46 Mon Sep 17 00:00:00 2001 From: Hubert Lu <55214931+hubertlu-tw@users.noreply.github.com> Date: Tue, 5 Jul 2022 22:28:15 -0700 Subject: [PATCH] Leverage vectorized load/write for SkipLayerNorm (#11803) * First attempt for half2 vectorized memory access in SkipLayerNorm * Add some functions for debugging * Clean up the code * Clean up the code * Generalize the vectorized kernels with aligned_vector and remove cudaDeviceProp * Add a unit test for a larger input size * Fix some Lint C++ warnings * Use ILP = 4 for the vectorized kernels * Rewrite the vectorized kernel and templatize ComputeSkipLayerNorm * Use conditional operator for input_v * Refactor LaunchSkipLayerNormKernel and replace the original SkipLayerNormKernelSmall with the vectorized kernel * Clean some comments and rename the layernorm function * Use ComputeSkipLayerNorm to replace LaunchSkipLayerNormKernel * Resolve a Lint C++ warning * Fix SkipLayerNormBatch1_Float16_vec output data --- .../contrib_ops/cuda/bert/layer_norm.cuh | 33 ++- .../contrib_ops/cuda/bert/skip_layer_norm.cc | 18 +- .../cuda/bert/skip_layer_norm_impl.cu | 223 +++++++++++------- .../cuda/bert/skip_layer_norm_impl.h | 16 +- .../test/contrib_ops/skiplayernorm_op_test.cc | 92 ++++++++ 5 files changed, 279 insertions(+), 103 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/bert/layer_norm.cuh b/onnxruntime/contrib_ops/cuda/bert/layer_norm.cuh index 54bc93467b..7d6cea8c09 100644 --- a/onnxruntime/contrib_ops/cuda/bert/layer_norm.cuh +++ b/onnxruntime/contrib_ops/cuda/bert/layer_norm.cuh @@ -106,20 +106,31 @@ __device__ inline void LayerNorm( } } -template -__device__ inline void LayerNormSmall(const T val, const cub::KeyValuePair& thread_data, const int ld, const int idx, - const T* beta, const T* gamma, const T epsilon, T* output) { +template +__device__ inline void LayerNormSmall(const T* input_v, const cub::KeyValuePair& thread_data, + const int ld, const int idx, const T* beta, const T* gamma, + const T epsilon, T* output) { // Assuming thread_data is already divided by ld // Small settings: the block covers the leading dimension TPB >= ld. The input // value is available in a register - + using VecT = aligned_vector; using BlockReduce = cub::BlockReduce, TPB>; __shared__ typename BlockReduce::TempStorage temp_storage; __shared__ T mu; // mean __shared__ T rsigma; // 1 / std.dev. + T beta_v[ILP], gamma_v[ILP], output_v[ILP]; + + if (beta != nullptr) { + VecT* beta_val = reinterpret_cast(&beta_v); + *beta_val = *reinterpret_cast(&beta[threadIdx.x * ILP]); + } + VecT* gamma_val = reinterpret_cast(&gamma_v); + *gamma_val = *reinterpret_cast(&gamma[threadIdx.x * ILP]); + + VecT* output_val = reinterpret_cast(&output_v); KeyValuePairSum pair_sum; - const auto sum_kv = BlockReduce(temp_storage).Reduce(thread_data, pair_sum); + const cub::KeyValuePair sum_kv = BlockReduce(temp_storage).Reduce(thread_data, pair_sum); if (threadIdx.x == 0) { mu = sum_kv.key; @@ -127,13 +138,17 @@ __device__ inline void LayerNormSmall(const T val, const cub::KeyValuePair } __syncthreads(); - if (threadIdx.x < ld) { - const T g(gamma[threadIdx.x]); - const T b = (nullptr == beta) ? (T)0 : beta[threadIdx.x]; - output[idx] = g * (val - mu) * rsigma + b; + if (ILP * threadIdx.x < ld) { + #pragma unroll + for (int i = 0; i < ILP; i++) { + output_v[i] = (beta != nullptr) ? gamma_v[i] * (input_v[i] - mu) * rsigma + beta_v[i] : + gamma_v[i] * (input_v[i] - mu) * rsigma; + } + *(reinterpret_cast(&output[idx])) = *reinterpret_cast(&output_v[0]); } } } // namespace cuda } // namespace contrib } // namespace onnxruntime + diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc index 1baaaa69e4..a87437dc75 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc @@ -94,18 +94,19 @@ Status SkipLayerNorm::ComputeInternal(OpKernelContext* ctx) const { int hidden_size = static_cast(input_dims[2]); int64_t element_count = input_dims[0] * sequence_length * hidden_size; size_t element_size = sizeof(T); + typedef typename ToCudaType::MappedType CudaT; - if (!LaunchSkipLayerNormKernel( + if (!LaunchSkipLayerNormKernel( Stream(), - output->template MutableData(), - input->template Data(), - skip->template Data(), - gamma->template Data(), - beta != nullptr ? beta->template Data() : nullptr, - bias != nullptr ? bias->template Data() : nullptr, + reinterpret_cast(output->template MutableData()), + reinterpret_cast(input->template Data()), + reinterpret_cast(skip->template Data()), + reinterpret_cast(gamma->template Data()), + (beta != nullptr) ? reinterpret_cast(beta->template Data()) : nullptr, + (bias != nullptr) ? reinterpret_cast(bias->template Data()) : nullptr, epsilon_, hidden_size, - static_cast(element_count), //TODO: check range + static_cast(element_count), element_size)) { // Get last error to reset it to cudaSuccess. CUDA_CALL(cudaGetLastError()); @@ -118,3 +119,4 @@ Status SkipLayerNorm::ComputeInternal(OpKernelContext* ctx) const { } //namespace cuda } // namespace contrib } // namespace onnxruntime + diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu index 4501186946..76c82f7656 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu @@ -1,7 +1,7 @@ /* 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"); @@ -20,6 +20,13 @@ 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 @@ -28,31 +35,22 @@ namespace onnxruntime { namespace contrib { namespace cuda { -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) { - const T reverse_ld = T(1.f / ld); - const int offset = blockIdx.x * ld; +template +T maybe2half(float x); - KeyValuePairSum pair_sum; - // reduce x and x^2 - cub::KeyValuePair thread_data(0, 0); - const int idx = offset + threadIdx.x; - T val = 0; +template<> +float maybe2half(float x) { + return x; +} - if (threadIdx.x < ld) { - val = (bias == nullptr) ? input[idx] + skip[idx] : input[idx] + skip[idx] + bias[threadIdx.x]; - const T rldval = reverse_ld * val; - thread_data = pair_sum(thread_data, cub::KeyValuePair(rldval, rldval * val)); - } - - LayerNormSmall(val, thread_data, ld, idx, beta, gamma, epsilon, output); +template<> +half maybe2half(float x) { + return __float2half_rn(x); } template __global__ void SkipLayerNormKernel( - const int ld, const T* input, const T* skip, const T* beta, const T* gamma, const T* bias, + const int ld, const T* input, const T* skip, const T* beta, const T* gamma, const T* bias, const T epsilon, T* output) { const T reverse_ld = T(1.f / ld); const int offset = blockIdx.x * ld; @@ -72,72 +70,139 @@ __global__ void SkipLayerNormKernel( LayerNorm(thread_data, ld, offset, beta, gamma, epsilon, output); } -template -bool ComputeSkipLayerNorm( - cudaStream_t stream, const int ld, const int n, const T* input, const T* skip, - const T* beta, const T* gamma, const T* bias, const T epsilon, T* output) { - // this must be true because n is the total size of the tensor - assert(n % ld == 0); - const int grid_size = n / ld; +// 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, bool hasBias) { + const T rld = T(1.f / ld); + const int idx = blockIdx.x * ld + threadIdx.x * ILP; // grid_size = n / ld - if (ld <= 32) { - constexpr int block_size = 32; - SkipLayerNormKernelSmall - <<>>(ld, input, skip, beta, gamma, bias, epsilon, output); - } else if (ld <= 128) { - constexpr int block_size = 128; - SkipLayerNormKernelSmall - <<>>(ld, input, skip, beta, gamma, bias, epsilon, output); - } else if (ld == 384) { - constexpr int block_size = 384; - SkipLayerNormKernelSmall - <<>>(ld, input, skip, beta, gamma, bias, epsilon, output); + using VecT = aligned_vector; + __shared__ T mu; // mean + __shared__ T rsigma; // 1 / std.dev. + + T input_v[ILP], skip_v[ILP], bias_v[ILP], 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]; + const T rldval = rld * input_v[i]; + rldval_sum += rldval; + rldvalsq_sum += rldval * input_v[i]; + } + thread_data = cub::KeyValuePair(rldval_sum, rldvalsq_sum); + } + LayerNormSmall(input_v, thread_data, ld, idx, beta, gamma, epsilon, output); +} + +template +bool LaunchSkipLayerNormKernel( + cudaStream_t stream, T* 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; + if (0 == (ld % 4)) { + const int grid_size = element_count / ld; + if (ld <= 32) { + constexpr int block_size = 32; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 64) { + constexpr int block_size = 64 / 2; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 128) { + constexpr int block_size = 128 / 4; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 384) { + constexpr int block_size = 384 / 4; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 768) { + constexpr int block_size = 768 / 4; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 1024) { + constexpr int block_size = 1024 / 4; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else { + constexpr int block_size = 256; + SkipLayerNormKernel + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output); + } } else { - constexpr int block_size = 256; - SkipLayerNormKernel<<>>(ld, input, skip, beta, gamma, bias, epsilon, output); + const int grid_size = element_count / ld; + if (ld <= 32) { + constexpr int block_size = 32; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 64) { + constexpr int block_size = 64; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld <= 128) { + constexpr int block_size = 128; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else if (ld == 384) { + constexpr int block_size = 384; + SkipLayerNormKernelSmall + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output, hasBias); + } else { + constexpr int block_size = 256; + SkipLayerNormKernel + <<>>(ld, input, skip, beta, gamma, bias, + maybe2half(epsilon), output); + } } return CUDA_CALL(cudaPeekAtLastError()); } -bool LaunchSkipLayerNormKernel( - cudaStream_t stream, - void* output, - const void* input, - const void* skip, - const void* gamma, - const void* beta, - const void* bias, - float epsilon, - int hidden_size, - int element_count, - size_t element_size) { - if (element_size == 2) { - return ComputeSkipLayerNorm( - stream, - hidden_size, - element_count, - reinterpret_cast(input), - reinterpret_cast(skip), - reinterpret_cast(beta), - reinterpret_cast(gamma), - reinterpret_cast(bias), - __float2half_rn(epsilon), - reinterpret_cast(output)); - } else { - return ComputeSkipLayerNorm( - stream, - hidden_size, - element_count, - reinterpret_cast(input), - reinterpret_cast(skip), - reinterpret_cast(beta), - reinterpret_cast(gamma), - reinterpret_cast(bias), - epsilon, - reinterpret_cast(output)); - } -} +template bool LaunchSkipLayerNormKernel(cudaStream_t stream, float* output, const float* input, + const float* skip, const float* gamma, const float* beta, + const float* bias, float epsilon, const int ld, + const int element_count, size_t element_size); + +template bool LaunchSkipLayerNormKernel(cudaStream_t stream, half* output, const half* input, + const half* skip, const half* gamma, const half* beta, + const half* bias, float epsilon, const int ld, + const int element_count, size_t element_size); } // 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 index 0148231f2b..ace02b63fd 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.h +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.h @@ -7,20 +7,22 @@ namespace onnxruntime { namespace contrib { namespace cuda { +template bool LaunchSkipLayerNormKernel( cudaStream_t stream, - void* output, // output tensor - const void* input, // input tensor - const void* skip, // skip tensor - const void* gamma, // Layer normalization gamma tensor - const void* beta, // Layer normalization beta tensor - const void* bias, // Layer normalization beta tensor + T* output, // output tensor + 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 // element size of input tensor + size_t element_size ); } // namespace cuda } // namespace contrib } // namespace onnxruntime + diff --git a/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc b/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc index 009201f130..f48be8b8c8 100644 --- a/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc +++ b/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc @@ -183,6 +183,98 @@ TEST(SkipLayerNormTest, SkipLayerNormBatch1_Float16) { true); } +TEST(SkipLayerNormTest, SkipLayerNormBatch1_Float16_vec) { + int batch_size = 1; + int sequence_length = 2; + int hidden_size = 64; + + std::vector input_data = { + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 1 + -0.8f, -0.5f, 2.0f, 1.f, 0.5f, 0.2f, 0.3f, 0.2f, // 2 + 0.8f, -0.5f, 0.0f, 1.f, -0.5f, 0.2f, 0.3f, 0.6f, // 3 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.1f, 0.3f, -0.3f, // 4 + 0.8f, -3.5f, 0.9f, 1.f, 0.5f, 0.2f, 0.2f, -0.6f, // 5 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 6 + 0.9f, -0.5f, 0.8f, 2.f, 0.3f, 0.3f, 0.3f, -0.6f, // 7 + 0.8f, -0.8f, 3.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 8 + 0.8f, -0.5f, 0.1f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 9 + 0.8f, -1.5f, 0.0f, 6.f, 0.5f, 0.2f, 0.3f, -0.6f, // 10 + 0.8f, -0.5f, 0.0f, 2.f, 0.5f, 0.2f, 0.3f, -0.6f, // 11 + 0.8f, -0.2f, 7.0f, 1.f, -0.2f, 0.2f, 0.3f, 0.6f, // 12 + 0.8f, -0.5f, 0.0f, 1.f, 0.6f, 0.2f, 0.3f, -0.6f, // 13 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.3f, 0.3f, -0.6f, // 14 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, -0.4f, 0.6f, // 15 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, 0.1f}; // 16 + + std::vector skip_data = { + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 1 + -0.8f, -0.5f, 2.0f, 1.f, 0.5f, 0.2f, 0.3f, 0.2f, // 2 + 0.8f, -0.5f, 0.0f, 1.f, -0.5f, 0.2f, 0.3f, 0.6f, // 3 + 0.8f, -0.5f, 0.0f, 3.f, 0.5f, 0.1f, 0.3f, -0.4f, // 4 + 0.8f, -3.5f, 2.9f, -0.f, 0.5f, 0.2f, 0.2f, 0.6f, // 5 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, -0.2f, 0.3f, 0.6f, // 6 + 0.9f, -0.5f, 0.8f, 2.f, 0.3f, 0.3f, 0.3f, -0.6f, // 7 + 0.8f, -1.8f, 3.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 8 + 0.8f, -0.5f, 0.1f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 9 + 0.8f, -1.5f, 0.0f, 6.f, 0.5f, 0.2f, -1.2f, 0.6f, // 10 + 0.8f, -3.5f, 0.0f, 2.f, -0.9f, 0.2f, 0.3f, 0.6f, // 11 + 0.8f, -0.2f, 7.0f, 0.f, -0.2f, 0.2f, 0.3f, 0.6f, // 12 + 0.8f, -0.5f, 4.0f, 1.f, 1.6f, 0.2f, 1.3f, -0.6f, // 13 + 0.8f, -0.5f, 0.1f, 1.f, 0.5f, 0.3f, 0.3f, -0.6f, // 14 + 0.8f, -0.5f, 1.0f, 0.f, 0.5f, 2.2f, -0.4f, 0.6f, // 15 + 0.8f, -0.5f, 0.2f, 1.f, 0.5f, 0.2f, 0.3f, 0.1f}; // 16 + + std::vector gamma_data = { + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 4.3f, -0.6f, // 1 + -0.8f, -3.5f, 2.0f, 1.f, 0.2f, 0.2f, 0.3f, 0.2f, // 2 + 0.8f, -0.5f, 0.0f, 1.f, -0.5f, 0.2f, 0.3f, 0.6f, // 3 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.1f, 0.3f, -0.3f, // 4 + 0.2f, -3.5f, 0.9f, -2.f, 0.5f, 1.2f, 0.2f, 0.6f, // 5 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 3.3f, -0.6f, // 6 + 0.9f, -0.5f, -0.8f, 2.f, 0.3f, 0.3f, 0.3f, 0.6f, // 7 + 0.1f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, 0.1f}; // 8 + + std::vector beta_data = { + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.6f, // 1 + -0.8f, -0.5f, 2.0f, 0.f, 0.5f, 0.2f, 4.9f, 0.2f, // 2 + 0.2f, -0.5f, 0.0f, 1.f, -0.5f, 0.2f, 0.3f, 0.6f, // 3 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, -0.3f, // 4 + 0.1f, -3.5f, 4.9f, 0.f, 0.5f, 0.2f, 0.2f, -0.6f, // 5 + 0.8f, -1.5f, 0.0f, 3.f, 0.5f, 0.7f, 0.8f, -0.6f, // 6 + 0.9f, -0.5f, 0.8f, 0.f, 0.3f, 0.3f, 0.3f, -0.6f, // 7 + 0.8f, -0.5f, 0.0f, 1.f, 0.5f, 0.2f, 0.3f, 0.1f}; // 8 + + std::vector output_data = { + 1.2490234f, -0.044372559f, 0.f, 1.7890625f, 0.61132812f, 0.1763916f, 0.28100586f, 0.014961243f, + 0.20117188f, 2.6894531f, 5.8476562f, 0.78955078f, 0.54443359f, 0.1763916f, 4.8984375f, 0.1763916f, + 0.64941406f, -0.044372559f, 0.f, 1.7890625f, -0.044372559f, 0.1763916f, 0.29882812f, 0.80175781f, + 1.2490234f, -0.044372559f, 0.f, 2.9238281f, 0.61132812f, 0.17687988f, 0.29882812f, -0.077514648f, + 0.21240234f, 11.59375f, 6.5273438f, -0.44458008f, 0.61132812f, 0.058441162f, 0.1763916f, -0.80664062f, + 1.2490234f, -1.0439453f, 0.f, 3.7890625f, 0.61132812f, 0.63134766f, 0.78515625f, -0.39331055f, + 1.5078125f, -0.044372559f, 0.3503418f, 3.8457031f, 0.29882812f, 0.29882812f, 0.29882812f, -1.2148438f, + 0.85595703f, 0.40893555f, 0.f, 1.7890625f, 0.61132812f, 0.1763916f, 0.29882812f, -0.0025119781f, + 1.0097656f, -0.12133789f, 0.f, 1.4189453f, 0.51367188f, 0.1583252f, -0.25830078f, -0.098876953f, + -1.0097656f, 4.8945312f, 1.2695312f, 4.3398438f, 0.50537109f, 0.1583252f, 4.6835938f, 0.12695312f, + 0.40966797f, 0.46655273f, 0.f, 2.203125f, -0.23901367f, 0.1583252f, 0.26123047f, 0.38110352f, + 1.0097656f, -0.23901367f, 0.f, 1.0273438f, 0.23901367f, 0.17907715f, 0.26123047f, -0.33178711f, + 0.15234375f, -0.84863281f, 5.9804688f, -0.83789062f, 0.74853516f, -0.050079346f, 0.25244141f, -1.1015625f, + 1.0097656f, -1.1210938f, 0.f, 3.4179688f, 0.51367188f, 0.67431641f, 0.37133789f, -0.098876953f, + 1.1357422f, -0.12133789f, 0.77832031f, 0.053985596f, 0.30810547f, 0.47265625f, 0.096557617f, -0.53662109f, + 0.82617188f, -0.12133789f, 0.f, 1.4189453f, 0.51367188f, 0.1583252f, 0.26123047f, 0.071289062f}; + + RunTest(input_data, + skip_data, + gamma_data, + beta_data, + std::vector(), + output_data, + epsilon_, + batch_size, + sequence_length, + hidden_size, + true); +} + TEST(SkipLayerNormTest, SkipLayerNormBatch1_NoBeta) { int batch_size = 1; int sequence_length = 2;