Kernel函数
如异构编程QuickStartDemoACL.cce所示,Kernel函数是由__global__关键字标注的特殊函数,是设备函数的入口,由Host侧通过特定的异构调用语法<<<>>>执行。一个cce文件可包含若干个Kernel函数,Host代码既可以和Device代码共存于同一文件中,也可以分开存放,基本符合常规C++函数调用习惯。
// moduleA.cce
// Kernel 函数
__global__ [aicore] void foo(/* Args */...) {}
__global__ [aicore] void bar(/* Args */...) {}
RetCode_t FooAPI(/* Args */ ...) {
foo<<<blockDim, nullptr, stream>>>(...);
}
// moduleB.cce
extern __global__ [aicore] void bar(/*Args*/ ...);
extern RetCode_t FooAPI(...);
int main() {
bar<<<...>>>(...);
FooAPI(...);
}
Kernel函数(含设备侧函数)有以下使用约束:
- 不支持Kernel函数调用Kernel函数。
- 不支持递归调用。
- 不支持动态内存分配。
- 不支持C++标准库。
- 不支持某些基础数据类型,如double。
- 部分支持C++11、C++14、C++17特性。
- AICore运行环境不包含操作系统,简而言之,除了上述明确不支持的特性外,设备侧编程也不支持依赖OS、IO、中断等相关的特性。
父主题: CCE Intrinsic介绍