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介绍