GPU并行编程:如何启动多维块-技术开发专区

GPU并行编程:如何启动多维块

作者：黄永兵译编辑：胡铭娅 2011-07-13 00:30 IT168网站原创

　　【IT168 专稿】欢迎继续阅读本系列文章的第5部分，在这篇文章中，我将介绍如何在GPU(Grid)上启动多维块。我们将创建和上一篇文章一样的程序，但这一次我们要显示二维数组块，每个块显示一个计算的值。

　　系列文章：

　　GPU并行编程：内核及函数的实现

　　GPU并行编程：创建一个CUDA应用程序

　　GPU并行编程：熟练使用CUDA C语言

▲图 1 网格(Grid)

　　这些块的工作方式和我们在这篇文章中看到的其它块是一样的，但因为它们是二维的，你可以将它们看作是一个坐标系统，在x和y轴上有许多块，从本质上来说，和我们前面做的事情是一样的，但这次我们要使用多维索引。

　　我们怎么做呢?首先，我们需要使用CUDA C库中的一个关键词定义我们的变量：

　　dim3 multiBlockArray(X,Y);

　　你可能要问，为什么是dim3呢?因为将来CUDA C也可能会支持三维数组，但现在它只是留作备用的，因此当你创建数组时，你只要指定x和y轴的维度，然后第三轴自动被设置为1。

　　实现我们的实验方案

　　首先，包括stdio.h和定义块数组大小：

　　#include <stdio.h>
　　#define BLOCKS 10
　　Next, we create our main-function.
　　int main( void )
　　{

　　然后定义一个二维数组，一个从GPU拷贝出或拷贝到GPU指针，以及我们的dim3变量：

　　int hostArray[BLOCKS][BLOCKS];
　　int *deviceArray;

　　接下来，我们给数组分配在设备上需要的内存，正如你所看到的，我们要处理一个二维数组，使用BLOCKS*BLOCKS分配内存：

　　cudaMalloc( (void**)&deviceArray, BLOCKS * BLOCKS * sizeof(int) );
　　cudaMemcpy( deviceArray,
　　hostArray, BLOCKS * BLOCKS * sizeof(int),
　　cudaMemcpyHostToDevice );

　　获得设备上的空间后，我们开始启动内核，在GPU上执行计算任务。

　　generateArray<<<multiBlockArray,1>>>( deviceArray );

　　这里唯一的不同是，我们传递了multiBlockArray参数，它就是我们前面提到的我们想要运行的块数。

　　接下来我们拷贝GPU上的数组到主机，这样我们才能显示它：

　　cudaMemcpy( hostArray,
　　deviceArray,
　　BLOCKS * BLOCKS * sizeof(int),
　　cudaMemcpyDeviceToHost );
　　for (int i=0; i<BLOCKS; i++)
　　{
　　 printf( “Thread ID running: %d”, hostArray[0][i] );
　　 for (int j=1; j<BLOCKS; j++)
　　 {
　　 printf( ” %d”, hostArray[j][i] );
　　 }
　　 printf( “\n” );
　　}
　　cudaFree( deviceArray );

　　这里也没什么要说的，除了我们现在也从设备拷贝BLOCKS*BLOCKS外，我们也要遍历每个块，打印出它们的内容。

　　最后不要忘了main函数的结束括号“}”。

　　现在，我们添加内核：

　　__global__ void generateArray( int *hostArray )
　　{
　　 int ThreadIndex = blockIdx.x + blockIdx.y * gridDim.x;
　　 hostArray[ThreadIndex] = ThreadIndex;
　　}

　　看起来和前面的完全一样，blockId.x用于获得在x维度上工作的块，blockId.y用于获得在y维度上工作的块，girdDim是我们网格维度的最大长度，.x表示x轴，.y表示y轴，和创建块数组指定的数字是一样的。

　　如果运行这个例子，你将看到类似下图所示的结果：

▲图 2 运行结果

　　正如你看到的，有很多不同的情景使用多维索引更好，另一个例子是，当你处理二维图像时，你可以在相同的坐标系统为每个像素创建一个块，如下图所示。

　　最后，完整的代码如下：

　　#include <stdio.h>
　　#define BLOCKS 10
　　__global__ void generateArray( int *hostArray )
　　{
　　 int ThreadIndex = blockIdx.x + blockIdx.y * BLOCKS;
　　 hostArray[ThreadIndex] = ThreadIndex;
　　}
　　int main( void )
　　{
　　 int hostArray[BLOCKS][BLOCKS];
　　 int *deviceArray;
　　 dim3 multiBlockArray(BLOCKS,BLOCKS);
　　 cudaMalloc( (void**)&deviceArray, BLOCKS * BLOCKS * sizeof(int) );
　　 cudaMemcpy( deviceArray,
　　 hostArray, BLOCKS * BLOCKS * sizeof(int),
　　 cudaMemcpyHostToDevice );
　　 generateArray<<<multiBlockArray,1>>>( deviceArray );
　　 cudaMemcpy( hostArray,
　　 deviceArray,
　　 BLOCKS * BLOCKS * sizeof(int),
　　 cudaMemcpyDeviceToHost );
　　 for (int i=0; i<BLOCKS; i++)
　　 {
　　 printf( “Thread ID running: %d”, hostArray[0][i] );
　　 for (int j=1; j<BLOCKS; j++)
　　 {
　　 printf( ” %d”, hostArray[j][i] );
　　 }
　　 printf( “\n” );
　　 }
　　 cudaFree( deviceArray );
　　 return 0;
　　}