下载掌阅APP，畅读海量书库
立即打开
调用任意由限定符__global__声明的函数时，需要指定一个执行配置（Execution Configuration）。该执行配置指定了与设备端执行函数相关的线程网格和线程块等配置信息。
通过在函数名和参数列表之间插入<<<Dg，Db，Ns，S>>>形式的表达式来指定执行配置（<<<和>>>是一个整体，中间不能有空格）。
● 参数Dg的类型为dim3，其指定了线程网格中三个维度的值。这三个维度的乘积Dg.x×Dg.y×Dg.z表示每个线程网格中线程块的个数。
● 参数Db的类型为dim3，其指定了线程块中三个维度的值。这三个维度的乘积Db.x×Db.y×Db.z表示每个线程块中线程的个数。
● 参数Ns的类型为size_t，其指定了本次调用中每个线程块在共享内存中动态分配的字节数量。例如，Ns为1024表示每个线程块在共享内存中动态分配了1024字节，Ns分配的内存是共享内存中除静态分配的内存之外的内存。该内存可被external array声明的shared数组使用。参数Ns是可选的，如果不声明Ns的值，则Ns默认为0。
● 参数S的类型为mcStream_t，其指定了相关的流（参考第5.3节）。S是可选参数，其默认值为0。
当声明了一个__global__函数__global__void kernel（int*arg），则可以使用如下方式调用。
这里省略了参数S的值，将其设置为默认值0。同理，可以省略参数Ns的值，如果省略，则Ns为0，例如
执行配置的参数会在函数的参数计算之前完成计算，例如
其中，i的初始值为1，会优先计算执行配置中的++i，计算后i值为2，函数参数值也为2。变量j亦同样处理。因此，上面的调用等价于
如果执行配置中的参数Dg、Db或Ns超出设备允许的最大值，则调用失败。
函数执行空间限定符主要用于区分当前定义或声明的函数需要在主机上还是在设备上运行，进而约束了该函数可以被主机端或者设备端的其他函数调用。以下分别介绍__global__、__device__、__host__等函数执行空间限定符。
（1）用函数执行空间限定符__global__修饰的函数一般称为核函数，其特性如下：该函数的函数体实际运行在设备端；该函数可以被主机端的函数调用；该函数可以被设备端的函数调用。核函数的返回类型必须为void，且不能成为类（Class）的成员函数。对核函数的调用必须指定其执行配置（参考第4.3.1节），且核函数是异步调用的，其返回可能先于函数体运行完成。
（2）用函数执行空间限定符__device__修饰的函数一般称为设备端函数。设备端函数有如下特性：该函数的函数体实际运行在设备端；该函数可以被设备端的函数调用。需要注意的是，__global__和__device__不能同时修饰同一个函数。
（3）用函数执行空间限定符__host__修饰的函数一般被称为主机端函数，其特性如下：该函数的函数体实际运行在主机端；该函数可以被主机端的函数调用。如果一个函数没有任何函数执行空间限定符进行修饰，则默认该函数为主机端函数。例如系统声明或用户声明的一些函数
这些函数没有任何的函数执行空间限定符，因此，这些函数均被视为主机端函数。
需要注意的是，__global__和__host__不能同时修饰同一个函数，__host__和__device__可以同时修饰同一个函数（这表示该函数既可以运行在主机端也可以运行在设备端）。
使用宏__MACA_ARCH__可以区分主机端和设备端不同的代码路径，例如
如果代码中产生了“跨执行空间”的调用行为，则其结果是未定义的。
变量存储空间限定符用于描述变量在设备上的内存位置，以下分别介绍__device__、__constant__、__shared__、__managed__等变量存储空间限定符。
（1）变量存储空间限定符__device__用于修饰需要存储在设备端内存中的变量。使用__device__修饰的变量具有如下特性：该变量存储于全局内存空间；其生命周期持续到创建该变量的上下文被销毁为止；每个设备都有一个不重复的对象；可以被单个线程网格中的所有线程访问，也可以通过运行时库函数（如mcGetSymbolAddress、mcGetSymbolSize、mcMemcpyToSymbol等）访问。
（2）变量存储空间限定符__constant__用于修饰需要存储在设备端内存中的变量，且该变量一旦定义就不能被修改。使用__constant__修饰的变量具有如下特性：该变量存储在constant内存空间中；其生命周期持续到创建该变量的上下文被销毁为止；每个设备都有一个不重复的对象；可以被单个线程网格中的所有线程访问，也可以通过运行时库函数（如mcGetSymbolAddress、mcGetSymbolSize、mcMemcpyToSymbol等）访问。
（3）变量存储空间限定符__shared__用于修饰需要存储在片上共享内存中的变量。使用__shared__修饰的变量具有如下特性：该变量存储在线程块所分配到的共享内存空间中；该变量的生命周期持续到该线程块任务执行结束为止；每个线程块都拥有一个不同的shared对象；该变量只能被所在的线程块中的线程访问；该变量没有固定的地址。
声明一个共享内存数组可用如下代码。
该数组的大小受执行配置（默认为0）控制。以该方式声明的数组，在共享内存中有相同的起始地址，且数组中的元素需要根据偏移量进行访问。
在共享内存中，若想达到下述数组初始化的效果
可以按照下述方式对数组进行声明和初始化。
需要注意的是，指针需要根据它所指向的类型对齐，因此，下述代码将无法正常工作。
上述代码中，array1没有4字节对齐。内置向量类型的对齐要求可参见第4.2.1节。
（4）变量存储空间限定符__managed__用于修饰主机端和设备端均可访问和修改的变量。使用__managed__修饰的变量具有如下特性：可以同时被设备端和主机端的代码引用，它的地址可以从设备端或者主机端函数获取，也可以从设备端或主机端函数中直接读取或写入；该变量拥有与当前应用程序相同的生命周期。
MXMACA提供了丰富的内置向量类型（见表4-3），包括不同的维度和不同的数据类型，其在主机端和设备端均可使用。可以用构造函数make_<type>构建一个数组，使用x、y、z访问不同的分量。
表4-3 内置向量类型
续表
dim3是基于uint3的整型向量类型，相当于由3个unsigned int型组成的结构体，有3个数据成员unsigned int x、unsigned int y和unsigned int z，可被用于指定3个维度。dim3可用于一维、二维或三维的索引来标识线程，以构成一维、二维或三维线程块。在定义dim3类型的变量时，未指定的数据成员会被初始化为1。
除内置向量类型外，MXMACA还提供了一些内置变量，可用来指定线程网格和线程块的大小和线程的索引，可在核函数或设备函数内使用。以下分别介绍threadIdx、blockIdx、gridDim、blockDim、waveSize等内置变量。
（1）threadIdx被用于表示线程在线程块内的索引，其数据类型为uint3，threadIdx.x、threadIdx.y、threadIdx.z分别表示线程在x、y、z维度上的索引。
（2）blockIdx被用于表示线程块在线程网格中的索引，其数据类型为dim3，blockIdx.x、blockIdx.y、blockIdx.z分别表示线程块在x、y、z维度上的索引。
（3）gridDim被用于表示线程网格的尺寸，其数据类型为dim3，gridDim.x、gridDim.y、gridDim.z分别表示线程网格在x、y、z维度上的尺寸。
（4）blockDim被用于表示线程块的尺寸，其数据类型为dim3，blockDim.x、blockDim.y、blockDim.z分别表示线程块在x、y、z维度上的尺寸。
（5）waveSize用于表示线程束的大小，其数据类型为int，在MXMACA中默认为64。其使用示例如下。
向量运算单元用于求解例如 D = A × B + C 的矩阵乘加问题。这些操作需要一个线程束中所有的线程共同完成。这些操作所处的条件分支必须相同，否则其执行结果将是未定义行为（Undefined Behavior）。以下分别介绍fragment、load_matrix_sync、store_matrix_sync、fill_fragment、mma_sync等向量运算单元。
（1）fragment的类型声明如下。fragment描述了矩阵的分片信息，该类型是线程束中所有线程共同存储矩阵的数据结构。
（2）load_matrix_sync的函数原型如下，其被用于从内存中加载待计算的矩阵。
（3）store_matrix_sync的函数原型如下，其被用于将计算后的矩阵存储到内存中。
（4）fill_fragment的函数原型如下，其根据入参v的值遵循fragment规格填入矩阵分片。
（5）mma_sync的函数原型如下，其是以线程束为单位的矩阵乘加执行函数。
当线程束中的所有线程均执行到mma_sync语句时，执行矩阵运算 D = A × B + C 。该函数支持 C ＇ = A × B + C 的原地计算。默认支持的向量矩阵形状表见表4-4。
表4-4 默认支持的向量矩阵形状表
续表
