Ascend C自定义算子开发的相关概念。

编程模型

• SPMD模型：单程序多数据的并行计算模型，通过__global__ __aicore__限定符定义核函数，示例：

  extern "C" __global__ __aicore__ void my_kernel(GM_ADDR a, GM_ADDR b, GM_ADDR c) {
      // 核函数实现
  }

• 核函数：设备侧算子实现的入口点，示例中的核函数my_kernel接受三个全局内存地址作为参数。

矢量编程

• 算子类实现：基于矢量编程范式，实现算子类，包括数据搬运和计算操作，示例：

• class VectorAdd {
  public:
      __aicore__ inline void Init() {
          // 初始化操作
      }
      __aicore__ inline void Process() {
          // 处理逻辑
      }
  private:
      TPipe pipe;
      // 其他私有成员
  private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        
  
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        
    }
  };

• 核函数调用算子类：在核函数中创建算子类的实例并调用其方法，示例：

  extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) {
      VectorAdd op;
      op.Init();
      op.Process();
  }

矩阵编程（高阶API）

矩阵编程（高阶API）

在Ascend C自定义算子开发中，矩阵编程是一个关键部分，尤其是使用高阶API进行矩阵乘法（MatMul）操作。以下是关于矩阵编程（高阶API）的详细总结：

基础知识

• MatMul概述：MatMul是基础的矩阵运算，表示为C = A * B + bias，其中A和B是输入矩阵，C是输出矩阵，bias是添加到结果的偏置项。
• 数据流：MatMul的数据流涉及将数据从全局内存（Global Memory）搬运到局部内存（Local Memory），进行计算，并将结果搬运回全局内存。
• 数据格式：MatMul操作中涉及到的数据格式主要有ND（普通格式）和NZ（特殊格式），其中NZ格式是为满足AICore的高性能计算需求而设计。

Tiling策略

• 多核切分：为了实现多核并行计算，需要将矩阵数据切分到不同的核上。切分策略可以是仅切分M、N轴或同时切分M、N、K轴。
• 核内切分：由于Local Memory的限制，需要将数据进一步切分，以便在核内进行多次计算。

使用Matmul高阶API

• 创建Matmul对象：使用模板定义Matmul对象，传入A、B、C和Bias的参数类型信息。
• 初始化操作：在核函数中注册Matmul对象，并设置系统workspace。
• 设置矩阵和偏置：使用SetTensorA、SetTensorB和SetBias方法设置Matmul的输入矩阵和偏置。
• 执行计算：使用Iterate或IterateAll方法执行矩阵乘法计算。Iterate提供灵活的迭代控制，而IterateAll简化了循环迭代。
• 结束操作：使用End方法结束Matmul操作。

示例代码

// 定义Matmul类型
typedef MatmulType<TPosition::GM, CubeFormat::ND, half> aType;
// ... 其他类型定义 ...

// 创建Matmul实例
Matmul<aType, bType, cType, biasType> mm;

// 初始化并设置输入
REGIST_MATMUL_OBJ(&pipe, GetSysWorkSpacePtr(), mm);
mm.Init(&tiling);
mm.SetTensorA(gm_a);
mm.SetTensorB(gm_b);
mm.SetBias(gm_bias);

// 执行计算
while (mm.Iterate()) {
    mm.GetTensorC(gm_c);
}

// 结束Matmul操作
mm.End();

注意事项

• 在使用Matmul高阶API时，需要确保系统workspace已经设置，以便API可以正确执行。
• 在多核场景中，需要通过SetDim接口设置参与计算的核数。
• 对于非对齐场景，需要在kernel侧进行特殊的尾块处理，以确保计算的正确性。

融合算子编程

• 融合算子数据流：通过融合算子的数据流优化性能，示例中展示Cube输出到Vector输入：

  // Cube计算结果输出到Vector
  CO2 -> VECIN

• 融合算子编程范式：基于Matmul高阶API，实现融合算子的编程范式，示例：

  template<typename aType, typename bType, typename cType, typename biasType>
  __aicore__ inline void MatmulLeakyKernel::Process() {
      // 初始化Matmul对象
      // 进行MatMul计算
      // 将结果搬运到Vector核上
      // 进行Vector矢量计算
      // 将输出结果搬运到Global Memory上
  }

算子开发

• Kernel直调工程：直接使用device指针进行kernel调用，示例：

  int64_t shape_1[] = {1024, 1024};
  void *data[3];
  ACLRT_LAUNCH_KERNEL(my_add)(blockDim, stream, data[0], data[1], data[2]);

• 自定义算子工程：支持多种调用方式，包括单算子API调用，示例：

  size_t workspace_size = 0;
  aclOpExecutor *handle;
  aclTensor *tensors[3];
  // 创建aclTensor并执行算子
  aclnnMyAddGetWorkspaceSize(tensors[0], tensors[1], tensors[2], &workspace_size, &handle);
  void *workspace;
  if (aclrtMalloc(&workspace, workspace_size, ACL_MEM_MALLOC_HUGE_FIRST) != ACL_SUCCESS) {
      printf("Malloc device memory failed\n");
  }
  aclnnMyAdd(workspace, workspace_size, handle, stream);

算子调试调优

• 孪生调试：Ascend C提供孪生调试功能，可以在CPU和NPU上进行调试，示例：

  #ifdef ASCENDC_CPU_DEBUG
      // CPU调试代码
  #else
      // NPU调试代码
  #endif

• double buffer：使用double buffer机制优化数据搬运和计算并行性，示例：

  pipe.InitBuffer(inQueueX, 2, 256); // 申请两块内存，实现double buffer

其他

• 精度问题：如果核函数运行验证时存在精度问题，可以通过CPU域调试、gdb调试或printf打印来定位问题。
• 内存分配失败：如果出现AllocTensor/FreeTensor失败，可能是因为违反了队列缓冲区数量的限制。

Reference

• https://www.hiascend.com/document/detail/zh/CANNCommunityEdition/80RC2alpha003/devguide/opdevg/ascendcopdevg/atlas_ascendc_10_0019.html