cuda编程模型

host和device：
- host：即CPU，CPU所关联的内存就叫host memory
- device：即GPU，GPU内的内存就叫device memory
- 运行CUDA程序主要有三步：1）host-to-device transfer：将数据从host memory拷到device memory；2）加载GPU程序并执行，将数据缓存到on-chip上读取更快；3）device-to-host transfer：将device memory中的结果拷到host memory
【逻辑上】CUDA kernel：就是在GPU上执行的函数
- start with a __global__ declaration specifier
- a kernel is executed as a grid of blocks of threads
【物理上】一个CUDA block在一个SM（streaming multiprocessor）上执行；一个SM可以运行多个并发的CUDA block；一个kernel在一个deivce上执行，一个device上可以同时运行多个kernels。如下图是逻辑和硬件资源的映射关系：
- 其中，SM的基本执单元是包含32个线程的线程数，所以block大小一般设置为32的倍数
- 每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管一个warp中的线程同时从同一程序地址执行，但可能具有不同的行为。比如遇到了分支结构，一些线程可能进入这个分支，但另外一些可能并不用执行，只能进入死等（等待warp内其他未执行完的线程执行完毕），因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。
CUDA中的index问题
- 几个built-in变量：gridDim , blockDim用于指示自己的维度；blockIdx, threadIdx用于指示自己在上一层中是第几个。如下图的例子中：
  - gridDim.x=3（表示grid的x维度有3个blocks），gridDim.y=2（grid的y维度有2个blocks）
  - blockDim.x=4（表示block的x维度有4个threads），blockDim.y=3（表示block的y维度有3个threads）
  - blockIdx.x=0, blockIdx.y=1 表示的是 block(0, 1)
  - threadIdx.x=2, threadIdx.y=1表示的是 thread(2, 1)
- thread indexing: 为每个thread分配一个唯一的id
  - 1D grid of 1D blocks: threadId = blockIdx.x * blockDim.x + threadIdx.x
  - 1D grid of 2D blocks: threadId = blockIdx.x * blockDim.x * blockDim * y + threadIdx.y * blockDim.x + threadIdx.x;
  - 2D grid of 1D blocks: threadId = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x
    - blockId = blockIdx.y * gridDim.x + blockIdx.x
    - threadId = blockId * blockDim.x + threadIdx.x
  - 2D grid of 2D blocks: threadId = blockIdx.y * gridDim.x * blockDim.x * blockDim.y + blockIdx.x * blockDim.x * blockDim.y + theadIdx.y * blockDim.x + threadIdx.x
  - 总结：二维下，[小Id]就是[小Idx.y]*[大Dim.x]+[小Idx.x]，即若位于第n行，前面就有n*上一级列数个，然后再加上是第几列（因为索引值从零开始，所以直接加即可）
GPU存储

上图中：
L1/SMEM：指L1 cache/Shared memory，是每个SM独有的。Shared memory可以由用户写代码进行数据的读写控制，L1则不行；
Read only：只读缓存；
L2 Cache：所有SM都可以访问
Global Memory：全局内存，是所有线程都能访问的内存，也是和CPU内存进行数据传递的地方。通常说的显存就是global memory

每个SM都拥有自己的shared memory，而这些SM们都位于同一块芯片上（on the same chip），这块芯片通过PCB电路连接内存芯片（DRAM）
SP（cuda core、流处理器，一个thread占用一个SP）对shared memory的访问属于片上访问，可以立刻获得数据
SP对内存芯片（DRAM）的访问需要通过请求内存控制器等一系列操作，然后才能得到数据。

下图是一个速度问题：其中s、t、u是本地内存（local memory）中的变量，a、b、c是shared memory中的变量，所以t=s最快
__syncthreads()：
- 确保这行代码之前，同一个block内的所有线程都完成了各自的工作（如将数据从全局内存加载到了共享内存中）。只同步一个线程块中的线程，其他线程块不受影响