CANN学习资源开源仓的算子调用快速入门

算子API调用：调用C语言实现的单算子API接口完成算子计算，设计为两段式接口。

aclnnStatus aclxxXxxGetWorkspaceSize(const aclTensor *src, ..., aclTensor *out, ..., uint64_t *workspaceSize, aclOpExecutor **executor);
aclnnStatus aclxxXxx(void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, aclrtStream stream);

以CANN内置abs算子为例，第一段接口原型与参数：

// 第一段接口
aclnnStatus aclnnAbsGetWorkspaceSize(const aclTensor *self,  aclTensor *out,  uint64_t *workspaceSize,  aclOpExecutor  **executor)

第二段接口原型与参数：

// 第二段接口
aclnnStatus aclnnAbs( void *workspace, uint64_t workspaceSize, aclOpExecutor *executor, const aclrtStream stream)

以下是代码备注。首先，头文件引入

• acl.h是ACL框架的汇总头文件，内部已通过#include引入所有核心子头文件，引入该文件即可调用ACL的全部核心功能，是单算子API调用的固定引入头文件；//同自定义算子开发快速入门
• aclnn_abs.h是CANN算子库中Abs算子的专属头文件，内含Abs算子两段式API的接口定义；开发不同算子的API调用时，需替换为对应算子的专属头文件。//不同。自定义算子开发快速入门引入的是：“kernel_operator.h”

然后下面2步和自定义算子开发快速入门（下面简称前例）是一样的。

1. aclInit()
2. aclrtSetDevice() + aclrtCreateStream()

然后在Device侧申请内存存放算子的输入和输出数据，通常申请的内存指针数量与算子的输入输出张量数量保持一致。通过aclrtMalloc接口实现，在Device侧分配指定大小的线性内存，返回已分配内存的指针，且内存首地址默认64字节对齐；需要注意的是，aclrtMalloc分配内存时会做额外的字节对齐处理，会将用户传入的申请大小size向上对齐为32字节的整数倍，再额外多分配32字节。这里定义的输入为单个形状为 [32, 32] 的 float32 ，输出为同形状、同数据类型。

然后构造算子的输入数据，通常是通过读取文件或者从其他算子输出中获取，这里模拟数据，从host的inputData.data()拷贝到device。而不像前例，专门使用了aclrtMallocHost。哦不太对，aclrtMallocHost是专门为输出而使用的。输入和这里也是类似的。

std::vector<float> inputData(INPUT_DIM0 * INPUT_DIM1, -1.0f); //32*32

void FillWithNegativeRandomSimple() {
     for (size_t i = 0; i < inputData.size(); ++i) {
        inputData[i] = -static_cast<float>(i + 1);
    }
}

所以这2步和前例也基本一样：

3. aclrtMalloc() 分配设备内存

4. aclrtMemcpy() 拷贝数据到设备

在然后，就是重大区别了。前例使用内核调用符 add<<<blockDim, nullptr, stream>>>(xDevice, yDevice, zDevice, tiling);直接调用自定义核函数，传递设备内存指针和tiling结构体。这里却要从构造输入输出Tensor开始，并与tiling毫不相干。Tensor是算子计算数据的容器，包含了数据内容和Tensor属性信息。

由于单算子API调用时输入输出都是Tensor，因此需要先构造输入输出Tensor。这里需要使用aclCreateTensor接口根据已有的Device内存（selfDevice和outDevice）和Tensor属性信息构造Tensor。

aclTensor *inputTensor = nullptr;
aclTensor *outputTensor = nullptr;

std::vector<int64_t> inputShape = {32, 32};
std::vector<int64_t> outputShape = {32, 32};

int ConstructingTensors() {
  inputTensor = aclCreateTensor(inputShape.data(), inputShape.size(), aclDataType::ACL_FLOAT, nullptr, 0, aclFormat::ACL_FORMAT_ND, inputShape.data(),
                              inputShape.size(), selfDevice);
  outputTensor = aclCreateTensor(outputShape.data(), outputShape.size(), aclDataType::ACL_FLOAT, nullptr, 0, aclFormat::ACL_FORMAT_ND, outputShape.data(),
                              outputShape.size(), outDevice);
  return 0;
}

在aclCreateTensor() 创建Tensor对象后，还需要显式的aclnnAbsGetWorkspaceSize() 获取workspace大小，然后再aclrtMalloc() 申请workspace内存。在此之后才可以像前例调用核函数一样，调用第二段接口aclnnAbs() 执行算子运算。

再往后的步骤，就与前例大同小异了：

9. aclrtSynchronizeStream()
10. aclrtMemcpy() 拷贝结果回主机(Device侧的数据无法直接打印或者保存文件)
11. aclrtFree() 释放资源（包括workspace）
12. aclFinalize()

最后是编译和运行

g++ -std=c++11 -fPIC -O0 -g -Wall  \
-I$ASCEND_TOOLKIT_HOME/include \
-I$ASCEND_TOOLKIT_HOME/include/aclnn \
Sources/test_abs.cpp -o Sources/opapi_test \
$ASCEND_TOOLKIT_HOME/lib64/libascendcl.so $ASCEND_TOOLKIT_HOME/lib64/libnnopbase.so $ASCEND_TOOLKIT_HOME/lib64/libopapi_math.so

./Sources/opapi_test

附完整代码：

#include <iostream>
#include <vector>
#include "acl/acl.h"
#include "aclnnop/aclnn_abs.h"

#define CHECK_RET(cond, return_expr) \
  do {                               \
    if (!(cond)) {                   \
      return_expr;                   \
    }                                \
  } while (0)

#define LOG_PRINT(message, ...)     \
  do {                              \
    printf(message, ##__VA_ARGS__); \
  } while (0)

int InitACL() {
  auto ret = aclInit(nullptr);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret);
  return 0;
}

int32_t deviceId = 0;
aclrtStream stream;

int InitResource() {
  auto ret = aclrtSetDevice(deviceId);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret);
  ret = aclrtCreateStream(&stream);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret);
  return 0;
}

// 定义输入输出参数
constexpr uint32_t INPUT_DIM0 = 32;
constexpr uint32_t INPUT_DIM1 = 32;
constexpr float VALUE_X = 0.0f;
constexpr float VALUE_Y = 0.0f;

void* selfDevice = nullptr;
void* outDevice = nullptr;

uint32_t dataSize = INPUT_DIM0 * INPUT_DIM1 * sizeof(float);

int AllocateMemory() {

  auto ret = aclrtMalloc(&selfDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for x failed. ERROR: %d\n", ret));

  ret = aclrtMalloc(&outDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for z failed. ERROR: %d\n", ret));
  return 0;
}

std::vector<float> inputData(INPUT_DIM0 * INPUT_DIM1, -1.0f);


void FillWithNegativeRandomSimple() {
     for (size_t i = 0; i < inputData.size(); ++i) {
        inputData[i] = -static_cast<float>(i + 1);
    }
}

int TransferData() {
  auto ret = aclrtMemcpy(selfDevice, dataSize, inputData.data(), dataSize, ACL_MEMCPY_HOST_TO_DEVICE);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy H2D for x failed. ERROR: %d\n", ret); return ret);

  return 0;
}
aclTensor *inputTensor = nullptr;

aclTensor *outputTensor = nullptr;

std::vector<int64_t> inputShape = {32, 32};

std::vector<int64_t> outputShape = {32, 32};
int ConstructingTensors() {
  inputTensor = aclCreateTensor(inputShape.data(), inputShape.size(), aclDataType::ACL_FLOAT, nullptr, 0, aclFormat::ACL_FORMAT_ND, inputShape.data(),
                              inputShape.size(), selfDevice);
  outputTensor = aclCreateTensor(outputShape.data(), outputShape.size(), aclDataType::ACL_FLOAT, nullptr, 0, aclFormat::ACL_FORMAT_ND, outputShape.data(),
                              outputShape.size(), outDevice);
  return 0;
}

void* workspace = nullptr;
uint64_t  workspaceSize = 0;
aclOpExecutor* executor;
int CalculateAndAllocateWorkspace() {
  auto ret = aclnnAbsGetWorkspaceSize(inputTensor, outputTensor, &workspaceSize, &executor);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAbsGetWorkspaceSize failed. ERROR: %d\n", ret); return ret);

  if (workspaceSize > 0) {
    ret = aclrtMalloc(&workspace, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST);
    CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc for workspace failed. ERROR: %d\n", ret); return ret);
  }

  return 0;
}

int ExecuteOperator() {
  auto ret = aclnnAbs(workspace, workspaceSize, executor, stream);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnAbs failed. ERROR: %d\n", ret); return ret);
  return 0;
}

int SynchronizeStream() {
  // 等待流执行完成
  auto ret = aclrtSynchronizeStream(stream);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret);

  return 0;
}

int ProcessOutput() {
  std::vector<float> outHost(INPUT_DIM0 * INPUT_DIM1);
  uint32_t dataSize = INPUT_DIM0 * INPUT_DIM1 * sizeof(float);

  auto ret = aclrtMemcpy(outHost.data(), dataSize, outDevice, dataSize, ACL_MEMCPY_DEVICE_TO_HOST);
  CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy D2H for z failed. ERROR: %d\n", ret); return ret);

  bool resultCorrect = true;
  for (size_t i = 0; i < outHost.size(); ++i) {
    printf("abs(%f) =  is: %f\n", inputData[i], outHost[i]);
  }
  return 0;
}

int ReleaseResource() {
  aclDestroyTensor(inputTensor);
  aclDestroyTensor(outputTensor);

  // 7. 释放device资源
  aclrtFree(selfDevice);
  aclrtFree(outDevice);
  if (workspaceSize > 0) {
    aclrtFree(workspace);
  }
  aclrtDestroyStream(stream);
  aclrtResetDevice(deviceId);

  return 0;
}
int FinalizeACL(){
    aclFinalize();
    return 0;
}
int main() {
    InitACL();
    InitResource();
    AllocateMemory();
    FillWithNegativeRandomSimple();
    TransferData();
    ConstructingTensors();
    CalculateAndAllocateWorkspace();
    ExecuteOperator();
    SynchronizeStream();
    ProcessOutput();
    ReleaseResource();
    FinalizeACL();
}