huawei
/
mindspore2022
0
0
Fork 53
Code Issues Pull Requests 29 Projects Releases Wiki Activity

mindspore代码评注-等风也等你 #28

Open
darrenlu wants to merge 8 commits from darrenlu/mindspore2022:master into master
pull from: darrenlu/mindspore2022:master
merge into: huawei:master
Conversation 0 Commits 8 Files Changed 16 +446 -98
16 changed files with 446 additions and 98 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
gitup.sh Executable file
View File

@ -0,0 +1,5 @@
#!/bin/bash
git pull origin master
git add .
git commit -m "addd"
git push origin master

30
mindspore/ccsrc/common/graph_kernel/substitute_dropout.cc
View File

@ -37,29 +37,38 @@ inline const PrimitivePtr kPrimGkDropout = std::make_shared<Primitive>("GkDropou
namespace graphkernel {
using opt::CheckCNodeInputSize;
using opt::kDropoutInputTensorNum;
// 静态成员变量，用于保存随机数生成器的种子，初始化为当前时间戳
int64_t DropoutExpander::seed_ = time(nullptr);
// 预处理函数，用于处理Dropout节点之前的操作
AnfNodePtr DropoutExpander::PreProcess(const FuncGraphPtr &func_graph, const AnfNodePtr &node) {
// 检查节点和节点类型是否有效
MS_EXCEPTION_IF_NULL(node);
CNodePtr cnode = node->cast<CNodePtr>();
MS_EXCEPTION_IF_NULL(cnode);
// 检查Dropout节点输入的数量是否正确
CheckCNodeInputSize(cnode, kDropoutInputTensorNum);
// 获取输入节点的设备上的形状
auto shape = AnfAlgo::GetInputDeviceShape(cnode, 0);
ShapeVector shape_i64;
(void)std::transform(shape.begin(), shape.end(), std::back_inserter(shape_i64), SizeToLong);
// Get seed from original dropout's attrs, rather than set seed by time.
// Only seed0 and seed1 are all equal to 0, then set seed = time.
// 从Dropout节点的属性中获取种子值
auto node_prim = GetCNodePrimitive(node);
MS_EXCEPTION_IF_NULL(node_prim);
int64_t seed = GetValue<int64_t>(node_prim->GetAttr("Seed0"));
// 如果Seed0和Seed1都等于0，则使用全局种子值seed_
if (seed == 0) {
seed = GetValue<int64_t>(node_prim->GetAttr("Seed1"));
if (seed == 0) {
seed = seed_++;
}
}
// Create a uniform_real kernel to generate random value.
// 创建一个uniform_real节点，用于生成随机值
auto tensor = std::make_shared<tensor::Tensor>(kNumberTypeInt64, ShapeVector(1, SizeToLong(shape.size())),
static_cast<void *>(&shape[0]), kNumberTypeInt64);
AnfNodePtrList uniform_real_input = {NewValueNode(prim::kPrimCudnnUniformReal), NewValueNode(tensor)};
@ -69,7 +78,8 @@ AnfNodePtr DropoutExpander::PreProcess(const FuncGraphPtr &func_graph, const Anf
SetNodeAttrSafely("seed", MakeValue(seed), uniform_real_node);
common::AnfAlgo::SetNodeAttr("seed2", MakeValue(static_cast<int64_t>(0)), uniform_real_node);
uniform_real_node->set_abstract(std::make_shared<abstract::AbstractTensor>(kFloat32, shape_i64));
// Set kernel_info for uniform_real node
// 为uniform_real节点设置kernel_info
auto uniform_real_kernel_info_builder = std::make_shared<kernel::KernelBuildInfo::KernelBuildInfoBuilder>();
uniform_real_kernel_info_builder->SetInputsFormat({kOpFormat_DEFAULT});
uniform_real_kernel_info_builder->SetInputsDeviceType({kNumberTypeInt32});
@ -79,23 +89,29 @@ AnfNodePtr DropoutExpander::PreProcess(const FuncGraphPtr &func_graph, const Anf
uniform_real_kernel_info_builder->SetProcessor(kernel::Processor::CUDA);
AnfAlgo::SetSelectKernelBuildInfo(uniform_real_kernel_info_builder->Build(), uniform_real_node.get());
// Create a GKDropout node with uniform_real as its second input.
// 创建一个GKDropout节点，其第二个输入为uniform_real节点
AnfNodePtrList gkdropout_inputs = {NewValueNode(prim::kPrimGkDropout), cnode->input(1), uniform_real_node};
auto new_dropout_node = func_graph->NewCNode(gkdropout_inputs);
SetNodeAttrSafely("keep_prob", MakeValue(common::AnfAlgo::GetNodeAttr<float>(cnode, "keep_prob")), new_dropout_node);
// the output info is unchanged.
// 输出信息与原始节点相同
new_dropout_node->set_abstract(node->abstract());
auto old_kernel_info = AnfAlgo::GetSelectKernelBuildInfo(node);
auto dropout_kernel_info_builder = std::make_shared<kernel::KernelBuildInfo::KernelBuildInfoBuilder>(old_kernel_info);
dropout_kernel_info_builder->SetInputsFormat({old_kernel_info->GetInputFormat(0), kOpFormat_DEFAULT});
dropout_kernel_info_builder->SetInputsDeviceType({old_kernel_info->GetInputDeviceType(0), kNumberTypeFloat32});
AnfAlgo::SetSelectKernelBuildInfo(dropout_kernel_info_builder->Build(), new_dropout_node.get());
return new_dropout_node;
}
// 运行函数，用于处理Dropout节点
AnfNodePtr DropoutExpander::Run(const AnfNodePtr &node) {
// 预处理Dropout节点
auto gkdropout_node = PreProcess(node->func_graph(), node);
// 调用父类的Run方法处理GKDropout节点
return PyExpander::Run(gkdropout_node);
}
} // namespace graphkernel
} // namespace mindspore

10
mindspore/ccsrc/common/graph_kernel/substitute_dropout.h
View File

@ -19,15 +19,25 @@
#include "common/graph_kernel/adapter/graph_kernel_expander_with_py.h"
namespace mindspore::graphkernel {
// 定义一个名为DropoutExpander的类，它继承自PyExpander类
class DropoutExpander : public PyExpander {
public:
// 默认构造函数
DropoutExpander() = default;
// 虚析构函数
virtual ~DropoutExpander() = default;
// 重写PyExpander类的Run方法，接受一个AnfNodePtr类型的参数node
AnfNodePtr Run(const AnfNodePtr &node) override;
private:
// 私有成员函数，用于在处理节点之前进行预处理
AnfNodePtr PreProcess(const FuncGraphPtr &, const AnfNodePtr &);
// 静态成员变量，用于保存随机数生成器的种子
static int64_t seed_;
};
} // namespace mindspore::graphkernel
#endif // MINDSPORE_CCSRC_BACKEND_OPTIMIZER_GRAPH_KERNEL_SUBSTITUTE_DROPOUT_H_

23
mindspore/ccsrc/plugin/device/cpu/kernel/mkldnn/batch_norm_cpu_kernel.h
View File

@ -21,19 +21,29 @@
namespace mindspore {
namespace kernel {
// 定义一个名为BatchNormCpuKernelMod的类，它继承自MKLCpuKernelMod类
class BatchNormCpuKernelMod : public MKLCpuKernelMod {
public:
// 默认构造函数
BatchNormCpuKernelMod() = default;
// 虚析构函数
~BatchNormCpuKernelMod() override = default;
// 初始化函数，用于解析CNode节点并设置内核参数
void InitKernel(const CNodePtr &kernel_node) override;
// 启动CPU核函数，执行Batch Normalization 操作
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
const std::vector<AddressPtr> &outputs) override;
protected:
// 初始化输入和输出数据大小
void InitInputOutputSize(const CNodePtr &kernel_node) override;
// 获取操作支持的数据类型
std::vector<KernelAttr> GetOpSupport() override {
// 定义支持的输入和输出数据类型列表
static std::vector<KernelAttr> support_list = {KernelAttr()
.AddInputAttr(kNumberTypeFloat32)
.AddInputAttr(kNumberTypeFloat32)
@ -49,13 +59,14 @@ class BatchNormCpuKernelMod : public MKLCpuKernelMod {
}
private:
bool is_train{false};
float momentum{0.9};
size_t batch_size{0};
size_t channel{0};
size_t hw_size{0};
size_t nhw_size{0};
bool is_train{false}; // 是否处于训练模式
float momentum{0.9}; // 动量参数
size_t batch_size{0}; // 批量大小
size_t channel{0}; // 通道数
size_t hw_size{0}; // 单个样本的特征维度大小
size_t nhw_size{0}; // 批次中所有样本的特征维度总大小
};
} // namespace kernel
} // namespace mindspore

32
mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/detection_post_process_fp32.c
View File

@ -166,15 +166,17 @@ int DetectionPostProcessFast(const int num_boxes, const int num_classes_with_bg,
}
return NNACL_OK;
}
int DetectionPostProcessRegular(const int num_boxes, const int num_classes_with_bg, const float *input_scores,
float *output_boxes, float *output_classes, float *output_scores, float *output_num,
void (*PartialArgSort)(const float *, int *, int, int),
const DetectionPostProcessParameter *param) {
// 检查输入参数是否为NULL
if (input_scores == NULL || output_boxes == NULL || output_classes == NULL || output_scores == NULL ||
output_num == NULL || param == NULL || PartialArgSort == NULL) {
return NNACL_NULL_PTR;
}
// 提取输入参数
const int first_class_index = num_classes_with_bg - (int)(param->num_classes_);
float *decoded_boxes = (float *)param->decoded_boxes_;
int *selected = (int *)param->selected_;
@ -184,35 +186,54 @@ int DetectionPostProcessRegular(const int num_boxes, const int num_classes_with_
int *all_indexes = (int *)(param->all_class_indexes_);
int all_classes_sorted_num = 0;
int all_classes_output_num = 0;
// 循环处理每个类别
for (int j = first_class_index; j < num_classes_with_bg; ++j) {
// process single class
// 处理单个类别的分数
for (int i = 0; i < num_boxes; ++i) {
scores[i] = input_scores[i * num_classes_with_bg + j];
}
// 执行非极大值抑制 (NMS) 操作，选择一些边界框
int selected_num =
NmsSingleClass(num_boxes, decoded_boxes, param->detections_per_class_, scores, selected, PartialArgSort, param);
// 复制已排序的索引
for (int i = 0; i < all_classes_sorted_num; ++i) {
indexes[i] = all_indexes[i];
all_indexes[i] = i;
}
// process all classes
// 处理所有类别
for (int i = 0; i < selected_num; ++i) {
indexes[all_classes_sorted_num] = selected[i] * num_classes_with_bg + j;
all_indexes[all_classes_sorted_num] = all_classes_sorted_num;
all_scores[all_classes_sorted_num++] = scores[selected[i]];
}
// 限制输出数量为param->max_detections_
all_classes_output_num =
all_classes_sorted_num < param->max_detections_ ? all_classes_sorted_num : param->max_detections_;
// 对所有分数进行排序
PartialArgSort(all_scores, all_indexes, all_classes_output_num, all_classes_sorted_num);
// 更新输出分数和索引
for (int i = 0; i < all_classes_output_num; ++i) {
scores[i] = all_scores[all_indexes[i]];
all_indexes[i] = indexes[all_indexes[i]];
}
// 更新所有分数
for (int i = 0; i < all_classes_output_num; ++i) {
all_scores[i] = scores[i];
}
// 更新已排序的类别数量
all_classes_sorted_num = all_classes_output_num;
}
// 将结果写入输出参数
for (int i = 0; i < param->max_detections_ * param->max_classes_per_detection_; ++i) {
if (i < all_classes_output_num) {
NNACL_CHECK_ZERO_RETURN_ERR(num_classes_with_bg);
@ -222,6 +243,7 @@ int DetectionPostProcessRegular(const int num_boxes, const int num_classes_with_
output_classes[i] = (float)class_index;
output_scores[i] = all_scores[i];
} else {
// 如果没有更多的结果，将输出设置为0
((BboxCorner *)(output_boxes) + i)->ymin = 0;
((BboxCorner *)(output_boxes) + i)->xmin = 0;
((BboxCorner *)(output_boxes) + i)->ymax = 0;
@ -230,6 +252,10 @@ int DetectionPostProcessRegular(const int num_boxes, const int num_classes_with_
output_scores[i] = 0.0f;
}
}
// 更新输出对象的数量
*output_num = (float)all_classes_output_num;
// 返回操作成功的标志
return NNACL_OK;
}

72
mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/pooling_fp32.c
View File

@ -45,20 +45,34 @@
MS_ST_F32(block_size, dst_c_ptr, tmp_avg); \
} \
} while (0)
// 函数名：AvgPoolingBatch
// 参数说明：
// - src_b_ptr: 输入数据的指针，指向当前批次的源数据
// - dst_b_ptr: 输出数据的指针，指向当前批次的目标数据
// - pooling_param: 池化操作的参数结构体，包含输入、输出尺寸等信息
// - task_id: 当前任务的线程 ID，用于并行处理
// - minf: 数据的最小值
// - maxf: 数据的最大值
// 返回值：操作是否成功的标志，NNACL_OK 表示成功
int AvgPoolingBatch(const float *src_b_ptr, float *dst_b_ptr, const PoolingParameter *pooling_param, int task_id,
float minf, float maxf) {
// 从参数结构体中提取输入和输出尺寸信息
int win_w = pooling_param->window_w_, win_h = pooling_param->window_h_;
int in_w = pooling_param->input_w_, in_h = pooling_param->input_h_;
int output_w = pooling_param->output_w_, output_h = pooling_param->output_h_;
int channel = pooling_param->input_channel_;
int out_plane = output_w * output_h;
int out_tile_count = UP_DIV(out_plane, TILE_NUM);
// 检查输出宽度是否为零，如果是，则返回错误
NNACL_CHECK_ZERO_RETURN_ERR(output_w);
// 循环遍历输出矩阵中的每个元素块（称为瓦片）
for (int thread_id = task_id; thread_id < out_tile_count; thread_id += pooling_param->thread_num_) {
int cal_start_index = thread_id * TILE_NUM;
int real_cal_num = (out_plane - cal_start_index) > TILE_NUM ? TILE_NUM : (out_plane - cal_start_index);
// 遍历当前瓦片中的每个元素
for (int i = 0; i < real_cal_num; i++) {
int index = cal_start_index + i;
int out_w_index = index % output_w;
@ -66,24 +80,30 @@ int AvgPoolingBatch(const float *src_b_ptr, float *dst_b_ptr, const PoolingParam
int in_w_index = out_w_index * pooling_param->stride_w_ - pooling_param->pad_l_;
int in_h_index = out_h_index * pooling_param->stride_h_ - pooling_param->pad_u_;
// 获取源数据和目标数据的指针
const float *src_plane_ptr = src_b_ptr;
float *dst_plane_ptr = dst_b_ptr + index * channel;
// 计算有效的窗口范围，确保不超出输入数据的边界
int real_win_h_start = MSMAX(0, -in_h_index);
int real_win_h_end = MSMIN(win_h, in_h - in_h_index);
int real_win_w_start = MSMAX(0, -in_w_index);
int real_win_w_end = MSMIN(win_w, in_w - in_w_index);
int ci = 0;
// 使用 SIMD 指令运行平均池化操作的核心计算
MS_SIMD_RUN_NO_SCALAR(SimdFp32AvgPoolingBatchCoreCalc, src_plane_ptr, channel, dst_plane_ptr, ci,
real_win_h_start, real_win_h_end, real_win_w_start, real_win_w_end, in_h_index, in_w,
in_w_index);
// 对于每个通道，执行平均池化操作
for (; ci < channel; ci++) {
const float *src_c_ptr = src_plane_ptr + ci;
float *dst_c_ptr = dst_plane_ptr + ci;
float tmp_avg = 0;
int real_count = 0;
// 在窗口范围内计算平均值
for (int h = real_win_h_start; h < real_win_h_end; h++) {
for (int w = real_win_w_start; w < real_win_w_end; w++) {
const float *src_win_ptr = src_c_ptr + ((in_h_index + h) * in_w + in_w_index + w) * channel;
@ -91,6 +111,8 @@ int AvgPoolingBatch(const float *src_b_ptr, float *dst_b_ptr, const PoolingParam
++real_count;
} // win_w loop
} // win_h loop
// 确保 real_count 不为零，然后计算平均值并将其限制在 [minf, maxf] 范围内
MS_CHECK_TRUE_RET(real_count != 0, NNACL_ERR);
tmp_avg = tmp_avg / (float)real_count;
tmp_avg = fmaxf(tmp_avg, minf);
@ -99,11 +121,22 @@ int AvgPoolingBatch(const float *src_b_ptr, float *dst_b_ptr, const PoolingParam
} // channel_res loop
} // real_cal_num loop
} // out_plane loop
// 返回操作成功标志
return NNACL_OK;
}
// 函数名：AvgPooling
// 参数说明：
// - input_ptr: 输入数据的指针，指向整个输入批次的数据
// - output_ptr: 输出数据的指针，指向整个输出批次的数据
// - pooling_param: 池化操作的参数结构体，包含输入、输出尺寸等信息
// - task_id: 当前任务的线程 ID，用于并行处理
// - minf: 数据的最小值
// - maxf: 数据的最大值
// 返回值：操作是否成功的标志，NNACL_OK 表示成功
int AvgPooling(const float *input_ptr, float *output_ptr, const PoolingParameter *pooling_param, int task_id,
float minf, float maxf) {
// 从参数结构体中提取输入和输出尺寸信息以及通道数、批次大小等信息
int in_w = pooling_param->input_w_;
int in_h = pooling_param->input_h_;
int output_w = pooling_param->output_w_;
@ -111,17 +144,26 @@ int AvgPooling(const float *input_ptr, float *output_ptr, const PoolingParameter
int channel = pooling_param->input_channel_;
int output_batch = pooling_param->output_batch_;
// 循环遍历每个批次的数据
for (int batch = 0; batch < output_batch; batch++) {
// 获取当前批次的输入数据和输出数据的指针
const float *src_b_ptr = input_ptr + batch * in_h * in_w * channel;
float *dst_b_ptr = output_ptr + batch * output_h * output_w * channel;
// 调用 AvgPoolingBatch 函数处理当前批次的数据
int ret = AvgPoolingBatch(src_b_ptr, dst_b_ptr, pooling_param, task_id, minf, maxf);
// 检查处理是否成功，如果失败则返回错误码
if (ret != NNACL_OK) {
return ret;
}
}
// 返回操作成功标志
return NNACL_OK;
}
int AvgPoolingFromNC4HW4ToNHWCLessC(const float *src_b_ptr, float *dst_b_ptr, const PoolingParameter *pooling_param,
int task_id, float minf, float maxf) {
int in_w = pooling_param->input_w_, in_h = pooling_param->input_h_;
@ -382,19 +424,34 @@ int AvgPoolingFromNC4HW4ToNHWC(const float *input_ptr, float *output_ptr, const
} \
} while (0)
// 函数名：MaxPoolingBatch
// 参数说明：
// - src_b_ptr: 输入数据的指针，指向当前批次的源数据
// - dst_b_ptr: 输出数据的指针，指向当前批次的目标数据
// - pooling_param: 池化操作的参数结构体，包含输入、输出尺寸等信息
// - task_id: 当前任务的线程 ID，用于并行处理
// - minf: 数据的最小值
// - maxf: 数据的最大值
// 返回值：操作是否成功的标志，NNACL_OK 表示成功
int MaxPoolingBatch(const float *src_b_ptr, float *dst_b_ptr, const PoolingParameter *pooling_param, int task_id,
float minf, float maxf) {
// 从参数结构体中提取输入和输出尺寸信息
int in_w = pooling_param->input_w_, in_h = pooling_param->input_h_;
int win_w = pooling_param->window_w_, win_h = pooling_param->window_h_;
int output_w = pooling_param->output_w_, output_h = pooling_param->output_h_;
int channel = pooling_param->input_channel_;
int out_plane = output_w * output_h;
int out_tile_count = UP_DIV(out_plane, TILE_NUM);
// 检查输出宽度是否为零，如果是，则返回错误
NNACL_CHECK_ZERO_RETURN_ERR(output_w);
// 循环遍历输出矩阵中的每个元素块（称为瓦片）
for (int thread_id = task_id; thread_id < out_tile_count; thread_id += pooling_param->thread_num_) {
int cal_start_index = thread_id * TILE_NUM;
int real_cal_num = (out_plane - cal_start_index) > TILE_NUM ? TILE_NUM : (out_plane - cal_start_index);
// 遍历当前瓦片中的每个元素
for (int i = 0; i < real_cal_num; i++) {
int index = cal_start_index + i;
int out_w_index = index % output_w;
@ -402,38 +459,49 @@ int MaxPoolingBatch(const float *src_b_ptr, float *dst_b_ptr, const PoolingParam
int in_w_index = out_w_index * pooling_param->stride_w_ - pooling_param->pad_l_;
int in_h_index = out_h_index * pooling_param->stride_h_ - pooling_param->pad_u_;
// 获取源数据和目标数据的指针
const float *src_plane_ptr = src_b_ptr;
float *dst_plane_ptr = dst_b_ptr + index * channel;
// 计算有效的窗口范围，确保不超出输入数据的边界
int real_win_h_start = MSMAX(0, -in_h_index);
int real_win_h_end = MSMIN(win_h, in_h - in_h_index);
int real_win_w_start = MSMAX(0, -in_w_index);
int real_win_w_end = MSMIN(win_w, in_w - in_w_index);
int ci = 0;
// 使用 SIMD 指令运行最大池化操作的核心计算
MS_SIMD_RUN_NO_SCALAR(SimdFp32MaxPoolingBatchCoreCalc, src_plane_ptr, channel, dst_plane_ptr, ci,
real_win_h_start, real_win_h_end, real_win_w_start, real_win_w_end, in_h_index, in_w,
in_w_index);
// 对于每个通道，执行最大池化操作
for (; ci < channel; ci++) {
float *dst_c_ptr = dst_plane_ptr + ci;
const float *src_c_ptr = src_plane_ptr + ci;
float tmp_max = -FLT_MAX;
// 在窗口范围内查找最大值
for (int kh = real_win_h_start; kh < real_win_h_end; kh++) {
for (int kw = real_win_w_start; kw < real_win_w_end; kw++) {
const float *src_win_ptr = src_c_ptr + ((in_h_index + kh) * in_w + in_w_index + kw) * channel;
tmp_max = fmaxf(tmp_max, src_win_ptr[0]);
} // win_w loop
} // win_h loop
// 限制最大值的范围在 [minf, maxf] 内，并将结果存储在目标数据中
tmp_max = fmaxf(tmp_max, minf);
tmp_max = fminf(tmp_max, maxf);
dst_c_ptr[0] = tmp_max;
} // channel_res loop
} // real_cal_num loop
} // out_plane loop
// 返回操作成功标志
return NNACL_OK;
}
int MaxPooling(const float *input_ptr, float *output_ptr, const PoolingParameter *pooling_param, int task_id,
float minf, float maxf) {
int in_w = pooling_param->input_w_;

5
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu
View File

@ -28,20 +28,25 @@ __device__ __forceinline__ half SqrtFunc(half input) {
}
template <typename T, typename S, typename G>
// 定义一个 CUDA 核函数，用于执行 Adagrad 算法的更新操作
__global__ void ApplyAdagradKernel(const size_t size,
const bool update_slots,
const S *learning_rate,
const G *gradient,
T *variable,
T *accumulation) {
// 使用线程索引迭代处理数据，每个线程处理一个元素
for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < size; i += gridDim.x * blockDim.x) {
// 如果 update_slots 为真，则更新梯度积累
if (update_slots) {
accumulation[i] += gradient[i] * gradient[i];
}
// 计算变量的更新值，根据 Adagrad 公式进行更新
variable[i] -= learning_rate[0] * gradient[i] / SqrtFunc(accumulation[i]);
}
}
template <>
__global__ void ApplyAdagradKernel(const size_t size,
const bool update_slots,

10
mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cu
View File

@ -17,16 +17,24 @@
#include <iostream>
#include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cuh"
// 定义一个 CUDA 核函数，用于执行 RmsProp 算法的参数更新
template <typename T>
__global__ void RmsPropKernel(const T* learning_rate, const T decay, const T momentum, const T epsilon, T* variable,
T* mean_square, T*moment, T* gradients, const size_t size) {
T* mean_square, T* moment, T* gradients, const size_t size) {
// 使用线程索引迭代处理数据，每个线程处理一个元素
for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < (size); i += blockDim.x * gridDim.x) {
// 计算均方(mean_square)梯度的加权平均，用于自适应调整学习率
mean_square[i] = decay * mean_square[i] + (1.0 - decay) * gradients[i] * gradients[i];
// 计算动量(moment)项，用于加速收敛，同时应用学习率和均方根逆的调整
moment[i] = momentum * moment[i] + learning_rate[0] * rsqrt(mean_square[i] + epsilon) * gradients[i];
// 更新模型参数 variable，减去动量项
variable[i] -= moment[i];
}
}
template <typename T>
void RmsProp(const T* learning_rate, const T decay, const T momentum, const T epsilon,
T* variable, T* mean_square, T* moment, T* gradients, const size_t size, cudaStream_t cuda_stream) {

26
mindspore/ccsrc/plugin/device/gpu/kernel/nn/adagrad_gpu_kernel.h
View File

@ -26,8 +26,10 @@
namespace mindspore {
namespace kernel {
template <typename T, typename S, typename G>
// 定义一个名为AdagradGpuKernelMod的类，它继承自DeprecatedNativeGpuKernelMod类
class AdagradGpuKernelMod : public DeprecatedNativeGpuKernelMod {
public:
// 默认构造函数，初始化成员变量
AdagradGpuKernelMod()
: variable_size_(0),
accumulation_size_(0),
@ -37,22 +39,30 @@ class AdagradGpuKernelMod : public DeprecatedNativeGpuKernelMod {
is_null_input_(false),
kernel_name_("ApplyAdagrad") {}
// 虚析构函数
~AdagradGpuKernelMod() override = default;
// 启动GPU核函数，执行Adagrad算法的参数更新
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &,
const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
// 如果输入为空，直接返回
if (is_null_input_) {
return true;
}
// 获取输入和输出地址
T *variable = GetDeviceAddress<T>(inputs, 0);
T *accumulation = GetDeviceAddress<T>(inputs, 1);
S *learning_rate = GetDeviceAddress<S>(inputs, 2);
G *gradient = GetDeviceAddress<G>(inputs, 3);
T *variable_out = GetDeviceAddress<T>(outputs, 0);
T *accumulation_out = GetDeviceAddress<T>(outputs, 1);
// 调用ApplyAdagrad函数来执行参数更新
ApplyAdagrad(inputs[0]->size / sizeof(T), update_slots, learning_rate, gradient, variable, accumulation,
reinterpret_cast<cudaStream_t>(stream_ptr));
// 使用CUDA内存复制函数将更新后的数据拷贝回输出地址
CHECK_CUDA_RET_WITH_EXCEPT(kernel_node_,
cudaMemcpyAsync(variable_out, variable, variable_size_, cudaMemcpyDeviceToDevice,
reinterpret_cast<cudaStream_t>(stream_ptr)),
@ -65,29 +75,43 @@ class AdagradGpuKernelMod : public DeprecatedNativeGpuKernelMod {
return true;
}
// 初始化函数，用于解析CNode节点和初始化成员变量
bool Init(const CNodePtr &kernel_node) override {
// 获取CNode节点的名称
kernel_name_ = common::AnfAlgo::GetCNodeName(kernel_node);
// 获取输入数据的数量和是否更新slots的标志
size_t input_num = common::AnfAlgo::GetInputTensorNum(kernel_node);
update_slots = common::AnfAlgo::GetNodeAttr<bool>(kernel_node, "update_slots");
kernel_node_ = kernel_node;
// 检查输入数据数量是否为4
if (input_num != 4) {
MS_LOG(EXCEPTION) << "For '" << kernel_name_ << "', the number of inputs should be 4, but got " << input_num;
}
// 初始化各种数据的大小
variable_size_ = sizeof(T);
accumulation_size_ = sizeof(T);
learning_rate_size_ = sizeof(S);
gradient_size_ = sizeof(G);
// 获取输入数据的形状
auto variable_shape = common::AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 0);
auto accumulation_shape = common::AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 1);
auto gradient_shape = common::AnfAlgo::GetPrevNodeOutputInferShape(kernel_node, 3);
// 检查输入数据是否为空
is_null_input_ = CHECK_SHAPE_NULL(variable_shape, kernel_name_, "var") ||
CHECK_SHAPE_NULL(accumulation_shape, kernel_name_, "accum") ||
CHECK_SHAPE_NULL(gradient_shape, kernel_name_, "grad");
if (is_null_input_) {
InitSizeLists();
return true;
}
// 计算各输入数据的总大小
for (size_t i = 0; i < variable_shape.size(); i++) {
variable_size_ *= variable_shape[i];
}
@ -104,6 +128,7 @@ class AdagradGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
protected:
// 初始化输入和输出数据大小列表
void InitSizeLists() override {
input_size_list_.push_back(variable_size_);
input_size_list_.push_back(accumulation_size_);
@ -122,6 +147,7 @@ class AdagradGpuKernelMod : public DeprecatedNativeGpuKernelMod {
bool is_null_input_;
std::string kernel_name_;
};
} // namespace kernel
} // namespace mindspore

30
mindspore/ccsrc/plugin/device/gpu/kernel/nn/batch_norm_gpu_kernel.h
View File

@ -29,18 +29,26 @@ namespace kernel {
constexpr size_t CUDNN_BATCHNORM_OPS_BN_ADD_ACTIVATION_INPUT_NUM = 6;
constexpr size_t NO_CUDNN_BATCHNORM_OPS_BN_ADD_ACTIVATION_INPUT_NUM = 5;
template <typename T>
// 定义一个名为BatchNormGpuKernelMod的类，它继承自DeprecatedNativeGpuKernelMod类
class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
public:
// 构造函数，在构造对象时重置资源
BatchNormGpuKernelMod() { ResetResource(); }
// 析构函数，在销毁对象时销毁资源
~BatchNormGpuKernelMod() override { DestroyResource(); }
// 启动GPU核函数，执行批量归一化（Batch Normalization）操作
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &workspace,
const std::vector<AddressPtr> &outputs, void *stream_ptr) override {
// 未使用的变量
VARIABLE_NOT_USED(workspace);
VARIABLE_NOT_USED(stream_ptr);
if (is_null_input_) {
return true;
}
// 获取输入数据地址
auto x = GetDeviceAddress<T>(inputs, 0);
auto scale = GetDeviceAddress<float>(inputs, 1);
auto bias = GetDeviceAddress<float>(inputs, 2);
@ -51,15 +59,19 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
z = GetPossiblyNullDeviceAddress<T>(inputs, 5);
}
// 获取输出数据地址
auto y = GetDeviceAddress<T>(outputs, 0);
T *workspace_addr = GetPossiblyNullDeviceAddress<T>(workspace, 0);
const float alpha = 1;
const float beta = 0;
if (is_train_) {
// 训练模式下执行批量归一化
auto reserve_addr = GetPossiblyNullDeviceAddress<float>(outputs, 2);
auto save_mean = GetDeviceAddress<float>(outputs, 3);
auto save_variance = GetDeviceAddress<float>(outputs, 4);
// 调用cuDNN函数执行批量归一化
CHECK_CUDNN_RET_WITH_EXCEPT(
kernel_node_,
cudnnBatchNormalizationForwardTrainingEx(
@ -68,6 +80,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
workspace_addr, workspace_size_, reserve_addr, reserve_size_),
"Kernel launch failed");
} else {
// 推理模式下执行批量归一化
CHECK_CUDNN_RET_WITH_EXCEPT(kernel_node_,
cudnnBatchNormalizationForwardInference(
handle_, mode_, &alpha, &beta, x_desc_, x, y_desc_, y, scale_bias_mean_var_desc_,
@ -77,10 +90,13 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
return true;
}
// 初始化函数，用于解析CNode节点并设置内核参数
bool Init(const CNodePtr &kernel_node) override {
kernel_node_ = kernel_node;
MS_EXCEPTION_IF_NULL(kernel_node);
std::string kernel_name = common::AnfAlgo::GetCNodeName(kernel_node);
// 根据节点名称确定批量归一化操作类型
if (kernel_name == kBatchNorm) {
bn_ops_ = CUDNN_BATCHNORM_OPS_BN;
} else if (kernel_name == kBatchNormWithActivation) {
@ -92,12 +108,18 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
<< kBatchNormWithAddAndActivation << ", but got " << kernel_name;
}
// 初始化资源
InitResource();
// 获取节点的训练模式、epsilon和momentum属性
is_train_ = GetAttr<bool>(kernel_node, "is_training");
epsilon_ = GetAttr<float>(kernel_node, "epsilon");
exp_avg_factor_ = GetAttr<float>(kernel_node, "momentum");
// 获取cuDNN数据类型
cudnn_data_type_ = GetCudnnDataType(TypeIdLabel(AnfAlgo::GetInputDeviceDataType(kernel_node, 0)));
// 检查输入的数量和形状
size_t input_num = common::AnfAlgo::GetInputTensorNum(kernel_node);
if (bn_ops_ == CUDNN_BATCHNORM_OPS_BN_ADD_ACTIVATION) {
if (input_num != CUDNN_BATCHNORM_OPS_BN_ADD_ACTIVATION_INPUT_NUM) {
@ -111,6 +133,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
}
// 获取输入的形状和格式
auto shape = AnfAlgo::GetInputDeviceShape(kernel_node, 0);
if (shape.size() != 4 && shape.size() != 2) {
MS_LOG(EXCEPTION) << "For '" << kernel_name << "', the dimension of input should be 2 or 4, but got "
@ -141,6 +164,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
return true;
}
// 重置资源
void ResetResource() noexcept override {
input_x_size_ = 0;
input_z_size_ = 0;
@ -166,6 +190,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
workspace_size_list_.clear();
}
// 销毁资源
void DestroyResource() noexcept override {
CHECK_CUDNN_RET_WITH_ERROR(kernel_node_, cudnnDestroyTensorDescriptor(x_desc_), "Destroy x desc failed");
CHECK_CUDNN_RET_WITH_ERROR(kernel_node_, cudnnDestroyTensorDescriptor(y_desc_), "Destroy y desc failed");
@ -182,6 +207,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
protected:
// 初始化资源
void InitResource() override {
handle_ = device::gpu::GPUDeviceManager::GetInstance().GetCudnnHandle();
CHECK_CUDNN_RET_WITH_EXCEPT(kernel_node_, cudnnCreateTensorDescriptor(&x_desc_), "Create x desc failed");
@ -198,6 +224,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
}
// 初始化输入输出大小列表
void InitSizeLists() override {
if (!is_null_input_) {
CHECK_CUDNN_RET_WITH_EXCEPT(kernel_node_, cudnnGetTensorSizeInBytes(x_desc_, &input_x_size_),
@ -242,6 +269,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
private:
// 设置cuDNN张量描述符
void SetTensorDescriptor(const std::string &format, const std::vector<size_t> &shape) {
cudnnTensorFormat_t cudnn_format;
int batch, channel, height, width;
@ -292,6 +320,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
}
// 成员变量声明
size_t input_x_size_;
size_t input_z_size_;
size_t para_size_;
@ -313,6 +342,7 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod {
cudnnHandle_t handle_;
cudnnDataType_t cudnn_data_type_;
};
} // namespace kernel
} // namespace mindspore
#endif // MINDSPORE_CCSRC_BACKEND_KERNEL_COMPILER_GPU_NN_FUSED_BATCH_NORM_EX_GPU_KERNEL_H_

49
mindspore/ccsrc/plugin/device/gpu/kernel/nn/rmsprop_gpu_kernel.h
View File

@ -25,27 +25,38 @@
namespace mindspore {
namespace kernel {
template <typename T>
// 定义一个名为RMSPropGpuKernelMod的类，它继承自DeprecatedNativeGpuKernelMod类
class RMSPropGpuKernelMod : public DeprecatedNativeGpuKernelMod {
public:
// 默认构造函数，初始化成员变量
RMSPropGpuKernelMod()
: size_(1), use_center_(false), is_null_input_(false), decay_(0.0), momentum_(0.9), epsilon_(1e-12) {}
// 虚析构函数
~RMSPropGpuKernelMod() override = default;
// 启动GPU核函数，执行RMSProp或Centered RMSProp算法的更新操作
bool Launch(const std::vector<AddressPtr> &inputs, const std::vector<AddressPtr> &,
const std::vector<AddressPtr> &outputs, void *stream) override {
// 如果输入为空，直接返回
if (is_null_input_) {
return true;
}
// 根据是否使用Centered RMSProp选择执行相应的操作
if (!use_center_) {
// 获取输入地址
T *variable = GetDeviceAddress<T>(inputs, 0);
T *mean_square = GetDeviceAddress<T>(inputs, 1);
T *moment = GetDeviceAddress<T>(inputs, 2);
T *learning_rate = GetDeviceAddress<T>(inputs, 3);
T *gradients = GetDeviceAddress<T>(inputs, 4);
// 调用RMSProp函数来执行参数更新
RmsProp(learning_rate, decay_, momentum_, epsilon_, variable, mean_square, moment, gradients, size_,
reinterpret_cast<cudaStream_t>(stream));
} else {
// 获取输入地址
T *variable = GetDeviceAddress<T>(inputs, 0);
T *mean_gradients = GetDeviceAddress<T>(inputs, 1);
T *mean_square = GetDeviceAddress<T>(inputs, 2);
@ -56,37 +67,50 @@ class RMSPropGpuKernelMod : public DeprecatedNativeGpuKernelMod {
T *momentum = GetDeviceAddress<T>(inputs, 7);
T *epsilon = GetDeviceAddress<T>(inputs, 8);
// 调用RMSPropCenter函数来执行参数更新
RmsPropCenter(learning_rate, decay, momentum, epsilon, variable, mean_gradients, mean_square, moment, gradients,
size_, reinterpret_cast<cudaStream_t>(stream));
}
return true;
}
// 初始化函数，用于解析CNode节点和初始化成员变量
bool Init(const CNodePtr &kernel_node) override {
auto node_name = common::AnfAlgo::GetCNodeName(kernel_node);
kernel_node_ = kernel_node;
// 根据节点名称判断是否使用Centered RMSProp
if (node_name == "ApplyCenteredRMSProp") {
use_center_ = true;
}
// 如果使用RMSProp，从CNode节点获取rho、momentum和epsilon属性值
if (node_name == "ApplyRMSProp") {
decay_ = GetAttr<float>(kernel_node, "rho");
momentum_ = GetAttr<float>(kernel_node, "momentum");
epsilon_ = GetAttr<float>(kernel_node, "epsilon");
}
// 获取输入数据的形状
auto input_shape = common::AnfAlgo::GetOutputInferShape(kernel_node, 0);
// 检查输入数据是否为空
is_null_input_ = CHECK_SHAPE_NULL(input_shape, node_name, "var");
if (is_null_input_) {
InitSizeLists();
return true;
}
for (auto &dim : input_shape) {
size_ *= dim;
// 计算输入数据的总大小
if (!is_null_input_) {
for (auto &dim : input_shape) {
size_ *= dim;
}
}
// 初始化输入和输出数据大小列表
InitSizeLists();
return true;
}
protected:
// 初始化输入和输出数据大小列表
void InitSizeLists() override {
size_t input_size = size_ * sizeof(T);
if (!use_center_) {
@ -111,13 +135,14 @@ class RMSPropGpuKernelMod : public DeprecatedNativeGpuKernelMod {
}
private:
size_t size_;
bool use_center_;
bool is_null_input_;
float decay_;
float momentum_;
float epsilon_;
size_t size_; // 数据总大小
bool use_center_; // 是否使用Centered RMSProp
bool is_null_input_; // 输入数据是否为空
float decay_; // RMSProp中的rho参数
float momentum_; // RMSProp中的momentum参数
float epsilon_; // RMSProp中的epsilon参数
};
} // namespace kernel
} // namespace mindspore

48
mindspore/core/ops/max_pool.h
View File

@ -28,54 +28,46 @@
namespace mindspore {
namespace ops {
// 声明一个常量字符串 kNameMaxPool 来存储操作的名称 "MaxPool"
constexpr auto kNameMaxPool = "MaxPool";
/// \brief Max pooling operation. Refer to Python API @ref mindspore.ops.MaxPool for more details.
// 定义 MaxPool 类，继承自 BaseOperator 类
// 该类表示最大池化操作
class MIND_API MaxPool : public BaseOperator {
public:
MIND_API_BASE_MEMBER(MaxPool);
/// \brief Constructor.
// 构造函数，用于创建 MaxPool 类的实例
MaxPool() : BaseOperator(kNameMaxPool) { InitIOName({"x"}, {"output"}); }
explicit MaxPool(const std::string k_name) : BaseOperator(k_name) { InitIOName({"x"}, {"output"}); }
/// \brief Init. Refer to the parameters of Python API @ref mindspore.ops.MaxPool for the inputs.
// 初始化函数，设置最大池化操作的参数
void Init(const std::vector<int64_t> &kernel_size = {1}, const std::vector<int64_t> &stride = {1},
const PadMode &pad_mode = VALID, const Format &format = NCHW,
const std::vector<int64_t> &pad = {0, 0, 0, 0}, const RoundMode &round_mode = FLOOR);
/// \brief Set pad_mode.
// 设置 pad_mode 参数
void set_pad_mode(const PadMode &pad_mode);
/// \brief Set kernel_size.
// 设置 kernel_size 参数
void set_kernel_size(const std::vector<int64_t> &kernel_size);
/// \brief Set strides.
// 设置 strides 参数
void set_strides(const std::vector<int64_t> &strides);
/// \brief Set format.
// 设置 format 参数
void set_format(const Format &format);
/// \brief Set pad.
// 设置 pad 参数
void set_pad(const std::vector<int64_t> &pad);
/// \brief Set round_mode.
// 设置 round_mode 参数
void set_round_mode(const RoundMode &round_mode);
/// \brief Get kernel_size.
///
/// \return kernel_size.
// 获取 kernel_size 参数
std::vector<int64_t> get_kernel_size() const;
/// \brief Get strides.
///
/// \return strides.
// 获取 strides 参数
std::vector<int64_t> get_strides() const;
/// \brief Get pad_mode.
///
/// \return pad_mode.
// 获取 pad_mode 参数
PadMode get_pad_mode() const;
/// \brief Get format.
///
/// \return format.
// 获取 format 参数
Format get_format() const;
/// \brief Get pad.
///
/// \return pad.
// 获取 pad 参数
std::vector<int64_t> get_pad() const;
/// \brief Get round_mode.
///
/// \return round_mode.
// 获取 round_mode 参数
RoundMode get_round_mode() const;
};

25
mindspore/lite/experimental/kernel/convolution_fp32.cc
View File

@ -18,43 +18,60 @@
#include "src/common/tensor_util.h"
namespace mindspore::kernel {
// 构造函数，用于初始化 ConvolutionCPUFp32 类的实例
ConvolutionCPUFp32::ConvolutionCPUFp32(OpParameter *parameter, std::vector<lite::Tensor *> in_tensors,
std::vector<lite::Tensor *> out_tensors, const lite::Context *ctx)
: LiteKernel(parameter, in_tensors, out_tensors, ctx) {
// 遍历输入张量，为每个输入张量创建相应的 TensorC 结构
for (size_t i = 0; i < in_tensors.size(); i++) {
in[i] = reinterpret_cast<TensorC *>(malloc(sizeof(TensorC)));
Tensor2TensorC(in_tensors[i], in[i]);
}
// 为输出张量创建相应的 TensorC 结构
out[0] = reinterpret_cast<TensorC *>(malloc(sizeof(TensorC)));
Tensor2TensorC(out_tensors[0], out[0]);
}
// 析构函数，用于释放资源
ConvolutionCPUFp32::~ConvolutionCPUFp32() {
if (kernel == nullptr) {
return;
}
// 释放 kernel 资源
kernel->release(kernel);
free(kernel);
// 释放输入张量的资源
for (size_t i = 0; i < in_tensors_.size(); i++) {
free(in[i]);
}
// 释放输出张量的资源
free(out[0]);
}
// 准备函数，用于初始化 kernel 并调整其大小
int ConvolutionCPUFp32::Prepare() {
// 创建 kernel 对象
kernel = CreateKernel(op_parameter_, in, in_tensors_.size(), out, 1);
if (kernel == nullptr) {
return -1;
}
// 调整 kernel 的大小以适应输入和输出张量的尺寸
auto ret = kernel->resize(kernel, in, in_tensors_.size(), out, 1);
if (ret != NNACL_OK) {
return ret;
}
// 准备 kernel
return kernel->prepare(kernel);
}
int ConvolutionCPUFp32::Run() { return kernel->compute(kernel); }
int ConvolutionCPUFp32::ReSize() { return kernel->resize(kernel, in, in_tensors_.size(), out, 1); }
// 运行函数，执行卷积计算
int ConvolutionCPUFp32::Run() {
return kernel->compute(kernel);
}
// 重新调整函数，用于调整 kernel 的大小以适应输入和输出张量的尺寸
int ConvolutionCPUFp32::ReSize() {
return kernel->resize(kernel, in, in_tensors_.size(), out, 1);
}
} // namespace mindspore::kernel

74
mindspore/lite/src/runtime/kernel/arm/base/convolution_base.cc
View File

@ -26,31 +26,44 @@ using mindspore::lite::RET_OK;
using mindspore::schema::ActivationType;
namespace mindspore::kernel {
// 用于分配带对齐要求的内存块的函数
void *ConvolutionBaseCPUKernel::MallocAlignedData(size_t alignment, size_t size) {
// 检查分配的内存块大小是否超出最大允许大小
MS_CHECK_TRUE_RET(size + alignment < MAX_MALLOC_SIZE, nullptr);
// 分配内存块，大小为 size + alignment
auto ptr = malloc(size + alignment);
if (ptr == nullptr) {
MS_LOG(ERROR) << "MallocAlignedData failed!";
return nullptr;
}
// 计算对齐后的指针地址
uintptr_t aligned_ptr = (reinterpret_cast<uintptr_t>(ptr) + alignment - 1) & (~(alignment - 1));
// 在 addr_map 中记录原始指针与对齐后指针的映射关系
addr_map[aligned_ptr] = ptr;
// 返回对齐后的指针
return reinterpret_cast<void *>(aligned_ptr);
}
// 用于释放带对齐要求的内存块的函数
void ConvolutionBaseCPUKernel::FreeAlignedData(void **ptr) {
if (*ptr != nullptr && addr_map[reinterpret_cast<uintptr_t>(*ptr)] != nullptr) {
// 通过对齐后的指针查找原始指针，并释放原始指针对应的内存块
free(addr_map[reinterpret_cast<uintptr_t>(*ptr)]);
// 将对应关系从 addr_map 中移除
addr_map[reinterpret_cast<uintptr_t>(*ptr)] = nullptr;
// 将传入的指针置为空
*ptr = nullptr;
}
}
// 析构函数，用于释放内存和资源
ConvolutionBaseCPUKernel::~ConvolutionBaseCPUKernel() {
if (addr_map.find(reinterpret_cast<uintptr_t>(packed_weight_)) != addr_map.end()) {
// 如果 packed_weight_ 对应的内存块需要对齐，调用 FreeAlignedData 释放
FreeAlignedData(reinterpret_cast<void **>(&packed_weight_));
} else if (!op_parameter_->is_train_session_) {
#ifdef SERVER_INFERENCE
// 在非训练会话模式下，如果 packed_weight_ 不需要对齐且使用了 MALLOC 方式分配内存，则直接释放
if (packed_weight_ != nullptr && weight_is_packed_ == lite::MALLOC) {
#endif
free(packed_weight_);
@ -60,26 +73,32 @@ ConvolutionBaseCPUKernel::~ConvolutionBaseCPUKernel() {
#endif
}
if (addr_map.find(reinterpret_cast<uintptr_t>(bias_data_)) != addr_map.end()) {
// 如果 bias_data_ 对应的内存块需要对齐，调用 FreeAlignedData 释放
FreeAlignedData(reinterpret_cast<void **>(&bias_data_));
} else if (bias_data_ != nullptr) {
// 如果 bias_data_ 不需要对齐，直接释放内存
free(bias_data_);
bias_data_ = nullptr;
}
}
// 用于释放量化参数的函数
void ConvolutionBaseCPUKernel::FreeQuantParam() {
if (conv_quant_arg_ == nullptr) {
return;
}
if (conv_quant_arg_->real_multiplier_ != nullptr) {
// 释放 real_multiplier_ 内存
free(conv_quant_arg_->real_multiplier_);
conv_quant_arg_->real_multiplier_ = nullptr;
}
if (conv_quant_arg_->left_shift_ != nullptr) {
// 释放 left_shift_ 内存
free(conv_quant_arg_->left_shift_);
conv_quant_arg_->left_shift_ = nullptr;
}
if (conv_quant_arg_->right_shift_ != nullptr) {
// 释放 right_shift_ 内存
free(conv_quant_arg_->right_shift_);
conv_quant_arg_->right_shift_ = nullptr;
}
@ -149,42 +168,63 @@ bool ConvolutionBaseCPUKernel::CheckParamsValid() const {
}
return true;
}
// 初始化卷积操作的权重和偏置
int ConvolutionBaseCPUKernel::InitConvWeightBias() {
// 如果当前会话是训练会话，更新原始权重和偏置
if (op_parameter_->is_train_session_) {
UpdateOriginWeightAndBias();
}
// 获取权重张量
auto weight_tensor = in_tensors_.at(kWeightIndex);
CHECK_NULL_RETURN(weight_tensor);
// 获取权重张量的形状
auto shape = weight_tensor->shape();
// 检查权重张量是否包含未知维度（-1），如果是，将在运行时初始化权重和偏置
if (std::find(shape.begin(), shape.end(), -1) != shape.end()) {
MS_LOG(WARNING) << "The shape of weight tensor is not ready, the weight and bias would be inited in runtime.";
return RET_OK;
}
// 分配内存用于权重和偏置数据
if (MallocWeightBiasData() != RET_OK) {
MS_LOG(ERROR) << "Malloc data for bias and weight failed.";
return RET_ERROR;
}
// 如果输入张量数量为2，表示存在偏置项
if (in_tensors_.size() == kInputSize2) {
auto bias_tensor = in_tensors_.at(kBiasIndex);
CHECK_NULL_RETURN(bias_tensor);
// 检查偏置张量的大小是否为0
MS_CHECK_FALSE(bias_tensor->Size() == 0, RET_ERROR);
// 检查原始偏置数据是否为nullptr
if (origin_bias_ == nullptr) {
MS_LOG(ERROR) << "Convolution op " << this->name() << " bias data is nullptr.";
return RET_ERROR;
}
// 将原始偏置数据复制到分配的偏置数据中
memcpy(bias_data_, origin_bias_, bias_tensor->Size());
} else {
// 如果输入张量数量为1，表示没有偏置项
MS_ASSERT(in_tensors_.size() == kInputSize1);
}
// 如果不是训练会话并且运行在服务器推理模式下
if (!op_parameter_->is_train_session_) {
#ifdef SERVER_INFERENCE
// 如果权重已经打包且处于服务器推理模式，则不进行权重打包
if (weight_is_packed_ == lite::PACKED) {
MS_LOG(DEBUG) << "not do weight pack.";
return RET_OK;
}
#endif
// 如果原始权重数据不为空，则进行权重打包
if (origin_weight_ != nullptr) {
PackWeight();
} else {
@ -194,26 +234,38 @@ int ConvolutionBaseCPUKernel::InitConvWeightBias() {
}
return RET_OK;
}
// 重新打包卷积操作的权重数据
int ConvolutionBaseCPUKernel::RepackWeight() {
// 如果原始权重数据和输入权重张量的数据都为空，报错并返回错误
if (origin_weight_ == nullptr && in_tensors_.at(kWeightIndex)->data() == nullptr) {
MS_LOG(ERROR) << "Convolution op " << this->name() << " weight data is nullptr.";
return RET_ERROR;
}
// 如果原始权重数据为空，将其设置为输入权重张量的数据
origin_weight_ = origin_weight_ != nullptr ? origin_weight_ : in_tensors_.at(kWeightIndex)->data();
// 如果打包后的权重数据为空并且初始化权重和偏置失败，报错并返回错误
if (packed_weight_ == nullptr && InitConvWeightBias() != RET_OK) {
MS_LOG(ERROR) << "Malloc data for bias and weight failed.";
return RET_ERROR;
}
// 如果需要重新打包或者当前会话为训练会话
if (IsRepack() || (op_parameter_->is_train_session_)) {
if (op_parameter_->is_train_session_) {
// 如果当前会话为训练会话，将打包后的权重数据指针设置为工作空间的指针，并初始化为零
packed_weight_ = reinterpret_cast<float *>(workspace());
memset(packed_weight_, 0, workspace_size());
} else {
// 如果不是训练会话，将重新打包标志设置为false
is_repack_ = false;
}
// 打包权重数据
PackWeight();
}
return RET_OK;
}
@ -426,39 +478,54 @@ void ConvolutionBaseCPUKernel::SetRoundingAndMultipilerMode() {
}
}
}
// 设置卷积操作的量化参数
int ConvolutionBaseCPUKernel::SetQuantParam() {
// 分配量化参数内存
auto ret = MallocQuantParam();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Malloc quant param failed.";
return ret;
}
// 设置输入张量的量化参数
ret = SetInputTensorQuantParam();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Set Input Tensor Quant Param Failed.";
return ret;
}
// 设置滤波器张量的量化参数
ret = SetFilterTensorQuantParam();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Set Filter Tensor Quant Param Failed.";
return ret;
}
// 设置输出张量的量化参数
ret = SetOutputTensorQuantParam();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Set Output Tensor Quant Param Failed.";
return ret;
}
// 设置是否为每通道量化
ret = SetIfPerChannel();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Set if per tensor channel failed.";
return ret;
}
// 设置舍入模式和乘法器模式
SetRoundingAndMultipilerMode();
// 设置量化乘法器
ret = SetQuantMultiplier();
if (ret != RET_OK) {
MS_LOG(ERROR) << "Set Quant Multiplier Failed.";
return ret;
}
// 检查激活函数类型，计算量化后的激活范围
bool relu = conv_param_->act_type_ == ActType_Relu;
bool relu6 = conv_param_->act_type_ == ActType_Relu6;
CalculateActivationRangeQuantized(relu, relu6, conv_param_->conv_quant_arg_.output_quant_args_[0].zp_,
@ -468,6 +535,7 @@ int ConvolutionBaseCPUKernel::SetQuantParam() {
return RET_OK;
}
void ConvolutionBaseCPUKernel::UpdateOriginWeightAndBias() {
if (in_tensors_.at(kWeightIndex)->data() != nullptr) {
origin_weight_ = in_tensors_.at(kWeightIndex)->data();

48
mindspore/python/mindspore/ops/composite/base.py
View File

@ -434,6 +434,15 @@ class _Grad(GradOperation_):
def __init__(self, get_by_list=False, sens_param=False, get_by_position=False):
"""Initialize _Grad."""
"""
参数:
- get_by_list: 是否按列表获取梯度，默认为 False
- sens_param: 是否使用灵敏度参数，默认为 False
- get_by_position: 是否按位置获取梯度，默认为 False
- has_aux: 是否包含辅助输出，默认为 False
- get_value: 是否获取值，默认为 False
- return_ids: 是否返回标识符，默认为 False
"""
if not isinstance(get_by_position, bool):
raise TypeError(f"For '_Grad', the 'get_by_position' should be bool, "
f"but got {type(get_by_position).__name__}")
@ -453,6 +462,14 @@ class _Grad(GradOperation_):
self.grad_position = None
def __call__(self, fn, weights=None, grad_position=0):
"""
调用 _Grad 类实例以生成梯度函数。
参数：
- fn: 输入的函数
- weights: 权重参数，默认为 None
- grad_position: 梯度位置，默认为 0
"""
if self.grad_fn is not None and self.fn == fn and self.grad_position == grad_position:
return self.grad_fn
grad_ = _Grad(self.get_by_list, self.sens_param, self.get_by_position)
@ -507,25 +524,32 @@ class _Grad(GradOperation_):
def _pynative_forward_run(self, grad, args, kwargs, fn):
""" Pynative forward runs to build grad graph. """
new_kwargs = kwargs
new_kwargs = kwargs # 将传入的关键字参数 kwargs 赋值给新变量 new_kwargs
if self.sens_param:
# 如果 self.sens_param 为 True，表示函数需要灵敏度参数
if 'sens' in kwargs.keys():
new_kwargs = kwargs.copy()
new_kwargs.pop('sens')
# 如果 'sens' 存在于 kwargs 的键中
new_kwargs = kwargs.copy() # 复制 kwargs 到新的字典 new_kwargs
new_kwargs.pop('sens') # 从 new_kwargs 中删除 'sens' 键
else:
args = args[:-1]
args = args[:-1] # 否则，移除参数列表 args 的最后一个参数
if isinstance(fn, FunctionType):
# 如果 fn 是 FunctionType 类型的对象（通常表示普通函数或方法）
if not _pynative_executor.check_run(grad, fn, *args, **new_kwargs):
_pynative_executor.set_grad_flag(True)
_pynative_executor.new_graph(fn, *args, **new_kwargs)
outputs = fn(*args, **new_kwargs)
_pynative_executor.end_graph(fn, outputs, *args, **new_kwargs)
# 如果还没有运行过与 grad 关联的 fn 函数
_pynative_executor.set_grad_flag(True) # 设置梯度标志为 True，表示进入梯度计算模式
_pynative_executor.new_graph(fn, *args, **new_kwargs) # 创建新的计算图
outputs = fn(*args, **new_kwargs) # 执行 fn 函数并获取其输出
_pynative_executor.end_graph(fn, outputs, *args, **new_kwargs) # 结束计算图的构建
else:
# Check if fn has run already.
# 如果 fn 不是 FunctionType 类型的对象（可能是 nn.Cell 或其他类型的对象）
# 检查 fn 是否已经运行过，如果没有，则设置梯度标志为 True，运行 fn，然后将梯度标志设置为 False
if not _pynative_executor.check_run(grad, fn, *args, **new_kwargs):
fn.set_grad()
fn(*args, **new_kwargs)
fn.set_grad(False)
fn.set_grad() # 设置 fn 为梯度计算模式
fn(*args, **new_kwargs) # 运行 fn 函数
fn.set_grad(False) # 将 fn 设置为非梯度计算模式
class _Vmap(VmapOperation_):

57
tests/perf_test/lenet.py
View File

@ -39,13 +39,14 @@ def weight_variable():
class LeNet5(nn.Cell):
"""
Lenet network
LeNet network
Args:
num_class (int): Num classes. Default: 10.
Returns:
Tensor, output tensor
Examples:
>>> LeNet(num_class=10)
@ -53,26 +54,42 @@ class LeNet5(nn.Cell):
def __init__(self, num_class=10, channel=1):
super(LeNet5, self).__init__()
self.num_class = num_class
self.conv1 = conv(channel, 6, 5)
self.conv2 = conv(6, 16, 5)
self.fc1 = fc_with_initialize(16 * 5 * 5, 120)
self.fc2 = fc_with_initialize(120, 84)
self.fc3 = fc_with_initialize(84, self.num_class)
# 定义卷积层1，输入通道数为channel，输出通道数为6，卷积核大小为5
self.conv1 = nn.Conv2d(channel, 6, kernel_size=5)
# 定义卷积层2，输入通道数为6，输出通道数为16，卷积核大小为5
self.conv2 = nn.Conv2d(6, 16, kernel_size=5)
# 定义全连接层1，输入特征大小为16*5*5，输出特征数量为120
self.fc1 = nn.Dense(16 * 5 * 5, 120)
# 定义全连接层2，输入特征数量为120，输出特征数量为84
self.fc2 = nn.Dense(120, 84)
# 定义全连接层3，输入特征数量为84，输出特征数量为num_class
self.fc3 = nn.Dense(84, self.num_class)
self.relu = nn.ReLU()
# 定义最大池化层，池化核大小为2，步长为2
self.max_pool2d = nn.MaxPool2d(kernel_size=2, stride=2)
# 定义展平层，用于将多维的输入展平成一维
self.flatten = nn.Flatten()
def construct(self, x):
x = self.conv1(x)
x = self.relu(x)
x = self.max_pool2d(x)
x = self.conv2(x)
x = self.relu(x)
x = self.max_pool2d(x)
x = self.flatten(x)
x = self.fc1(x)
x = self.relu(x)
x = self.fc2(x)
x = self.relu(x)
x = self.fc3(x)
return x
# 前向传播函数定义
x = self.conv1(x) # 卷积层1
x = self.relu(x) # ReLU激活函数
x = self.max_pool2d(x) # 池化层1
x = self.conv2(x) # 卷积层2
x = self.relu(x) # ReLU激活函数
x = self.max_pool2d(x) # 池化层2
x = self.flatten(x) # 展平层
x = self.fc1(x) # 全连接层1
x = self.relu(x) # ReLU激活函数
x = self.fc2(x) # 全连接层2
x = self.relu(x) # ReLU激活函数
x = self.fc3(x) # 全连接层3
return x