一、多线程
(1)设计理念
CPU是面向延迟的核心,尽可能快完成计算任务
GPU是面向吞吐量的核心,通过多个线程实现在某一段时间内的高吞吐率
打个比喻,前者是一架飞机,后者是一千辆牛车
可以看到,CPU的Cache更大,通过复杂的控制逻辑实现性能优化,而GPU的控制逻辑相对简单,通过多线程来实现高吞吐率
(2)多线程
CPU中的SIMD
✓所有内核同时执行相同的指令,但数据不同
✓类似于CRAY超级计算机上的矢量计算
✓例如:SSE4,AVX指令集
SMX中的SIMT
✓多线程CUDA核心
✓每个SMX上的线程执行组共享相同指令
✓细粒度并行
✓自然用于图形处理和大量科学计算
✓SIMT也是多核芯片的自然选择简化每个核心
线程:具有自己的PC和数据的指令流
✓拥有专用寄存器,专用内存,程序计数器和线程执行状态
✓线程级并行(TLP):利用线程之间固有的并行性
多线程
✓多线程共享功能单元1 处理器通过重叠
✓处理器必须复制每个线程的独立状态,例如,寄存器文件的单独副本,单独的PC
✓通常,用于快速线程切换的硬件
✓要共享的内存
✓解决内存访问停顿
(3)GPU多线程
GPU是一种多线程处理器,例如GTX8800支持96个线程
取值→译码→执行→内存→回写
显然,内存读写是最耗时的阶段,在这个停滞时期进行其他线程的计算可以提高性能
二、CUDA抽象
CPU诞生于2007年,是一种并行计算平台和编程模型,可通过利用GPU的强大功能显着提高计算性能
CUDA作为并行计算平台
✓语言:基于C的CUDA C,带有一些扩展(广泛的C ++支持)
✓编辑器:Eclipse / VisualStudio
✓编译器:nvcc
✓SDK:CUDA工具包,库,样本
✓Profiler&Debugger:Nsight
CUDA作为编程模型
✓GPU硬件的软件抽象
✓独立于操作系统,CPU,Nvidia GPU
CUDA虚拟化物理硬件
✓线程是虚拟化的CUDA核心(寄存器,PC,状态)
✓block是虚拟化流式多处理器(SM)(线程,共享内存)
无需抢占计划到物理硬件上
✓线程/块启动并运行到完成/暂停
✓块应该是独立的
(1)并发线程的层次结构
第一个参数是block数量,第二个参数是每个block的线程数
线程对应于CUDA核,block对应于SM,而Grid对应于多个SM
(2)协作线程的共享内存模型
每个线程有自己的私有内存和寄存器,同一个block内有共享内存,grid间有全局内存
(3)轻量级同步
全局同步
完成一个kernel再进行下一个
线程同步
kernel内部完成一个线程再进行下一个
竞争
原子操作
atomic{Add, Sub, Exch, Min, Max, Inc, Dec, CAS, And, Or, Xor}
原子操作可以避免竞争,但是执行速度略慢
三、Kernel执行
(1)Kernel函数
使用__global__修饰,返回值为void
__device__和__host__可以一起使用
KernelFunc<<<gridDim, blockDim, nSMem,iStream>>>(args);
gridDim是网格的维度
blockDim是每个块的维度
其中这些参数可以是dim3类型,要注意下面这些维度限制
(2)线程索引
对于一维情况:
blockIdx.x * blockDim.x + threadIdx.x
对于二维情况:
i = blockIdx.x * blockDim.x + threadIdx.x
j = blockIdx.y * blockDim.y + threadIdx.y
idx = j * gridDim.x * blockDim.x + i
如果数据无法恰好比配,则有上面的情况。
对于三位情况图示如下,推导类似于二维情况。
(3)设备占用率
✓每个多处理器的活动warp数与同时在多处理器上可以激活的最大warp数之比
✓分母是一个常数,仅取决于设备的计算能力
✓该表达式的分子是一个函数与以下内容有关:
•计算能力
•每个块的线程数
•每个线程的寄存器数
•每个块的共享内存和共享内存配置
四、线程束调度
在较低级别,在GPU内:
✓执行内核的每个块在SMX上执行,一个SM最大由16个block共享
✓如果块的数量超过SMX的数量,那么如果有足够的寄存器,则每个SMX上将一次运行多个块和共享内存,其他人将在队列中等待并稍后执行
✓一个块中的所有线程都可以访问本地共享内存,但无法看到其他块正在做什么(即使它们在同一个SMX上)
✓不保证块执行的顺序
block映射到物理GPU上按照每32个线程一组,即一个线程束,几乎是同时运行同时结束,可以看到warp0处在执行状态
当进行访存时warp0处于挂起状态,同时warp1进入执行状态,此时warp1如果也需要访存也将挂起
当4个warp都挂起时,硬件此时会空闲
当某个warp访存完成后,则会进入执行阶段
五、CUDA开发工具和生态圈
(1)软件架构
从图中可以发现,CUDA应用有多层结构,而CUDA编程本身其实是基于其中“CUDA Runtime”层次。
(2)编译过程
从图中可以看到,编写完.cu文件后,nvcc编译成GPU与CPU运行的代码,然后合并,经过Host编译器生成Host执行文件,然后运行,并通过CUDA驱动。运行涉及到CUDA的一些库,cudart.lib和cuda.lib。
(3)NVCC
常用命令行选项:
-g:产生可调式代码,调试模式下必须的
-G:产生可调式的设备代码
-lineinfo:为设备代码产生行number信息
-o outfile:指定输出文件的位置和名称
-include xxx.h:指定预处理和编译时预先需要包含的头文件
-l yyy.lib:指定链接时需要的库文件
-arch sm_??:指定GPU架构
(4)CUDA软件生态
数值分析工具:Matlab, ArrayFire, Mathematica, LabView, Jacket
GPU加速库:Thrust, cuDNN, cuFFT, cuBLAS
语言和接口:CUDA Toolkit, OpenACC, CUDA Fortran, PyCUDA, PGI compilers
性能分析工具:Nsight, NVIDIA Visual Profiler,nvprof
Debug工具: Nsight, CUDA-GDB, CUDA-MEMCHECK