CUDA 线程 分层执行

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>

//Defining number of elements in Array
#define N	50000
//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) 
{
	//Getting block index of current kernel

	int tid = threadIdx.x + blockIdx.x * blockDim.x;	
	while (tid < N)
	{
		d_c[tid] = d_a[tid] + d_b[tid];
		tid += blockDim.x * gridDim.x;
	}		
}

int main() 
{
	//Defining host arrays
	int h_a[N], h_b[N], h_c[N];
	//Defining device pointers
	int *d_a, *d_b, *d_c;
	// allocate the memory
	cudaMalloc((void**)&d_a, N * sizeof(int));
	cudaMalloc((void**)&d_b, N * sizeof(int));
	cudaMalloc((void**)&d_c, N * sizeof(int));
	//Initializing Arrays
	for (int i = 0; i < N; i++) 
	{
		h_a[i] = 2 * i*i;
		h_b[i] = i;
	}
	// Copy input arrays from host to device memory
	cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
	//Calling kernels with N blocks and one thread per block, passing device pointers as parameters
	gpuAdd << <512, 512 >> >(d_a, d_b, d_c);
	//Copy result back to host memory from device memory
	cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
	cudaDeviceSynchronize();
	int Correct = 1;
	printf("Vector addition on GPU \n");
	//Printing result on console
	for (int i = 0; i < N; i++) 
	{
		if ((h_a[i] + h_b[i] != h_c[i]))
		{
			Correct = 0;
		}
	}
	if (Correct == 1)
	{
		printf("GPU has computed Sum Correctly\n");
	}
	else
	{
		printf("There is an Error in GPU Computation\n");
	}
	//Free up memory
	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);
	return 0;
}

本内核的代码和上一章我们写过的那个很相似。但是有两处不同:一处是计算初始的tid的时候，另一处则是添加了while循环部分。计算初始的tid的变化，是因为我们现在是启动多个块，每个里面有多个线程，直接看成ID的结构，多个块横排排列，每个块里面有N个线程，那么自然计算tid的时候是用当前块的ID 当前块里面的线程数量＋当前线程在块中的ID，即 tid=blockldx.x(当前块的ID) * blockDim.x (当前块里面的线程数量)+threadIdx.x （当前线程在块中的ID)。而 while部分每次增加现有的线程数量（因为你没有启动到N)，直到达到N。这就如同你有一个卡，一次最多只能启动100个块，每个块里有7个线程，也就是一次最多能启动700个线程。但N的规模是8 000，远远超过700怎么办?答案是直接启动K个(K≥700 )，这样就能安全启动。然后里面添加一个while循环,这700个线程第一次处理[0,699)，第二次处理[700,1 400)，第三次处理[1 400,2 100)……直到这8 000个元素都被处理完。这就是我们本例中看到的代码。初始化时候的tid =threadIdx.x+blockDim.x * blockldx.x，每次 while循环的时候tid += blockDim.xgridDim.x。