darrenlu
/
mindspore2022
forked from huawei/mindspore2022
0
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

!30559 modify gpu opt adapter dynamic shape 

Merge pull request !30559 from changzherui/mod_gpu_opt
This commit is contained in:
i-robot 2022-03-10 03:30:43 +00:00 committed by Gitee
parent 662a7f8408 47ef477fef
commit c04faad4ad
No known key found for this signature in database
GPG Key ID: 173E9B9CA92EEF8F
18 changed files with 364 additions and 22 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

65
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/one_hot_gpu_kernel.cc
View File

@ -47,5 +47,70 @@ MS_REG_GPU_KERNEL_TWO(OneHot,
.AddInputAttr(kNumberTypeFloat16)
.AddOutputAttr(kNumberTypeFloat16),
OneHotFwdGpuKernelMod, half, int64_t)
// dynamic shape
MS_REG_GPU_KERNEL_TWO(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeFloat32)
.AddOutputAttr(kNumberTypeFloat32),
OneHotFwdGpuKernelMod, float, int)
MS_REG_GPU_KERNEL_TWO(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeFloat16)
.AddOutputAttr(kNumberTypeFloat16),
OneHotFwdGpuKernelMod, half, int)
MS_REG_GPU_KERNEL_THREE(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeFloat32)
.AddOutputAttr(kNumberTypeFloat32),
OneHotFwdGpuKernelMod, float, int64_t, int)
MS_REG_GPU_KERNEL_THREE(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeFloat16)
.AddOutputAttr(kNumberTypeFloat16),
OneHotFwdGpuKernelMod, half, int64_t, int)
MS_REG_GPU_KERNEL_THREE(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeFloat32)
.AddOutputAttr(kNumberTypeFloat32),
OneHotFwdGpuKernelMod, float, int, int64_t)
MS_REG_GPU_KERNEL_THREE(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeFloat16)
.AddOutputAttr(kNumberTypeFloat16),
OneHotFwdGpuKernelMod, half, int, int64_t)
MS_REG_GPU_KERNEL_TWO(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeFloat32)
.AddOutputAttr(kNumberTypeFloat32),
OneHotFwdGpuKernelMod, float, int64_t)
MS_REG_GPU_KERNEL_TWO(OneHot,
KernelAttr()
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeFloat16)
.AddOutputAttr(kNumberTypeFloat16),
OneHotFwdGpuKernelMod, half, int64_t)
} // namespace kernel
} // namespace mindspore

34
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/one_hot_gpu_kernel.h
View File

@ -24,7 +24,8 @@
namespace mindspore {
namespace kernel {
template <typename T, typename S>
constexpr int DynamicInputNum = 4;
template <typename T, typename S, typename G = int>
class OneHotFwdGpuKernelMod : public NativeGpuKernelMod {
public:
OneHotFwdGpuKernelMod()
@ -37,9 +38,15 @@ class OneHotFwdGpuKernelMod : public NativeGpuKernelMod {
return true;
}
VARIABLE_NOT_USED(workspace);
size_t on_value_idx = 1;
size_t off_value_idx = 2;
if (is_dynamic_shape_) {
on_value_idx++;
off_value_idx++;
}
const S *indices = GetDeviceAddress<S>(inputs, 0);
const T *on_value = GetDeviceAddress<T>(inputs, 1);
const T *off_value = GetDeviceAddress<T>(inputs, 2);
const T *on_value = GetDeviceAddress<T>(inputs, on_value_idx);
const T *off_value = GetDeviceAddress<T>(inputs, off_value_idx);
T *output = GetDeviceAddress<T>(outputs, 0);
OneHot(indices, depth_, on_value, off_value, left_dim_size_, right_dim_size_, output,
reinterpret_cast<cudaStream_t>(stream_ptr));
@ -51,6 +58,10 @@ class OneHotFwdGpuKernelMod : public NativeGpuKernelMod {
int64_t axis = GetAttr<int64_t>(kernel_node, "axis");
auto input_shape = common::AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
auto output_shape = common::AnfAlgo::GetOutputInferShape(kernel_node, 0);
size_t input_num = common::AnfAlgo::GetInputTensorNum(kernel_node);
if (input_num == DynamicInputNum) {
is_dynamic_shape_ = true;
}
is_null_input_ =
CHECK_SHAPE_NULL(input_shape, kernel_name, "input") || CHECK_SHAPE_NULL(output_shape, kernel_name, "output");
if (is_null_input_) {
@ -91,11 +102,27 @@ class OneHotFwdGpuKernelMod : public NativeGpuKernelMod {
InitSizeLists();
return true;
}
void ResetResource() noexcept override {
is_dynamic_shape_ = false;
input_size_ = 1;
output_size_ = 1;
depth_ = 0;
left_dim_size_ = 1;
right_dim_size_ = 1;
is_null_input_ = false;
input_size_list_.clear();
output_size_list_.clear();
workspace_size_list_.clear();
is_dynamic_shape_ = false;
}
protected:
void InitSizeLists() override {
// inputs: indices, depth
input_size_list_.push_back((input_size_ + 1) * sizeof(S));
if (is_dynamic_shape_) {
input_size_list_.push_back(sizeof(int64_t));
}
output_size_list_.push_back(output_size_ * sizeof(T));
}
@ -103,6 +130,7 @@ class OneHotFwdGpuKernelMod : public NativeGpuKernelMod {
size_t input_size_;
size_t output_size_;
bool is_dynamic_shape_ = false;
size_t depth_;
size_t left_dim_size_;
size_t right_dim_size_;

2
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/oneslike_gpu_kernel.h
View File

@ -25,7 +25,7 @@ namespace kernel {
template <typename T>
class OnesLikeGpuKernelMod : public NativeGpuKernelMod {
public:
OnesLikeGpuKernelMod() : input_size_(0), output_size_(0), is_null_input_(false) { ResetResource(); }
OnesLikeGpuKernelMod() : input_size_(0), output_size_(0), is_null_input_(false) {}
~OnesLikeGpuKernelMod() override = default;
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &,

12
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/pack_gpu_kernel.h
View File

@ -106,6 +106,18 @@ class PackFwdGpuKernelMod : public NativeGpuKernelMod {
return true;
}
void ResetResource() noexcept override {
axis_ = 0;
is_null_input_ = false;
input_num_ = 1;
output_size_ = 0;
dims_behind_axis_ = 1;
inputs_host_ = nullptr;
input_size_list_.clear();
output_size_list_.clear();
workspace_size_list_.clear();
}
protected:
void InitSizeLists() override {}

12
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/strided_slice_gpu_common.h
View File

@ -19,6 +19,8 @@
#include <vector>
#include <algorithm>
#include "plugin/device/gpu/kernel/gpu_kernel.h"
#include "plugin/device/gpu/kernel/gpu_kernel_factory.h"
#include "backend/common/session/anf_runtime_algorithm.h"
#include "include/common/utils/anfalgo.h"
#include "kernel/common_utils.h"
@ -31,10 +33,12 @@ class StridedSliceGpuCommon {
StridedSliceGpuCommon() : null_output_(false) {}
~StridedSliceGpuCommon() = default;
void CollectInfo(const CNodePtr &kernel_node) {
begin_ = common::AnfAlgo::GetNodeAttr<std::vector<int64_t>>(kernel_node, kAttrBegin);
end_ = common::AnfAlgo::GetNodeAttr<std::vector<int64_t>>(kernel_node, kAttrEnd);
strides_ = common::AnfAlgo::GetNodeAttr<std::vector<int64_t>>(kernel_node, kAttrStrides);
void CollectInfo(const CNodePtr &kernel_node, bool is_dynamic_attr_ = false) {
if (!is_dynamic_attr_) {
begin_ = common::AnfAlgo::GetNodeAttr<std::vector<int64_t>>(kernel_node, kAttrBegin);
end_ = common::AnfAlgo::GetNodeAttr<std::vector<int64_t>>(kernel_node, kAttrEnd);
strides_ = common::AnfAlgo::GetNodeAttr<std::vector<int64_t>>(kernel_node, kAttrStrides);
}
FillEmptyDims(kernel_node, &begin_, &end_, &strides_, &input_shape_);
ParseStrideSliceMasks(kernel_node, &begin_, &end_, &strides_, input_shape_);
FillOutputDim();

96
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/strided_slice_gpu_kernel.cc
View File

@ -42,5 +42,101 @@ MS_REG_GPU_KERNEL_ONE(StridedSlice, KernelAttr().AddInputAttr(kNumberTypeUInt8).
StridedSliceGpuKernelMod, uchar)
MS_REG_GPU_KERNEL_ONE(StridedSlice, KernelAttr().AddInputAttr(kNumberTypeBool).AddOutputAttr(kNumberTypeBool),
StridedSliceGpuKernelMod, bool)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeFloat64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeFloat64),
StridedSliceGpuKernelMod, double, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeFloat32),
StridedSliceGpuKernelMod, float, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeFloat16),
StridedSliceGpuKernelMod, half, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeInt64),
StridedSliceGpuKernelMod, int64_t, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeInt32)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeInt32),
StridedSliceGpuKernelMod, int, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeInt16)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeInt16),
StridedSliceGpuKernelMod, short, int64_t) // NOLINT
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeInt8)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeInt8),
StridedSliceGpuKernelMod, int8_t, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeUInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeUInt64),
StridedSliceGpuKernelMod, uint64_t, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeUInt32)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeUInt32),
StridedSliceGpuKernelMod, uint32_t, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeUInt16)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeUInt16),
StridedSliceGpuKernelMod, uint16_t, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeUInt8)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeUInt8),
StridedSliceGpuKernelMod, uchar, int64_t)
MS_REG_GPU_KERNEL_TWO(StridedSlice,
KernelAttr()
.AddInputAttr(kNumberTypeBool)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeBool),
StridedSliceGpuKernelMod, bool, int64_t)
} // namespace kernel
} // namespace mindspore

31
mindspore/ccsrc/plugin/device/gpu/kernel/arrays/strided_slice_gpu_kernel.h
View File

@ -27,7 +27,8 @@
namespace mindspore {
namespace kernel {
template <typename T>
constexpr int DynamicInputNum = 4;
template <typename T, typename S = int64_t>
class StridedSliceGpuKernelMod : public NativeGpuKernelMod, public StridedSliceGpuCommon {
public:
StridedSliceGpuKernelMod() = default;
@ -48,6 +49,10 @@ class StridedSliceGpuKernelMod : public NativeGpuKernelMod, public StridedSliceG
bool Init(const CNodePtr &kernel_node) override {
auto kernel_name = common::AnfAlgo::GetCNodeName(kernel_node);
size_t input_num = common::AnfAlgo::GetInputTensorNum(kernel_node);
if (input_num == DynamicInputNum) {
is_dynamic_attr_ = true;
}
input_shape_ = common::AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
kernel_node_ = kernel_node;
null_output_ = CHECK_SHAPE_NULL(input_shape_, kernel_name, "input");
@ -59,11 +64,25 @@ class StridedSliceGpuKernelMod : public NativeGpuKernelMod, public StridedSliceG
MS_LOG(EXCEPTION) << "For '" << kernel_name << "', the dimension of input cannot be greater than " << MAX_DIMS
<< ", but got " << input_shape_.size();
}
CollectInfo(kernel_node);
if (!is_dynamic_attr_) {
GetDynamicAttrIntValue(kernel_node, kBeginIndex_, &begin_);
GetDynamicAttrIntValue(kernel_node, kEndIndex_, &end_);
GetDynamicAttrIntValue(kernel_node, kStrideIndex_, &strides_);
}
CollectInfo(kernel_node, is_dynamic_attr_);
InitSizeLists();
return true;
}
void ResetResource() noexcept override {
ResetSizeLists();
begin_.clear();
end_.clear();
strides_.clear();
input_shape_.clear();
output_shape_.clear();
is_null_input_ = false;
is_dynamic_attr_ = false;
}
protected:
void InitSizeLists() override {
@ -79,6 +98,12 @@ class StridedSliceGpuKernelMod : public NativeGpuKernelMod, public StridedSliceG
}
output_size_list_.push_back(size1);
}
bool is_null_input_{false};
bool is_dynamic_attr_{false};
bool get_dynamic_attr_value_{false};
static constexpr size_t kBeginIndex_{1};
static constexpr size_t kEndIndex_{2};
static constexpr size_t kStrideIndex_{3};
};
} // namespace kernel
} // namespace mindspore

8
mindspore/ccsrc/plugin/device/gpu/kernel/math/assign_add_gpu_kernel.h
View File

@ -69,6 +69,14 @@ class AssignAddFwdGpuKernelMod : public NativeGpuKernelMod {
return true;
}
void ResetResource() noexcept override {
is_null_input_ = false;
input_size_ = 0;
input_size_list_.clear();
output_size_list_.clear();
workspace_size_list_.clear();
}
protected:
void InitSizeLists() override {
input_size_list_.push_back(input_size_);

4
mindspore/ccsrc/plugin/device/gpu/kernel/nn/bias_add_grad_gpu_kenel.h
View File

@ -44,9 +44,7 @@ class BiasAddGradGpuKernelMod : public NativeGpuKernelMod {
cudnn_compute_format_(CUDNN_TENSOR_NCHW),
dy_desc_(nullptr),
db_desc_(nullptr),
op_desc_(nullptr) {
ResetResource();
}
op_desc_(nullptr) {}
~BiasAddGradGpuKernelMod() override { DestroyResource(); }
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,

15
mindspore/ccsrc/plugin/device/gpu/kernel/nn/conv2d_grad_filter_gpu_kernel.cc
View File

@ -26,5 +26,20 @@ MS_REG_GPU_KERNEL_ONE(
Conv2DBackpropFilter,
KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeFloat16).AddOutputAttr(kNumberTypeFloat16),
ConvGradFilterBkwGpuKernelMod, half)
MS_REG_GPU_KERNEL_TWO(Conv2DBackpropFilter,
KernelAttr()
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeFloat32),
ConvGradFilterBkwGpuKernelMod, float, int64_t)
MS_REG_GPU_KERNEL_TWO(Conv2DBackpropFilter,
KernelAttr()
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeFloat16)
.AddInputAttr(kNumberTypeInt64)
.AddOutputAttr(kNumberTypeFloat16),
ConvGradFilterBkwGpuKernelMod, half, int64_t)
} // namespace kernel
} // namespace mindspore

2
mindspore/ccsrc/plugin/device/gpu/kernel/nn/conv2d_grad_filter_gpu_kernel.h
View File

@ -46,7 +46,7 @@ constexpr size_t k2DDilationSize = 4;
constexpr size_t kHeight2DDilationIndex = 2;
constexpr size_t kWidth2DDilationIndex = 3;
template <typename T>
template <typename T, typename S = int64_t>
class ConvGradFilterBkwGpuKernelMod : public NativeGpuKernelMod {
public:
ConvGradFilterBkwGpuKernelMod()

4
mindspore/ccsrc/plugin/device/gpu/kernel/nn/ctcloss_gpu_kernel.h
View File

@ -83,9 +83,7 @@ class CtcLossGpuKernelMod : public NativeGpuKernelMod {
batch_label(0),
label_value_with_blank(nullptr),
log_alpha_b(nullptr),
log_beta_b(nullptr) {
ResetResource();
}
log_beta_b(nullptr) {}
~CtcLossGpuKernelMod() override = default;
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,

11
mindspore/ccsrc/plugin/device/gpu/kernel/nn/dropout_grad_kernel.h
View File

@ -73,6 +73,17 @@ class DropoutGradBwdGpuKernelMod : public NativeGpuKernelMod {
return true;
}
void ResetResource() noexcept override {
cudnn_handle_ = nullptr;
is_null_input_ = false;
kernel_name_ = "DropoutGrad";
num_count_ = 0;
keep_prob_ = 0.0;
input_size_list_.clear();
output_size_list_.clear();
workspace_size_list_.clear();
}
protected:
void InitResource() override { cudnn_handle_ = device::gpu::GPUDeviceManager::GetInstance().GetCudnnHandle(); }
void InitSizeLists() override {

62
mindspore/ccsrc/plugin/device/gpu/kernel/nn/flatten_gpu_kernel.cc
View File

@ -178,6 +178,68 @@ MS_REG_GPU_KERNEL_TWO(
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeBool).AddInputAttr(kNumberTypeInt64).AddOutputAttr(kNumberTypeBool),
FlattenFwdGpuKernelMod, bool, int64_t)
// float64
MS_REG_GPU_KERNEL_TWO(
Reshape,
KernelAttr().AddInputAttr(kNumberTypeFloat64).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeFloat64),
FlattenFwdGpuKernelMod, double, int32_t)
// float32
MS_REG_GPU_KERNEL_TWO(
Reshape,
KernelAttr().AddInputAttr(kNumberTypeFloat32).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeFloat32),
FlattenFwdGpuKernelMod, float, int32_t)
// float16
MS_REG_GPU_KERNEL_TWO(
Reshape,
KernelAttr().AddInputAttr(kNumberTypeFloat16).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeFloat16),
FlattenFwdGpuKernelMod, half, int32_t)
// int64
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeInt64).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeInt64),
FlattenFwdGpuKernelMod, int64_t, int32_t)
// int32
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeInt32).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeInt32),
FlattenFwdGpuKernelMod, int32_t, int32_t)
// int16
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeInt16).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeInt16),
FlattenFwdGpuKernelMod, int16_t, int32_t)
// int8
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeInt8).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeInt8),
FlattenFwdGpuKernelMod, char, int32_t)
// uint64
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeUInt64).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeUInt64),
FlattenFwdGpuKernelMod, uint64_t, int32_t)
// uint32
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeUInt32).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeUInt32),
FlattenFwdGpuKernelMod, uint, int32_t)
// uint16
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeUInt16).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeUInt16),
FlattenFwdGpuKernelMod, uint16_t, int32_t)
// uint8
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeUInt8).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeUInt8),
FlattenFwdGpuKernelMod, uchar, int32_t)
// bool
MS_REG_GPU_KERNEL_TWO(
Reshape, KernelAttr().AddInputAttr(kNumberTypeBool).AddInputAttr(kNumberTypeInt32).AddOutputAttr(kNumberTypeBool),
FlattenFwdGpuKernelMod, bool, int32_t)
} // namespace
} // namespace kernel
} // namespace mindspore

2
mindspore/ccsrc/plugin/device/gpu/kernel/nn/layer_norm_gpu_kernel.h
View File

@ -27,7 +27,7 @@ namespace kernel {
template <typename T>
class LayerNormGpuKernelMod : public NativeGpuKernelMod {
public:
LayerNormGpuKernelMod() : input_row_(1), input_col_(1), param_dim_(1), is_null_input_(false) { ResetResource(); }
LayerNormGpuKernelMod() : input_row_(1), input_col_(1), param_dim_(1), is_null_input_(false) {}
~LayerNormGpuKernelMod() override = default;
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &,

10
mindspore/ccsrc/plugin/device/gpu/kernel/nn/layer_norm_grad_gpu_kernel.h
View File

@ -90,6 +90,16 @@ class LayerNormGradGpuKernelMod : public NativeGpuKernelMod {
return true;
}
void ResetResource() noexcept override {
input_row_ = 1;
input_col_ = 1;
param_dim_ = 1;
is_null_input_ = false;
input_size_list_.clear();
output_size_list_.clear();
workspace_size_list_.clear();
}
protected:
void InitSizeLists() override {
input_size_list_.push_back(input_row_ * input_col_ * sizeof(T));

12
mindspore/ccsrc/plugin/device/gpu/kernel/nn/layer_norm_grad_grad_gpu_kernel.h
View File

@ -109,6 +109,18 @@ class LayerNormGradGradGpuKernelMod : public NativeGpuKernelMod {
return true;
}
void ResetResource() noexcept override {
input_row_ = 1;
input_col_ = 1;
param_dim_ = 1;
input_size_ = 1;
is_null_input_ = false;
epsilon_ = 1e-12;
input_size_list_.clear();
output_size_list_.clear();
workspace_size_list_.clear();
}
protected:
void InitSizeLists() override {
input_size_ = input_row_ * input_col_ * sizeof(T);

4
mindspore/ccsrc/plugin/device/gpu/kernel/nn/softmax_gpu_kernel.h
View File

@ -47,9 +47,7 @@ class SoftmaxGpuKernelMod : public NativeGpuKernelMod {
batch_size_(0),
channel_size_(0),
height_(0),
width_(0) {
ResetResource();
}
width_(0) {}
~SoftmaxGpuKernelMod() override { DestroyResource(); }
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,