GPU

GPU是一种SIMD(Single Instruction Multiple Data)架构, 成百上千个核在同一时间最好能做同样的事情。

GPU的工作计算量大,但没什么技术含量,而且要重复很多很多次。
GPU用很多简单的计算单元去完成大量的计算任务,纯粹的人海战术。这种策略基于一个前提,就是小学生A和小学生B的工作是互相独立的。
GPU的运算速度取决于雇了多少小学生,CPU的运算速度取决于请了多么厉害的教授。教授处理复杂任务的能力是碾压小学生的,但是对于没那么复杂的任务,还是顶不住人多。
现在的GPU也能做一些稍微复杂的工作了,相当于升级成初中生高中生的水平。但还需要CPU来把数据喂到嘴边才能开始干活,究竟还是靠CPU来管的。
程序员将挖矿和**密码这种事情分成小学生都能做的简单任务。
GPU
其中绿色的是计算单元,橙红色的是存储单元,橙黄色的是控制单元。
1:CPU中负责逻辑算数的部分并不多;而GPU整个就是一个庞大的计算阵列(包括alu和shader填充),GPU大致分为3d setup和流计算两个部分,流计算是属于可编程的,那么GPU正是利用OpenCL、CUDA这些API让流计算的部分来解决非图形问题,比如流体模拟、爆炸等等这些适合大规模密集型运算的工作。
2:GPU对Cache的依赖比CPU小。缓存的目的不是保存后面需要访问的数据的,这点和CPU不同,而是为thread提高服务的。如果有很多线程需要访问同一个相同的数据,缓存会合并这些访问,然后再去访问dram
3:GPU中的SP就是简单的核心。但是CPU在core的设计上是比GPU复杂的多。GPU因为设计问题天生适合大规模SIMT/SIMD运算方式。

GPU通用计算:用GPU来处理一些原本CPU可以处理的计算。
浮点运算:例如高分辨率三角形渲染,材质渲染,光影效果,抗锯齿,物理引擎等等
图形学之外的领域专家开始注意到GPU与众不同的计算能力,开始尝试把GPU用于通用计算(即GPGPU)。GPU开始在通用计算领域得到广泛应用,包括:数值分析,海量数据处理,金融分析等等。
通用计算规范:

  • NVIDIA独家主导的CUDA,只能使用NVIDIA的显卡实现;
  • 微软主导的DirectCompute,A卡N卡乃至Intel核显都能使用;
  • 苹果提出、多家厂商支持的开放性规范,OpenCL,能够对不同架构CPU、GPU等硬件提供支持。

在操作系统中,进程是资源分配的基本单元,而线程是CPU时间调度的基本单元(这里假设只有1个CPU)。

在CUDA程序设计中,线程就是执行CUDA程序的最小单元。在GPU上每个线程都会运行一次核函数。

GPU上线程没有优先级概念,所有线程机会均等。
线程状态只有等待资源和执行两种状态,如果资源未就绪,那么就等待;一旦就绪,立即执行。
当GPU资源很充裕时,所有线程都是并发执行的,这样加速效果很接近理论加速比;
而GPU资源少于总线程个数时,有一部分线程就会等待前面执行的线程释放资源,从而变为串行化执行。

块并行相当于操作系统中多进程的情况。CUDA有线程组(线程块)的概念,将一组线程组织到一起,共同分配一部分资源,然后内部调度执行。线程块与线程块之间,毫无瓜葛。这有利于做更粗粒度的并行。

Grid-Block-Thread结构:
一组线程并行组织为一个block
一组block并行组织为一个Grid(线程格)
一组Grid并行组织为一个流
GPU

线程并行和块并行的区别:

线程并行是细粒度并行,调度效率高;块并行是粗粒度并行,每次调度都要重新分配资源。

将一个大问题分解为几个小规模问题,将每个小规模问题用一个线程块实现,线程块内可以采用细粒度的线程并行,而块之间为粗粒度并行,这样可以充分利用硬件资源,降低线程并行的计算复杂度。

sp(streaming processor): 最基本的处理单元, GPU进行并行计算是很多个sp同时做处理
sm(streaming multiprocessor):多个sp加上其他资源(存储资源,共享内存,寄储器等)组成一个sm
warp:GPU执行程序时的调度单位,目前cuda的warp的大小为32,同在一个warp的线程,以不同数据资源执行相同的指令。

流可以实现在一个设备上运行多个核函数。前面的块并行,线程并行运行的核函数都是相同的。
流并行,可以执行不同的核函数,也可以实现对同一个核函数传递不同的参数,实现任务级别的并行。

如一组数求和,求最大(小)值,各个线程不再是相互独立的,而是产生一定关联,线程2可能会用到线程1的结果,这时需要利用线程通信技术。

线程通信在CUDA中有三种实现方式:

  1. Shared Memory(共享存储器)
    一个SM(流多处理器)中包含若干个SP(流处理器),一部分高速Cache,寄存器组,共享内存等
    GPU
    Shared Memory由这M个SP共同占有。一个SM中所有SP在同一时间执行同一代码。

当线程A需要线程B计算的结果作为输入时,需要确保线程B已经将结果写入共享内存中,然后线程A再从共享内存中读出。
当某个线程执行到__syncthreads()函数时,进入等待状态,直到同一线程块中所有线程都执行到这个函数为止,然后线程进入运行状态。

  1. 原子操作;
    不同Block中的线程采用原子操作或主机介入。