Cambricon · Chuancysun · Jul 23, 2024 · Apr 30, 2024 · May 24, 2024 · May 31, 2024
diff --git a/docs/design_docs/sgetrf/1.png b/docs/design_docs/sgetrf/1.png
diff --git a/docs/design_docs/sgetrf/10.png b/docs/design_docs/sgetrf/10.png
diff --git a/docs/design_docs/sgetrf/11.png b/docs/design_docs/sgetrf/11.png
diff --git a/docs/design_docs/sgetrf/12.png b/docs/design_docs/sgetrf/12.png
diff --git a/docs/design_docs/sgetrf/13.png b/docs/design_docs/sgetrf/13.png
diff --git a/docs/design_docs/sgetrf/14.png b/docs/design_docs/sgetrf/14.png
diff --git a/docs/design_docs/sgetrf/15.png b/docs/design_docs/sgetrf/15.png
diff --git a/docs/design_docs/sgetrf/16.png b/docs/design_docs/sgetrf/16.png
diff --git a/docs/design_docs/sgetrf/17.png b/docs/design_docs/sgetrf/17.png
diff --git a/docs/design_docs/sgetrf/18.png b/docs/design_docs/sgetrf/18.png
diff --git a/docs/design_docs/sgetrf/19.png b/docs/design_docs/sgetrf/19.png
diff --git a/docs/design_docs/sgetrf/2.png b/docs/design_docs/sgetrf/2.png
diff --git a/docs/design_docs/sgetrf/20.png b/docs/design_docs/sgetrf/20.png
diff --git a/docs/design_docs/sgetrf/21.png b/docs/design_docs/sgetrf/21.png
diff --git a/docs/design_docs/sgetrf/22.png b/docs/design_docs/sgetrf/22.png
diff --git a/docs/design_docs/sgetrf/23.png b/docs/design_docs/sgetrf/23.png
diff --git a/docs/design_docs/sgetrf/24.png b/docs/design_docs/sgetrf/24.png
diff --git a/docs/design_docs/sgetrf/25.png b/docs/design_docs/sgetrf/25.png
diff --git a/docs/design_docs/sgetrf/26.png b/docs/design_docs/sgetrf/26.png
diff --git a/docs/design_docs/sgetrf/3.png b/docs/design_docs/sgetrf/3.png
diff --git a/docs/design_docs/sgetrf/4.png b/docs/design_docs/sgetrf/4.png
diff --git a/docs/design_docs/sgetrf/5.png b/docs/design_docs/sgetrf/5.png
diff --git a/docs/design_docs/sgetrf/6.png b/docs/design_docs/sgetrf/6.png
diff --git a/docs/design_docs/sgetrf/7.png b/docs/design_docs/sgetrf/7.png
diff --git a/docs/design_docs/sgetrf/8.png b/docs/design_docs/sgetrf/8.png
diff --git a/docs/design_docs/sgetrf/9.png b/docs/design_docs/sgetrf/9.png
diff --git a/docs/design_docs/sgetrf/sgetrf.md b/docs/design_docs/sgetrf/sgetrf.md
diff --git a/docs/user_guide/9_operators/index.rst b/docs/user_guide/9_operators/index.rst
@@ -793,3 +793,10 @@ mluOpLgamma
 
 - ``x`` 为输入张量。
 
+.. Sgetrf2::
+
+mluOpSgetrf2
+---------------
+执行 LU 分解，将一个矩阵分解为一个下三角矩阵（L）和一个上三角矩阵（U），参数``mode``用来指定是否进行选主元操作。
+
+该算子包含7个输入：handle 为操作句柄，x_desc 与 x 分别描述并提供输入矩阵的信息；两个输出：y_desc 与 y 分别描述并存储输出矩阵的信息；此外，还包含一个参数 mode，用于指定是否进行选主元，值为0表示选择非主元模式，ipiv表示置换矩阵，以及一个 workspace 用于临时存储计算过程中的数据。
-该算子包含7个输入：handle 为操作句柄，x_desc 与 x 分别描述并提供输入矩阵的信息；两个输出：y_desc 与 y 分别描述并存储输出矩阵的信息；此外，还包含一个参数 mode，用于指定是否进行选主元，值为0表示选择非主元模式，ipiv表示置换矩阵，以及一个 workspace 用于临时存储计算过程中的数据。
+该算子包含7个输入：其中，``handle`` 为操作句柄，``x_desc`` 与 ``x`` 分别描述并提供输入矩阵的信息；两个输出：``y_desc`` 与 ``y`` 分别描述并存储输出矩阵的信息；此外，还包含一个参数 ``mode`` 用于指定是否进行选主元操作，值为0时，表示选择非主元模式，``ipiv`` 表示置换矩阵，以及 ``workspace`` 用于临时存储计算过程中的数据。
-该算子包含7个输入：handle 为操作句柄，x_desc 与 x 分别描述并提供输入矩阵的信息；两个输出：y_desc 与 y 分别描述并存储输出矩阵的信息；此外，还包含一个参数 mode，用于指定是否进行选主元，值为0表示选择非主元模式，ipiv表示置换矩阵，以及一个 workspace 用于临时存储计算过程中的数据。
+该算子包含7个输入：其中，``handle`` 为操作句柄，``x_desc`` 与 ``x`` 分别描述并提供输入矩阵的信息；两个输出：``y_desc`` 与 ``y`` 分别描述并存储输出矩阵的信息；此外，还包含一个参数 ``mode`` 用于指定是否进行选主元操作，值为0时，表示选择非主元模式，``ipiv`` 表示置换矩阵，以及 ``workspace`` 用于临时存储计算过程中的数据。
diff --git a/kernels/sgetrf2/mycnrtMemcpy2D_union1.mlu b/kernels/sgetrf2/mycnrtMemcpy2D_union1.mlu
@@ -0,0 +1,61 @@
+/*************************************************************************
+ * Copyright (C) [2022] by Cambricon, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the
+ * "Software"), to deal in the Software without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Software, and to
+ * permit persons to whom the Software is furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+ *************************************************************************/
+#include "sgetrf2.h"
+#include "sgetrf2_utils.h"
+#include "core/logging.h"
+#include "kernels/debug.h"
+#include "kernels/kernel.h"
+
+__mlu_entry__ void MLUKernelMyCnrtMemcpy2D(int batch, int m, int n, float *dA,
+                                           int ldda, int stride_a, float *dB,
+                                           int lddb, int stride_b, int mode) {
+  int id, batch_id, tx, taskdim;
+
+  if (batch > 1) {
+    id = taskId;
+    batch_id = id / 4;
+    if (batch_id >= batch) return;
+    tx = taskId % 4;
+    dA += batch_id * stride_a;
+    dB += batch_id * stride_b;
+    taskdim = TaskUnion1;
+  } else {
+    id = taskId;
+    batch_id = 0;
+    taskdim = taskDim;
+    tx = taskId;
+  }
+  if (tx == 0)
+    __memcpy(dB, dA, n * sizeof(float), GDRAM2GDRAM, lddb * sizeof(float),
+             ldda * sizeof(float), m - 1);
+}
+
+mluOpStatus_t MLUOP_WIN_API KernelMyCnrtMemcpy2D(
+    cnrtDim3_t k_dim, cnrtFunctionType_t k_type, cnrtQueue_t queue,
+    mluOpDataType_t d_type, int batch, int m, int n, float *dA, int ldda,
+    int stride_a, float *dB, int lddb, int stride_b, int mode) {
+  KERNEL_CHECK(MLUKernelMyCnrtMemcpy2D<<<k_dim, k_type, queue>>>(
+      batch, m, n, dA, ldda, stride_a, dB, lddb, stride_b, mode));
+
+  return MLUOP_STATUS_SUCCESS;
+}