AI CPU编程

AI CPU核函数定义

在进行AI CPU编程时，与AI Core类似，同样需要定义设备侧函数入口（即核函数），该函数必须通过__aicpu__标识符进行声明，并且需与__global__标识符联合使用以表明其只能被Host侧调用。AI CPU的Device侧实现文件需要以.aicpu为后缀（或者以.cpp为后缀，在编译时增加-x aicpu选项）。该实现文件中包括上面介绍的核函数以及AI CPU普通函数定义，AI CPU普通函数无需添加执行空间标识符。

如下是一个AI CPU“Hello World”程序的示例，hello_world.aicpu文件内容如下：

1
2
3
4
5
6
7
8

// 调用printf接口需要包含的头文件
#include "aicpu_api.h"

__global__ __aicpu__ uint32_t hello_world(void *args)
{
    AscendC::printf("Hello World!!!\n");
    return 0;
}

编程时需要遵循如下规范：

• __aicpu__ __global__函数不能是void返回类型，并且入参只能是一个指针。
• __aicpu__ __global__函数不能是类的成员函数，也不能存在于匿名空间下。
• 尽管AI CPU的Kernel函数有返回值，但该返回值仅用于Runtime组件报告运行状态，开发者无需编写返回逻辑，也无法使用该返回值。因此，对于用户而言，AI CPU Kernel函数等同于void类型，不能作为右值使用。

AI CPU核函数调用

1

hello_world<<<numBlocks, nullptr, stream>>>(&args, sizeof(KernelArgs));

在编写调用代码时需要遵循如下规范：

• __aicpu__ __global__函数不能在.asc文件中进行定义，只能声明，且需要使用extern。
• Host侧调用__global__ __aicpu__函数时必须使用<<<>>>异构调用语法，输入的函数入参在入参指针的基础上需要输入从指针中读取的数据大小。
• 在Host侧使用内核调用符<<<...>>>调用AI Core与AI CPU算子时不能使用同一条stream。

加载和运行算子时，需要使用Runtime API，完成运行时管理和配置，详细内容请参考算子运行。AI CPU算子的编译请参考AI CPU算子编译。

AI CPU模板核函数

template<typename T, int BUFF_SIZE>
__global__ __aicpu__ uint32_t hello_world(void *args)
{
    AscendC::printf("Hello World!!!\n");
    AscendC::printf("buffer_size is %d\n", BUFF_SIZE);
    return 0;
}
template __global__ __aicpu__ uint32_t hello_world<KernelArgs, 4096>(void *args);

并在.asc文件中新增模板核函数实例化的extern声明：

template<typename T, int BUFF_SIZE>
extern __global__ __aicpu__ uint32_t hello_world(void *args);

template extern __global__ __aicpu__ uint32_t hello_world<KernelArgs, 4096>(void *args);

更多进阶用法

更多AI CPU API的使用方法请参考AI CPU API。