mirror of
https://github.com/saymrwulf/onnxruntime.git
synced 2026-07-10 17:37:14 +00:00
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
This commit is contained in:
Hubert Lu
GitHub
parent
7b8f45dd60
commit
835ecb264d
5 changed files with 279 additions and 103 deletions
|
|
@ -106,20 +106,31 @@ __device__ inline void LayerNorm(
|
|||
}
|
||||
}
|
||||
|
||||
template <typename T, int TPB>
|
||||
__device__ inline void LayerNormSmall(const T val, const cub::KeyValuePair<T, T>& thread_data, const int ld, const int idx,
|
||||
const T* beta, const T* gamma, const T epsilon, T* output) {
|
||||
template <typename T, int TPB, int ILP>
|
||||
__device__ inline void LayerNormSmall(const T* input_v, const cub::KeyValuePair<T, T>& 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<T, ILP>;
|
||||
using BlockReduce = cub::BlockReduce<cub::KeyValuePair<T, T>, 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<VecT*>(&beta_v);
|
||||
*beta_val = *reinterpret_cast<const VecT*>(&beta[threadIdx.x * ILP]);
|
||||
}
|
||||
VecT* gamma_val = reinterpret_cast<VecT*>(&gamma_v);
|
||||
*gamma_val = *reinterpret_cast<const VecT*>(&gamma[threadIdx.x * ILP]);
|
||||
|
||||
VecT* output_val = reinterpret_cast<VecT*>(&output_v);
|
||||
|
||||
KeyValuePairSum pair_sum;
|
||||
const auto sum_kv = BlockReduce(temp_storage).Reduce(thread_data, pair_sum);
|
||||
const cub::KeyValuePair<T, T> 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<T, T>
|
|||
}
|
||||
__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<VecT*>(&output[idx])) = *reinterpret_cast<VecT*>(&output_v[0]);
|
||||
}
|
||||
}
|
||||
|
||||
} // namespace cuda
|
||||
} // namespace contrib
|
||||
} // namespace onnxruntime
|
||||
|
||||
|
|
|
|||
|
|
@ -94,18 +94,19 @@ Status SkipLayerNorm<T>::ComputeInternal(OpKernelContext* ctx) const {
|
|||
int hidden_size = static_cast<int>(input_dims[2]);
|
||||
int64_t element_count = input_dims[0] * sequence_length * hidden_size;
|
||||
size_t element_size = sizeof(T);
|
||||
typedef typename ToCudaType<T>::MappedType CudaT;
|
||||
|
||||
if (!LaunchSkipLayerNormKernel(
|
||||
if (!LaunchSkipLayerNormKernel<CudaT>(
|
||||
Stream(),
|
||||
output->template MutableData<T>(),
|
||||
input->template Data<T>(),
|
||||
skip->template Data<T>(),
|
||||
gamma->template Data<T>(),
|
||||
beta != nullptr ? beta->template Data<T>() : nullptr,
|
||||
bias != nullptr ? bias->template Data<T>() : nullptr,
|
||||
reinterpret_cast<CudaT*>(output->template MutableData<T>()),
|
||||
reinterpret_cast<const CudaT*>(input->template Data<T>()),
|
||||
reinterpret_cast<const CudaT*>(skip->template Data<T>()),
|
||||
reinterpret_cast<const CudaT*>(gamma->template Data<T>()),
|
||||
(beta != nullptr) ? reinterpret_cast<const CudaT*>(beta->template Data<T>()) : nullptr,
|
||||
(bias != nullptr) ? reinterpret_cast<const CudaT*>(bias->template Data<T>()) : nullptr,
|
||||
epsilon_,
|
||||
hidden_size,
|
||||
static_cast<int>(element_count), //TODO: check range
|
||||
static_cast<int>(element_count),
|
||||
element_size)) {
|
||||
// Get last error to reset it to cudaSuccess.
|
||||
CUDA_CALL(cudaGetLastError());
|
||||
|
|
@ -118,3 +119,4 @@ Status SkipLayerNorm<T>::ComputeInternal(OpKernelContext* ctx) const {
|
|||
} //namespace cuda
|
||||
} // namespace contrib
|
||||
} // namespace onnxruntime
|
||||
|
||||
|
|
|
|||
|
|
@ -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 <cuda_fp16.h>
|
||||
|
|
@ -28,31 +35,22 @@ namespace onnxruntime {
|
|||
namespace contrib {
|
||||
namespace cuda {
|
||||
|
||||
template <typename T, unsigned TPB>
|
||||
__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<typename T>
|
||||
T maybe2half(float x);
|
||||
|
||||
KeyValuePairSum pair_sum;
|
||||
// reduce x and x^2
|
||||
cub::KeyValuePair<T, T> 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<T, T>(rldval, rldval * val));
|
||||
}
|
||||
|
||||
LayerNormSmall<T, TPB>(val, thread_data, ld, idx, beta, gamma, epsilon, output);
|
||||
template<>
|
||||
half maybe2half(float x) {
|
||||
return __float2half_rn(x);
|
||||
}
|
||||
|
||||
template <typename T, unsigned TPB>
|
||||
__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<T, TPB>(thread_data, ld, offset, beta, gamma, epsilon, output);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
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 <typename T, unsigned TPB, int ILP>
|
||||
__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<T, block_size>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias, epsilon, output);
|
||||
} else if (ld <= 128) {
|
||||
constexpr int block_size = 128;
|
||||
SkipLayerNormKernelSmall<T, block_size>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias, epsilon, output);
|
||||
} else if (ld == 384) {
|
||||
constexpr int block_size = 384;
|
||||
SkipLayerNormKernelSmall<T, block_size>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias, epsilon, output);
|
||||
using VecT = aligned_vector<T, ILP>;
|
||||
__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<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];
|
||||
const T rldval = rld * input_v[i];
|
||||
rldval_sum += rldval;
|
||||
rldvalsq_sum += rldval * input_v[i];
|
||||
}
|
||||
thread_data = cub::KeyValuePair<T, T>(rldval_sum, rldvalsq_sum);
|
||||
}
|
||||
LayerNormSmall<T, TPB, ILP>(input_v, thread_data, ld, idx, beta, gamma, epsilon, output);
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
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<T, block_size, 1>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 64) {
|
||||
constexpr int block_size = 64 / 2;
|
||||
SkipLayerNormKernelSmall<T, block_size, 2>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 128) {
|
||||
constexpr int block_size = 128 / 4;
|
||||
SkipLayerNormKernelSmall<T, block_size, 4>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 384) {
|
||||
constexpr int block_size = 384 / 4;
|
||||
SkipLayerNormKernelSmall<T, block_size, 4>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 768) {
|
||||
constexpr int block_size = 768 / 4;
|
||||
SkipLayerNormKernelSmall<T, block_size, 4>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 1024) {
|
||||
constexpr int block_size = 1024 / 4;
|
||||
SkipLayerNormKernelSmall<T, block_size, 4>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else {
|
||||
constexpr int block_size = 256;
|
||||
SkipLayerNormKernel<T, block_size>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output);
|
||||
}
|
||||
} else {
|
||||
constexpr int block_size = 256;
|
||||
SkipLayerNormKernel<T, block_size><<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias, epsilon, output);
|
||||
const int grid_size = element_count / ld;
|
||||
if (ld <= 32) {
|
||||
constexpr int block_size = 32;
|
||||
SkipLayerNormKernelSmall<T, block_size, 1>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 64) {
|
||||
constexpr int block_size = 64;
|
||||
SkipLayerNormKernelSmall<T, block_size, 1>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld <= 128) {
|
||||
constexpr int block_size = 128;
|
||||
SkipLayerNormKernelSmall<T, block_size, 1>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else if (ld == 384) {
|
||||
constexpr int block_size = 384;
|
||||
SkipLayerNormKernelSmall<T, block_size, 1>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(epsilon), output, hasBias);
|
||||
} else {
|
||||
constexpr int block_size = 256;
|
||||
SkipLayerNormKernel<T, block_size>
|
||||
<<<grid_size, block_size, 0, stream>>>(ld, input, skip, beta, gamma, bias,
|
||||
maybe2half<T>(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<const half*>(input),
|
||||
reinterpret_cast<const half*>(skip),
|
||||
reinterpret_cast<const half*>(beta),
|
||||
reinterpret_cast<const half*>(gamma),
|
||||
reinterpret_cast<const half*>(bias),
|
||||
__float2half_rn(epsilon),
|
||||
reinterpret_cast<half*>(output));
|
||||
} else {
|
||||
return ComputeSkipLayerNorm(
|
||||
stream,
|
||||
hidden_size,
|
||||
element_count,
|
||||
reinterpret_cast<const float*>(input),
|
||||
reinterpret_cast<const float*>(skip),
|
||||
reinterpret_cast<const float*>(beta),
|
||||
reinterpret_cast<const float*>(gamma),
|
||||
reinterpret_cast<const float*>(bias),
|
||||
epsilon,
|
||||
reinterpret_cast<float*>(output));
|
||||
}
|
||||
}
|
||||
template bool LaunchSkipLayerNormKernel<float>(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<half>(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
|
||||
|
||||
|
||||
|
|
|
|||
|
|
@ -7,20 +7,22 @@ namespace onnxruntime {
|
|||
namespace contrib {
|
||||
namespace cuda {
|
||||
|
||||
template <typename T>
|
||||
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
|
||||
|
||||
|
|
|
|||
|
|
@ -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<float> 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<float> 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<float> 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<float> 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<float> 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<float>(),
|
||||
output_data,
|
||||
epsilon_,
|
||||
batch_size,
|
||||
sequence_length,
|
||||
hidden_size,
|
||||
true);
|
||||
}
|
||||
|
||||
TEST(SkipLayerNormTest, SkipLayerNormBatch1_NoBeta) {
|
||||
int batch_size = 1;
|
||||
int sequence_length = 2;
|
||||
|
|
|
|||
Loading…
Reference in a new issue