From 336624806e7acd3d3defb1b581b13341735cfdf6 Mon Sep 17 00:00:00 2001 From: Weixing Zhang Date: Thu, 23 Apr 2020 10:12:55 -0700 Subject: [PATCH] Simplify and clean code (#3655) 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 --- .../contrib_ops/cuda/bert/attention.cc | 1 - onnxruntime/contrib_ops/cuda/bert/attention.h | 2 +- .../contrib_ops/cuda/bert/embed_layer_norm.cc | 1 - .../contrib_ops/cuda/bert/fast_gelu.cc | 2 +- .../contrib_ops/cuda/bert/skip_layer_norm.cc | 2 +- onnxruntime/contrib_ops/cuda/layer_norm.cc | 2 +- .../contrib_ops/cuda/layer_norm_impl.cu | 28 ++++++++++--------- .../contrib_ops/cuda/tensor/image_scaler.h | 2 +- onnxruntime/core/providers/cuda/math/gemm.cc | 2 +- .../core/providers/cuda/math/softmax_impl.cuh | 1 - onnxruntime/core/providers/cuda/nn/shrink.h | 2 +- .../core/providers/cuda/tensor/concat.h | 2 +- .../core/providers/cuda/tensor/gather.h | 2 +- .../training_ops/cuda/communication/recv.h | 1 - .../training_ops/cuda/communication/send.h | 2 +- .../training_ops/cuda/math/isfinite.h | 2 +- .../cuda/math/mixed_precision_scale.h | 2 +- .../training_ops/cuda/nn/layer_norm.h | 1 - .../training_ops/cuda/nn/layer_norm_impl.cu | 8 +++--- .../training_ops/cuda/optimizer/adam.h | 1 - .../cuda/optimizer/gradient_control.h | 1 - .../training_ops/cuda/optimizer/lamb.h | 1 - .../training_ops/cuda/optimizer/sg.h | 1 - 23 files changed, 31 insertions(+), 38 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/bert/attention.cc b/onnxruntime/contrib_ops/cuda/bert/attention.cc index 6135bf07c1..a6b4cfb8b7 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention.cc +++ b/onnxruntime/contrib_ops/cuda/bert/attention.cc @@ -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" diff --git a/onnxruntime/contrib_ops/cuda/bert/attention.h b/onnxruntime/contrib_ops/cuda/bert/attention.h index 9ec55f1793..b32ca9ee35 100644 --- a/onnxruntime/contrib_ops/cuda/bert/attention.h +++ b/onnxruntime/contrib_ops/cuda/bert/attention.h @@ -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 { diff --git a/onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.cc b/onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.cc index 8df41e9aab..d873971e8c 100644 --- a/onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.cc +++ b/onnxruntime/contrib_ops/cuda/bert/embed_layer_norm.cc @@ -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" diff --git a/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc b/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc index 1d9fe64336..0c965cd468 100644 --- a/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc +++ b/onnxruntime/contrib_ops/cuda/bert/fast_gelu.cc @@ -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" diff --git a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc index 1eeaeb773d..61a1d64ca1 100644 --- a/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc +++ b/onnxruntime/contrib_ops/cuda/bert/skip_layer_norm.cc @@ -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" diff --git a/onnxruntime/contrib_ops/cuda/layer_norm.cc b/onnxruntime/contrib_ops/cuda/layer_norm.cc index 5bb3b237ab..bd6d14eef0 100644 --- a/onnxruntime/contrib_ops/cuda/layer_norm.cc +++ b/onnxruntime/contrib_ops/cuda/layer_norm.cc @@ -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 { diff --git a/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu b/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu index d6867949d0..747a2ff70e 100644 --- a/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu +++ b/onnxruntime/contrib_ops/cuda/layer_norm_impl.cu @@ -107,13 +107,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; - 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(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(curr.x), mu, sigma2, count); + cuWelfordOnlineSum(static_cast(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 shared; U* buf = shared.getPointer(); U mu, sigma2; diff --git a/onnxruntime/contrib_ops/cuda/tensor/image_scaler.h b/onnxruntime/contrib_ops/cuda/tensor/image_scaler.h index 70f6590f62..de431d45f5 100644 --- a/onnxruntime/contrib_ops/cuda/tensor/image_scaler.h +++ b/onnxruntime/contrib_ops/cuda/tensor/image_scaler.h @@ -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 { diff --git a/onnxruntime/core/providers/cuda/math/gemm.cc b/onnxruntime/core/providers/cuda/math/gemm.cc index 87e500cc19..21b771fdbf 100644 --- a/onnxruntime/core/providers/cuda/math/gemm.cc +++ b/onnxruntime/core/providers/cuda/math/gemm.cc @@ -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 { diff --git a/onnxruntime/core/providers/cuda/math/softmax_impl.cuh b/onnxruntime/core/providers/cuda/math/softmax_impl.cuh index b9c07da7f2..19508f8538 100644 --- a/onnxruntime/core/providers/cuda/math/softmax_impl.cuh +++ b/onnxruntime/core/providers/cuda/math/softmax_impl.cuh @@ -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 { diff --git a/onnxruntime/core/providers/cuda/nn/shrink.h b/onnxruntime/core/providers/cuda/nn/shrink.h index 68fd27d00d..850dd9781e 100644 --- a/onnxruntime/core/providers/cuda/nn/shrink.h +++ b/onnxruntime/core/providers/cuda/nn/shrink.h @@ -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 { diff --git a/onnxruntime/core/providers/cuda/tensor/concat.h b/onnxruntime/core/providers/cuda/tensor/concat.h index 7c542709bc..9cefcafb59 100644 --- a/onnxruntime/core/providers/cuda/tensor/concat.h +++ b/onnxruntime/core/providers/cuda/tensor/concat.h @@ -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; }; diff --git a/onnxruntime/core/providers/cuda/tensor/gather.h b/onnxruntime/core/providers/cuda/tensor/gather.h index bc7e2508f2..917bff8fc4 100644 --- a/onnxruntime/core/providers/cuda/tensor/gather.h +++ b/onnxruntime/core/providers/cuda/tensor/gather.h @@ -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; }; diff --git a/orttraining/orttraining/training_ops/cuda/communication/recv.h b/orttraining/orttraining/training_ops/cuda/communication/recv.h index abfc8ca03a..0d1a812038 100644 --- a/orttraining/orttraining/training_ops/cuda/communication/recv.h +++ b/orttraining/orttraining/training_ops/cuda/communication/recv.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/communication/send.h b/orttraining/orttraining/training_ops/cuda/communication/send.h index 3350c519ed..878fee48d7 100644 --- a/orttraining/orttraining/training_ops/cuda/communication/send.h +++ b/orttraining/orttraining/training_ops/cuda/communication/send.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/math/isfinite.h b/orttraining/orttraining/training_ops/cuda/math/isfinite.h index ea11c1d43b..c5073f81cf 100644 --- a/orttraining/orttraining/training_ops/cuda/math/isfinite.h +++ b/orttraining/orttraining/training_ops/cuda/math/isfinite.h @@ -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; diff --git a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h index 64e8d05224..7e92dc1b17 100644 --- a/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h +++ b/orttraining/orttraining/training_ops/cuda/math/mixed_precision_scale.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/nn/layer_norm.h b/orttraining/orttraining/training_ops/cuda/nn/layer_norm.h index c1cc447ace..fd21b09ba4 100644 --- a/orttraining/orttraining/training_ops/cuda/nn/layer_norm.h +++ b/orttraining/orttraining/training_ops/cuda/nn/layer_norm.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/nn/layer_norm_impl.cu b/orttraining/orttraining/training_ops/cuda/nn/layer_norm_impl.cu index e8ba8a0e6c..7f85c5676f 100644 --- a/orttraining/orttraining/training_ops/cuda/nn/layer_norm_impl.cu +++ b/orttraining/orttraining/training_ops/cuda/nn/layer_norm_impl.cu @@ -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(curr.x), mu, sigma2, count); + cuWelfordOnlineSum(static_cast(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 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]; diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/adam.h b/orttraining/orttraining/training_ops/cuda/optimizer/adam.h index fcfa617dbc..a35625885b 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/adam.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/adam.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h index bd4d2fd3de..bf8c12d51a 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/gradient_control.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h b/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h index bbb34ba208..55a828a94f 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/lamb.h @@ -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 { diff --git a/orttraining/orttraining/training_ops/cuda/optimizer/sg.h b/orttraining/orttraining/training_ops/cuda/optimizer/sg.h index 91ffbbbab9..a58d98e1f9 100644 --- a/orttraining/orttraining/training_ops/cuda/optimizer/sg.h +++ b/orttraining/orttraining/training_ops/cuda/optimizer/sg.h @@ -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 {