文章目录

通过前面博文的介绍我们已经对CUDA有一定了解了，相信大家还记得第一篇博文中，代码的矢量和简单实例，在这里我将对其进行丰富，并详细介绍其中的细节。

代码如下：

__global__ void add( int *a, int *b, int *c){
	int tid = blockIdx.x;
	if( tid < N )
		c[ tid ] = a[ tid ] + b[tid];
}

int  main (void) {
	int a[N], b[N],c[N];
	int *dev_a, *dev_b, *dev_c;
	
	cudaMalloc( (void **) &dev_a, N * sizeof(int));
	cudaMalloc( (void**) &dev_b, N* sizeof(int));
	cudaMalloc((void **) &dev_c, N *sizeof(int));
	
	for(int i = 0; i<N; i++ ) {
		a[i] = -i;
		b[i] = i * i;
	}
	
	cudaMemcpy( dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy( dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);
	
	add<<<N, 1>>>(dev_a, dev_b, dev_c);
	
	cudaMemcpy(c, dev_c , N * sizeof(int), cudaMemcpyDeviceToHost);
	
	cudaFree( dev_a );
	cudaFree( dev_b );
	cudaFree( dev_c );
	
	return 0;
}

代码说明：

1. main() 中的通用模式：

   • 调用cudaMalloc() 在设备上为三个数组分配内存；
   • 为避免内存泄露，在使用完GPU内存后通过cudaFree()释放它们。
   • 通过cudaMemcpy()将输入数据复制到设备中，同时指定参数cudaMemcpyHostToDevice;
   • 通过<<<…>>>语法，在主机代码main()中执行add()的设备代码

2. add()中的通用模式：

   • 编写一个在设备上执行的函数add()。使用C语言，并在函数名前添加一个__global__.
   • 启动方式：add<<<N,1>>>(dev_a, dev_b, dev_c)；
     • N表示将有N个线程块在GPU上运行，1表示每个线程块中只包含一个线程。
     • 即GPU将运行核函数的N个副本

3. 如何知道当前正在运行哪个线程块？

   • 注意核函数中, 语句 int tid = blockIdx.x;由于blockIdx为内置变量，表示当前执行设备代码的线程块索引。
   • blockIdx.x 表示线程块的第一个索引值，blockIdx.y分别表示第二，且最多两个块索引。（线程索引为三维）

4. 为什么需要判断tid是否<N ?

   • 核函数中假设tid总是小于N，可此实例中一旦tid>N 则会造成内存非法访问。
   • 任意长度的矢量相加，假设Add<<<128,128>>>

__global__ void add (int *a, int *b, int *c){
	int tid = threadIdx.x + blockIdx.x * blockDim.x;
	While(tid < N){
		c[tid] = a[tid] + b[tid];
		Tid +=blockDim.x * gridDim.x    //网格维度 * 块维度
	}
}

1. 实例中的cuda函数：

   cudaError_t cudaMalloc( void** devPtr，size_t count )
cudaError_t cudaFree (void* devPtr)
cudaError_t cudaMemcpy( void* dst，const void* src，size_t count，enum cudaMemcpyKind kind )