在前面的例子中,只是简单的介绍如何在GPU中执行GPU代码,介绍了核函数及CUDA的一些接口函数,并未涉及到并行计算的概念。这次总结一下CUDA是如何做并行计算的吧,我们还是举简单的例子,以两个向量相加来解释线程(thread)是如何运行的吧。
一、SUMMING VECTORS
主函数如下:
1.我们使用cudaMalloc()分配三个GPU内存,输入数据dev_a,dev_b和输出dev_c
2.我们使用cudaFree()函数来释放内存。
3.我们使用cudaMemcpy()函数将数据从host拷贝到device,或者将数据从device拷贝到host,分别使用标志cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost。
4.通过主函数调用核函数add<<<N,1>>>( dev_a, dev_b, dev_c)。其中尖括号内部参数1代表Blocks(线程块)数目,参数2代表Threads(线程)数目,这两个参数都可以是二维或者三维的。然后将device的指针传到核函数当中去。事实上,在上述核函数中,并行计算会以如下方式并行。
二、CUDA内建变量
在并行运算时,会出现网格(grid),线程块(block),线程(thread)三个常见的概念。一个网格下包含一个或多个线程块,一个线程块下包含多个线程。在使用时 block和grid都可以用三维的向量表示,其中block向量的元素是thread,grid的元素是block。当然线程块内的线程数不是无限的,如先前的G80一个线程块内线程数最多为512个,Fermi架构下的线程块的线程数增加到1024个,Kepler架构下的线程数达到了2048个。如下图所示:
接下来我们看看核函数实现的内容:
如同之前我们使用CUDA 关键字__global__,在这里出现了新的blockIdx.x,这是CUDA的内建变量。因为我们之前传入的add<<<N,1>>>(),其中N和1都是以为的,所以只需要对一维进行索引,这里的blockIdx代表对N的索引。我们在看看几个其他的内建变量:
1.threadIdx(.x/.y/.z代表几维索引):线程所在block中各个维度上的线程号。
2.blockIdx(.x/.y/.z代表几维索引):块所在grid中各个维度上的块号。
3.blockDim(.x/.y/.z代表各维度上block的大小):block的大小即block中线程的数量,blockDim.x代表块中x轴上的线程数量, blockDim.y代表块中y轴上的线程数量,blockDim.z代表块中z轴上的线程数量。
4.gridDim(.x/.y/.z代表个维度上grid的大小):grid的大小即grid中block的数量,gridDim.x代表grid中x轴上块的数量,gridDim.y代 表grid中y轴上块的数量,gridDim.z代表grid中z轴上块的数量。
如图所示:我们以二维为例,比如我们有一个1000 *1000的矩阵,我们将矩阵进行分块,每块大小为100*100,那么矩阵就可以分成10*10的块,在这里blockDim.x = blockDim.y = 100,gridDim.x = gridDim.y = 10。
参考文献:
[1] Addison.Wesley.CUDA.By.Example.Jul.2010.ISBN.0131387685