这是本人21年读书时学习CUDA基础知识保留的一些笔记，学习时的内容出处和图片来源不记得了，仅作为个人记录！
CUDA编程关键术语：
- host : cpu
- device : GPU
- host memory : cpu 内存
- device memory : gpu onboard显存
- kernels : 调用CPU上的在GPU执行的函数
- device function : 只能在GPU上执行的函数
安装Numba库：
$ conda install numba
检查一下CUDA和Numba是否安装成功
from numba import cuda print(cuda.gpus)
如果上述步骤没有问题，可以得到结果：<Managed Device 0>…
一般使用CUDA_VISIBLE_DEVICES这个环境变量来选择某张卡。如选择5号GPU卡运行你的程序。
CUDA_VISIBLE_DEVICES='5' python example.py
GPU计算流程
- 初始化，并将必要的数据拷贝到GPU设备的显存上
- CPU调用GPU函数，启动GPU多个核心同时进行计算。
- CPU与GPU异步计算。
- 将GPU计算结果拷贝回主机端，得到计算结果。
与传统CPU对比
from numba import cuda
def cpu_print():
print("print by cpu.")
@cuda.jit
def gpu_print():
# GPU核函数
print("print by gpu.")
def main():
gpu_print[1, 2]()
cuda.synchronize()
cpu_print()
if __name__ == "__main__":
main()
- 使用
from numba import cuda引入cuda库 - 在GPU函数上添加
@cuda.jit装饰符，表示该函数是一个在GPU设备上运行的函数，GPU函数又被称为核函数。 - 主函数调用GPU核函数时，需要添加如
[1, 2]这样的执行配置，这个配置是在告知GPU以多大的并行粒度同时进行计算。gpu_print[1, 2]()表示同时开启2个线程并行地执行gpu_print函数，函数将被并行地执行2次。下文会深入探讨如何设置执行配置。 - GPU核函数的启动方式是异步的：启动GPU函数后，CPU不会等待GPU函数执行完毕才执行下一行代码。必要时，需要调用
cuda.synchronize()，告知CPU等待GPU执行完核函数后，再进行CPU端后续计算。这个过程被称为同步，也就是GPU执行流程图中的红线部分。如果不调用cuda.synchronize()函数，执行结果也将改变，"print by cpu.将先被打印。虽然GPU函数在前，但是程序并没有等待GPU函数执行完，而是继续执行后面的cpu_print函数，由于CPU调用GPU有一定的延迟，反而后面的cpu_print先被执行，因此cpu_print的结果先被打印了出来。
Thread层次结构
-
CUDA将核函数所定义的运算称为线程（Thread），多个线程组成一个块（Block），多个块组成网格（Grid）。这样一个grid可以定义成千上万个线程，也就解决了并行执行上万次操作的问题。例如，把前面的程序改为并行执行8次：可以用2个block，每个block中有4个thread。原来的代码可以改为
gpu_print[2, 4]()，其中方括号中第一个数字表示整个grid有多少个block，方括号中第二个数字表示一个block有多少个thread。 -
实际上，线程（thread）是一个编程上的软件概念。从硬件来看，thread运行在一个CUDA核心上，多个thread组成的block运行在Streaming Multiprocessor（SM的概念详见本系列第一篇文章）多个block可以运行在一个SM上，但是一个额block只能运行在一个SM上，多个block组成的grid运行在一个GPU显卡上。
CUDA提供了一系列内置变量，以记录thread和block的大小及索引下标。以[2, 4]这样的配置为例：
blockDim.x变量表示block的大小是4，即每个block有4个thread，threadIdx.x变量是一个从0到blockDim.x - 1（4-1=3）的索引下标，记录这是第几个thread；`- gridDim.x`变量表示grid的大小是2，即每个grid有2个block，
blockIdx.x变量是一个从0到gridDim.x - 1（2-1=1）的索引下标，记录这是第几个block。
- 某个thread在整个grid中的位置编号为：
threadIdx.x + blockIdx.x * blockDim.x。
from numba import cuda
def cpu_print(N):
for i in range(0, N):
print(i)
@cuda.jit
def gpu_print(N):
idx = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x
if (idx < N):
print(idx)
def main():
print("gpu print:")
gpu_print[2, 4](8)
cuda.synchronize()
print("cpu print:")
cpu_print(8)
if __name__ == "__main__":
main() #结果相同
-
这里的GPU函数在每个CUDA thread中打印了当前thread的编号，起到了CPU函数for循环同样的作用。因为for循环中的计算内容互相不依赖，也就是说，某次循环只是专心做自己的事情，循环第i次不影响循环第j次的计算，所以这样互相不依赖的for循环非常适合放到CUDAthread里做并行计算。 在实际使用中，我们一般将CPU代码中互相不依赖的的for循环适当替换成CUDA代码。
-
这份代码打印了8个数字，核函数有一个参数N，N = 8，假如我们只想打印5个数字呢？当前的执行配置共2 * 4 = 8个线程，线程数8与要执行的次数5不匹配，不过我们已经在代码里写好了if (idx < N)的判断语句，判断会帮我们过滤不需要的计算。我们只需要把N = 5传递给gpu_print函数中就好，CUDA仍然会启动8个thread，但是大于等于N的thread不进行计算。注意，当线程数与计算次数不一致时，一定要使用这样的判断语句，以保证某个线程的计算不会影响其他线程的数据。
Block大小的设置
-
不同的执行配置会影响GPU程序的速度，一般需要多次调试才能找到较好的执行配置，在实际编程中，执行配置
[gridDim, blockDim]应参考下面的方法：- block运行在SM上，不同硬件架构（Turing、Volta、Pascal…）的CUDA核心数不同，一般需要根据当前硬件来设置block的大小
blockDim（执行配置中第二个参数）。一个block中的thread数最好是32、128、256的倍数。注意，限于当前硬件的设计，block大小不能超过1024。 - grid的大小
gridDim（执行配置中第一个参数），即一个grid中block的个数可以由总次数N除以blockDim，并向上取整。
例如，我们想并行启动1000个thread，可以将blockDim设置为128，1000 ÷ 128 = 7.8，向上取整为8。使用时，执行配置可以写成gpuWork[8, 128]()，CUDA共启动8 * 128 = 1024个thread，实际计算时只使用前1000个thread，多余的24个thread不进行计算。
- block运行在SM上，不同硬件架构（Turing、Volta、Pascal…）的CUDA核心数不同，一般需要根据当前硬件来设置block的大小
注意，这几个变量比较容易混淆，再次明确一下：blockDim是block中thread的个数(如果是二维blockDim.x是列数，blockDim.y是行数)，一个block中的threadIdx最大不超过blockDim；gridDim是grid中block的个数，一个grid中的blockIdx最大不超过gridDim。
以上讨论中，block和grid大小均是一维，实际编程使用的执行配置常常更复杂，block和grid的大小可以设置为二维甚至三维，如下图所示。
- 例如Thread(2,0)的位置计算，threadIdx.x = 2,threadIdx.y=0,blockIdx.x = 1,blockIdx.y = 1,
blockDim.x = 4,blockDim.y = 3，计算得该线程在grid中所有线程的索引为
Thread_x = blockIdx.x* blockDim.x+threadIdx.x = 6
Thread_y = blockIdx.y* blockDim.y+threadIdx.y = 3
内存分配
cuda.device_array()：在设备上分配一个空向量，类似于numpy.empty()cuda.to_device()：将主机的数据拷贝到设备
ary = np.arange(10)
device_ary = cuda.to_device(ary)
cuda.copy_to_host()：将设备的数据拷贝回主机
host_ary = device_ary.copy_to_host()
- Globel Memory : 最常用，传输到gpu的数据基本都在这里（显存）所有thread共享（相对于其他memory相对慢）
- Constant Memory：常量内存
- Texture Memory：纹理内存
- Shared Memory ：block中所有thread共享
- local Memory ：thread独自使用
- registers : 寄存器 thread独自使用