saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-30 20:18:08 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Simplify and clean code (#3655)

Browse source 
1. It is not necessary to include cudnn_common.h for kernels which are not implemented with CUDNN.
2. Minor change in layer norm kernel to simplify the code and resolve building warning.

Co-authored-by: Weixing Zhang <wezhan@microsoft.com>
This commit is contained in:
Weixing Zhang 2020-04-23 10:12:55 -07:00 committed by GitHub
parent 125f68f305
commit 336624806e
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
23 changed files with 31 additions and 38 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
onnxruntime/contrib_ops/cuda/bert/attention.cc
View file

@ -3,7 +3,6 @@
#include "attention.h"
#include "core/framework/tensorprotoutils.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
#include "attention_impl.h"

2
onnxruntime/contrib_ops/cuda/bert/attention.h
View file

@ -5,7 +5,7 @@
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
#include "contrib_ops/cpu/bert/attention.h"
namespace onnxruntime {

1
onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.cc
View file

@ -2,7 +2,6 @@
// Licensed under the MIT License.
#include "core/providers/common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/framework/tensorprotoutils.h"
#include "onnx/defs/tensor_proto_util.h"
#include "contrib_ops/cpu/bert/embed_layer_norm_helper.h"

2
onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc
View file

@ -2,7 +2,7 @@
// Licensed under the MIT License.
#include "core/providers/common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/framework/tensorprotoutils.h"
#include "fast_gelu.h"
#include "fast_gelu_impl.h"

2
onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc
View file

@ -2,7 +2,7 @@
// Licensed under the MIT License.
#include "core/providers/common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/framework/tensorprotoutils.h"
#include "onnx/defs/tensor_proto_util.h"
#include "skip_layer_norm.h"

2
onnxruntime/contrib_ops/cuda/layer_norm.cc
View file

@ -5,7 +5,7 @@
#include "layer_norm_impl.h"
#include "core/providers/common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
namespace contrib {

28
onnxruntime/contrib_ops/cuda/layer_norm_impl.cu
View file

@ -107,13 +107,14 @@ __device__ void cuWelfordMuSigma2(
cuWelfordOnlineSum<U>(curr, mu, sigma2, count);
}
// intra-warp reductions
for (int l = 0; l <= 4; ++l) {
int srcLaneB = (threadIdx.x + (1 << l)) & 31;
U muB = WARP_SHFL(mu, srcLaneB);
U countB = WARP_SHFL(count, srcLaneB);
U sigma2B = WARP_SHFL(sigma2, srcLaneB);
#pragma unroll
for (int stride = GPU_WARP_SIZE / 2; stride > 0; stride /= 2) {
U muB = WARP_SHFL_DOWN(mu, stride);
U countB = WARP_SHFL_DOWN(count, stride);
U sigma2B = WARP_SHFL_DOWN(sigma2, stride);
cuChanOnlineSum<U>(muB, sigma2B, countB, mu, sigma2, count);
}
// threadIdx.x == 0 has correct values for each warp
// inter-warp reductions
if (blockDim.y > 1) {
@ -192,8 +193,8 @@ __device__ void cuWelfordMuSigma2(
for (; l + 7 < n2; l += 8 * numx) {
for (int k = 0; k < 8; k += 2) {
float2 curr = __half22float2(*((__half2*)(lvals + l + k)));
cuWelfordOnlineSum(curr.x, mu, sigma2, count);
cuWelfordOnlineSum(curr.y, mu, sigma2, count);
cuWelfordOnlineSum(static_cast<float>(curr.x), mu, sigma2, count);
cuWelfordOnlineSum(static_cast<float>(curr.y), mu, sigma2, count);
}
}
for (; l < n2; ++l) {
@ -201,13 +202,14 @@ __device__ void cuWelfordMuSigma2(
cuWelfordOnlineSum(curr, mu, sigma2, count);
}
// intra-warp reductions
for (int l = 0; l <= 4; ++l) {
int srcLaneB = (threadIdx.x + (1 << l)) & 31;
float muB = WARP_SHFL(mu, srcLaneB);
float countB = WARP_SHFL(count, srcLaneB);
float sigma2B = WARP_SHFL(sigma2, srcLaneB);
#pragma unroll
for (int stride = GPU_WARP_SIZE / 2; stride > 0; stride /= 2) {
float muB = WARP_SHFL_DOWN(mu, stride);
float countB = WARP_SHFL_DOWN(count, stride);
float sigma2B = WARP_SHFL_DOWN(sigma2, stride);
cuChanOnlineSum(muB, sigma2B, countB, mu, sigma2, count);
}
// threadIdx.x == 0 has correct values for each warp
// inter-warp reductions
if (blockDim.y > 1) {
@ -310,7 +312,7 @@ __global__ void cuApplyLayerNorm(
// 1) blockDim.x == GPU_WARP_SIZE
// 2) Tensors are contiguous
//
for (auto i1 = blockIdx.y; i1 < n1; i1 += gridDim.y) {
for (int i1 = blockIdx.y; i1 < n1; i1 += gridDim.y) {
SharedMemory<U> shared;
U* buf = shared.getPointer();
U mu, sigma2;

2
onnxruntime/contrib_ops/cuda/tensor/image_scaler.h
View file

@ -5,7 +5,7 @@
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
namespace contrib {

2
onnxruntime/core/providers/cuda/math/gemm.cc
View file

@ -3,7 +3,7 @@
#include "gemm.h"
#include "core/providers/cpu/math/gemm_helper.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/shared_inc/fpgeneric.h"
namespace onnxruntime {

1
onnxruntime/core/providers/cuda/math/softmax_impl.cuh
View file

@ -17,7 +17,6 @@
// The code below is mostly copied from Pytorch PersistentSoftmax.cuh
#pragma once
#include "core/providers/cuda/cu_inc/common.cuh"
namespace onnxruntime {

2
onnxruntime/core/providers/cuda/nn/shrink.h
View file

@ -4,7 +4,7 @@
#pragma once
#include "gsl/gsl"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
namespace cuda {

2
onnxruntime/core/providers/cuda/tensor/concat.h
View file

@ -11,7 +11,7 @@ namespace cuda {
class Concat final : public CudaKernel, public ConcatBase {
public:
Concat(const OpKernelInfo& info) : ConcatBase(info), CudaKernel(info) {}
Concat(const OpKernelInfo& info) : CudaKernel(info), ConcatBase(info) {}
Status ComputeInternal(OpKernelContext* context) const override;
};

2
onnxruntime/core/providers/cuda/tensor/gather.h
View file

@ -11,7 +11,7 @@ namespace cuda {
class Gather final : public CudaKernel, public GatherBase {
public:
Gather(const OpKernelInfo& info) : GatherBase(info), CudaKernel(info) {}
Gather(const OpKernelInfo& info) : CudaKernel(info), GatherBase(info) {}
Status ComputeInternal(OpKernelContext* context) const override;
};

1
orttraining/orttraining/training_ops/cuda/communication/recv.h
View file

@ -6,7 +6,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
namespace onnxruntime {
namespace cuda {

2
orttraining/orttraining/training_ops/cuda/communication/send.h
View file

@ -6,7 +6,7 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
namespace onnxruntime {
namespace cuda {

2
orttraining/orttraining/training_ops/cuda/math/isfinite.h
View file

@ -4,7 +4,7 @@
#pragma once
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/multi_tensor/common.cuh"
constexpr int PARALLEL_LOADS = 4;

2
orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h
View file

@ -4,7 +4,7 @@
#pragma once
#include "core/common/common.h"
#include "core/framework/op_kernel.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/cuda_common.h"
namespace onnxruntime {
namespace cuda {

1
orttraining/orttraining/training_ops/cuda/nn/layer_norm.h
View file

@ -1,7 +1,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
namespace onnxruntime {
namespace cuda {

8
orttraining/orttraining/training_ops/cuda/nn/layer_norm_impl.cu
View file

@ -190,8 +190,8 @@ __device__ void cuWelfordMuSigma2(
for (; l + 7 < n2; l += 8 * numx) {
for (int k = 0; k < 8; k += 2) {
float2 curr = __half22float2(*((__half2*)(lvals + l + k)));
cuWelfordOnlineSum(curr.x, mu, sigma2, count);
cuWelfordOnlineSum(curr.y, mu, sigma2, count);
cuWelfordOnlineSum(static_cast<float>(curr.x), mu, sigma2, count);
cuWelfordOnlineSum(static_cast<float>(curr.y), mu, sigma2, count);
}
}
for (; l < n2; ++l) {
@ -308,7 +308,7 @@ __global__ void cuApplyLayerNorm(
// 1) blockDim.x == GPU_WARP_SIZE
// 2) Tensors are contiguous
//
for (auto i1 = blockIdx.y; i1 < n1; i1 += gridDim.y) {
for (int i1 = blockIdx.y; i1 < n1; i1 += gridDim.y) {
SharedMemory<U> shared;
U* buf = shared.getPointer();
U mu, sigma2;
@ -576,7 +576,7 @@ __global__ void cuComputeGradInput(
const U* __restrict__ invvar,
const T* gamma,
T* grad_input) {
for (auto i1 = blockIdx.y; i1 < n1; i1 += gridDim.y) {
for (int i1 = blockIdx.y; i1 < n1; i1 += gridDim.y) {
U sum_loss1 = U(0);
U sum_loss2 = U(0);
const U c_mean = mean[i1];

1
orttraining/orttraining/training_ops/cuda/optimizer/adam.h
View file

@ -4,7 +4,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
namespace onnxruntime {
namespace cuda {

1
orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h
View file

@ -4,7 +4,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
namespace onnxruntime {
namespace cuda {

1
orttraining/orttraining/training_ops/cuda/optimizer/lamb.h
View file

@ -4,7 +4,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
#include "core/providers/cuda/multi_tensor/common.cuh"
namespace onnxruntime {

1
orttraining/orttraining/training_ops/cuda/optimizer/sg.h
View file

@ -4,7 +4,6 @@
#pragma once
#include "core/common/common.h"
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cuda/cudnn_common.h"
namespace onnxruntime {
namespace cuda {