TIK C++介绍
TIK C++
是一种使用**C/C++
** 作为前端语言的算子开发工具,通过四层接口抽象、并行编程范式、孪生调试等技术,极大提高算子开发效率,助力Al开发者低成本完成算子开发和模型调优部署
使用TIK C++开发自定义算子的优势:
- C/C++原语言编程
- 编程模型屏蘞硬件差异,编程范式提高开发效率
- 多层级API封装,从简单到灵活,兼顾易用与高效
- 孪生调试,CPU侧模拟NPU侧的行为,可先在CPU侧调试
核函数
核函数 (Kernel Function)
是TIK C++算子设备侧的入口。TIK C++允许用户使用核函数这种C/C++函数的语法扩展来管理设备仅的运行代码,用户在核西数中实现算子透辑的编写,例如自定义算子类及其成员函数以实现该算子的所有功能。核函数是主机侧和设备侧连接的桥梁
TIK C++和CUDA的差异
核函数是直接在设备侧执行的代码。在核函数中,需要为在一个核上执行的代码规定要进行的数据访问和计算操作,当核函数被调用时,多个核将并行执行同一个计算任务
编写核函数
函数类型限定符
除了需要按照C/C++函数声明的方式定义核函数之外,还要为核函数加上额外的函数类型限定符,包含含__global__
和__aicore__
使用 __global__
函数类型限定符来标识它是一个核函数,可以被 <<<...>>>
调用:使用 __aicore__
函数类型限定符来标识该核函数在设备侧Al Core上执行:
函数类型限定符 | 执行 | 调用 | 备注 |
---|---|---|---|
__global__ | 在设备侧执行 | 由<<<...>>> 来调用 | 必须为void 返回值类型 |
__aicore__ | 在设备侧执行 | 仅从设备端调用 | - |
函数类型限定符
变量类型限定符
为了方便:指针入参变量统一的类型定义为 __gm__ uint8_t
用户可统一使用*uint8_t
*类型的指针,并在使用时转化为实际的指针类型;亦可直接传入实际的指针类型
变量类型限定符 | 内存空间 | 意义 |
---|---|---|
__gm__ | 驻留在Global Memory 上 | 表明该指针变量指向Global Memory 上某处内存地址 |
变量类型限定符
其他规则
- 必须具有void返回类型
- 使用
extern "C"
- 仅支持入参为指针类型或C/C++内置数据类型
(Primitive Data Types)
,如:half* s0
、float* s1
、int32_t c