Huawei_Technology
/
mindspore
Not watched
1
0
Fork 0
Code
Releases 13 Wiki evaluate Activity
Issues 0 Pull Requests 0 Datasets Model Cloudbrain HPC
Browse Source

!6330 GPU fix reviewbot : large function 

Merge pull request !6330 from VectorSL/pooling
tags/v1.0.0
mindspore-ci-bot Gitee 5 years ago
parent
33b4886ca9 01ef6db9e7
commit
f6522ef5bc
2 changed files with 72 additions and 72 deletions
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. +35
    -32
      mindspore/ccsrc/backend/kernel_compiler/gpu/nn/pooling_gpu_kernel.h
  2. +37
    -40
      mindspore/ccsrc/backend/kernel_compiler/gpu/nn/pooling_grad_gpu_kernel.h

+ 35
- 32
mindspore/ccsrc/backend/kernel_compiler/gpu/nn/pooling_gpu_kernel.h View File

@@ -103,38 +103,8 @@ class PoolingGpuFwdKernel : public GpuKernel {
CHECK_CUDNN_RET_WITH_EXCEPT(
cudnnSetTensorNdDescriptor(output_descriptor_, cudnn_data_type_, nbDims, dimAout, strideAout),
"cudnnSetTensor4dDescriptor failed");
auto window = GetValue<std::vector<int>>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("ksize"));
int window_height = window[2];
int window_width = window[3];
stride_ = GetValue<std::vector<int>>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("strides"));
SetPoolingMode(kernel_node);
int windowDimA[2] = {window_height, window_width};
int paddingA[2] = {0, 0};
int strideA[2] = {stride_[2], stride_[3]};
if (pad_mode_ == kSamePadModeUpperCase || pad_mode_ == kSamePadModeLowerCase) {
pad_height_ =
std::max<int>(0, (((old_height_ / stride_[2]) * stride_[2] == old_height_ ? (old_height_ / stride_[2])
: (old_height_ / stride_[2]) + 1) -
1) *
stride_[2] +
window_height - old_height_);
pad_width_ =
std::max<int>(0, (((old_width_ / stride_[3]) * stride_[3] == old_width_ ? (old_width_ / stride_[3])
: (old_width_ / stride_[3]) + 1) -
1) *
stride_[3] +
window_width - old_width_);
pad_top_ = pad_height_ / 2;
pad_left_ = pad_width_ / 2;
paddingA[0] = pad_top_;
paddingA[1] = pad_left_;
} else {
pad_height_ = 0;
pad_width_ = 0;
}
CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetPoolingNdDescriptor(pooling_descriptor_, pooling_mode_, CUDNN_NOT_PROPAGATE_NAN,
2, windowDimA, paddingA, strideA),
"cudnnSetPoolingNdDescriptor failed");
SetPad(kernel_node);
InitSizeLists();
return true;
}
@@ -172,7 +142,6 @@ class PoolingGpuFwdKernel : public GpuKernel {
}

void SetPoolingMode(const CNodePtr &kernel_node) {
pad_mode_ = GetValue<std::string>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("padding"));
mode_ = AnfAlgo::GetCNodeName(kernel_node);
if (mode_ == "AvgPool") {
pooling_mode_ = CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING;
@@ -182,6 +151,40 @@ class PoolingGpuFwdKernel : public GpuKernel {
pad_value_ = kSignedMinFloat;
}
}
void SetPad(const CNodePtr &kernel_node) {
pad_mode_ = GetValue<std::string>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("padding"));
auto window = GetValue<std::vector<int>>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("ksize"));
int window_height = window[2];
int window_width = window[3];
stride_ = GetValue<std::vector<int>>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("strides"));
int windowDimA[2] = {window_height, window_width};
int paddingA[2] = {0, 0};
int strideA[2] = {stride_[2], stride_[3]};
if (pad_mode_ == kSamePadModeUpperCase || pad_mode_ == kSamePadModeLowerCase) {
pad_height_ =
std::max<int>(0, (((old_height_ / stride_[2]) * stride_[2] == old_height_ ? (old_height_ / stride_[2])
: (old_height_ / stride_[2]) + 1) -
1) *
stride_[2] +
window_height - old_height_);
pad_width_ =
std::max<int>(0, (((old_width_ / stride_[3]) * stride_[3] == old_width_ ? (old_width_ / stride_[3])
: (old_width_ / stride_[3]) + 1) -
1) *
stride_[3] +
window_width - old_width_);
pad_top_ = pad_height_ / 2;
pad_left_ = pad_width_ / 2;
paddingA[0] = pad_top_;
paddingA[1] = pad_left_;
} else {
pad_height_ = 0;
pad_width_ = 0;
}
CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetPoolingNdDescriptor(pooling_descriptor_, pooling_mode_, CUDNN_NOT_PROPAGATE_NAN,
2, windowDimA, paddingA, strideA),
"cudnnSetPoolingNdDescriptor failed");
}
void DestroyResource() noexcept {
CHECK_CUDNN_RET_WITH_ERROR(cudnnDestroyPoolingDescriptor(pooling_descriptor_),
"cudnnDestroyPoolingDescriptor failed");


+ 37
- 40
mindspore/ccsrc/backend/kernel_compiler/gpu/nn/pooling_grad_gpu_kernel.h View File

@@ -81,10 +81,6 @@ class PoolingGradGpuKernel : public GpuKernel {
if (!CheckParam(kernel_node)) {
return false;
}
auto window = GetAttr<std::vector<int>>(kernel_node, "ksize");
int window_height = window[2];
int window_width = window[3];
SetPoolingMode(kernel_node);
auto input_shape = AnfAlgo::GetInputDeviceShape(kernel_node, 0);
auto input_mask = AnfAlgo::GetInputDeviceShape(kernel_node, 1);
auto dout_shape = AnfAlgo::GetInputDeviceShape(kernel_node, 2);
@@ -97,9 +93,6 @@ class PoolingGradGpuKernel : public GpuKernel {
return true;
}
SetNCHW(input_shape, &n_, &c_, &old_height_, &old_width_, data_format_);
int windowDimA[2] = {window_height, window_width};
int paddingA[2] = {0, 0};
int strideA[2] = {stride_[2], stride_[3]};
const int nbDims = 4;
int dimA[4];
int strideAin[4];
@@ -126,32 +119,8 @@ class PoolingGradGpuKernel : public GpuKernel {
"cudnnSetTensor4dDescriptor failed");
CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetTensorNdDescriptor(x_descriptor_, cudnn_data_type_, nbDims, dimA, strideAin),
"cudnnSetTensor4dDescriptor failed");
if (kSamePadModeUpperCase == pad_mode_ || kSamePadModeLowerCase == pad_mode_) {
pad_height_ =
std::max<int>(0, (((old_height_ / stride_[2]) * stride_[2] == old_height_ ? (old_height_ / stride_[2])
: (old_height_ / stride_[2]) + 1) -
1) *
stride_[2] +
window_height - old_height_);
pad_width_ =
std::max<int>(0, (((old_width_ / stride_[3]) * stride_[3] == old_width_ ? (old_width_ / stride_[3])
: (old_width_ / stride_[3]) + 1) -
1) *
stride_[3] +
window_width - old_width_);
pad_top_ = pad_height_ / 2;
pad_left_ = pad_width_ / 2;
paddingA[0] = pad_top_;
paddingA[1] = pad_left_;
} else {
if (pad_mode_ == kValidPadModeUpperCase || pad_mode_ == kValidPadModeLowerCase) {
pad_height_ = 0;
pad_width_ = 0;
}
}
CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetPoolingNdDescriptor(pooling_descriptor_, pooling_mode_, CUDNN_NOT_PROPAGATE_NAN,
2, windowDimA, paddingA, strideA),
"cudnnSetPoolingNdDescriptor failed");
SetPoolingMode(kernel_node);
SetPad(kernel_node);
InitSizeLists();
return true;
}
@@ -198,15 +167,43 @@ class PoolingGradGpuKernel : public GpuKernel {
}
return true;
}
void SetPad(const std::vector<size_t> &input_shape, const int &window_height, const int &window_width) {
n_ = SizeToInt(input_shape[0]);
c_ = SizeToInt(input_shape[1]);
old_height_ = SizeToInt(input_shape[2]);
old_width_ = SizeToInt(input_shape[3]);
}
void SetPoolingMode(const CNodePtr &kernel_node) {
void SetPad(const CNodePtr &kernel_node) {
pad_mode_ = GetAttr<std::string>(kernel_node, "padding");
stride_ = GetAttr<std::vector<int>>(kernel_node, "strides");
auto window = GetAttr<std::vector<int>>(kernel_node, "ksize");
int window_height = window[2];
int window_width = window[3];
int windowDimA[2] = {window_height, window_width};
int paddingA[2] = {0, 0};
int strideA[2] = {stride_[2], stride_[3]};
if (kSamePadModeUpperCase == pad_mode_ || kSamePadModeLowerCase == pad_mode_) {
pad_height_ =
std::max<int>(0, (((old_height_ / stride_[2]) * stride_[2] == old_height_ ? (old_height_ / stride_[2])
: (old_height_ / stride_[2]) + 1) -
1) *
stride_[2] +
window_height - old_height_);
pad_width_ =
std::max<int>(0, (((old_width_ / stride_[3]) * stride_[3] == old_width_ ? (old_width_ / stride_[3])
: (old_width_ / stride_[3]) + 1) -
1) *
stride_[3] +
window_width - old_width_);
pad_top_ = pad_height_ / 2;
pad_left_ = pad_width_ / 2;
paddingA[0] = pad_top_;
paddingA[1] = pad_left_;
} else {
if (pad_mode_ == kValidPadModeUpperCase || pad_mode_ == kValidPadModeLowerCase) {
pad_height_ = 0;
pad_width_ = 0;
}
}
CHECK_CUDNN_RET_WITH_EXCEPT(cudnnSetPoolingNdDescriptor(pooling_descriptor_, pooling_mode_, CUDNN_NOT_PROPAGATE_NAN,
2, windowDimA, paddingA, strideA),
"cudnnSetPoolingNdDescriptor failed");
}
void SetPoolingMode(const CNodePtr &kernel_node) {
cudnn_data_type_ = GetCudnnDataType(TypeIdLabel(AnfAlgo::GetInputDeviceDataType(kernel_node, 0)));
mode_ = AnfAlgo::GetCNodeName(kernel_node);
if (mode_ == "AvgPoolGradGpu") {


Write Preview
Loading…
Cancel
Save