You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@mxnet.apache.org by di...@apache.org on 2020/08/18 22:17:25 UTC
[incubator-mxnet] branch master updated: Fix setting cudnn bias
stride (#18905)
This is an automated email from the ASF dual-hosted git repository.
dickjc123 pushed a commit to branch master
in repository https://gitbox.apache.org/repos/asf/incubator-mxnet.git
The following commit(s) were added to refs/heads/master by this push:
new 32994bb Fix setting cudnn bias stride (#18905)
32994bb is described below
commit 32994bbad18445a5bd790d795055fc14264cb0cf
Author: Przemyslaw Tredak <pt...@nvidia.com>
AuthorDate: Tue Aug 18 15:15:50 2020 -0700
Fix setting cudnn bias stride (#18905)
---
src/operator/nn/cudnn/cudnn_convolution-inl.h | 27 +++++++++++++------------
src/operator/nn/cudnn/cudnn_deconvolution-inl.h | 27 +++++++++++++------------
2 files changed, 28 insertions(+), 26 deletions(-)
diff --git a/src/operator/nn/cudnn/cudnn_convolution-inl.h b/src/operator/nn/cudnn/cudnn_convolution-inl.h
index 056f93b..c5beb8a 100644
--- a/src/operator/nn/cudnn/cudnn_convolution-inl.h
+++ b/src/operator/nn/cudnn/cudnn_convolution-inl.h
@@ -212,16 +212,6 @@ class CuDNNConvolutionOp {
typename DataType<DType>::ScaleType alpha = 1.0f;
typename DataType<DType>::ScaleType beta = 0.0f;
typename DataType<DType>::ScaleType beta_add = 1.0f;
- if (!param_.no_bias && (req[conv::kBias] != kNullOp)) {
- Tensor<gpu, 1, DType> gbias = in_grad[conv::kBias].get<gpu, 1, DType>(s);
- CUDNN_CALL(cudnnConvolutionBackwardBias(s->dnn_handle_,
- &alpha,
- out_desc_,
- grad_ptr,
- req[conv::kBias] == kAddTo ? &beta_add : &beta,
- bias_desc_,
- gbias.dptr_));
- }
if (req[conv::kWeight] != kNullOp) {
CHECK_EQ(add_to_weight_, req[conv::kWeight] == kAddTo);
CUDNN_CALL(cudnnConvolutionBackwardFilter(s->dnn_handle_,
@@ -238,6 +228,16 @@ class CuDNNConvolutionOp {
filter_desc_,
gwmat_ptr));
}
+ if (!param_.no_bias && (req[conv::kBias] != kNullOp)) {
+ Tensor<gpu, 1, DType> gbias = in_grad[conv::kBias].get<gpu, 1, DType>(s);
+ CUDNN_CALL(cudnnConvolutionBackwardBias(s->dnn_handle_,
+ &alpha,
+ out_desc_,
+ grad_ptr,
+ req[conv::kBias] == kAddTo ? &beta_add : &beta,
+ bias_desc_,
+ gbias.dptr_));
+ }
if (req[conv::kData] != kNullOp) {
CUDNN_CALL(cudnnConvolutionBackwardData(s_dgrad.GetStream()->dnn_handle_,
&alpha,
@@ -459,13 +459,14 @@ class CuDNNConvolutionOp {
if (!param_.no_bias) {
mxnet::TShape bias = in_shape[conv::kBias];
+ int bias_dim = static_cast<int>(bias[0]);
std::vector<int> bias_shape = {1,
- static_cast<int>(bias[0]),
+ bias_dim,
1, 1};
- std::vector<int> bias_stride = {static_cast<int>(bias[0]), 1, 1, 1};
+ std::vector<int> bias_stride = {bias_dim, 1, bias_dim, bias_dim};
if (param_.kernel.ndim() == 3) {
bias_shape.push_back(1);
- bias_stride.push_back(1);
+ bias_stride.push_back(bias_dim);
}
CUDNN_CALL(cudnnSetTensorNdDescriptor(bias_desc_,
dtype_,
diff --git a/src/operator/nn/cudnn/cudnn_deconvolution-inl.h b/src/operator/nn/cudnn/cudnn_deconvolution-inl.h
index b701883..4f02511 100644
--- a/src/operator/nn/cudnn/cudnn_deconvolution-inl.h
+++ b/src/operator/nn/cudnn/cudnn_deconvolution-inl.h
@@ -201,16 +201,6 @@ class CuDNNDeconvolutionOp {
req[deconv::kData] == kAddTo ? 1.0f : 0.0f;
typename DataType<DType>::ScaleType weight_beta =
req[deconv::kWeight] == kAddTo ? 1.0f : 0.0f;
- if (!param_.no_bias && (req[deconv::kBias] != kNullOp)) {
- Tensor<gpu, 1, DType> gbias = in_grad[deconv::kBias].get<gpu, 1, DType>(s);
- CUDNN_CALL(cudnnConvolutionBackwardBias(s->dnn_handle_,
- &alpha,
- out_desc_,
- grad_ptr + out_offset_ * g,
- &bias_beta,
- bias_desc_,
- gbias.dptr_ + bias_offset_ * g));
- }
if (req[deconv::kWeight] != kNullOp) {
CHECK_EQ(add_to_weight_, req[deconv::kWeight] == kAddTo);
CUDNN_CALL(cudnnConvolutionBackwardFilter(
@@ -228,6 +218,16 @@ class CuDNNDeconvolutionOp {
filter_desc_,
gwmat_ptr + weight_offset_ * g));
}
+ if (!param_.no_bias && (req[deconv::kBias] != kNullOp)) {
+ Tensor<gpu, 1, DType> gbias = in_grad[deconv::kBias].get<gpu, 1, DType>(s);
+ CUDNN_CALL(cudnnConvolutionBackwardBias(s->dnn_handle_,
+ &alpha,
+ out_desc_,
+ grad_ptr + out_offset_ * g,
+ &bias_beta,
+ bias_desc_,
+ gbias.dptr_ + bias_offset_ * g));
+ }
if (req[deconv::kData] != kNullOp) {
CUDNN_CALL(cudnnConvolutionForward(s->dnn_handle_,
&alpha,
@@ -460,13 +460,14 @@ class CuDNNDeconvolutionOp {
if (!param_.no_bias) {
mxnet::TShape bias = in_shape[deconv::kBias];
bias_offset_ = bias[0] / param_.num_group;
+ int bias_dim = static_cast<int>(bias_offset_);
std::vector<int> bias_shape = {1,
- static_cast<int>(bias[0] / param_.num_group),
+ bias_dim,
1, 1};
- std::vector<int> bias_stride = {static_cast<int>(bias_offset_), 1, 1, 1};
+ std::vector<int> bias_stride = {bias_dim, 1, bias_dim, bias_dim};
if (param_.kernel.ndim() == 3) {
bias_shape.push_back(1);
- bias_stride.push_back(1);
+ bias_stride.push_back(bias_dim);
}
CUDNN_CALL(cudnnSetTensorNdDescriptor(bias_desc_,
dtype_,