购买
下载掌阅APP,畅读海量书库
立即打开
畅读海量书库
扫码下载掌阅APP

3.5 核函数

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)显示为异步行为。 8w/Iijw4QdwxC7BurupdbPiajN4l555oS5opHieUy7j3qaiFDiZJ6uWcYsBy8nox

点击中间区域
呼出菜单
上一章
目录
下一章
×