CUDA

GPU上线程没有优先级概念，所有线程机会均等，线程状态只有等待资源和执行两种状态，如果资源未就绪，那么就等待；一旦就绪，立即执行

线程并行是细粒度并行，调度效率高；块并行是粗粒度并行，每次调度都要重新分配资源，有时资源只有一份，那么所有线程块都只能排成一队，串行执行。

流可以实现在一个设备上运行多个核函数。前面的块并行也好，线程并行也好，运行的核函数都是相同的（代码一样，传递参数也一样）。而流并行(grid级并行)，可以执行不同的核函数，也可以实现对同一个核函数传递不同的参数，实现任务级别的并行。

实际应用中只有向量相加、向量对应点乘这类才会有如此高的并行度，其余的没有如此高并行度，所以线程间需要通信。

Block内部同步机制可以用CUDA内置函数：__syncthreads()；当某个线程执行到该函数时，进入等待状态，直到同一线程块（Block）中所有线程都执行到这个函数为止，即一个__syncthreads()相当于一个线程同步点，确保一个Block中所有线程都达到同步，然后线程进入运行状态
 

不同Block中的线程不能通过共享内存、同步进行通信，而应采用原子操作或主机介入，对于原子操作，如果感兴趣可以翻阅《GPU高性能编程CUDA实战》第九章“原子性”。

有了这个思路，结合上节看到的SM结构图，看到有一片存储器叫做Shared Memory，它位于SM内部，处理时访问速度相当快（差不多每个时钟周期读一次），而全局存储器读一次需要耗费几十甚至上百个时钟周期。于是，我们就制定A计划如下：

线程，有时被称为轻量进程，是程序执行流的最小单元。一个标准的线程由线程ID，当前指令指针(PC），寄存器集合和堆栈组成。线程是进程中的一个实体，是被系统独立调度和分派的基本单位，线程不拥有私有的系统资源，但它可与同属一个进程的其它线程共享进程所拥有的全部资源。一

参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的shared memory大小，单位为byte。不需要动态分配时该值为0或省略不写，所以在l1 cache设计中需要考虑。

 

分析：其实在cuda上线程去读取内存时，是以线程束的方式去读取的，一般来说一个线程束有32个线程，如果这32个线程读取的数据具有连续性，则cuda会合并其访问，即只发起一次内存访问，在计算功能集比较高的平台上，由于有缓存的存在，这样也会使得缓存命中率提高，提高访存效率。但访存合并仅限于同一个block中，所以当我们改变block中的线程格局时，当达到32*16和64*8等时会比16*32访存效率高，从而提高程序性能，但由于缓存大小的影响，访存合并的最大字节有限，所以当线程格局太多时反而会影响程序性能。

线程-》PE映射：

for (i=0,i<=256,i++)

  c[i] = a[i]+b[i];

直接映射到256个线程进行处理

for (i=0, i<=256,i++)

  c = a + c

利用寄存器及pipeline的处理，也可映射到256个线程上进行处理，具体需要想想怎么一个过程。

即使简单的 w=a+b,也可能由于a,b各自位宽比较大超过alu的运算bit，而被拆分到不同的thread进行运算。

 

warp是SM的基本执行单元。一个warp包含32个并行thread，这32个thread执行于SMIT模式。也就是说所有thread执行同一条指令，并且每个thread会使用各自的data执行该指令。

线程的执行过程：

图中的Residency Slots中包含很多slot，每个slot代表一个warp，空的表示目前还没有部署warp。而部署了warp的slots一共有三个状态，绿色表示active，黄色表示ready，可以执行了，红色表示阻塞；active的warp接下来会在执行单元上执行，如图右侧所示，所有的32个work-item同时并行执行。Ready的会在下一个执行周期被调度执行；阻塞的则是因为读写等原因进入该状态。

下图是USC中流水线示意图，其中包含4个warp的调度。Warp0首先被执行，warp0会一直执行到它进入阻塞状态，例如读写全局存储，此时调度器会停止调度warp0，开始执行warp1；因为warp中的所有工作项执行相同的kernel代码，因此就有相同的特性，例如同时进入阻塞；在warp2进入阻塞状态时，warp0读写结束，进入read状态；最后在调度器调度完warp3后，重新开始调度warp0。这样并发执行可以实现对内存访问延迟的隐藏。因此在编程实现中一般使用较大的工作组，来实现warp切换对内存访问的延迟（当然，这不是绝对的，在实际中还要考虑寄存器等资源的消耗情况）。

GPU与CPU内核是不同的, 少了三级缓存, 分支预测等等. 但是增加了ALU(PE)的数量, 扩大了上下文存储池(Pool of context storge).上下文会对应实际的warp

所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等，因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。

如何决定kernal程序分成多少个block?

Block <-> SM
Thread执行 <-> CUDA Cores
Thread数据 <-> Register/Local Memory

同一Grid下的不同Block会被分发到不同的SM上执行。SM上可能同时存在多个Block被执行，它们不一定来自同一个kernel函数。每个Thread中的局域变量被映射到SM的寄存器上，而Thread的执行则由CUDA cores来完成。

SM上可以同时存在多少个Block？这由硬件资源的消耗决定：每个SM会占用一定数量的寄存器和Shared Memory，因此SM上同时存活的Block数目不应当超过这些硬件资源的限制。由于SM上可以同时有来自不同kernel的Block存在，因此有时候即便SM上剩余资源不足以再容纳一个kernel A的Block，但却仍可能容纳下一个kernel B的Block.
 

PC的奇偶thread是为了解决 warp divergence? 以提高效率?

https://www.jianshu.com/p/87cf95b1faa0

一个kernel的所有线程其实在物理层是不一定同时并发的。所以kernel的grid和block的配置不同

cam 设计？text cache

如图, GPU硬件的一个核心组件是SM, SM是英文名是Streaming Multiprocessor, 翻译过来就是流式多处理器. SM的核心组件包括CUDA核心(其实就是ALU, 如上图绿色小块就是一个CUDA核心), 共享内存, 寄存器等, SM可以并发地执行数百个线程, 并发能力就取决于SM所拥有的资源数. 当一个kernel被执行时, 它的gird中的线程块被分配到SM上, 一个线程块只能在一个SM上被调度. SM一般可以调度多个线程块, 这要看SM本身的能力. 那么有可能一个kernel的各个线程块被分配多个SM, 所以grid只是逻辑层, 而SM才是执行的物理层.