昇腾论坛

使用Ascend C开发的算子运行在AI Core上，AI Core是昇腾NPU硬件平台的计算核心，NPU内部有多个AI Core。Ascend C编程过程中会将需要处理的数据拆分同时在多个AI Core上运行，从而获取更高的性能。多个AI Core共享相同的指令代码，每个核上的运行实例唯一的区别是block_idx不同，开发者只需要关注单核上的处理程序,也就是核函数。

Ascend C提供流水线式的编程范式，基于编程范式可以快速搭建算子实现的代码框架，实现流水并行。流水线并行的概念和工业生产中的流水线是类似的，任务1完成对某片数据的处理后，将其加入到通信队列，任务2空闲时就会从队列中取出数据继续处理；可以类比为生产流水线中的工人只完成某一项固定工序，完成后就交由下一项工序负责人继续处理。

Ascend C编程范式是一种流水线式的编程范式，把算子核内的处理程序，分成多个流水任务：“搬入、计算、搬出”，通过队列（Queue）完成任务间通信和同步，并通过统一的内存管理模块（Pipe）管理任务间通信内存。开发者只需聚焦实现“搬入、计算、搬出”内容。

基于NPU域算子的调用接口编写程序，通过毕昇编译器编译后运行，可以完成算子NPU域的运行验证；基于CPU域算子的调用接口编写程序，通过标准的GCC编译器进行编译后运行，并通过GDB通用调试工具进行单步调试，精准验证程序执行流程是否符合预期。孪生调试的能力，大大提升了算子的调试效率。

开发流程：

CANN软件安装、环境变量配置：安装CANN软件后，使用CANN运行用户编译、运行时，需要以CANN运行用户登录环境，执行source ${install_path}/set_env.sh命令设置环境变量，其中${install_path}为CANN软件的安装目录。

进行核函数的定义，并在核函数中调用算子类的Init和Process函数。

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)

{

KernelAdd op;

op.Init(x, y, z);

op.Process();

}

内置宏__CCE_KT_TEST__ 是区分运行CPU模式或NPU模式逻辑的标志，在同一个main函数中通过对__CCE_KT_TEST__宏定义的判断来区分CPU和NPU侧的运行程序。

int32_t main(int32_t argc, char* argv[])

{

size_t inputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half

size_t outputByteSize = 8 * 2048 * sizeof(uint16_t); // uint16_t represent half

uint32_t blockDim = 8;

#ifdef __CCE_KT_TEST__

// 用于CPU调试的调用程序



#else

// NPU侧运行算子的调用程序

#endif

return 0;

}

CPU运行验证：

uint8_t* x = (uint8_t*)AscendC::GmAlloc(inputByteSize);

uint8_t* y = (uint8_t*)AscendC::GmAlloc(inputByteSize);

uint8_t* z = (uint8_t*)AscendC::GmAlloc(outputByteSize);

ReadFile("./input/input_x.bin", inputByteSize, x, inputByteSize);

ReadFile("./input/input_y.bin", inputByteSize, y, inputByteSize);

AscendC::SetKernelMode(KernelMode::AIV_MODE);

ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug

WriteFile("./output/output_z.bin", z, outputByteSize);

AscendC::GmFree((void *)x);

AscendC::GmFree((void *)y);

AscendC::GmFree((void *)z);

NPU侧运行验证：

// AscendCL初始化

CHECK_ACL(aclInit(nullptr));

// 创建Context绑定设备

aclrtContext context;

int32_t deviceId = 0;

CHECK_ACL(aclrtSetDevice(deviceId));

CHECK_ACL(aclrtCreateContext(&context, deviceId));

aclrtStream stream = nullptr;

CHECK_ACL(aclrtCreateStream(&stream));

// 分配Host内存，并进行数据初始化

uint8_t *xHost, *yHost, *zHost;

uint8_t *xDevice, *yDevice, *zDevice;

CHECK_ACL(aclrtMallocHost((void**)(&xHost), inputByteSize));

CHECK_ACL(aclrtMallocHost((void**)(&yHost), inputByteSize));

CHECK_ACL(aclrtMallocHost((void**)(&zHost), outputByteSize));

// 分配Device内存，并将数据从Host上拷贝到Device上

CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

CHECK_ACL(aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));

ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);

ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize);

CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));

// 用内核调用符<<<>>>调用核函数完成指定的运算

add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice);

CHECK_ACL(aclrtSynchronizeStream(stream));

// 将Device上的运算结果拷贝回Host

CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));

WriteFile("./output/output_z.bin", zHost, outputByteSize);

// 释放申请的资源

CHECK_ACL(aclrtFree(xDevice));

CHECK_ACL(aclrtFree(yDevice));

CHECK_ACL(aclrtFree(zDevice));

CHECK_ACL(aclrtFreeHost(xHost));

CHECK_ACL(aclrtFreeHost(yHost));

CHECK_ACL(aclrtFreeHost(zHost));

CHECK_ACL(aclrtDestroyStream(stream));

CHECK_ACL(aclrtDestroyContext(context));

CHECK_ACL(aclrtResetDevice(deviceId));

CHECK_ACL(aclFinalize());

CPU模式下执行如下命令：bash run.sh add_custom ascend910 AiCore cpu

NPU模式下执行如下命令：bash run.sh add_custom ascend910 AiCore npu

分别得到

关于昇腾

新闻与活动

交流与资讯

支持与服务

开源社区