CUDA学习，第一个kernel函数及代码讲解

前一篇CUDA学习，我们已经完成了编程环境的配置，现在我们继续深入去了解CUDA编程。本博文分为三个部分，第一部分给出一个代码示例，第二部分对代码进行讲解，第三部分根据这个例子介绍如何部署和发起一个kernel函数。

一、代码示例

#include< stdio.h>  
#include "cuda_runtime.h"  
#include "device_launch_parameters.h"  
#include 
   
   
    
    
#include 
    
    
     
     

// Simple utility function to check for CUDA runtime errors
void checkCUDAError(const char *msg);

// Part 3 of 5: implement the kernel
__global__ void myFirstKernel(int *d_a)
{
	int i=blockIdx.x*blockDim.x+threadIdx.x;
	d_a[i]=1000 * blockIdx.x + threadIdx.x;
}


// Program main

int main( int argc, char** argv) 
{
    // pointer for host memory
    int *h_a;

    // pointer for device memory
    int *d_a;

    // define grid and block size
    int numBlocks = 8;
    int numThreadsPerBlock = 8;

    // Part 1 of 5: allocate host and device memory
    size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
    h_a = (int *) malloc(memSize);
    cudaMalloc((void**)&d_a,memSize);

    // Part 2 of 5: configure and launch kernel
    dim3 dimGrid(numBlocks,1,1);
    dim3 dimBlock(numThreadsPerBlock,1,1);
    myFirstKernel<<
     
     
      
      >>(d_a);

    // block until the device has completed
    cudaThreadSynchronize();

    // check if kernel execution generated an error
    checkCUDAError("kernel execution");

    // Part 4 of 5: device to host copy
    cudaMemcpy(h_a,d_a,memSize,cudaMemcpyDeviceToHost);

    // Check for any CUDA errors
    checkCUDAError("cudaMemcpy");

    // Part 5 of 5: verify the data returned to the host is correct
    for (int i = 0; i 
      
      
     
     
    
    
   
   

二、代码解说

void checkCUDAError(const char *msg);

申明一个函数，用于检测CUDA运行中是否出错。

__global__ void myFirstKernel(int *d_a)
{
	int i=blockIdx.x*blockDim.x+threadIdx.x;
	d_a[i]=1000 * blockIdx.x + threadIdx.x;
}

kernel函数，blockIdx.x表示block在x方向的索引号，blockDim.x表示block在x方向的维度，threadIdx.x表示thread在x方向的索引号。

这里也许你会问，为什么在x方向？难道还有其他方向？

对的，grid可以是一维、二维，block可以是一维、二维和三维的。一个grid里包含多个block,一个block里也包含多个thread，可参考下图。

从而，i是每个thread的索引号，也许你会问为什么不直接是threadIdx.x呢？

因为每个block里的threadIdx.x都是从0到blockDim.x（假设block是一维的），那么不同的block里threadIdx.x会出现相同的值，我们就不知道该调用那个thread来执行，因为用threadIdx.x+blockIdx.x*blockDim.x表示每个thread的索引号，这样就是唯一的了。举个例子，假设grid和block都是一维的，blockDim.x=8，threadIdx.x从0到7，blockIdx.x也是从0到7，那么i就是0,1,2,3,4,5,6,7,8,9,10,11...63，从而保证了每个thread索引号的唯一性。

// pointer for host memory
    int *h_a;

    // pointer for device memory
    int *d_a;

    // define grid and block size
    int numBlocks = 8;
    int numThreadsPerBlock = 8;

申请两个指针*h_a和*d_a，分别指向host内存和device内存，host是指主程序cpu内存，device是指gpu内存（global memory）。并定义一个grid里block数（即numBlocks）和每个block里thread数numThreadsPerBlock。

size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
h_a = (int *) malloc(memSize);
cudaMalloc((void**)&d_a,memSize);

分配host和device内存，cudaMalloc((void**)&d_a,memSize)是给device中d_a分配memSize字节的内存，为什么前面有(void**)&呢？因为这个内存分配要通过cpu传给gpu。

dim3 dimGrid(numBlocks,1,1);
dim3 dimBlock(numThreadsPerBlock,1,1);
myFirstKernel<<
    
    
     
     >>(d_a);

    
    

部署并发起kernel函数，kernel函数是跑在gpu上的那段程序，即我们之前申明的myFirstKernel函数。部署：dimGrid(numBlocks,1,1)，一维的grid,里面包含numBlocks个block；dimBlock(numThreadsPerBlock,1,1)，grid里每个block都是一维的，每个block有numThreadsPerBlock个thread。发起kernel函数：myFirstKernel<<<dimGrid,dimBlock>>>(d_a)，第一个是函数名，括号里是kernel的部署，后面一个是函数的参数。

cudaThreadSynchronize();

checkCUDAError("kernel execution");

同步并检查kernel函数运行是否出错。为什么要同步呢？因为每个thread运行的时间是不一样的，只有等所有线程都跑完了，我们才做下一件事。这会造成运行的性能降低，但是是必要的。

 cudaMemcpy(h_a,d_a,memSize,cudaMemcpyDeviceToHost);

cudaMemcpy函数完成数据在host和device之间的传输，第一个参数是传输的目标，第二个参数是传输的源数据，第三个参数传输的数据量，第四个参数是传输的方向，这里是从device传到host。

 checkCUDAError("cudaMemcpy");

检查cudaMemcpy函数运行是否出错。

for (int i = 0; i 
    
    

检查从device传回的数据是否正确。

cudaFree(d_a);

free(h_a);

释放host和device内存。

void checkCUDAError(const char *msg)
{
    cudaError_t err = cudaGetLastError();
    if( cudaSuccess != err) 
    {
        fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
        exit(-1);
    }                         
}

检查CUDA函数是否运行正确的函数。

三、部署和发起一个kernel函数

如上述，对kernel函数先申明，在函数体中实现线程的算法，即：

__global__ void myFirstKernel(int *d_a)
{
	int i=blockIdx.x*blockDim.x+threadIdx.x;
	d_a[i]=1000 * blockIdx.x + threadIdx.x;
}

然后部署和发起kernel函数，即：

    // define grid and block size
    int numBlocks = 8;
    int numThreadsPerBlock = 8;

    // Part 1 of 5: allocate host and device memory
    size_t memSize = numBlocks * numThreadsPerBlock * sizeof(int);
    h_a = (int *) malloc(memSize);
    cudaMalloc((void**)&d_a,memSize);

    // Part 2 of 5: configure and launch kernel
    dim3 dimGrid(numBlocks,1,1);
    dim3 dimBlock(numThreadsPerBlock,1,1);
    myFirstKernel<<
     
     
      
      >>(d_a);

     
     

综上，我们完成代码的讲解，部署和发起一个kernel函数。想必你现在应该对CUDA有了较为深入的了解了！是吧？