From 7c1c7c665fa3e88c724495f860affc77e2144335 Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Tue, 19 Sep 2023 15:45:56 +0800 Subject: [PATCH 1/8] grad_ --- .../python/mindspore/ops/composite/base.py | 48 ++++++++++++++----- 1 file changed, 36 insertions(+), 12 deletions(-) diff --git a/mindspore/python/mindspore/ops/composite/base.py b/mindspore/python/mindspore/ops/composite/base.py index 495323ede99..005ac7c9cc0 100644 --- a/mindspore/python/mindspore/ops/composite/base.py +++ b/mindspore/python/mindspore/ops/composite/base.py @@ -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_): -- 2.34.1 From c00c8e3a7348eec072c3b9f9985ec073b74db955 Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Sun, 24 Sep 2023 12:07:47 +0800 Subject: [PATCH 2/8] LeNet5 --- tests/perf_test/lenet.py | 57 ++++++++++++++++++++++++++-------------- 1 file changed, 37 insertions(+), 20 deletions(-) diff --git a/tests/perf_test/lenet.py b/tests/perf_test/lenet.py index 3864315dba3..2af4ec8fbf5 100644 --- a/tests/perf_test/lenet.py +++ b/tests/perf_test/lenet.py @@ -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 \ No newline at end of file -- 2.34.1 From 4704690e43269555353abb2a1f29064c44c2e325 Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Tue, 10 Oct 2023 12:44:39 +0800 Subject: [PATCH 3/8] regularization --- .../common/graph_kernel/substitute_dropout.h | 10 ++++++ .../nnacl/fp32/detection_post_process_fp32.c | 32 +++++++++++++++++-- 2 files changed, 39 insertions(+), 3 deletions(-) diff --git a/mindspore/ccsrc/common/graph_kernel/substitute_dropout.h b/mindspore/ccsrc/common/graph_kernel/substitute_dropout.h index 53103d29958..b6bd3c0a79a 100644 --- a/mindspore/ccsrc/common/graph_kernel/substitute_dropout.h +++ b/mindspore/ccsrc/common/graph_kernel/substitute_dropout.h @@ -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_ diff --git a/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/detection_post_process_fp32.c b/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/detection_post_process_fp32.c index f5555c0fb32..dd2c09df858 100644 --- a/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/detection_post_process_fp32.c +++ b/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/detection_post_process_fp32.c @@ -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; } -- 2.34.1 From 88f3bf423e4f898c8137ad46c987d83f85a7ab74 Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Tue, 10 Oct 2023 13:10:16 +0800 Subject: [PATCH 4/8] dropout --- .../common/graph_kernel/substitute_dropout.cc | 30 ++++++++++++++----- 1 file changed, 23 insertions(+), 7 deletions(-) diff --git a/mindspore/ccsrc/common/graph_kernel/substitute_dropout.cc b/mindspore/ccsrc/common/graph_kernel/substitute_dropout.cc index 5edb1b4a1a0..7d8d07fc957 100644 --- a/mindspore/ccsrc/common/graph_kernel/substitute_dropout.cc +++ b/mindspore/ccsrc/common/graph_kernel/substitute_dropout.cc @@ -37,29 +37,38 @@ inline const PrimitivePtr kPrimGkDropout = std::make_shared("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(); 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(node_prim->GetAttr("Seed0")); + + // 如果Seed0和Seed1都等于0，则使用全局种子值seed_ if (seed == 0) { seed = GetValue(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(kNumberTypeInt64, ShapeVector(1, SizeToLong(shape.size())), static_cast(&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(0)), uniform_real_node); uniform_real_node->set_abstract(std::make_shared(kFloat32, shape_i64)); - // Set kernel_info for uniform_real node + + // 为uniform_real节点设置kernel_info auto uniform_real_kernel_info_builder = std::make_shared(); 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(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(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 -- 2.34.1 From 82778f3a37be81f6adcc541fa5e8d7ec0f3e747c Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Tue, 10 Oct 2023 15:45:29 +0800 Subject: [PATCH 5/8] addd --- gitup.sh | 5 +++++ .../device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu | 5 +++++ 2 files changed, 10 insertions(+) create mode 100755 gitup.sh diff --git a/gitup.sh b/gitup.sh new file mode 100755 index 00000000000..aefc43a39c9 --- /dev/null +++ b/gitup.sh @@ -0,0 +1,5 @@ +#!/bin/bash +git pull origin master +git add . +git commit -m "addd" +git push origin master diff --git a/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu b/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu index 765b73ac229..af84eb8c7df 100644 --- a/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu +++ b/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/adagrad_impl.cu @@ -28,20 +28,25 @@ __device__ __forceinline__ half SqrtFunc(half input) { } template +// 定义一个 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, -- 2.34.1 From 353ec770491a4d40c36d183c3f7729f3a7ebd5ba Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Tue, 10 Oct 2023 15:54:05 +0800 Subject: [PATCH 6/8] addd --- .../kernel/cuda_impl/cuda_ops/rmsprop_impl.cu | 10 +++- .../device/gpu/kernel/nn/adagrad_gpu_kernel.h | 26 ++++++++++ .../device/gpu/kernel/nn/rmsprop_gpu_kernel.h | 49 ++++++++++++++----- 3 files changed, 72 insertions(+), 13 deletions(-) diff --git a/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cu b/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cu index ee6f93f0f73..3e0c51d2a39 100644 --- a/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cu +++ b/mindspore/ccsrc/plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cu @@ -17,16 +17,24 @@ #include #include "plugin/device/gpu/kernel/cuda_impl/cuda_ops/rmsprop_impl.cuh" +// 定义一个 CUDA 核函数，用于执行 RmsProp 算法的参数更新 template __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 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) { diff --git a/mindspore/ccsrc/plugin/device/gpu/kernel/nn/adagrad_gpu_kernel.h b/mindspore/ccsrc/plugin/device/gpu/kernel/nn/adagrad_gpu_kernel.h index c31d78d79a2..227e6df5f5f 100644 --- a/mindspore/ccsrc/plugin/device/gpu/kernel/nn/adagrad_gpu_kernel.h +++ b/mindspore/ccsrc/plugin/device/gpu/kernel/nn/adagrad_gpu_kernel.h @@ -26,8 +26,10 @@ namespace mindspore { namespace kernel { template +// 定义一个名为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 &inputs, const std::vector &, const std::vector &outputs, void *stream_ptr) override { + // 如果输入为空，直接返回 if (is_null_input_) { return true; } + + // 获取输入和输出地址 T *variable = GetDeviceAddress(inputs, 0); T *accumulation = GetDeviceAddress(inputs, 1); S *learning_rate = GetDeviceAddress(inputs, 2); G *gradient = GetDeviceAddress(inputs, 3); T *variable_out = GetDeviceAddress(outputs, 0); T *accumulation_out = GetDeviceAddress(outputs, 1); + + // 调用ApplyAdagrad函数来执行参数更新 ApplyAdagrad(inputs[0]->size / sizeof(T), update_slots, learning_rate, gradient, variable, accumulation, reinterpret_cast(stream_ptr)); + // 使用CUDA内存复制函数将更新后的数据拷贝回输出地址 CHECK_CUDA_RET_WITH_EXCEPT(kernel_node_, cudaMemcpyAsync(variable_out, variable, variable_size_, cudaMemcpyDeviceToDevice, reinterpret_cast(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(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 diff --git a/mindspore/ccsrc/plugin/device/gpu/kernel/nn/rmsprop_gpu_kernel.h b/mindspore/ccsrc/plugin/device/gpu/kernel/nn/rmsprop_gpu_kernel.h index 40e4969912e..09bd38cdbc3 100644 --- a/mindspore/ccsrc/plugin/device/gpu/kernel/nn/rmsprop_gpu_kernel.h +++ b/mindspore/ccsrc/plugin/device/gpu/kernel/nn/rmsprop_gpu_kernel.h @@ -25,27 +25,38 @@ namespace mindspore { namespace kernel { template +// 定义一个名为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 &inputs, const std::vector &, const std::vector &outputs, void *stream) override { + // 如果输入为空，直接返回 if (is_null_input_) { return true; } + + // 根据是否使用Centered RMSProp选择执行相应的操作 if (!use_center_) { + // 获取输入地址 T *variable = GetDeviceAddress(inputs, 0); T *mean_square = GetDeviceAddress(inputs, 1); T *moment = GetDeviceAddress(inputs, 2); T *learning_rate = GetDeviceAddress(inputs, 3); T *gradients = GetDeviceAddress(inputs, 4); + // 调用RMSProp函数来执行参数更新 RmsProp(learning_rate, decay_, momentum_, epsilon_, variable, mean_square, moment, gradients, size_, reinterpret_cast(stream)); } else { + // 获取输入地址 T *variable = GetDeviceAddress(inputs, 0); T *mean_gradients = GetDeviceAddress(inputs, 1); T *mean_square = GetDeviceAddress(inputs, 2); @@ -56,37 +67,50 @@ class RMSPropGpuKernelMod : public DeprecatedNativeGpuKernelMod { T *momentum = GetDeviceAddress(inputs, 7); T *epsilon = GetDeviceAddress(inputs, 8); + // 调用RMSPropCenter函数来执行参数更新 RmsPropCenter(learning_rate, decay, momentum, epsilon, variable, mean_gradients, mean_square, moment, gradients, size_, reinterpret_cast(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(kernel_node, "rho"); momentum_ = GetAttr(kernel_node, "momentum"); epsilon_ = GetAttr(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 -- 2.34.1 From 2c321a710c3e51fc1f856a01be41435a2925ad90 Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Tue, 10 Oct 2023 23:44:55 +0800 Subject: [PATCH 7/8] addd --- .../cpu/kernel/mkldnn/batch_norm_cpu_kernel.h | 23 ++++++++++---- .../gpu/kernel/nn/batch_norm_gpu_kernel.h | 30 +++++++++++++++++++ 2 files changed, 47 insertions(+), 6 deletions(-) diff --git a/mindspore/ccsrc/plugin/device/cpu/kernel/mkldnn/batch_norm_cpu_kernel.h b/mindspore/ccsrc/plugin/device/cpu/kernel/mkldnn/batch_norm_cpu_kernel.h index 38243f0b0be..3a421991682 100644 --- a/mindspore/ccsrc/plugin/device/cpu/kernel/mkldnn/batch_norm_cpu_kernel.h +++ b/mindspore/ccsrc/plugin/device/cpu/kernel/mkldnn/batch_norm_cpu_kernel.h @@ -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 &inputs, const std::vector &workspace, const std::vector &outputs) override; protected: + // 初始化输入和输出数据大小 void InitInputOutputSize(const CNodePtr &kernel_node) override; + + // 获取操作支持的数据类型 std::vector GetOpSupport() override { + // 定义支持的输入和输出数据类型列表 static std::vector 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 diff --git a/mindspore/ccsrc/plugin/device/gpu/kernel/nn/batch_norm_gpu_kernel.h b/mindspore/ccsrc/plugin/device/gpu/kernel/nn/batch_norm_gpu_kernel.h index d2a70ece76b..7887cc42882 100644 --- a/mindspore/ccsrc/plugin/device/gpu/kernel/nn/batch_norm_gpu_kernel.h +++ b/mindspore/ccsrc/plugin/device/gpu/kernel/nn/batch_norm_gpu_kernel.h @@ -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 +// 定义一个名为BatchNormGpuKernelMod的类，它继承自DeprecatedNativeGpuKernelMod类 class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod { public: + // 构造函数，在构造对象时重置资源 BatchNormGpuKernelMod() { ResetResource(); } + + // 析构函数，在销毁对象时销毁资源 ~BatchNormGpuKernelMod() override { DestroyResource(); } + // 启动GPU核函数，执行批量归一化（Batch Normalization）操作 bool Launch(const std::vector &inputs, const std::vector &workspace, const std::vector &outputs, void *stream_ptr) override { + // 未使用的变量 VARIABLE_NOT_USED(workspace); VARIABLE_NOT_USED(stream_ptr); if (is_null_input_) { return true; } + + // 获取输入数据地址 auto x = GetDeviceAddress(inputs, 0); auto scale = GetDeviceAddress(inputs, 1); auto bias = GetDeviceAddress(inputs, 2); @@ -51,15 +59,19 @@ class BatchNormGpuKernelMod : public DeprecatedNativeGpuKernelMod { z = GetPossiblyNullDeviceAddress(inputs, 5); } + // 获取输出数据地址 auto y = GetDeviceAddress(outputs, 0); T *workspace_addr = GetPossiblyNullDeviceAddress(workspace, 0); const float alpha = 1; const float beta = 0; if (is_train_) { + // 训练模式下执行批量归一化 auto reserve_addr = GetPossiblyNullDeviceAddress(outputs, 2); auto save_mean = GetDeviceAddress(outputs, 3); auto save_variance = GetDeviceAddress(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(kernel_node, "is_training"); epsilon_ = GetAttr(kernel_node, "epsilon"); exp_avg_factor_ = GetAttr(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 &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_ -- 2.34.1 From 7381e119276663d3c3ce1569fd5c87090ebaa200 Mon Sep 17 00:00:00 2001 From: darrenlu666888 <2945034270@qq.com> Date: Wed, 11 Oct 2023 01:11:21 +0800 Subject: [PATCH 8/8] addd --- .../cpu/kernel/nnacl/fp32/pooling_fp32.c | 72 +++++++++++++++++- mindspore/core/ops/max_pool.h | 48 +++++------- .../experimental/kernel/convolution_fp32.cc | 25 ++++++- .../kernel/arm/base/convolution_base.cc | 74 ++++++++++++++++++- 4 files changed, 182 insertions(+), 37 deletions(-) diff --git a/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/pooling_fp32.c b/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/pooling_fp32.c index 35cff1a47fe..83e3c214b52 100644 --- a/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/pooling_fp32.c +++ b/mindspore/ccsrc/plugin/device/cpu/kernel/nnacl/fp32/pooling_fp32.c @@ -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_; diff --git a/mindspore/core/ops/max_pool.h b/mindspore/core/ops/max_pool.h index c7500283536..8d106ec250e 100644 --- a/mindspore/core/ops/max_pool.h +++ b/mindspore/core/ops/max_pool.h @@ -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 &kernel_size = {1}, const std::vector &stride = {1}, const PadMode &pad_mode = VALID, const Format &format = NCHW, const std::vector &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 &kernel_size); - /// \brief Set strides. + // 设置 strides 参数 void set_strides(const std::vector &strides); - /// \brief Set format. + // 设置 format 参数 void set_format(const Format &format); - /// \brief Set pad. + // 设置 pad 参数 void set_pad(const std::vector &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 get_kernel_size() const; - /// \brief Get strides. - /// - /// \return strides. + // 获取 strides 参数 std::vector 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 get_pad() const; - /// \brief Get round_mode. - /// - /// \return round_mode. + // 获取 round_mode 参数 RoundMode get_round_mode() const; }; diff --git a/mindspore/lite/experimental/kernel/convolution_fp32.cc b/mindspore/lite/experimental/kernel/convolution_fp32.cc index 84da4db58f0..132b9510b58 100644 --- a/mindspore/lite/experimental/kernel/convolution_fp32.cc +++ b/mindspore/lite/experimental/kernel/convolution_fp32.cc @@ -18,43 +18,60 @@ #include "src/common/tensor_util.h" namespace mindspore::kernel { +// 构造函数，用于初始化 ConvolutionCPUFp32 类的实例 ConvolutionCPUFp32::ConvolutionCPUFp32(OpParameter *parameter, std::vector in_tensors, std::vector 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(malloc(sizeof(TensorC))); Tensor2TensorC(in_tensors[i], in[i]); } + // 为输出张量创建相应的 TensorC 结构 out[0] = reinterpret_cast(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 diff --git a/mindspore/lite/src/runtime/kernel/arm/base/convolution_base.cc b/mindspore/lite/src/runtime/kernel/arm/base/convolution_base.cc index 9ef3b8591a4..8235125670c 100644 --- a/mindspore/lite/src/runtime/kernel/arm/base/convolution_base.cc +++ b/mindspore/lite/src/runtime/kernel/arm/base/convolution_base.cc @@ -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(ptr) + alignment - 1) & (~(alignment - 1)); + // 在 addr_map 中记录原始指针与对齐后指针的映射关系 addr_map[aligned_ptr] = ptr; + // 返回对齐后的指针 return reinterpret_cast(aligned_ptr); } +// 用于释放带对齐要求的内存块的函数 void ConvolutionBaseCPUKernel::FreeAlignedData(void **ptr) { if (*ptr != nullptr && addr_map[reinterpret_cast(*ptr)] != nullptr) { + // 通过对齐后的指针查找原始指针，并释放原始指针对应的内存块 free(addr_map[reinterpret_cast(*ptr)]); + // 将对应关系从 addr_map 中移除 addr_map[reinterpret_cast(*ptr)] = nullptr; + // 将传入的指针置为空 *ptr = nullptr; } } +// 析构函数，用于释放内存和资源 ConvolutionBaseCPUKernel::~ConvolutionBaseCPUKernel() { if (addr_map.find(reinterpret_cast(packed_weight_)) != addr_map.end()) { + // 如果 packed_weight_ 对应的内存块需要对齐，调用 FreeAlignedData 释放 FreeAlignedData(reinterpret_cast(&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(bias_data_)) != addr_map.end()) { + // 如果 bias_data_ 对应的内存块需要对齐，调用 FreeAlignedData 释放 FreeAlignedData(reinterpret_cast(&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(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(); -- 2.34.1