diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc index 53da91420e..d2f7d974be 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc @@ -2,8 +2,8 @@ // Licensed under the MIT License. #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 { @@ -106,25 +106,29 @@ Status SkipLayerNorm::ComputeInternal(OpKernelContext* ctx) const } } - int sequence_length = static_cast(input_dims[1]); - 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; + int sequence_length = gsl::narrow_cast(input_dims[1]); + int hidden_size = gsl::narrow_cast(input_dims[2]); + int row_count = gsl::narrow_cast(input_dims[0] * sequence_length); - return LaunchSkipLayerNormKernel( + typedef typename ToCudaType::MappedType CudaT; + HostApplyLayerNorm( + GetDeviceProp(), Stream(ctx), - 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); + 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(); } } // 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 deleted file mode 100644 index 0a815f7f49..0000000000 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.cu +++ /dev/null @@ -1,241 +0,0 @@ -/* - 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 deleted file mode 100644 index da2894928c..0000000000 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm_impl.h +++ /dev/null @@ -1,28 +0,0 @@ -// 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/nn/layer_norm.cc b/onnxruntime/core/providers/cuda/nn/layer_norm.cc index 24a6b196c1..7dd10f9c29 100644 --- a/onnxruntime/core/providers/cuda/nn/layer_norm.cc +++ b/onnxruntime/core/providers/cuda/nn/layer_norm.cc @@ -95,6 +95,7 @@ Status LayerNorm::ComputeInternal(OpKernelContext* ctx) con HostApplyLayerNorm(GetDeviceProp(), Stream(ctx), Y_data, mean_data, inv_var_data, X_data, n1, n2, epsilon_, scale_data, bias_data); + CUDA_RETURN_IF_ERROR(cudaGetLastError()); return Status::OK(); } diff --git a/onnxruntime/core/providers/cuda/nn/layer_norm_impl.cu b/onnxruntime/core/providers/cuda/nn/layer_norm_impl.cu index 717b994001..7f9d41298d 100644 --- a/onnxruntime/core/providers/cuda/nn/layer_norm_impl.cu +++ b/onnxruntime/core/providers/cuda/nn/layer_norm_impl.cu @@ -1,18 +1,18 @@ /** -* Copyright (c) 2016-present, Facebook, Inc. -* -* 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) 2016-present, Facebook, Inc. + * + * 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) 2017, NVIDIA CORPORATION. All rights reserved. @@ -85,7 +85,9 @@ __device__ void cuWelfordMuSigma2( const int i1, U& mu, U& sigma2, - U* buf) { + U* buf, + const T* __restrict__ skip, + const T* __restrict__ bias) { // Assumptions: // 1) blockDim.x == GPU_WARP_SIZE // 2) Tensor is contiguous @@ -102,15 +104,34 @@ __device__ void cuWelfordMuSigma2( const int numx = blockDim.x * blockDim.y; const int thrx = threadIdx.x + threadIdx.y * blockDim.x; const T* lvals = vals + i1 * n2; + const T* skip_vals = (skip != NULL) ? skip + i1 * n2 : NULL; int l = 4 * thrx; for (; l + 3 < n2; l += 4 * numx) { for (int k = 0; k < 4; ++k) { U curr = static_cast(lvals[l + k]); + + if (bias != NULL) { + curr += static_cast(bias[l + k]); + } + + if (skip_vals != NULL) { + curr += static_cast(skip_vals[l + k]); + } + cuWelfordOnlineSum(curr, mu, sigma2, count); } } for (; l < n2; ++l) { U curr = static_cast(lvals[l]); + + if (bias != NULL) { + curr += static_cast(bias[l]); + } + + if (skip_vals != NULL) { + curr += static_cast(skip_vals[l]); + } + cuWelfordOnlineSum(curr, mu, sigma2, count); } // intra-warp reductions @@ -314,7 +335,10 @@ __global__ void cuApplyLayerNorm( const int n2, const U epsilon, const V* __restrict__ gamma, - const V* __restrict__ beta) { + const V* __restrict__ beta, + const T* __restrict__ skip, + const T* __restrict__ bias, + T* __restrict__ skip_input_bias_add_output) { // Assumptions: // 1) blockDim.x == GPU_WARP_SIZE // 2) Tensors are contiguous @@ -323,20 +347,35 @@ __global__ void cuApplyLayerNorm( SharedMemory shared; U* buf = shared.getPointer(); U mu, sigma2; - cuWelfordMuSigma2(vals, n1, n2, i1, mu, sigma2, buf); + cuWelfordMuSigma2(vals, n1, n2, i1, mu, sigma2, buf, skip, bias); const T* lvals = vals + i1 * n2; + const T* skip_vals = (skip != NULL) ? skip + i1 * n2 : NULL; V* ovals = output_vals + i1 * n2; + T* skip_input_bias_add_ovals = (skip_input_bias_add_output != NULL) ? skip_input_bias_add_output + i1 * n2 : NULL; U c_inv_std_dev = rsqrt(sigma2 + epsilon); const int numx = blockDim.x * blockDim.y; const int thrx = threadIdx.x + threadIdx.y * blockDim.x; for (int i = thrx; i < n2; i += numx) { U curr = static_cast(lvals[i]); - V gamma_i = (gamma != NULL) ? gamma[i] : (V)1; - V beta_i = (beta != NULL) ? beta[i] : (V)0; + + if (bias != NULL) { + curr += static_cast(bias[i]); + } + + if (skip_vals != NULL) { + curr += static_cast(skip_vals[i]); + } + + U gamma_i = (gamma != NULL) ? (U)gamma[i] : (U)1; + U beta_i = (beta != NULL) ? (U)beta[i] : (U)0; if (simplified) { - ovals[i] = gamma_i * static_cast(c_inv_std_dev * curr); + ovals[i] = static_cast(gamma_i * c_inv_std_dev * curr); } else { - ovals[i] = gamma_i * static_cast(c_inv_std_dev * (curr - mu)) + beta_i; + ovals[i] = static_cast(gamma_i * c_inv_std_dev * (curr - mu) + beta_i); + } + + if (skip_input_bias_add_ovals != NULL) { + skip_input_bias_add_ovals[i] = static_cast(curr); } } if (threadIdx.x == 0 && threadIdx.y == 0) { @@ -346,6 +385,17 @@ __global__ void cuApplyLayerNorm( } } +int32_t round_up_power_of_2(int32_t v) { + v--; + v |= v >> 1; + v |= v >> 2; + v |= v >> 4; + v |= v >> 8; + v |= v >> 16; + v++; + return v; +} + template void HostApplyLayerNorm( const cudaDeviceProp& prop, @@ -358,12 +408,20 @@ void HostApplyLayerNorm( int n2, double epsilon, const V* gamma, - const V* beta) { + const V* beta, + const T* skip, + const T* bias, + T* skip_input_bias_add_output) { const int maxGridY = prop.maxGridSize[1]; const int warp_size = prop.warpSize; ORT_ENFORCE(warp_size == GPU_WARP_SIZE_HOST); - dim3 threads(warp_size, 4, 1); + // Be careful for the logic on treads_y calc: + // * 4 is current implementation related, yet it does not using vectorized load + // * Do not using maxTreads as it will cause resource issue. + int threads_y = std::min(round_up_power_of_2((n2 + 4 * warp_size - 1) / (4 * warp_size)), + prop.maxThreadsPerBlock / (warp_size * 2)); + dim3 threads(warp_size, threads_y, 1); #ifdef __HIP_PLATFORM_HCC__ // Optimization for ROCm MI100 threads.y = 1; @@ -378,13 +436,15 @@ void HostApplyLayerNorm( input, n1, n2, U(epsilon), - gamma, beta); + gamma, beta, + skip, bias, skip_input_bias_add_output); } -#define LAYERNORM_LINEAR_IMPL(T, U, V, simplified) \ - template void HostApplyLayerNorm(const cudaDeviceProp& prop, cudaStream_t stream, V* output, \ - U* mean, U* inv_std_dev, const T* input, int n1, int n2, \ - double epsilon, const V* gamma, const V* beta); +#define LAYERNORM_LINEAR_IMPL(T, U, V, simplified) \ + template void HostApplyLayerNorm(const cudaDeviceProp& prop, cudaStream_t stream, V* output, \ + U* mean, U* inv_std_dev, const T* input, int n1, int n2, \ + double epsilon, const V* gamma, const V* beta, const T* skip, \ + const T* bias, T* skip_input_bias_add_output); LAYERNORM_LINEAR_IMPL(float, float, float, true) LAYERNORM_LINEAR_IMPL(half, float, half, true) diff --git a/onnxruntime/core/providers/cuda/nn/layer_norm_impl.h b/onnxruntime/core/providers/cuda/nn/layer_norm_impl.h index 8c904072c2..e3952eefae 100644 --- a/onnxruntime/core/providers/cuda/nn/layer_norm_impl.h +++ b/onnxruntime/core/providers/cuda/nn/layer_norm_impl.h @@ -40,7 +40,10 @@ void HostApplyLayerNorm( int n2, double epsilon, const V* gamma, - const V* beta); + const V* beta, + const T* skip = nullptr, + const T* bias = nullptr, + T* skip_input_bias_add_output = nullptr); } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/test/contrib_ops/layer_norm_test.cc b/onnxruntime/test/contrib_ops/layer_norm_test.cc index 5b6ae6b9c6..c6c82565be 100644 --- a/onnxruntime/test/contrib_ops/layer_norm_test.cc +++ b/onnxruntime/test/contrib_ops/layer_norm_test.cc @@ -111,6 +111,16 @@ TEST(CudaKernelTest, LayerNorm_LargeSizeTensor) { TestLayerNorm(X_dims, LAYER_NORM_OP, k_epsilon_default); } +TEST(CudaKernelTest, LayerNorm_4KTensor) { + std::vector X_dims{3, 10, 4096}; + TestLayerNorm(X_dims, LAYER_NORM_OP, k_epsilon_default); +} + +TEST(CudaKernelTest, LayerNorm_8KTensor) { + std::vector X_dims{3, 10, 8192}; + TestLayerNorm(X_dims, LAYER_NORM_OP, k_epsilon_default); +} + TEST(CudaKernelTest, LayerNorm_MidSizeTensor_NoBias) { std::vector X_dims{8, 80, 768}; constexpr int64_t axis = -1; diff --git a/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc b/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc index c501bd72a9..638b7565a3 100644 --- a/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc +++ b/onnxruntime/test/contrib_ops/skiplayernorm_op_test.cc @@ -280,23 +280,35 @@ TEST(SkipLayerNormTest, SkipLayerNormBatch1_Float16_vec) { 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 + // Update test data result which use internal fp32 calculation for fp16 input/parameters. + // Following pytorch code snippet are used to generate the result: (Not torch uses fp32 internal calculation for this) + // + // gamma_tensor = torch.tensor(gamma_data, dtype=torch.float32).reshape(hidden_size).to('cuda:0').to(torch.float16) + // beta_tensor = torch.tensor(beta_data, dtype=torch.float32).reshape(hidden_size).to('cuda:0').to(torch.float16) + // input_tensor = torch.tensor(input_data, dtype=torch.float32).reshape( + // batch_size, sequence_length, hidden_size).to('cuda:0').to(torch.float16) + // skip_tensor = torch.tensor(skip_data, dtype=torch.float32).reshape( + // batch_size, sequence_length, hidden_size).to('cuda:0').to(torch.float16) + // added_input = torch.add(input_tensor, skip_tensor) + // out32 = torch.layer_norm(added_input, [hidden_size], gamma_tensor, beta_tensor, eps=epsilon_).to(torch.float32) + // 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}; + 1.25000000f, -0.04403687f, 0.00000000f, 1.79003906f, 0.61132812f, 0.17639160f, 0.28125000f, 0.01530457f, + 0.20166016f, 2.69140625f, 5.84765625f, 0.78955078f, 0.54443359f, 0.17639160f, 4.89843750f, 0.17639160f, + 0.64990234f, -0.04403687f, 0.00000000f, 1.79003906f, -0.04403687f, 0.17639160f, 0.29882812f, 0.80175781f, + 1.25000000f, -0.04403687f, 0.00000000f, 2.92382812f, 0.61132812f, 0.17687988f, 0.29882812f, -0.07745361f, + 0.21240234f, 11.60156250f, 6.52734375f, -0.44482422f, 0.61132812f, 0.05844116f, 0.17639160f, -0.80712891f, + 1.25000000f, -1.04394531f, 0.00000000f, 3.78906250f, 0.61132812f, 0.63134766f, 0.78564453f, -0.39331055f, + 1.50878906f, -0.04403687f, 0.34985352f, 3.84765625f, 0.29882812f, 0.29882812f, 0.29882812f, -1.21582031f, + 0.85595703f, 0.40966797f, 0.00000000f, 1.79003906f, 0.61132812f, 0.17639160f, 0.29882812f, -0.00255013f, + 1.00976562f, -0.12152100f, 0.00000000f, 1.41894531f, 0.51367188f, 0.15832520f, -0.25805664f, -0.09875488f, + -1.00976562f, 4.89453125f, 1.26953125f, 4.33984375f, 0.50537109f, 0.15832520f, 4.68359375f, 0.12695312f, + 0.40942383f, 0.46655273f, 0.00000000f, 2.20312500f, -0.23913574f, 0.15832520f, 0.26123047f, 0.38110352f, + 1.00976562f, -0.23913574f, 0.00000000f, 1.02734375f, 0.23913574f, 0.17907715f, 0.26123047f, -0.33178711f, + 0.15234375f, -0.85058594f, 5.98046875f, -0.83789062f, 0.74853516f, -0.04998779f, 0.25244141f, -1.10156250f, + 1.00976562f, -1.12109375f, 0.00000000f, 3.41992188f, 0.51367188f, 0.67431641f, 0.37133789f, -0.09875488f, + 1.13574219f, -0.12152100f, 0.77832031f, 0.05398560f, 0.30810547f, 0.47265625f, 0.09643555f, -0.53662109f, + 0.82617188f, -0.12152100f, 0.00000000f, 1.41894531f, 0.51367188f, 0.15832520f, 0.26123047f, 0.07135010f}; RunTest(input_data, skip_data,