Ascend C算子开发（入门）—— 算子开发初体验

Ascend C算子开发（入门）—— 算子开发初体验

Host与Device

• Host指与Device相连接的x86服务器，ARM服务器，会利用Device提供的NN(Neural Network)计算能力完成任务。
• Device模块指安装了昇腾AI处理器的硬件设备，利用PCle接口与Host侧连接，提供NN计算能力。

以下图为例，理解传统计算机中Host端和Device端的概念。

而CANN框架中的Host和Device的概念不过是，将Device端的显卡部分替换成华为AI加速卡。

核心需要编写在AI Core上的代码。

核函数

什么是核函数

核函数是跑在AI Core上的，是Ascend C算子设备测的入口，用户使用核函数这种C/C++函数的语法扩展来管理设备测的运行代码，用户在核函数中实现算子逻辑的编写，例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧核设备侧连接的桥梁。

如何编写核函数？

核函数实现例子——Hello World

HelloWorld——实现(Device端)

#include "kernel_operator.h"
using namespace AscendC;
// 核函数定义
extern "C" __global__ __aicore__ void hello_world(){
  printf("Hello World!!!\n");
}
// 用于配置核函数
void hello_world_do(uint32_t blockDim, void* stream){
  hello_world<<<blockDim,nullptr,stream>>>();
}

• unit32_t是无符号32位整数类型。
• 在CUDA和Ascend C编程中，<<<>>>是用于启动GPU函数的一种特殊语法，称为执行配置(execution configuration)。这三个参数控制和函数的执行方式。具体如下；
  • blockDim：核函数启动的线程块的数量
  • nullptr：这个参数表示每个线程块使用的共享内存的大小(单位是字节)。当设置为nullptr或·0时(是默认值)，表示没有额外的共享内存分配。
  • stream用于指定核函数指定的流。

HelloWorld——调用(Host端)

hello_world_do(blockDim,stream);
aclrtSynchronizeStream(stream);// 阻塞，等待所有队列任务执行完成。

aclrtSynchronizeStream(stream)是在**CPU(主机端)执行的同步操作，它用于等待Device(设备端)**的AICore上的核函数执行完毕，以确保主机端能够获取正确的结果。

完整核函数泛讲

在核函数中完成初始化和数据处理，初始化阶段找到核需要计算的GlobalMemory，Process对每一段数据进行拷入、计算和拷出。

copyin阶段完成了将数据从Device端的显存移动到AICore中(从GlobalMemory到LocalMemory)，当将数据拷入时，先放入一个队列，计算时将数据从队列拿出来做计算并放到另一个队列中。

compute阶段完成计算

copyout阶段拷出数据