CUDA学习第一天： 基础概念扫盲

CUDA基础学习开始，以下内容是来自于各个博客和官方文档，如果我写不够全面，大家可以去往相应的参考链接进行学习。

参考链接：https://zhuanlan.zhihu.com/p/34587739
参考链接：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#application-compatibility
参考链接：https://www.cnblogs.com/skyfsm/p/9673960.html

CUDA编程模型是一个异构模型，需要CPU和GPU协同工作。

在CUDA中，host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。

CUDA程序中既包含host程序， 又包含device程序，他们分别在CPU和GPU上运行，同时， host与device之间可以进行通信，这样他们之间可以进行数据拷贝。

• 典型的CUDA程序的执行流程

  1.分配host内存，并进行数据初始化；---- CPU初始化

  2.分配device内存，并从host将数据拷贝到device上； — GPU初始化

  3.调用CUDA的核函数在device上完成指定的运算；----GPU并行运算（核心)

  4.将device上的运算结果拷贝到host上；— 将GPU结果传回CPU

  5.释放device和host上分配的内存。---- 初始化清空

• CUDA中的核函数 — Kernel

  1.kernel是在device上线程中并行执行的函数，核函数用__global__符号声明，在调用时需要用<<<grid, block>>>来指定kernel要执行的线程数量;

  2.在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。

• CPU 与 GPU的结构不一， 所以调用时需要区分host 与device上的代码

  在CUDA中是通过函数类型限定词开区别host和device上的函数，主要的三个函数类型限定词如下：

  • global：在device上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须是void，不支持可变参数参数，不能成为类成员函数。注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。

  • device：在device上执行，单仅可以从device中调用，不可以和__global__同时用;

  • host：在host上执行，仅可以从host上调用, 一般省略不写, 不可以和__global__同时用，但可和__device__，此时函数会在device和host都编译;

• 深刻理解Kernel

  • 网格 – grid

  GPU上有很多并行化的轻量级线程。kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次。

  • 线程块 – block

  网格可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次。

  • 数据类型定义 — dim3

  dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1; 对于图中结构（主要水平方向为x轴）。

定义的grid和block示例如下图所示，kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。

	dim3 grid(3, 2);
	dim3 block(5, 3);
	kernel_fun<<< grid, block >>>(prams...);

所以，一个线程需要两个内置的坐标变量**（blockIdx，threadIdx）**来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置。

• 线程块上的线程变量

  一个线程块上的线程是放在同一个流式多处理器（SM)上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数可达1024个。

  有时候，我们要知道一个线程在blcok中的全局ID，此时就必须还要知道block的组织结构。

  • 内置变量 – blockDim ： 获取线程块各个维度的大小

  对于一个2-dim的block（Dx, Dy）, 线程（x, y）的ID值为：（x+y*Dx）;

  对于一个3-dim的block（Dx, Dy, Dz）, 线程（x, y, z）的ID值为：（x+yDx+ zDx*Dy）;

  • 内置变量 – gridDim： 获取网格块各个维度的大小

• CUDA编程示例1

kernel的这种线程组织结构天然适合vector,matrix等运算，如我们将利用上图2-dim结构实现两个矩阵的加法，每个线程负责处理每个位置的两个元素相加，代码如下所示。线程块大小为(16, 16)，然后将N*N大小的矩阵均分为不同的线程块来执行加法运算。

	// Kernel定义
	__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) 
	{ 
	    int i = blockIdx.x * blockDim.x + threadIdx.x; 
	    int j = blockIdx.y * blockDim.y + threadIdx.y; 
	    if (i < N && j < N) 
		C[i][j] = A[i][j] + B[i][j]; 
	}
	int main() 
	{ 
	    ...
	    // Kernel 线程配置
	    dim3 threadsPerBlock(16, 16); 
	    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
	    // kernel调用
	    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); 
	    ...
	}