理论

• 并行和串行：是否有执行次序
• 一个内核启动一个网格，一个网格由多个线程块组成，同一线程块内包含多个线程，同一网格中共享相同的全局内存空间，同一线程块内的线程可以通过同步和共享内存实现协作，不同线程块内的线程不能协作。
• CPU是MIMD，而GPU是SIMT
• SM（流多处理器）相当于一个CPU核，但每次同时计算多个任务
• 任务执行：锁步思想lock-step，一个SM中的N个SP核，每个SP输入不同的数据但都执行相同的指令，即SPMD。SM以线程束为单位进行调度（半个线程束），当等待内存时被挂起，切换到另外一个线程束(warp)执行
  • 因为GPU有多个寄存器组，所以上下文切换代价小
  • 所以为了避免访问内存时SM闲置，每个线程块的线程数量要尽量多，避免小的线程块，但由于线程块只有当所有线程束执行完毕后，线程块才会从SM中撤走，线程越多，等待个别执行较慢的线程束的可能性也会增加，导致SM又因此被闲置
  • 因为是线程束为单位调度，所以线程块的线程数量尽量要是线程束的整数倍，否则最后一个线程束会有浪费
  • 不同架构，每个SM最大同时可以处理的线程块数量不同，如Ampere架构中计算力为8.6的RTX3060，最大支持16个线程块
  • 线程束大小(wrapSize)目前都是32
• 访存方式：SM会尽量合并同一个线程束的访存请求
  • 数据和处理器要保持紧密联系，数组的宽度最好是线程束大小的整数倍（可以补齐），线程块映射到数据需要注意内存访问连续
  • 尽量通过对齐且连续合并的访问方式，将内存请求合并，利用cudaMallocPitch对齐，避免数据结构在内存上交错分布，比如不要使用struct->a而是直接访问a[],b[],c[]这样分离的数组（这与CPU相反，因为CPU会缓存结构体所在的一行数据，所以使用struct->a更快）
  • 共享内存（用户可控的L1缓存）比全局内存更靠近SM，所以更快（1.5TB/s:190GB/s），可以将数据暂时读写在共享内存中，最后再一起合并写入到全局内存中，高低端的显卡中，全局内存速度差异极大，但共享内存的速度却差不多
  • 可以利用寄存器加速经常访问的变量（速度是共享内存的10倍），但是SM能够调度的最大线程块数量也受制于寄存器使用情况
  • 尽量避免对同一个全局变量的读写，同样的，共享内存要注意不要同时访问同一个存储体，因为存储体一个周期只能响应一个请求，可以按列访问共享内存来避免竞争
  • 常量内存有一个比较大的高速缓存，并且可以通过广播机制单周期内迅速分配到每个线程中
  • GPU加速的本质：通过多线程并行隐藏存储延迟和指令执行带来的延迟，能够相较于CPU更轻量级地切换上下文（更多的寄存器组），当并行足以掩盖存储延迟达到一个临界点后，可以考虑等大的存储事务，或者单个线程处理多个元素(ILP)
• 分支执行：GPU会执行所有分支，满足当前分支条件的线程会被激活，不满足的不会被激活
• 调度方式：线程块的分配是完全随机的，因此造成运算顺序的不确定性，导致浮点数运算结果有误

代码

• <<grid, block>>分别设置网格和线程块布局，blockIdx, threadIdx分别对应线程块和线程的索引

(blockIdx.x * blockDim.x) + threadIdx.x // idx
(blockIdx.y * blockDim.y) + threadIdx.y // idy
(gridDim.x * blockDim.x) * idy + idx // unique_id

• cudaMemcpy是隐式同步的，cudaDeviceSynchronize阻塞主机直到设备端全部结束
• cudaMallocPitch 分配对齐的内存
• __device只能从设备端调用, __host__只能从主机端调用, __global__都行
• _syncthreads同步所有线程，阻塞线程至线程块内所有线程都到达同步点
• __shared__ 共享内存， __constant__ 常量内存
• atomicAdd 原子加，注意最坏情况下，有可能会因为竞争导致程序完全串行化; atomicMin 原子取最小

工具

• 查看计算能力：nvidia-smi --query-gpu=compute_cap --format=csv
• 克隆仓库cuda-samples，并且构建运行Samples/1_Utilities/deviceQuery，即可获得硬件的详细信息

Detected 1 CUDA Capable device(s)

Device 0: "NVIDIA GeForce RTX 3060 Laptop GPU"
  CUDA Driver Version / Runtime Version          12.2 / 12.5
  CUDA Capability Major/Minor version number:    8.6
  Total amount of global memory:                 5938 MBytes (6226378752 bytes)
  (030) Multiprocessors, (128) CUDA Cores/MP:    3840 CUDA Cores
  GPU Max Clock rate:                            1425 MHz (1.42 GHz)
  Memory Clock rate:                             7001 Mhz
  Memory Bus Width:                              192-bit
  L2 Cache Size:                                 3145728 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total shared memory per multiprocessor:        102400 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  1536
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Managed Memory:                Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 12.2, CUDA Runtime Version = 12.5, NumDevs = 1
Result = PASS

• 性能分析
  • NsightCompute操作教程
  • 报错Cannot mix incompatible Qt library (5.15.3) with this library (5.15.2)：解决方法