CUDA学习第三天：Kernel+grid+block关系

1. 理一理前两天学到的概念之间的关系

• CUDA && GPU

  CUDA: NIVID的CPUs上的一个通用并行计算平台和编程模型；
  GPU: CPU+GPU的异构计算架构，CPU所在位置为主机端(host), 而GPU所在位置为设备端(device).

• SP && SM

  SP: 最基本的处理单元，streaming processor，也称为CUDA core。最后具体的指令和任务都是在SP上处理的。GPU进行并行计算，也就是很多个SP同时做处理。

  SM: 多个SP加上其他的一些资源组成一个streaming multiprocessor。也叫GPU大核，其他资源如：warp scheduler，register，shared memory等。SM可以看做GPU的心脏（对比CPU核心），register和shared memory是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的threads。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。
  

  SP是线程执行的硬件单位，SM中包含多个SP，一个GPU可以有多个SM（比如16个），最终一个GPU可能包含有上千个SP。

• Host && Device

  Host与Device的内存交互模型：
  

  明确Host与Device分别位于什么位置：
  

2. Kernel与网格（grid）与线程块（block）与线程（Thread）

在实际编写CUDA程序时，我们需要设计block等的大小，也就是合理安排线程的数量。

3. CUDA实例程序二

• 程序文件结构

  sample_01_plus.cu (cpp文件无法运行)

• CMakeLists添加执行文件

  cuda_add_executable(sample_01 src/sample_01_plus.cu)

• 程序源代码

		#include <iostream>
		#include <cuda.h>
		#include <cuda_runtime.h>
		#include <cuda_runtime_api.h>
		#include <device_launch_parameters.h>

		__global__ void addCuda(float* x, float* y, float* z, int n)
		{
			int index = threadIdx.x+blockIdx.x*blockDim.x;
	
			int stride =  blockDim.x*gridDim.x;
	
			for(int i=index; i<n; i+=stride)
			{
				z[i] = x[i] + y[i];
			}
		}

		int main()
		{
			int N =20;
	
			int nBytes = N*sizeof(float);
	
			float *x, *y, *z;
		/**********************************************
		 *  在device上分配内存函数：
		 * 
		 *  cudaError_t cudaMalloc(void** devPtr, size_t size);
		 * 
		 * 	devPtr: 指向所分配内存的指针，
		 * 
		 * *****************************************************************/
	
			x = (float*)malloc(nBytes);
			y = (float*)malloc(nBytes);
			z = (float*)malloc(nBytes);
	
			for(int i=0; i<N; i++)
			{
				x[i] = 10.0;
				y[i] = 20.0;
			}
	
			float *dx, *dy, *dz;
		/**********************************************
		 *  负责host与device之间数据通讯的cudaMemcpy函数：
		 * 
		 *  cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind)
		 * 
		 * 	其中src指向数据源，而dst是目标区域，count是复制的字节数，其中kind控制复制的方向：
		 * 
		 * 	cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice.
		 * 	如cudaMemcpyHostToDevice将host上数据拷贝到device上
		 * 
		 * *****************************************************************/
	
			cudaMalloc((void**)&dx, nBytes);
			cudaMalloc((void**)&dy, nBytes);
			cudaMalloc((void**)&dz, nBytes);
	
			// x-y为原始数据， 将其复制到gpu上计算；
	
			cudaMemcpy((void*)dx, (void*)x, nBytes, cudaMemcpyHostToDevice);
			cudaMemcpy((void*)dy, (void*)y, nBytes, cudaMemcpyHostToDevice);
	
			dim3 blockSize(256);
			dim3 gridsize((N + blockSize.x - 1) / blockSize.x);
	
			// 在gpu上计算的结果保存在dz上；
			addCuda<<<gridsize, blockSize>>>(dx, dy, dz, N);
	
			// dz为得到的结果，将其复制出来放到z上， 然后检测z与实际值的差距。
	
			cudaMemcpy((void*)z, (void*)dz, nBytes, cudaMemcpyDeviceToHost);
	
			float maxError = 0.0;
		    for (int i = 0; i < N; i++)
			maxError = fmax(maxError, fabs(z[i] - 30.0));
		    std::cout << "最大误差: " << maxError << std::endl;
	
	
			cudaFree(dx);
			cudaFree(dy);
			cudaFree(dz);
	
			free(x);
			free(y);
			free(z);
	
			return 0;
		}