forked from mindspore-Ecosystem/mindspore
!3831 CUDA - GPU MirrorPad New Op
Merge pull request !3831 from danishnxt/GPU_One
This commit is contained in:
Gitee
commit
8f17535045
|
|
@ -0,0 +1,182 @@
|
||||||
|
/**
|
||||||
|
* Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
*
|
||||||
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
* you may not use this file except in compliance with the License.
|
||||||
|
* You may obtain a copy of the License at
|
||||||
|
*
|
||||||
|
* http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
*
|
||||||
|
* Unless required by applicable law or agreed to in writing, software
|
||||||
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
* See the License for the specific language governing permissions and
|
||||||
|
* limitations under the License.
|
||||||
|
*/
|
||||||
|
#include <assert.h>
|
||||||
|
#include <stdio.h>
|
||||||
|
#include <stdint.h>
|
||||||
|
#include "backend/kernel_compiler/gpu/cuda_impl/mirror_pad_impl.cuh"
|
||||||
|
|
||||||
|
__inline__ __device__ bool range_check(int x, int y, int padded_width, int padded_height) {
|
||||||
|
// check for existence in current padded array
|
||||||
|
if (((x >= 0) && (x <= padded_width - 1)) && ((y >= 0) && (y <= padded_height - 1))) {
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
__global__ void MirrorPad(const size_t size, const T *input, const int num, const int channels, const int old_height,
|
||||||
|
const int old_width, const int padded_height, const int padded_width, const int padd_dim,
|
||||||
|
const int *paddings, int mode, T *output) {
|
||||||
|
int padd_offset = 4 * (padd_dim - 2);
|
||||||
|
int pad_left_ = paddings[padd_offset + 4];
|
||||||
|
int pad_top_ = paddings[padd_offset + 0];
|
||||||
|
|
||||||
|
// Create anchor points for old tensor positions inside new tensor
|
||||||
|
int ap1_x = pad_left_;
|
||||||
|
int ap1_y = pad_top_;
|
||||||
|
int ap2_x = pad_left_ + old_width - 1;
|
||||||
|
int ap2_y = pad_top_ + old_height - 1;
|
||||||
|
|
||||||
|
for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < size; pos += blockDim.x * gridDim.x) {
|
||||||
|
int block_num = (pos / padded_width) / padded_height;
|
||||||
|
const int padded_x = pos % padded_width;
|
||||||
|
const int padded_y = (pos / padded_width) % padded_height;
|
||||||
|
|
||||||
|
// distance to move from anchor point
|
||||||
|
int x_dist = 0;
|
||||||
|
int y_dist = 0;
|
||||||
|
|
||||||
|
// x,y value to mirror in new tenspr
|
||||||
|
int matchval_x_index = padded_x;
|
||||||
|
int matchval_y_index = padded_y;
|
||||||
|
|
||||||
|
if (padded_y - pad_top_ < 0 || padded_x - pad_left_ < 0 || padded_y - pad_top_ >= old_height ||
|
||||||
|
padded_x - pad_left_ >= old_width) {
|
||||||
|
if ((padded_x < ap1_x) || (padded_x > ap2_x)) {
|
||||||
|
x_dist = (padded_x < ap1_x) ? (ap1_x - padded_x) : (padded_x - ap2_x); // GEN DIST
|
||||||
|
matchval_x_index = (padded_x < ap1_x) ? (ap1_x + x_dist - mode) : (ap2_x - x_dist + mode);
|
||||||
|
}
|
||||||
|
if ((padded_y < ap1_y) || (padded_y > ap2_y)) {
|
||||||
|
y_dist = (padded_y < ap1_y) ? (ap1_y - padded_y) : (padded_y - ap2_y);
|
||||||
|
matchval_y_index = (padded_y < ap1_y) ? (ap1_y + y_dist - mode) : (ap2_y - y_dist + mode);
|
||||||
|
}
|
||||||
|
output[pos] =
|
||||||
|
input[(block_num * old_height + matchval_y_index - pad_top_) * old_width + matchval_x_index - pad_left_];
|
||||||
|
} else {
|
||||||
|
// existing values remain the same
|
||||||
|
output[pos] = input[(block_num * old_height + padded_y - pad_top_) * old_width + padded_x - pad_left_];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
__global__ void MirrorPadGrad(const size_t size, const T *dy, const int num, const int channels,
|
||||||
|
const int padded_height, const int padded_width, const int old_height,
|
||||||
|
const int old_width, const int padd_dim, const int *paddings, int mode, T *dx) {
|
||||||
|
int padd_offset = 4 * (padd_dim - 2);
|
||||||
|
int pad_left_ = paddings[padd_offset + 4];
|
||||||
|
int pad_top_ = paddings[padd_offset + 0];
|
||||||
|
|
||||||
|
// Create anchor points for positions in the dy array
|
||||||
|
int ap1_x = pad_left_;
|
||||||
|
int ap1_y = pad_top_;
|
||||||
|
int ap2_x = pad_left_ + old_width - 1;
|
||||||
|
int ap2_y = pad_top_ + old_height - 1;
|
||||||
|
|
||||||
|
int adjust = 0; // adjust dist from reflection axis for symmetric padding
|
||||||
|
if (mode == 1) {
|
||||||
|
adjust = 1;
|
||||||
|
}
|
||||||
|
|
||||||
|
for (size_t pos = blockIdx.x * blockDim.x + threadIdx.x; pos < size; pos += blockDim.x * gridDim.x) {
|
||||||
|
int block_num = (pos / old_width) / old_height;
|
||||||
|
|
||||||
|
// refer to indices of original values inside padded array
|
||||||
|
const int padded_x = (pos % old_width) + pad_left_;
|
||||||
|
const int padded_y = ((pos / old_width) % old_height) + pad_top_;
|
||||||
|
|
||||||
|
// copy positions own value into output
|
||||||
|
dx[pos] = dx[pos] + dy[(block_num * padded_height + padded_y) * padded_width + padded_x];
|
||||||
|
|
||||||
|
int x_dist_1 = (ap1_x - padded_x - adjust);
|
||||||
|
int y_dist_1 = (ap1_y - padded_y - adjust);
|
||||||
|
int x_dist_2 = (ap2_x - padded_x + adjust);
|
||||||
|
int y_dist_2 = (ap2_y - padded_y + adjust);
|
||||||
|
|
||||||
|
int axis_dist[] = {x_dist_1, x_dist_2, y_dist_1, y_dist_2};
|
||||||
|
int anch_point[] = {ap1_x, ap2_x, ap1_y, ap2_y};
|
||||||
|
bool x_axis_check[] = {true, true, false, false}; // true - update X , false - update Y
|
||||||
|
|
||||||
|
int temp_x = 0;
|
||||||
|
int temp_y = 0;
|
||||||
|
|
||||||
|
// mirroring in axis lines
|
||||||
|
for (int x = 0; x < 4; x++) {
|
||||||
|
if (axis_dist[x] != 0) {
|
||||||
|
if (x_axis_check[x]) {
|
||||||
|
temp_y = padded_y;
|
||||||
|
temp_x = anch_point[x] + axis_dist[x];
|
||||||
|
} else {
|
||||||
|
temp_x = padded_x;
|
||||||
|
temp_y = anch_point[x] + axis_dist[x];
|
||||||
|
}
|
||||||
|
if (range_check(temp_x, temp_y, padded_width, padded_height)) {
|
||||||
|
dx[pos] = dx[pos] + dy[(block_num * padded_height + temp_y) * padded_width + temp_x];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
// mirroring at corners
|
||||||
|
for (int x = 0; x < 2; x++) {
|
||||||
|
for (int y = 2; y < 4; y++) {
|
||||||
|
if ((axis_dist[x] != 0) && (axis_dist[y] != 0)) {
|
||||||
|
temp_x = anch_point[x] + axis_dist[x];
|
||||||
|
temp_y = anch_point[y] + axis_dist[y];
|
||||||
|
if (range_check(temp_x, temp_y, padded_width, padded_height)) {
|
||||||
|
dx[pos] = dx[pos] + dy[(block_num * padded_height + temp_y) * padded_width + temp_x];
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
}
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
void CalMirrorPad(const size_t size, const T *input, const int num, const int channels, const int old_height,
|
||||||
|
const int old_width, const int padded_height, const int padded_width, int padd_num,
|
||||||
|
const int *paddings, const int mode, T *output, cudaStream_t cuda_stream) {
|
||||||
|
MirrorPad<<<GET_BLOCKS(size), GET_THREADS, 0, cuda_stream>>>(
|
||||||
|
size, input, num, channels, old_height, old_width, padded_height, padded_width, padd_num, paddings, mode, output);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
void CalMirrorPadGrad(const size_t size, const T *dy, const int num, const int channels, const int padded_height,
|
||||||
|
const int padded_width, const int old_height, const int old_width, const int padd_dim,
|
||||||
|
const int *paddings, int mode, T *dx, cudaStream_t cuda_stream) {
|
||||||
|
MirrorPadGrad<<<GET_BLOCKS(size), GET_THREADS, 0, cuda_stream>>>(size, dy, num, channels, padded_height, padded_width,
|
||||||
|
old_height, old_width, padd_dim, paddings, mode, dx);
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
template void CalMirrorPad<float>(const size_t size, const float *input, const int num, const int channels,
|
||||||
|
const int old_height, const int old_width, const int padded_height,
|
||||||
|
const int padded_width, int padd_num, const int *paddings, int mode, float *output,
|
||||||
|
cudaStream_t cuda_stream);
|
||||||
|
template void CalMirrorPadGrad<float>(const size_t size, const float *dy, const int num, const int channels,
|
||||||
|
const int old_height, const int old_width, const int padded_height,
|
||||||
|
const int padded_width, const int padd_dim, const int *paddings, int mode,
|
||||||
|
float *dx, cudaStream_t cuda_stream);
|
||||||
|
template void CalMirrorPad<half>(const size_t size, const half *input, const int num, const int channels,
|
||||||
|
const int old_height, const int old_width, const int padded_height,
|
||||||
|
const int padded_width, int padd_num, const int *paddings, int mode, half *output,
|
||||||
|
cudaStream_t cuda_stream);
|
||||||
|
template void CalMirrorPadGrad<half>(const size_t size, const half *dy, const int num, const int channels,
|
||||||
|
const int old_height, const int old_width, const int padded_height,
|
||||||
|
const int padded_width, const int padd_dim, const int *paddings, int mode,
|
||||||
|
half *dx, cudaStream_t cuda_stream);
|
||||||
|
|
@ -0,0 +1,31 @@
|
||||||
|
/**
|
||||||
|
* Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
*
|
||||||
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
* you may not use this file except in compliance with the License.
|
||||||
|
* You may obtain a copy of the License at
|
||||||
|
*
|
||||||
|
* http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
*
|
||||||
|
* Unless required by applicable law or agreed to in writing, software
|
||||||
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
* See the License for the specific language governing permissions and
|
||||||
|
* limitations under the License.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_MIRROR_PAD_IMPL_H_
|
||||||
|
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_MIRROR_PAD_IMPL_H_
|
||||||
|
#include <cuda_runtime.h>
|
||||||
|
#include "runtime/device/gpu/cuda_common.h"
|
||||||
|
|
||||||
|
template <typename T>
|
||||||
|
void CalMirrorPad(const size_t size, const T *input, const int num, const int channels, const int old_height,
|
||||||
|
const int old_width, const int padded_height, const int padded_width, int padd_num,
|
||||||
|
const int *paddings, int mode, T *output, cudaStream_t cuda_stream);
|
||||||
|
template <typename T>
|
||||||
|
void CalMirrorPadGrad(const size_t size, const T *dy, const int num, const int channels, const int padded_height,
|
||||||
|
const int padded_width, const int old_height, const int old_width, const int padd_dim,
|
||||||
|
const int *paddings, int mode, T *dx, cudaStream_t cuda_stream);
|
||||||
|
|
||||||
|
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_CUDA_IMPL_MIRROR_PAD_IMPL_H_
|
||||||
|
|
@ -0,0 +1,30 @@
|
||||||
|
/**
|
||||||
|
* Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
*
|
||||||
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
* you may not use this file except in compliance with the License.
|
||||||
|
* You may obtain a copy of the License at
|
||||||
|
*
|
||||||
|
* http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
*
|
||||||
|
* Unless required by applicable law or agreed to in writing, software
|
||||||
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
* See the License for the specific language governing permissions and
|
||||||
|
* limitations under the License.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#include "backend/kernel_compiler/gpu/nn/mirror_pad_gpu_kernel.h"
|
||||||
|
|
||||||
|
namespace mindspore {
|
||||||
|
namespace kernel {
|
||||||
|
MS_REG_GPU_KERNEL_ONE(
|
||||||
|
MirrorPad,
|
||||||
|
KernelAttr().AddInputAttr(kNumberTypeFloat32).AddInputAttr(kNumberTypeInt64).AddOutputAttr(kNumberTypeFloat32),
|
||||||
|
MirrorPadGpuFwdKernel, float)
|
||||||
|
MS_REG_GPU_KERNEL_ONE(
|
||||||
|
MirrorPad,
|
||||||
|
KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeInt64).AddOutputAttr(kNumberTypeFloat16),
|
||||||
|
MirrorPadGpuFwdKernel, half)
|
||||||
|
} // namespace kernel
|
||||||
|
} // namespace mindspore
|
||||||
|
|
@ -0,0 +1,150 @@
|
||||||
|
/**
|
||||||
|
* Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
*
|
||||||
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
* you may not use this file except in compliance with the License.
|
||||||
|
* You may obtain a copy of the License at
|
||||||
|
*
|
||||||
|
* http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
*
|
||||||
|
* Unless required by applicable law or agreed to in writing, software
|
||||||
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
* See the License for the specific language governing permissions and
|
||||||
|
* limitations under the License.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_MIRROR_PAD_GPU_KERNEL_H_
|
||||||
|
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_MIRROR_PAD_GPU_KERNEL_H_
|
||||||
|
|
||||||
|
#include <iostream>
|
||||||
|
#include <vector>
|
||||||
|
#include <string>
|
||||||
|
#include "backend/kernel_compiler/gpu/gpu_kernel.h"
|
||||||
|
#include "backend/kernel_compiler/gpu/gpu_kernel_factory.h"
|
||||||
|
#include "backend/kernel_compiler/gpu/cuda_impl/mirror_pad_impl.cuh"
|
||||||
|
|
||||||
|
namespace mindspore {
|
||||||
|
namespace kernel {
|
||||||
|
template <typename T>
|
||||||
|
class MirrorPadGpuFwdKernel : public GpuKernel {
|
||||||
|
public:
|
||||||
|
MirrorPadGpuFwdKernel()
|
||||||
|
: num_input_(0), num_paddings_(0), mode_(0), input_size_(1), output_size_(1), workspace_size_(0) {}
|
||||||
|
~MirrorPadGpuFwdKernel() override = default;
|
||||||
|
|
||||||
|
const std::vector<size_t> &GetInputSizeList() const override { return input_size_list_; }
|
||||||
|
const std::vector<size_t> &GetOutputSizeList() const override { return output_size_list_; }
|
||||||
|
const std::vector<size_t> &GetWorkspaceSizeList() const override { return workspace_size_list_; }
|
||||||
|
|
||||||
|
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
|
||||||
|
const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
|
||||||
|
T *input = GetDeviceAddress<T>(inputs, 0);
|
||||||
|
int *paddings = GetDeviceAddress<int>(inputs, 1);
|
||||||
|
T *output = GetDeviceAddress<T>(outputs, 0);
|
||||||
|
|
||||||
|
size_t size = output_size_ / sizeof(T);
|
||||||
|
int dim_offset = output_shape_.size() - 2;
|
||||||
|
|
||||||
|
CalMirrorPad(size, input, input_shape_[0], input_shape_[1], input_shape_[2], input_shape_[3],
|
||||||
|
output_shape_[dim_offset + 0], output_shape_[dim_offset + 1], num_paddings_, paddings, mode_, output,
|
||||||
|
reinterpret_cast<cudaStream_t>(stream_ptr));
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
|
bool Init(const CNodePtr &kernel_node) override {
|
||||||
|
size_t input_num = AnfAlgo::GetInputTensorNum(kernel_node);
|
||||||
|
if (input_num != 2) {
|
||||||
|
MS_LOG(ERROR) << "Input number is " << input_num << ", but MirrorPad needs 2 input.";
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
// check number of output -> should be 1
|
||||||
|
size_t output_num = AnfAlgo::GetOutputTensorNum(kernel_node);
|
||||||
|
if (output_num != 1) {
|
||||||
|
MS_LOG(ERROR) << "Output number is " << output_num << ", but Pad needs 1 output.";
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
string mode = GetValue<string>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("mode"));
|
||||||
|
|
||||||
|
if (mode == "REFLECT") {
|
||||||
|
mode_ = 0; // reflected mirroring
|
||||||
|
} else {
|
||||||
|
mode_ = 1; // symmetric mirroring
|
||||||
|
}
|
||||||
|
|
||||||
|
auto input_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
|
||||||
|
// shape adjustement -> from 2d/3d to 4d to standardize
|
||||||
|
if (input_shape.size() == 4) {
|
||||||
|
} else if (input_shape.size() == 3) {
|
||||||
|
auto it = input_shape.begin();
|
||||||
|
input_shape.insert(it, 1); // batch padding
|
||||||
|
} else if (input_shape.size() == 2) {
|
||||||
|
auto it = input_shape.begin();
|
||||||
|
input_shape.insert(it, 2, 1); // channel padding
|
||||||
|
}
|
||||||
|
|
||||||
|
for (auto in_shape : input_shape) {
|
||||||
|
input_size_ *= in_shape;
|
||||||
|
input_shape_.push_back(in_shape);
|
||||||
|
}
|
||||||
|
num_input_ = input_size_;
|
||||||
|
input_size_ *= sizeof(T);
|
||||||
|
|
||||||
|
auto padding_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 1);
|
||||||
|
num_paddings_ = padding_shape[0];
|
||||||
|
input_size_ += 2 * num_paddings_ * sizeof(int);
|
||||||
|
|
||||||
|
output_size_ = sizeof(T);
|
||||||
|
auto output_shape = AnfAlgo::GetOutputInferShape(kernel_node, 0);
|
||||||
|
for (auto x : output_shape) {
|
||||||
|
output_size_ *= x;
|
||||||
|
output_shape_.push_back(x);
|
||||||
|
}
|
||||||
|
|
||||||
|
int max_width = input_shape_[3];
|
||||||
|
int max_height = input_shape_[2];
|
||||||
|
|
||||||
|
// basic error check for padding value
|
||||||
|
if (mode_ == 1) { // symmetric
|
||||||
|
max_width = max_width + (2 * max_width);
|
||||||
|
max_height = max_height + (2 * max_height);
|
||||||
|
} else { // reflect
|
||||||
|
max_width = max_width + (2 * (max_width - 1));
|
||||||
|
max_height = max_height + (2 * (max_height - 1));
|
||||||
|
}
|
||||||
|
|
||||||
|
if (output_shape_[(output_shape_.size() - 2) + 0] > max_width ||
|
||||||
|
output_shape_[(output_shape_.size() - 2) + 1] > max_width) {
|
||||||
|
MS_LOG(ERROR) << "ERROR: Padding value too high for input Tensor on 1 or more dims";
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
InitSizeLists();
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
|
protected:
|
||||||
|
void InitSizeLists() override {
|
||||||
|
input_size_list_.push_back(num_input_ * sizeof(T));
|
||||||
|
input_size_list_.push_back(2 * num_paddings_ * sizeof(int));
|
||||||
|
output_size_list_.push_back(output_size_);
|
||||||
|
}
|
||||||
|
|
||||||
|
private:
|
||||||
|
size_t num_input_;
|
||||||
|
int num_paddings_;
|
||||||
|
int mode_;
|
||||||
|
std::vector<int> input_shape_; // dims of the input data
|
||||||
|
std::vector<int> output_shape_; // dims of the output data
|
||||||
|
// default
|
||||||
|
size_t input_size_;
|
||||||
|
size_t output_size_;
|
||||||
|
size_t workspace_size_;
|
||||||
|
std::vector<size_t> input_size_list_;
|
||||||
|
std::vector<size_t> output_size_list_;
|
||||||
|
std::vector<size_t> workspace_size_list_;
|
||||||
|
};
|
||||||
|
} // namespace kernel
|
||||||
|
} // namespace mindspore
|
||||||
|
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_MIRROR_PAD_GPU_KERNEL_H_
|
||||||
|
|
@ -0,0 +1,30 @@
|
||||||
|
/**
|
||||||
|
* Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
*
|
||||||
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
* you may not use this file except in compliance with the License.
|
||||||
|
* You may obtain a copy of the License at
|
||||||
|
*
|
||||||
|
* http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
*
|
||||||
|
* Unless required by applicable law or agreed to in writing, software
|
||||||
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
* See the License for the specific language governing permissions and
|
||||||
|
* limitations under the License.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#include "backend/kernel_compiler/gpu/nn/mirror_pad_grad_gpu_kernel.h"
|
||||||
|
|
||||||
|
namespace mindspore {
|
||||||
|
namespace kernel {
|
||||||
|
MS_REG_GPU_KERNEL_ONE(
|
||||||
|
MirrorPadGrad,
|
||||||
|
KernelAttr().AddInputAttr(kNumberTypeFloat32).AddInputAttr(kNumberTypeInt64).AddOutputAttr(kNumberTypeFloat32),
|
||||||
|
MirrorPadGpuBackKernel, float)
|
||||||
|
MS_REG_GPU_KERNEL_ONE(
|
||||||
|
MirrorPadGrad,
|
||||||
|
KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeInt64).AddOutputAttr(kNumberTypeFloat16),
|
||||||
|
MirrorPadGpuBackKernel, half)
|
||||||
|
} // namespace kernel
|
||||||
|
} // namespace mindspore
|
||||||
|
|
@ -0,0 +1,150 @@
|
||||||
|
/**
|
||||||
|
* Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
*
|
||||||
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
* you may not use this file except in compliance with the License.
|
||||||
|
* You may obtain a copy of the License at
|
||||||
|
*
|
||||||
|
* http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
*
|
||||||
|
* Unless required by applicable law or agreed to in writing, software
|
||||||
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
* See the License for the specific language governing permissions and
|
||||||
|
* limitations under the License.
|
||||||
|
*/
|
||||||
|
|
||||||
|
#ifndef MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_MIRROR_PAD_GRAD_GPU_KERNEL_H_
|
||||||
|
#define MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_MIRROR_PAD_GRAD_GPU_KERNEL_H_
|
||||||
|
|
||||||
|
#include <iostream>
|
||||||
|
#include <vector>
|
||||||
|
#include <string>
|
||||||
|
#include "backend/kernel_compiler/gpu/gpu_kernel.h"
|
||||||
|
#include "backend/kernel_compiler/gpu/gpu_kernel_factory.h"
|
||||||
|
#include "backend/kernel_compiler/gpu/cuda_impl/mirror_pad_impl.cuh"
|
||||||
|
|
||||||
|
namespace mindspore {
|
||||||
|
namespace kernel {
|
||||||
|
template <typename T>
|
||||||
|
class MirrorPadGpuBackKernel : public GpuKernel {
|
||||||
|
public:
|
||||||
|
MirrorPadGpuBackKernel()
|
||||||
|
: num_input_(0), num_paddings_(0), mode_(0), input_size_(1), output_size_(1), workspace_size_(0) {}
|
||||||
|
~MirrorPadGpuBackKernel() override = default;
|
||||||
|
|
||||||
|
const std::vector<size_t> &GetInputSizeList() const override { return input_size_list_; }
|
||||||
|
const std::vector<size_t> &GetOutputSizeList() const override { return output_size_list_; }
|
||||||
|
const std::vector<size_t> &GetWorkspaceSizeList() const override { return workspace_size_list_; }
|
||||||
|
|
||||||
|
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
|
||||||
|
const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
|
||||||
|
T *input = GetDeviceAddress<T>(inputs, 0);
|
||||||
|
int *paddings = GetDeviceAddress<int>(inputs, 1);
|
||||||
|
T *output = GetDeviceAddress<T>(outputs, 0);
|
||||||
|
|
||||||
|
size_t size = output_size_ / sizeof(T);
|
||||||
|
int dim_offset = output_shape_.size() - 2;
|
||||||
|
|
||||||
|
CalMirrorPadGrad(size, input, input_shape_[0], input_shape_[1], input_shape_[2], input_shape_[3],
|
||||||
|
output_shape_[dim_offset + 0], output_shape_[dim_offset + 1], num_paddings_, paddings, mode_,
|
||||||
|
output, reinterpret_cast<cudaStream_t>(stream_ptr));
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
|
bool Init(const CNodePtr &kernel_node) override {
|
||||||
|
size_t input_num = AnfAlgo::GetInputTensorNum(kernel_node);
|
||||||
|
if (input_num != 2) {
|
||||||
|
MS_LOG(ERROR) << "Input number is " << input_num << ", but MirrorPadGrad needs 2 input.";
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
// check number of output -> should be 1
|
||||||
|
size_t output_num = AnfAlgo::GetOutputTensorNum(kernel_node);
|
||||||
|
if (output_num != 1) {
|
||||||
|
MS_LOG(ERROR) << "Output number is " << output_num << ", but MirrorPadGrad needs 1 output.";
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
string mode = GetValue<string>(AnfAlgo::GetCNodePrimitive(kernel_node)->GetAttr("mode"));
|
||||||
|
|
||||||
|
if (mode == "REFLECT") {
|
||||||
|
mode_ = 0; // reflected mirroring
|
||||||
|
} else {
|
||||||
|
mode_ = 1; // symmetric mirroring
|
||||||
|
}
|
||||||
|
|
||||||
|
auto input_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
|
||||||
|
// shape adjustement -> from 2d/3d to 4d to standardize
|
||||||
|
if (input_shape.size() == 4) {
|
||||||
|
} else if (input_shape.size() == 3) {
|
||||||
|
auto it = input_shape.begin();
|
||||||
|
input_shape.insert(it, 1); // batch padding
|
||||||
|
} else if (input_shape.size() == 2) {
|
||||||
|
auto it = input_shape.begin();
|
||||||
|
input_shape.insert(it, 2, 1); // channel padding
|
||||||
|
}
|
||||||
|
|
||||||
|
for (auto in_shape : input_shape) {
|
||||||
|
input_size_ *= in_shape;
|
||||||
|
input_shape_.push_back(in_shape);
|
||||||
|
}
|
||||||
|
num_input_ = input_size_;
|
||||||
|
input_size_ *= sizeof(T);
|
||||||
|
|
||||||
|
auto padding_shape = AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 1);
|
||||||
|
num_paddings_ = padding_shape[0];
|
||||||
|
input_size_ += +(2 * num_paddings_ * sizeof(int));
|
||||||
|
|
||||||
|
output_size_ = sizeof(T);
|
||||||
|
auto output_shape = AnfAlgo::GetOutputInferShape(kernel_node, 0);
|
||||||
|
for (auto x : output_shape) {
|
||||||
|
output_size_ *= x;
|
||||||
|
output_shape_.push_back(x);
|
||||||
|
}
|
||||||
|
|
||||||
|
int max_width = input_shape_[3];
|
||||||
|
int max_height = input_shape_[2];
|
||||||
|
|
||||||
|
// basic error check for padding value
|
||||||
|
if (mode_ == 1) { // symmetric
|
||||||
|
max_width = max_width + (2 * max_width);
|
||||||
|
max_height = max_height + (2 * max_height);
|
||||||
|
} else { // reflect
|
||||||
|
max_width = max_width + (2 * (max_width - 1));
|
||||||
|
max_height = max_height + (2 * (max_height - 1));
|
||||||
|
}
|
||||||
|
|
||||||
|
if (output_shape_[(output_shape_.size() - 2) + 0] > max_width ||
|
||||||
|
output_shape_[(output_shape_.size() - 2) + 1] > max_width) {
|
||||||
|
MS_LOG(ERROR) << "ERROR: Padding value too high for input Tensor on 1 or more DIMS";
|
||||||
|
return false;
|
||||||
|
}
|
||||||
|
|
||||||
|
InitSizeLists();
|
||||||
|
return true;
|
||||||
|
}
|
||||||
|
|
||||||
|
protected:
|
||||||
|
void InitSizeLists() override {
|
||||||
|
input_size_list_.push_back(num_input_ * sizeof(T));
|
||||||
|
input_size_list_.push_back(2 * num_paddings_ * sizeof(int));
|
||||||
|
output_size_list_.push_back(output_size_);
|
||||||
|
}
|
||||||
|
|
||||||
|
private:
|
||||||
|
size_t num_input_;
|
||||||
|
int num_paddings_;
|
||||||
|
int mode_;
|
||||||
|
std::vector<int> input_shape_; // dims of the input data
|
||||||
|
std::vector<int> output_shape_; // dims of the output data
|
||||||
|
// default
|
||||||
|
size_t input_size_;
|
||||||
|
size_t output_size_;
|
||||||
|
size_t workspace_size_;
|
||||||
|
std::vector<size_t> input_size_list_;
|
||||||
|
std::vector<size_t> output_size_list_;
|
||||||
|
std::vector<size_t> workspace_size_list_;
|
||||||
|
};
|
||||||
|
} // namespace kernel
|
||||||
|
} // namespace mindspore
|
||||||
|
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_MIRROR_PAD_GRAD_GPU_KERNEL_H_
|
||||||
|
|
@ -2761,12 +2761,17 @@ class MirrorPad(PrimitiveWithInfer):
|
||||||
paddings_value = paddings['value'].asnumpy()
|
paddings_value = paddings['value'].asnumpy()
|
||||||
paddings_size = paddings_value.size
|
paddings_size = paddings_value.size
|
||||||
validator.check_integer('paddings.shape', paddings_size, len(x_shape) * 2, Rel.EQ, self.name)
|
validator.check_integer('paddings.shape', paddings_size, len(x_shape) * 2, Rel.EQ, self.name)
|
||||||
if not np.all(paddings_size >= 0):
|
if not np.all(paddings_value >= 0):
|
||||||
raise ValueError('All elements of paddings must be >= 0.')
|
raise ValueError('All elements of paddings must be >= 0.')
|
||||||
|
adjust = 0
|
||||||
|
if self.mode == 'SYMMETRIC':
|
||||||
|
adjust = 1
|
||||||
|
for i in range(0, int(paddings_size / 2)):
|
||||||
|
if (paddings_value[i, 0] >= x_shape[i] + adjust) or (paddings_value[i, 1] >= x_shape[i] + adjust):
|
||||||
|
raise ValueError('At least one dim has too high a padding value for this input and mode')
|
||||||
y_shape = ()
|
y_shape = ()
|
||||||
for i in range(0, int(paddings_size / 2)):
|
for i in range(0, int(paddings_size / 2)):
|
||||||
y_shape += ((x_shape[i] + paddings_value[i, 0] + paddings_value[i, 1]),)
|
y_shape += ((x_shape[i] + paddings_value[i, 0] + paddings_value[i, 1]),)
|
||||||
|
|
||||||
return {'shape': y_shape,
|
return {'shape': y_shape,
|
||||||
'dtype': input_x['dtype'],
|
'dtype': input_x['dtype'],
|
||||||
'value': None}
|
'value': None}
|
||||||
|
|
|
||||||
|
|
@ -0,0 +1,88 @@
|
||||||
|
# Copyright 2020 Huawei Technologies Co., Ltd
|
||||||
|
#
|
||||||
|
# Licensed under the Apache License, Version 2.0 (the "License");
|
||||||
|
# you may not use this file except in compliance with the License.
|
||||||
|
# You may obtain a copy of the License at
|
||||||
|
#
|
||||||
|
# http://www.apache.org/licenses/LICENSE-2.0
|
||||||
|
#
|
||||||
|
# Unless required by applicable law or agreed to in writing, software
|
||||||
|
# distributed under the License is distributed on an "AS IS" BASIS,
|
||||||
|
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
||||||
|
# See the License for the specific language governing permissions and
|
||||||
|
# limitations under the License.
|
||||||
|
# ============================================================================
|
||||||
|
|
||||||
|
import pytest
|
||||||
|
import numpy as np
|
||||||
|
|
||||||
|
import mindspore
|
||||||
|
import mindspore.nn as nn
|
||||||
|
import mindspore.context as context
|
||||||
|
|
||||||
|
from mindspore import Tensor
|
||||||
|
from mindspore.ops.composite import GradOperation
|
||||||
|
|
||||||
|
@pytest.mark.level0
|
||||||
|
@pytest.mark.platform_x86_gpu_training
|
||||||
|
@pytest.mark.env_onecard
|
||||||
|
def test_mirror_pad():
|
||||||
|
context.set_context(mode=context.GRAPH_MODE, device_target="GPU")
|
||||||
|
|
||||||
|
test1_arr_in = [[[[1, 2, 3], [4, 5, 6], [7, 8, 9]]]]
|
||||||
|
test_1_paddings = ((0, 0), (0, 0), (1, 1), (2, 2))
|
||||||
|
test1_arr_exp = [[[[6, 5, 4, 5, 6, 5, 4], [3, 2, 1, 2, 3, 2, 1], [6, 5, 4, 5, 6, 5, 4],
|
||||||
|
[9, 8, 7, 8, 9, 8, 7], [6, 5, 4, 5, 6, 5, 4]]]]
|
||||||
|
|
||||||
|
test2_arr_in = [[[[1, 2, 3], [4, 5, 6], [7, 8, 9]]]]
|
||||||
|
test_2_paddings = ((0, 0), (0, 0), (1, 1), (2, 2))
|
||||||
|
test2_arr_exp = [[[[2, 1, 1, 2, 3, 3, 2], [2, 1, 1, 2, 3, 3, 2], [5, 4, 4, 5, 6, 6, 5],
|
||||||
|
[8, 7, 7, 8, 9, 9, 8], [8, 7, 7, 8, 9, 9, 8]]]]
|
||||||
|
|
||||||
|
reflectOp = nn.Pad(mode='REFLECT', paddings=test_1_paddings)
|
||||||
|
symmOp = nn.Pad(mode='SYMMETRIC', paddings=test_2_paddings)
|
||||||
|
|
||||||
|
x_test_1 = Tensor(np.array(test1_arr_in), dtype=mindspore.float32)
|
||||||
|
x_test_2 = Tensor(np.array(test2_arr_in), dtype=mindspore.float32)
|
||||||
|
|
||||||
|
y_test_1 = reflectOp(x_test_1).asnumpy()
|
||||||
|
y_test_2 = symmOp(x_test_2).asnumpy()
|
||||||
|
|
||||||
|
print(np.array(test1_arr_in))
|
||||||
|
print(y_test_1)
|
||||||
|
|
||||||
|
np.testing.assert_equal(np.array(test1_arr_exp), y_test_1)
|
||||||
|
np.testing.assert_equal(np.array(test2_arr_exp), y_test_2)
|
||||||
|
|
||||||
|
|
||||||
|
class Grad(nn.Cell):
|
||||||
|
def __init__(self, network):
|
||||||
|
super(Grad, self).__init__()
|
||||||
|
self.grad = GradOperation(name="get_all", get_all=True, sens_param=True)
|
||||||
|
self.network = network
|
||||||
|
def construct(self, input_, output_grad):
|
||||||
|
return self.grad(self.network)(input_, output_grad)
|
||||||
|
|
||||||
|
class Net(nn.Cell):
|
||||||
|
def __init__(self):
|
||||||
|
super(Net, self).__init__()
|
||||||
|
self.pad = nn.Pad(mode="REFLECT", paddings=((0, 0), (0, 0), (1, 0), (0, 2)))
|
||||||
|
def construct(self, x):
|
||||||
|
return self.pad(x)
|
||||||
|
|
||||||
|
|
||||||
|
@pytest.mark.level0
|
||||||
|
@pytest.mark.platform_x86_gpu_training
|
||||||
|
@pytest.mark.env_onecard
|
||||||
|
def test_mirror_pad_backprop():
|
||||||
|
context.set_context(mode=context.GRAPH_MODE, device_target="GPU")
|
||||||
|
test_arr_in = [[[[1, 2, 3], [4, 5, 6], [7, 8, 9]]]] # size -> 3*3
|
||||||
|
test_arr_in = Tensor(test_arr_in, dtype=mindspore.float32)
|
||||||
|
dy = (np.ones((1, 1, 4, 5)) * 0.1).astype(np.float32)
|
||||||
|
expected_dx = np.array([[[[0.2, 0.2, 0.1],
|
||||||
|
[0.4, 0.4, 0.2],
|
||||||
|
[0.2, 0.2, 0.1]]]])
|
||||||
|
net = Grad(Net())
|
||||||
|
dx = net(test_arr_in, Tensor(dy))
|
||||||
|
dx = dx[0].asnumpy()
|
||||||
|
np.testing.assert_array_almost_equal(dx, expected_dx)
|
||||||
Loading…
Reference in New Issue