在前面平台模型时,我们学过了,一个OpenCL平台由一个主机(host)和若干设备(device)构成。那么,在这样一个平台上,程序到底如何执行的呢?了解了这个,我们才能明白应该如何来编写OpenCL程序。今天学习“执行模型”,我们以“漫话”中的向量求和为例子,来了解一下OpenCL程序执行的问题,看看程序的执行与划分“主机”和“设备”这事儿有没有什么关系?当然,既然还是在讲“模型”,我们肯定还会顺便定义一些术语和概念,便于在以后学其它部分的时候理解。
一个平台被分为一个主机和设备是有用处的,至少它便于我们说明“执行模型”:一个OpenCL的程序分两部分,在不同的地方运行。一个部分,被称为"kernel"的,在“设备”上运行;另外一部分,我们就叫它主程序(host program)吧,它在“主机”上运行。简单的说,kernels完成真正的计算,比如向量的求和工作,而主程序呢,它定义了kernel的运行环境,并且管理这些kernel的执行。
以前面说的“向量求和”作为例子。一个传统的(比如说C)程序,它看起来是这个样子:
void host_vector_add(const float *A, const float *B, float *C) {
for ( int i = 0; i < VECTOR_SIZE; i++ ) {
C[i] = A[i] + B[i];
}
}
一个向量求和的典型OpenCL程序,它的kernel是这个样子的:
__kernel void vector_add(__global const float *A, __global const float *B, __global float *C) {
// Get the index of the current element to be processed
int i = get_global_id(0);
// Do the operation
C[i] = A[i] + B[i];
}
这个?它也太神奇了吧!它只讲了求和,没有讲向量,是不是?在OpenCL环境中,设备(比如显卡GPU)最擅长的是什么?是计算,比如这里的“求和”。它可不擅长什么判断、循环什么的。它擅长的,是同时让许多(比如几百、上千个)处理单元来做相同的计算。Kernel就是写给这些处理单元的,所以它只讲“求和”,至于需要多少个处理单元来进行求和,通常是“主程序”来控制的。主程序最主要的一句代码,我们可以先看看它的样子:
queue.enqueueNDRangeKernel(kernel, NullRange, NDRange(LIST_SIZE), NullRange);
就是其中的NDRange(LIST_SIZE)告诉OpenCL,我需要LIST_SIZE这么多个处理单元来进行计算。这么说,还是很神奇,对吧?的确很神奇。这个神奇,就在NDRange上。
NDRange是OpenCL的一个重要概念。OpenCL规范上说,"An NDRange is an N-dimensional index space". 这样说的话,又引入了如何解释"index space"的问题;这样又涉及到很多别的概念,比如kernel代码例子中出现的"global id"等等,这些概念性的东西,留着下次再说吧。这里先把NDRange说一说。NDRange,就暂时把它翻译为“N维区间”吧。“区间”是中学数学里学的概念,表示从某个数到另外一个数之间的一个范围,比如中学里常用[0,1)表示从0到1之间的所有实数,它包含0但是不包含1。[0,1)是一个一维区间的例子;而例子代码中的"NDRange(LIST_SIZE)"也是一个一维区间的例子,但是和[0,1)这个略有不同的是,NDRange既然表示的是"index space",它就只包含整数,它表示的是从0开始的LIST_SIZE这么多个整数,也就是0、1、2、...一直到(LIST_SIZE - 1)这些个整数。另外,NDRange既然叫做N维,顾名思义,它是对我们中学里学的一维区间的一个扩展,它还可以表示二维、三维的区间,比如NDRange(2,2),它表示一个平面(二维)区域内的整数点,共包括(0,0)、(0,1)、(1,0)、(1,1)四个点。同理NDRange(2,2,2)是三维的,它表示一个空间(立体)范围内的8个整数点。
回到向量求和的例子,回顾一下,OpenCL程序分为kernel和主程序两个部分来执行。kernel表明了计算单元应该进行的计算,而主程序使用NDRange来表明哪些计算单元需要参与计算,程序执行时,这些计算单元可以同时进行计算,所以我们马上就能得到向量求和结果。
分享:
4
喜欢
0
赠金笔