TIK2简介
TIK2是一种使用C/C++作为前端语言的编程框架,开发者可以使用TIK2提供的API编写自定义算子,并通过CCEC编译器将自定义算子编译成为可运行在昇腾AI处理器上的应用程序。
TIK2继承了TIK数据操作灵活的优点,除此之外,更多优势如下:
1、C/C++原语编程
2、编程模型屏蔽硬件差异,编程范式提高开发效率
3、多层级API封装,从简单到灵活,兼顾易用与高效
4、调试方法简单
编程模型
1、使用函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外,还要为核函数加上额外的函数类型限定符,包含__global__和__aicore__。
2、使用变量类型限定符
1)、指针入参变量统一的类型定义为__gm__ uint8_t*,这里统一使用uint8_t类型的指针,在后续的使用中需要将其转化为实际的指针类型;用户亦可直接传入实际的指针类型。
2)、使用__global__函数类型限定符来标识它是一个核函数,可以被<<<…>>>调用;使用__aicore__函数类型限定符来标识该核函数在设备端aicore上执行;参数列表中使用变量类型限定符__gm__来表明输入输出的指针变量指向Global Memory上的某处地址。
3、调用核函数
核函数使用内核调用符<<<…>>>这种语法形式,来规定核函数的执行配置:
kernel_name<<<blockDim, l2ctrl, stream>>>(argument list);
内核调用符仅可在NPU侧编译时调用,CPU侧编译无法识别该符号。4、
如:小王做算子调试时发现CPU域调试精度无问题,NPU域调试时编译报错,那么可能会出现这种场景的原因是(AC)
A.
小王调用了cmath函数实现算子功能
B.
小王在代码内添加了Add 计算API
C.
小王在代码内添加了std::cout打印
D.
小王在代码运算时使用了 + 号
解释:A:cmath函数不适用于NPU。
C: cpu可以使用打印,但是NPU不行。
执行配置由3个参数决定:
blockDim,规定了核函数将会在几个核上执行,blockDim的大小不能超过当前设备上核的配置个数。每个执行该核函数的核会被分配一个逻辑ID,表现为内置变量block_idx,可以在核函数的实现中直接使用;
l2ctrl,保留参数,暂时设置为固定值nullptr,开发者无需关注;
stream,类型为aclrtStream,stream是一个任务队列,应用程序通过stream来管理任务的并行。
4、内存管理
任务间数据传递使用到的内存统一由内存管理模块Pipe进行管理。如下图所示,Pipe作为片上内存管理者,通过InitBuffer接口对外提供Queue内存初始化功能,开发者可以通过该接口为指定的Queue分配内存。
5、算子分析Add
Add算子的数学表达式为: z = x + y
dstLocal–输出–目的操作数,类型为LocalTensor。
srcLocal–输入–源操作数,类型为LocalTensor
例如:小王想实现一个add算子,那么他可能会使用的代码为(AB)
A.
Add(zLocal, xLocal, yLocal, TILE_LENGTH);
B.
zLocal = xLocal + yLocal
C.
Adds(dstLocal, srcLocal, TILE_LENGTH);
D.
Sub(zLocal, xLocal, yLocal, TILE_LENGTH);
解释:B:根据Add计算规则 z=x+y。
A: 为调用。
注:其他规则
1.必须具有void返回类型。
2.使用extern “C”。
3.仅支持入参为指针或C/C++内置数据类型(Primitive data types),如:half* s0、float* s1、int32_t c。
注意事项
1、当前TIK2支持的AI处理器型号为昇腾310P AI处理器、昇腾910 AI处理器,其他型号暂不支持。
2、当前支持用户使用g++等C/C++编译器编译在cpu侧执行的TIK2算子,并使用gdb单步调试;
3、支持用户使用CCEC编译器编译在npu侧执行的TIK2算子,实现加速计算,暂不支持加载至网络模型中进行整网验证。
4、TIK2提供了多层级的0-3级API,随着级别增高,API使用的自由度降低,易用性增强。您可以根据需要选择合适的API,使用最通俗易懂的高级接口快速搭建算子逻辑,使用自由灵活的低级接口进行复杂的逻辑实现和性能调优。