GPU和CUDA的基本概念

source: 

  1. https://blog.csdn.net/jiangbo1017/article/details/53940428
  2. http://www.geeks3d.com/20100318/tips-what-is-a-texture-processor-cluster-or-tpc/
  3. http://www.voidcn.com/article/p-vwvreldo-v.html
  4. http://people.cs.umass.edu/~emery/classes/cmpsci691st/readings/Arch/gpu.pdf
  5. https://blog.csdn.net/ai_vivi/article/details/42235787

-----------------------

CUDA 软件编程相关的概念：

  • thread --> block --> grid

在利用cuda进行编程时，一个grid分为多个block，而一个block分为多个thread。其任务划分影响最后的执行效果。划分的依据是任务特性和GPU本身的硬件特性。GRID,BLOCK,THREAD是软件概念,而非硬件的概念。

both grid and block can be 3 dimensional.

从硬件角度讲，一个GPU由多个SM组成（当然还有其他部分），一个SM包含有多个SP（以及还有寄存器资源，shared memory资源，L1 cache，scheduler，SPU，LD/ST单元等等），1.x硬件，一个SM包含8个SP，2.0是32个，2.1是48个，3.0和3.5是192个。以及SP目前也称为CUDA CORE，而SM目前也称为MP，在KEPLER架构（SM3.0和3.5）下也称为SMX。

从软件角度讲，CUDA因为是SIMT(single instruction, multiple threads)的形式，grid，block，thread是thread的组织形式。最小的逻辑单位是一个thread，最小的硬件执行单位是thread warp（简称warp），若干个thread（典型值是128~512个）组成一个block，block被加载到SM上运行，多个block组成整体的grid，一个kernel对应一个grid。一个grid可以是一维，二维或者三维的，具体由GridDim定义。每一个block都有各自的ID, 如blockIdx.xyz。而block中的thread的分布由BlockDim定义，每个thread都有各自的ID, 如theadIdx.xzy。

这里为什么要有一个中间的层次block呢？这是因为CUDA通过这个概念，提供了细粒度(fine grained)的通信手段，因为block是加载在SM上运行的，所以可以利用SM提供的shared memory 和__syncthreads() 功能实现线程同步和通信，这带来了很多好处。而block之间，除了结束kernel之外是无法同步的，一般也不保证运行先后顺序，这是因为CUDA程序要保证在不同规模（不同SM数量）的GPU上都可以运行，必须具备规模的可扩展性，因此block之间不能有依赖。

GPU硬件相关的概念：

  • SP --> WARP --> SM --> TPC
  • SP: streaming processor, 最基本的处理单元, 最后具体的指令和任务都是在sp (also called CUDA core) 上处理的。GPU进行并行计算，也就是很多个sp同时做处理。现在SP的术语已经有点弱化了，而是直接使用thread来代替。一个SP对应一个thread。A SP includes several ALUs and FPUs. An ALU is an arithmetical and Logical Unit and a FPU is a Floating Point Unit. The SP is the real processing element that acts on vertex or pixel data.
  • WARP: warp是SM调度和执行的基础概念，同时也是一个硬件概念，注意到warp实际上是一个和硬件相关的概念，通常一个SM中的多个SP(thread)会分成几个warp(也就是SP在SM中是进行分组的，物理上进行的分组)，每一个WARP中在Tegra中是32个thread.这个WARP中的32个thread(sp)是一起工作的，执行相同的指令，如果没有这么多thread需要工作，那么这个WARP中的一些thread(sp)是不工作的 (a “warp” in NVIDIA hardware executes only one common instruction at a time on all threads in the work-group (since access to individual threads is through global SIMT instruc- tions), so full efficiency is only realized when all 32 threads in the warp agree on their execution path)。 每一个线程都有自己的寄存器内存和local memory，一个warp中的线程是同时执行的，也就是当进行并行计算时，线程数尽量为32的倍数，如果线程数不上32的倍数的话；假如是1，则warp会生成一个掩码，当一个指令控制器对一个warp单位的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个 进程会进入静默状态.
  • SM: streaming multiprocessor, 多个sp加上其他的一些资源组成一个SM. 其他资源也就是存储资源，共享内存，寄储器等。可见，一个SM中的所有SP是先分成warp的，是共享同一个memory和instruction unit. 每个SM通过使用两个特殊函数(Special Function Unit,SFU)单元进行超越函数和属性插值函数（根据顶点属性来对像素进行插值）计算。SFU用来执行超越函数、插值以及其他特殊运算.
  • TPC：texture/processor cluster, it's a groupd made of several SM and a texture, and some control logics.

  • Stream:

总而言之，一个kernel对应一个GRID，该GRID又包含若干个block，block内包含若干个thread。GRID跑在GPU上的时候，可能是独占一个GPU的，也可能是多个kernel并发占用一个GPU的（需要fermi及更新的GPU架构支持）。block是resident在SM上的，一个SM可能有一个或多个resident blocks，需要具体根据资源占用分析。thread以warp为单位被SM的scheduler 发射到SP或者其他单元，如SFU，LD/ST unit执行相关操作，需要等待的warp会被切出（依然是resident 状态），以空出执行单元给其他warps。

 

-----------------------

 

GPU的结构简单的可以分为：

  • 一个连接GPU和PCIe总线的主机接口；
  • 0~2个复制引擎；
  • 一个连接GPU与GPU内存设备的DRAM接口；
  • 一定数目的TPC或者GPC（纹理处理集群或图形处理集群)，每个包含一定的缓存和一些流处理器簇（SM）
  • example: Tesla 架构的SM的结构包括
    • 执行单元，用以执行32位整数和单、双精度浮点运算；
    • 特殊函数单元（SFU），用以计算log/exp，sin/cos，rcp/rsqrt的单精度近似值；
    • 一个线程束调度器，用以协调把指令分发到执行单元；
    • 一个常量缓存，用于广播式传送数据给SM；
    • 共享内存，用于线程之间的数据交换；
    • 纹理映射的专用硬件

 

posted @ 2018-05-14 01:15  icdane  阅读(278)  评论(0)    收藏  举报