addd

gitup.sh

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

mindspore/ccsrc/common/graph_kernel/substitute_dropout.cc

@@ -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

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_

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<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
 

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;
 }

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_;

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

@@ -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,

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

@@ -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) {

mindspore/ccsrc/plugin/device/gpu/kernel/nn/adagrad_gpu_kernel.h

@@ -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
 

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 <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_

mindspore/ccsrc/plugin/device/gpu/kernel/nn/rmsprop_gpu_kernel.h

@@ -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
 

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<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;
 };
 

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<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

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<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();

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_):

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