函数执行空间限定符（Function Execution Space Qualifier）指示函数是在Host侧执行还是在Device侧执行，以及能被调用的空间范围。
表 1 函数执行空间限定符概览
[object Object][object Object]
[object Object]__global__修饰的函数是核函数入口，有以下使用约束：
- 函数返回类型必须为void，不能是class、struct或者union的成员函数。
- 不支持递归调用。
- 对__global__函数的调用是异步的，调用后即返回Host侧的主机线程。
- 只能被Host侧函数调用，在Device上执行。
__aicore__修饰的函数只能在Device侧执行，只能被__global__函数，或者其他__aicore__函数调用。
__host__修饰的函数只能在Host侧被调用和执行。
使用内存空间限定符__ubuf__来表示动、静态内存，静态内存的大小在编译期是确定的，动态内存的大小在核函数执行时确定。
dim3[object Object][object Object]
用于指定和获取线程网格（grid）、线程块（block）在x、y、z维度上的内置结构体。
dim3由3个无符号整数组成，结构体定义为{dimx，dimy，dimz}，用于指定3个不同维度的大小，三维总数为dimx * dimy * dimz。开发者可以通过如下方式创建dim3结构。
[object Object]
当前提供了以下仅在Device上可用的dim3结构的内置变量：
gridDim[object Object][object Object]
内置全局变量，只能在核函数中使用，表示整个计算任务在各个维度上分别由多少个线程块构成。各个维度上线程块关系需满足gridDim.x * gridDim.y * gridDim.z <= 65535。
blockDim[object Object][object Object]
内置全局变量，在核函数中可以直接使用，用于获取线程块中配置的线程的三维层次结构，即启动VF时配置的dim3结构体实例值。blockDim.x，blockDim.y，blockDim.z分别表示线程块中三个维度的线程数。
blockIdx[object Object][object Object]
内置全局变量，只能在核函数中使用，用于获取块索引。表示当前线程所在的线程块在整个网格中的位置坐标。
- blockIdx.x的范围是[0, gridDim.x - 1]。
- blockIdx.y的范围是[0, gridDim.y - 1]。
- blockIdx.z的范围是[0, gridDim.z - 1]。
threadIdx[object Object][object Object]
内置全局变量，在核函数中可以直接使用，用于获取当前线程在线程块内部的索引。threadIdx.x，threadIdx.y，threadIdx.z分别表示当前线程在3个维度的索引，threadIdx.x的范围为[0, blockDim.x)，threadIdx.y的范围为[0, blockDim.y)，threadIdx.z的范围为[0, blockDim.z)。线程块内线程的索引与线程ID对应关系如下：
对于一维线程块，其线程ID为blockIdx.x * blockDim.x + threadIdx.x。
对于二维线程块，其线程ID为二维结构，其计算公式为：
[object Object]对于三维线程块，其线程ID为三维结构，其计算公式为：
[object Object]
当前提供了以下仅在Device上可用的int类型的内置变量：
warpSize
运行时变量，表示一个线程束（warp）中的线程数量，当前为固定值32。
目前提供了一系列适用于Device侧的数据类型，包括标量和短向量。短向量是由多个元素组成的简单向量。
表 2 标量数据类型
[object Object][object Object]
[object Object]短向量数据类型分为Vector X2、Vector X3、Vector X4，表示一个短向量变量有2、3、4个元素，当前支持的类型分布如下：
[object Object][object Object]
[object Object]表 3 短向量数据类型
[object Object][object Object]
[object Object]SIMT编程提供了一系列运算符，用于执行数学运算。以下是支持的运算符列表。
表 4 SIMT编程支持的运算符列表
[object Object][object Object]
[object Object]运算符使用示例如下所示：
在调用__global__限定符修饰的函数时必须指定执行配置。执行配置通过在函数名后带括号的参数列表之间插入，形如：
其中：
- grid_dim：int或dim3类型，用于指定网格（grid）的维度与规模，grid_dim.x * grid_dim.y * grid_dim.z等于启动的线程块总数。
- block_dim：int或dim3类型，用于指定每个线程块（block）的维度与规模，block_dim.x * block_dim.y * block_dim.z等于每个线程块包含的线程数。
- dynamic_mem_size：size_t类型，该参数指定除静态分配的内存外，本次调用为每个线程块动态分配的共享内存字节数。
- stream：aclrtStream类型指针，指定关联的流，用于维护异步操作的执行顺序。
以下示例展示了内核函数的声明与调用方式。
在执行函数之前，会先对上述配置参数进行校验。如果grid_dim或block_dim超出设备的最大允许规模，或dynamic_smem_bytes超过分配静态内存后剩余的可用共享内存，该函数将会执行失败。
在多线程并发执行时，每个线程使用较少的寄存器可以让更多的线程和线程块驻留在AI处理器上，从而提升性能。因此，编译器会采用启发式算法，将寄存器溢出（register spilling）和指令数量控制在最低水平，同时尽量减少寄存器的使用量。应用程序可以通过在__global__函数定义中使用__launch_bounds__()限定符来限制启动边界（launch bounds），提供附加信息辅助编译器优化这一过程，这属于可选配置。
__launch_bounds__(N) [object Object][object Object]
函数标记宏，在SIMT VF入口函数上可选配置，用于在编译期指定SIMT VF启动的最大线程数。若未配置__launch_bounds__，最大线程数默认为1024。参数N需要满足：
N >= dimx * dimy * dimz；dimx，dimy，dimz为表示线程的dim3结构体。
N的取值范围为1到2048。
最大线程数决定了每个线程可分配的寄存器数量，具体对应关系请见下表，寄存器用于存储线程中的局部变量，若局部变量的个数超出寄存器个数，容易出现栈溢出等问题。建议最大线程数与启动VF任务的dim3线程数保持一致。
表 5 __launch_bounds__的Thread数量与每个Thread可用寄存器数
[object Object][object Object]
[object Object]配置SIMT函数最大线程数为512，示例如下：
[object Object]