TIK C++算子执行的不同模式
TIK C++算子可用**CPU模式或NPU模式**执行
CPU模式: 算子功能调试用,可以模拟在NPU上的计算行为,不需要依赖昇腾设备
NPU模式: 算子功能/性能调试用,可以使用NPU的强大算力进行运算加速
使用方式
使用**内置宏 __CCE_KT_TEST__ **标识被宏包括的代码在特定的模式下编译
#ifdef __CCE_KT_TEST__ 表示在CPU模式下会编译该段代码
#ifndef __CCE_KT_TEST__ 表示在NPU模式下会编译该段代码
不同模式区别
CPU模式 | NPU模式 | |
|---|---|---|
| 依赖昇腾设备 | 不依赖 | 依赖 |
| 使用的编译器 | g++编译器 | 自研ccec编译器 |
| 使用场景 | 算子功能调试 | 算子功能/性能调试 |
| 推荐开发顺序 | 先 | 后 |
| 模式限制 | 只限于精度调试 | 调试较为复杂 |
编写思路
1. 运行CPU模式包含的头文件
#include “tikicpulib.h"
#define_aicore_
2. 运行NPU模式包含的头文件
#include"acl/acl.h"
#define_aicore [aicore]
3. 核函数的定义
extern "C" __global__ __aicore__ void Hellowor1d(__gm__ uint&_t*.foo){}
4. 设备侧执行逻辑
void Hellowor1d(__gm__ uint&_t*.foo){}
5. 主机侧执行逻辑
负责数据在主机侧内存的申请,主机到设备的拷贝,核函数执行同步和回收资源的工作
int32_t main(int32_t argc, char* argv[]){
size_t fooSize = 256;
uint32_t blockDim = 8;
#ifdef __CCE_KT_TEST__
内置宏
__CCE_KT_TEST__是区分运行CPU模式或NPU模式逻辑的标志
-
主机侧执行CPU模式逻辑
使用封装的执行宏
ICPU_RUN_KFuint8_t *foo = (uint8_t *)tik2::GmAl1oc(fooSize); ICPU_RUN_KF(Hellolorld, blockDim, foo); tik2::GmFree((void*)foo); -
主机侧执行NPU模式逻辑
使用内核调用符
<<<…>>>#else aclInit (nullptr); acirtStream stream = nullptr; aclrtCreatestream(&stream);uint8_t *fooDevice; aclrtMalloc((void**)&fooDevice, fooSize, ACL_MEM_MALLOC_HUGE_FIRST); Hellolorld<<<blockDim, nullptr, stream>>>(fooDevice); aclrtSynchronizestream(stream); aclrtFree(fooDevice); aclrtDestroystream(stream); aclFinalize(); ```
#endif
return 0;
}
编译程序
-
设置环境变量
install_path=/usr/local/Ascend/ascend-toolkit/latest source ${install_path}/bin/setenv.bash # 或 source ${install_path}/../set_env.sh
-
CPU模式的流程命令
(根据不同的芯片版本而异)
# 编译命令 g++ hello_world.cpp \ -D__CCE_KT_TEST__=1 -D__CCE_AICORE__=100 -D__DAV_C100__ \ -I${install_path}/tools/tikicpulib/lib/include \ -L${install_path}/toolkit/tools/simulator/Ascend910A/lib \ -L${install_path}/tools/tikicpulib/lib -ltikicpulib_stubreg \ -L${install_path}/tools/tikicpulib/lib/Ascend910 -lstdc++ \ -O2 -std=c++17 -o hello_world_cpu # 执行命令 export LD_LIBRARY_PATH=${install_path}/tools/tikicpulib/lib:$LD_LIBRARY_PATH ./hello_world_cpu
CPU模式
-
NPU模式的流程命令
(根据不同的芯片版本而异)
# 编译命令 ccec -x cce hello_world.cpp \ --cce-aicore-arch=dav-c100 \ -I${install_path}/acllib/include \ -L${install_path}/atc/lib64 \ -lruntime -lascendcl -lstdc++ \ -O2 -std=c++17 -o hello_world_npu # 执行命令 ./hello_world_npu
NPU模式