CUDA笔记之二：硬件与拓扑原理篇

这部分是一些枯燥的硬件知识的总结，但是对优化CUDA程序有着至关重要的作用，在后面的文章里，我将尽量结合实例来讲解这些东西

 

1 GPU硬件

i   GPU一个最小单元称为Streaming Processor(SP)，全流水线单事件无序微处理器，包含两个ALU和一个FPU，多组寄存器文件（register file，很多寄存器的组合），这个SP没有cache。事实上，现代GPU就是一组SP的array，即SPA。每一个SP执行一个thread

ii   多个SP组成Streaming Multiprocessor(SM)。每一个SM执行一个block。每个SM包含
     8个SP；
     2个special function unit(SFU)：这里面有4个FPU可以进行超越函数和插值计算
     MultiThreading Issue Unit：分发线程指令
     具有指令和常量缓存。
     包含shared memory

iii   Texture Processor Cluster(TPC) ：包含某些其他单元的一组SM

    硬件SPA体系：

2 Single-Program Multiple-Data （SPMD）模型（说明：以下硬件数据不通用，只做参考）

i   CPU以顺序结构执行代码，GPU以threads blocks组织并发执行的代码，即无数个threads同时执行
ii   回顾一下CUDA的概念：
     一个kernel程序执行在一个grid of threads blocks之中
     一个threads block是一批相互合作的threads：可以用过__syncthreads同步；通过shared memory共享变量，不同block的不能同步。
iii   Threads block声明：
      可以包含有1到512个并发线程，具有唯一的blockID，可以是1,2,3D
      同一个block中的线程执行同一个程序，不同的操作数，可以同步，每个线程具有唯一的ID
3 线程硬件原理
i  GPU通过Global block scheduler来调度block，根据硬件架构分配block到某一个SM。每个SM最多分配8个block，每个SM最多可接受768个thread（可以是一个block包含512个thread，也可以是3个block每个包含256个thread（3*256=768！））。同一个SM上面的block的尺寸必须相同。每个线程的调度与ID由该SM管理。
ii SM满负载工作效率最高！考虑某个Block，其尺寸可以为8*8,16*16,32*32
   8*8：每个block有64个线程，由于每个SM最多处理768个线程，因此需要768/64=12个block。但是由于SM最多8个block，因此一个SM实际执行的线程为8*64=512个线程。
   16*16：每个block有256个线程，SM可以同时接受三个block，3*256=768，满负载：）
    32*32：每个block有1024个线程，SM无法处理！  
iii  Block是独立执行的，每个Block内的threads是可协同的。
iv  每个线程由SM中的一个SP执行。当然，由于SM中仅有8个SP，768个线程是以warp为单位执行的，每个warp包含32个线程，这是基于线程指令的流水线特性完成的。Warp是SM基本调度单位，实际上，一个Warp是一个32路SIMD指令。基本单位是half-warp。
    如，SM满负载工作有768个线程，则共有768/32=24个warp，每一瞬时，只有一组warp在SM中执行。
    Warp全部线程是执行同一个指令，每个指令需要4个clock cycle，通过复杂的机制执行。
v  一个thread的一生：
    Grid在GPU上启动；block被分配到SM上；SM把线程组织为warp；SM调度执行warp；执行结束后释放资源；block继续被分配....
4 线程存储模型
i  Register and local memory：线程私有，对程序员透明。
    每个SM中有8192个register，分配给某些block，block内部的thread只能使用分配的寄存器。线程数多，每个线程使用的寄存器就少了。
ii shared memory：block内共享，动态分配。如__shared__ float region[N]。
  shared memory 存储器是被划分为16个小单元，与half-warp长度相同，称为bank，每个bank可以提供自己的地址服务。连续的32位word映射到连续的bank。
    对同一bank的同时访问称为bank conflict。尽量减少这种情形。

iii Global memory：没有缓存！容易称为性能瓶颈，是优化的关键！
    一个half-warp里面的16个线程对global memory的访问可以被coalesce成整块内存的访问，如果：
   数据长度为4,8或16bytes；地址连续；起始地址对齐；第N个线程访问第N个数据。
    Coalesce可以大大提升性能。

coalesced

uncoalesced

   Coalesced方法：如果所有线程读取同一地址，不妨使用constant memory；如果为不规则读取可以使用texture内存
   如果使用了某种结构体，其大小不是4 8 16的倍数，可以通过__align(X)强制对齐，X=4 8 16