下载掌阅APP，畅读海量书库
立即打开
MXMACA核函数调用是对C语言函数调用语句的延伸，<<<>>>运算符内是核函数的执行配置。
第3.4节介绍了MXMACA编程的线程层次结构，通过核函数执行配置可以指定线程在GPU上调度运行的方式。执行配置的值包含：第一个值是线程网格维度，即启动线程块的数目；第二个值是线程块的维度，即每个线程块的线程数量。通过指定线程网格和线程块的大小，可以配置核函数中线程的使用数量和线程的使用布局。
同一个线程块内的线程之间可以相互协作，不同线程块内的线程不能协作。对于一个给定的问题，可以使用不同的线程网格和线程块来组织线程。假设有32个数据元素，每线程块若设置为8个元素，则需要启动4个线程块。
二维线程网格的线程布局示例如图3-8所示。
图3-8 二维线程网格的线程布局示例
二维线程网格可以使用变量（blockIdx.x，blockIdx.y）来描述，二维线程块可以使用变量（threadIdx.x，threadIdx.y）来描述。由于数据在全局内存中是线性存储的，借助这些变量可以进行以下操作。
● 在线程网格中标识一个唯一的线程，一个GPU线程一般通过GPU的一个核函数进行处理。
● 建立线程块和数据元素之间的映射关系。各个线程块是并行执行的，线程块之间无法通信，也没有执行顺序。
以图3-7为例，该线程网格总的线程数量 N 可通过式（3-1）计算。
同时，通过blockIdx.x、blockIdx.y、blockIdx.z、threadIdx.x、threadIdx.y、threadIdx.z就可以完全定位一个线程的坐标位置。例如，将所有的线程排成一个序列，序列号为0，1，2，…， N ，可以按以下步骤找到当前的线程序列号Idx。
（1）按式（3-2）找到当前线程位于线程网格中的哪一个线程块（blockId）。
（2）再按式（3-3）找到当前线程位于线程块中的哪一个线程（threadId）。
（3）按式（3-4）计算一个线程块中共有多少个线程（ M ）。
（4）按式（3-5）求出当前的线程序列号Idx。
核函数调用和主机端线程是异步的，当核函数调用结束后，控制权立刻返回给主机端。可以使用下面的函数来强制主机端程序等待所有的核函数执行结束。
核函数是在设备端执行的代码，它负责定义与GPU线程的计算和管理相关的数据访问。当核函数被调用时，许多不同的GPU线程将会并行执行同一计算任务，其声明方式如下。
其中，__global__是一种函数类型限定符，函数类型限定符指定一个函数是在主机端执行还是在设备端执行，以及是可被主机端调用还是可被设备端调用。表3-2总结了MXMACA C/C++语言中不同的函数类型限定符，限定符__device__和__host__可以一起使用，这样，函数可以同时在主机端和设备端进行编译。
表3-2 MXMACA C/C++语言中不同的函数类型限定符
核函数通常有以下几个限制：（1）在常规的内存管理方式中只能访问设备内存，访问主机内存需要特殊的内存管理操作（详见第6.3节）；（2）必须有void返回类型；（3）不支持可变数量的参数；（4）不支持静态变量；（5）显示为异步行为。