cann/asc-devkit AI CPU编程

AI CPU编程【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkitAI CPU是位于Device侧ARM64架构的处理器其具备与AI Core相同的内存访问能力可直接访问Device侧内存资源也可以与Host侧的CPU一样进行类似的数据计算通常作为AI Core的补充主要承担非矩阵类、逻辑比较复杂的分支密集型计算。AI CPU的运行环境为基础的Linux环境编程时可使用libc库C标准库STL模板库等。其硬件架构图如下所示图1AI CPU硬件架构图本节介绍的AI CPU编程仅支持如下产品型号Ascend 950PR/Ascend 950DTAtlas A3 训练系列产品/Atlas A3 推理系列产品Atlas A2 训练系列产品/Atlas A2 推理系列产品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文件内容如下// 调用printf接口需要包含的头文件 #include aicpu_api.h __global__ __aicpu__ uint32_t hello_world(void *args) { AscendC::printf(Hello World!!!\n); return 0; }[!NOTE]说明 编程时需要遵循如下规范__aicpu__ __global__函数不能是void返回类型并且入参只能是一个指针。__aicpu__ __global__函数不能是类的成员函数也不能存在于匿名空间下。尽管AI CPU的Kernel函数有返回值但该返回值仅用于Runtime组件报告运行状态开发者无需编写返回逻辑也无法使用该返回值。因此对于用户而言AI CPU Kernel函数等同于void类型不能作为右值使用。AI CPU核函数调用AI CPU核函数的调用需要在.asc文件中进行和AI Core的算子调用类似同样使用语法。hello_worldnumBlocks, nullptr, stream(args, sizeof(KernelArgs));numBlocksAI CPU Device侧暂不支持分核逻辑因此Host侧调用多核无实际意义。建议设置为1。l2ctrl保留参数当前固定为nullptr开发者无需关注。stream类型为aclrtStreamstream用于维护一些异步操作的执行顺序确保按照应用程序中的代码调用顺序在Device上执行。stream创建等管理接口请参考《Runtime运行时API》。[!NOTE]说明 在编写调用代码时需要遵循如下规范__aicpu__ __global__函数不能在.asc文件中进行定义只能声明且需要使用extern。Host侧调用__global__ __aicpu__函数时必须使用异构调用语法输入的函数入参在入参指针的基础上需要输入从指针中读取的数据大小。在Host侧使用内核调用符...调用AI Core与AI CPU算子时不能使用同一条stream。加载和运行算子时需要使用Runtime API完成运行时管理和配置详细内容请参考算子运行。AI CPU算子的编译请参考AI CPU算子编译。AI CPU模板核函数若需要使用模板核函数则需要在.aicpu文件中给出模板核函数的实例化声明参考如下templatetypename 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_worldKernelArgs, 4096(void *args);并在.asc文件中新增模板核函数实例化的extern声明templatetypename T, int BUFF_SIZE extern __global__ __aicpu__ uint32_t hello_world(void *args); template extern __global__ __aicpu__ uint32_t hello_worldKernelArgs, 4096(void *args);更多进阶用法[!NOTE]说明 更多AI CPU API的使用方法请参考AI CPU API。【免费下载链接】asc-devkit本项目是CANN 推出的昇腾AI处理器专用的算子程序开发语言原生支持C和C标准规范主要由类库和语言扩展层构成提供多层级API满足多维场景算子开发诉求。项目地址: https://gitcode.com/cann/asc-devkit创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考