核函数
核函数的定义
Ascend C支持开发者自定义核函数来扩展C++，核函数在AI处理器上执行时实际是若干线程在并行执行，每个线程有独立的寄存器和栈，共同完成数据计算任务。
核函数的定义示例如下：
1 | __global__ void kernel_name(argument list) |
定义核函数时需要遵循以下规则：
- 使用函数类型限定符__global__， 标识它是一个核函数。
- 核函数必须具有void返回类型。
- 函数入参支持的类型如下：
- 基础数据类型，如int32_t、float等。
- 基础数据类型的指针类型，如int32_t*、float*等，实际上这些指针指向的是GlobalMemory内存。
核函数的调用
算子程序中的函数可以分为三类：host侧执行函数、核函数（device侧执行）、device侧执行函数（除核函数之外）。下图以Kernel直调算子开发方式为例，描述三者的调用关系：
- host侧执行函数可以调用其它host执行函数，也就是通用C/C++编程中的函数调用；也可以通过<<<...>>>调用核函数。
- 核函数可以调用除核函数之外的其它device侧执行函数。
- device侧执行函数（除核函数之外）使用类型限定符__aicore__标识，可以调用同类的其它device侧执行函数。
图1 算子程序中三种函数间的调用关系
Host侧通过内核调用符<<<...>>>的语法形式调用核函数，如下所示：
1 | kernel_name<<<numBlocks, threadsPerBlock, dynUBufSize, stream>>>(args...) |
执行配置由4个参数决定：
- numBlocks：为核函数配置的线程块的个数，即启用的核数, 支持int32_t和dim3类型；
- threadsPerBlock：每个线程块内并发执行的线程数量，支持int32_t和dim3类型；
- dynUBufSize：动态申请内存空间总大小，一般情况设置为0；
- stream：用于host侧和device侧的流同步。
父主题： AI Core SIMT编程