CUDA学习笔记1——线程组织、HelloWorld、数组相加、内存结构、核函数、设备函数

本文介绍了CUDA编程中的关键概念，如线程、块、网格、warp、GPU架构、内存类型（全局内存、共享内存、常量内存、纹理内存）、以及内存管理和数据同步。通过HelloWorld示例和数组相加的C++实现，展示了如何在CUDA中利用GPU进行高效计算。

1. CUDA线程组织与GPU架构

1.1 CUDA线程组织

1702220021616

一个核函数对应一个grid，grid里面有若干block，block里面有若干thread，32个thread称为一个warp。
下面这四段介绍copy自：理解CUDA中的thread,block,grid和warp - 知乎 (zhihu.com)

SM采用的SIMT(Single-Instruction, Multiple-Thread，单指令多线程)架构，warp(线程束)是最基本的执行单元，一个warp包含32个并行thread，这些thread以不同数据资源执行相同的指令。

当一个kernel被执行时，grid中的线程块被分配到SM上，一个线程块的thread只能在一个SM上调度，SM一般可以调度多个线程块，大量的thread可能被分到不同的SM上。每个thread拥有它自己的程序计数器和状态寄存器，并且用该线程自己的数据执行指令，这就是所谓的Single Instruction Multiple Thread(SIMT)。

一个CUDA core可以执行一个thread，一个SM的CUDA core会分成几个warp（即CUDA core在SM中分组)，由warp scheduler负责调度。尽管warp中的线程从同一程序地址，但可能具有不同的行为，比如分支结构，因为GPU规定warp中所有线程在同一周期执行相同的指令，warp发散会导致性能下降。一个SM同时并发的warp是有限的，因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器，所以SM的配置会影响其所支持的线程块和warp并发数量。

一个warp中的线程必然在同一个block中，如果block所含线程数目不是warp大小的整数倍，那么多出的那些thread所在的warp中，会剩余一些inactive的thread，也就是说，即使凑不够warp整数倍的thread，硬件也会为warp凑足，只不过那些thread是inactive状态，需要注意的是，即使这部分thread是inactive的，也会消耗SM资源。由于warp的大小一般为32，所以block所含的thread的大小一般要设置为32的倍数。

1.2 GPU架构

GPU

下图是一个GPU的架构图，里面有许多SM。

SM (Streaming MultiProcessor)

一个SM里面有两个SMP (SM Processing Block)。绿色的方块Core就是CUDA core；深绿色的方块LD/ST是内存操作load/store；要注意，CUDA core做的是单精度浮点数的运算，DP Unit做的是双精度浮点数的运算，二者数量比为2:1；绿色方块SFU指的是Special Function Unit，做一些特殊函数的计算，比如sin、cos等。

每个SM有自己的指令缓存、L1缓存、共享内存。每个SMP有自己的warp scheduler、Register File等。

CUDA core

1702257369899

GPU内存架构

Global Memory就是平常说的显存（GPU Memory），L1缓存指的就是共享内存（第五章会细说）。

每次shared memory (也就是L1)要去访问Global Memory时，就会先看看L2里面有没有，有的话，就不用去global memory了。

2. Hello World

GPU只是一个设备，要它工作的话，需要有一个主机来给它下达命令。这个主机就是CPU。

主机对设备的调用都是通过核函数（kernel function）来实现的。所以，一个简单的cuda程序的结构具有下面这种形式：

int main() {
   
   
	主机代码
	核函数的调用
	主机代码
	return 0;
}

核函数的前面必须有***__global__***来修饰，并且返回值类型必须是void，比如下面这样：

__global__ void hello_from_gpu() {
   
   
	printf("Hello World From the GPU!");
}

__global__ 和 void 的次序可以交换

上述核函数需要被主机调用才能发挥作用，如下：

#include<stdio.h>

__global__ void hello_from_gpu() {
   
   
	printf("Hello World From the GPU!");
}

int main() {
   
   
	hello_from_gpu<<<1, 1>>>();
	cudaDeviceSynchronize();
	return 0;
}

hello_from_gpu<<<1, 1>>>() : 主机在调用一个核函数时，需要指明给设备分配多少线程。第一个参数是grid_size，第二个参数是block_size。也就是说，第一个参数是block的数量，第二个参数是一个block中thread的数量。
cudaDeviceSynchronize() : 输出流先存放在缓冲区中，而缓冲区不会自动刷新，只有程序遇到某种同步操作时缓冲区才会刷新，这个函数可以同步主机与设备，进而能够促使缓冲区刷新。

3. 数组相加

#include<stdio.h>
#include<math.h>

const double a = 1.11;
const double b = 2.22;

__global__ void add(const double *x, const double *y, double *z);

int main