CUDA-线程同步-线程调度

线程同步（一个block内的同步）

一个block内的所有threads有时候是需要同步的（如使用shared memory优化的矩阵相乘中），方法是在kernel函数中的适当位置加上__syncthreads()。

当一个thread执行到__syncthreadas()时，这个thread会看它所在的block内的其他所有threads情况，如果发现还有其他threads没有执行到这个位置，则这个thread等待其他threads。直到block中所有active的threads都执行到此，接着向下执行。

注意这个同步只是这个block内的threads同步。而非是全局同步，CUDA中没有全局通过不的原因是，全局同步的系统开销会很大。

并行系统中负载均衡：要求线程的执行时间尽量接近。如果不均衡，在需要线程同步时，所有线程会等待最慢的那个thread，此时整体的速度就时最慢的那个thread的速度，其他速度快的threads 的优势完全没有了。

block与block之间是异步的，不存在相互等待（当）。而同一个warp内的threads天然同步。

注意线程同步，可能导致死锁，比如下情况：
1
2
3
4
5
6
if (func()){
__syncthreads();
}
else{
__syncthreads();
}
产生死锁，执行不同分支的threads相互等待，谁也等不到谁。
block间的同步

block与block之间是异步的，当前kernel执行结束后，block之间自然同步了。言外之意是kernel的执行时间是由最慢的哪个block决定，所以，上文强调负载均衡。

线程调度
1. SP和threads
  
  为什么GPU的threads数量远远多于物理执行单元（SP）。因为每个SM中与CPU一样也含有上下文空间，用于执行上下文切换，从而实现多线程。（？提升吞吐量？）
  
  这里有个SP和threads的关系。以GPU G80为例：
  
  G80 的硬件信息：16个SM，每个SM含有8个SP，（共有16x8=128个SP），每个SM最多驻扎768个threads，总共同时执行12288个threads。（所以可以通过每个SM中最多可以驻扎的threads数，除以每个SM中的SP数，就得到了）
  
  解释：
  - 16，表是芯片实际含有16个SM
  - 8，表示每个SM含有8个SP（Streaming Processors），真正执行指令的工人。
  - 12288，表示这个芯片上可以同时有12288个threads进行调度。调度不意味着一定要实际执行。
  - 128个SP，表示每个时钟周期内实际并行执行的指令流为128个。但总共有12288个指令流间不停的切换。切换的目的是达到延时隐藏效果。
    
    warp的调度是零开销的。因为warp的上下文是存在与物理空间中的，需要了这个warp干活时，程序切换到这个warp上去即可。
    
    warp中的所有threads执行相同的指令。当有分支时，由于warp内threads天然的同步，所以含有分支时的执行会有性能下降。
2. warp和SP
  
  每个warp含有32个threads，假如每个SM只有8个SP，此时一个SM如何调度一个warp?
  - 第一个周期内，有8个threads进入SP
  - 第二，三，四个周期SP各进入8个threads
  - 如此循环，直到所有指令执行完毕
  - 所以此情况一个SM调度一个warp，需要4个周期（4个周期只是调度完成，指令执行完成需要4的倍数个周期）
    
    现代GPU的SP数已经远远大于32了。所以不存在上述问题了。
3. 调度warp实现延时隐藏
  
  一个实例：有一个kernel含有
  - 一个对global memory的读操作，这个操作耗时200个时钟周期。
  - 4个独立的乘或加操作，一个乘或加操作耗时4个时钟周期。
    
    那么需要多少个warp才可以将对global memory访问的延迟隐藏掉?
    
    首先每个warp需要执行4个独立的乘或加操作，共耗时4x4=16个始终周期。要覆盖200个时钟周期，就需要200/16=12.5，即13个warp串行，才能有效隐藏对global memory的访问延时。
    
    回忆：每个SM一次只能执行一个warp。待确认。。。。