滚动新闻:
首页 > 客户服务 > 技术支持

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

  到现在为止,我们还没有真正触摸到了并行编程,这篇文章就是为此而写的。

  ▲图 1 并行中的内核

  回想一下我们之前在设备上使用kernelFunction<<<1,1>>>(..)执行一个函数的代码,我在那里还曾说过后面会细说,本文就详细介绍一下参数N1,<<>>,这里就是并行魔法发生地。

  N1是我们想并行运行的块数,如果我们调用kernelFunction<<<5,1>>>(..),这个函数将分成5个副本并行运行,每个副本称为一个块。

  接下来我们必须要做的事情是,使用一个索引让每个副本为解决方案的不同部分工作,如果所有线程做完全一样的事情,就没有必要并行计算了,幸运的是,CUDA内置了一个变量blockIdx可以用来跟踪每个块的运行。

  blockIdx是一个2D变量,包含x和y,你可以使用x或同时使用x和y,这取决于我们要解决什么问题,一个简单的例子是同时使用x和y处理2D图像,为x和y轴上的每个像素产生一个线程,你也可以只使用x,这里没有什么指导原则。

  现在,我们通过检查blockIdx.x知道线程运行的id,并且知道如何并行运行内核,让我们创建一个简单的例子吧。

  在这个例子中,我们将创建一个应用程序,完全以并行内核生成一个数组,这个数组将包含每个运行的线程的threadID,当线程结束后,我们使用printf将结果打印出来。

  实现内核

  我们从查看内核代码开始:

  __global__ void generateArray( int *hostArray )

  {

  int ThreadIndex = blockIdx.x;

  hostArray[ThreadIndex] = ThreadIndex;

  }

  我们在前面的文章已经说过,__global__会告诉主机这个函数将在设备上运行,它产生一个数组作为输出,将blockIdx.x存储在一个名为ThreadIndex的变量中,然后将这个值放在数组中正确的位置。

  blockIdx将为每个运行的块产生一个ID,从0开始,这样它将成为数组完美的索引。

  实现main()函数

  接下来,我们要看看我们的main函数,这里应该没有什么新东西:

  int main( void )

  {

  int hostArray[BLOCKS];

  int *deviceArray;

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

  cudaMemcpy( deviceArray,

  hostArray, BLOCKS * sizeof(int),

  cudaMemcpyHostToDevice );

  generateArray<<

  cudaMemcpy( hostArray,

  deviceArray,

  BLOCKS * sizeof(int),

  cudaMemcpyDeviceToHost );

  for (int i=0; i

  {

  printf( Thread ID running: %dn, hostArray[i] );

  }

  cudaFree( deviceArray );

  return 0;

  }

  首先,我们按BLOCKS大小创建一个数组,在设备上未数组分配空间,并调用:

  generateArray<<

  这个函数将在BLOCKS并行内核中运行,在一个调用中创建好全部数组。

  这个操作完成后,我们将结果从设备拷贝到主机,并将它打印在屏幕上,释放数组,最后退出。

  整个应用程序的源代码如下:

  #include

  #define BLOCKS 25

  __global__ void generateArray( int *hostArray )

  {

  int ThreadIndex = blockIdx.x;

  hostArray[ThreadIndex] = ThreadIndex;

  }

  int main( void )

  {

  int hostArray[BLOCKS];

  int *deviceArray;

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

  cudaMemcpy( deviceArray,

  hostArray, BLOCKS * sizeof(int),

  cudaMemcpyHostToDevice );

  generateArray<<

  cudaMemcpy( hostArray,

  deviceArray,

  BLOCKS * sizeof(int),

  cudaMemcpyDeviceToHost );

  for (int i=0; i

  {

  printf( Thread ID running: %dn, hostArray[i] );

  }

  cudaFree( deviceArray );

  return 0;

  }

  现在编译并运行这段代码,你将会看到像下面这样的输出:

  ▲图 2 程序运行输出结果

 

  恭喜,你已经使用CUDA成功创建了你的第一个并行应用程序!