22日上午，来自湖南大学的阳王东教授为本次暑期学校带来了《异构计算机系统上协同并行计算与应用》系列报告第二场：《CUDA体系结构与编程技术》。此次报告主要介绍了GPU的物理结构，基于CUDA框架的编程模式以及CUDA并行程序的优化策略。
首先，阳教授给大家简单介绍了GPU的发展历程。GPU的发展经历了三代变化，可编程性和应用范围都逐渐提高，具体见图1。
图1：GPU的发展历程
其中CUDA是计算统一设备构架（Compute Unified Device Architecture）的缩写，是一种通用并行计算平台和编程模型。CUDA使得开发者在GPU的并行计算引擎上，依然可以使用常用的编程语言C进行编程。
接下来，阳教授给大家简单介绍了CPU+GPU混合计算架构的三个层次：
1.最外层为CPU与GPU并立的串行执行结构。GPU作为CPU的协处理器完成图形处理或并行任务，两者都拥有自己的存储系统，CPU通过PCIE总线向GPU传输指令和数据，管理GPU的运行。
2.第二层为GPU内的多处理器（SM，Streaming Multiprocessor）与板载显存组成的并行执行结构。多处理器之间并行执行，并独立访问显存。
3.第三层为SM内部的多个CUDA核心（或称为流处理器SP，Streaming Processor）与共享内存、寄存器、指令单元等物理器件组成的SIMT（Single Instruction Multiple Threads）执行结构。
然后，阳教授以GeForceGTX 200为例（具体见图2），给大家详细介绍了GPU的物理结构。
1.GTX 200拥有10个TPC（TextureProcessor Cluster），每个TPC由3个SM（Streaming Multiprocessor）和一个纹理单元（Texture Unit）以及一些逻辑控制单元构成。
2.一个SM由8个流处理器SP（Streaming Processor）和一些板载显存构成，SM就是一个SIMD（SingleInstruction Multiple Data，单指令多数据流）单元，也是最小的指令分派单位。
3.SP是CUDA的最小运行单位（即CUDA Core），可运行一个逻辑上的线程。SP的执行延迟一般是4个周期，类似于执行管线长度有4级，所以每个SP可以执行4条相同的指令来充分消除这个延迟。
图2：GTX 200 GPU物理结构示意图
在介绍了GPU的物理结构之后，阳教授给大家继续讲解基于CUDA框架的编程模式。与硬件对应，CUDA框架将并发分为三个层次：Thread是最基本的并发单元，在SP上运行；Block由thread组成，在一个SM上运行；Grid由block组成，由GPU整体执行。
阳教授将CUDA编程称为“混杂编程”。所谓混杂，即CUDA编程综合了串行和并行两部分：CPU部分称为host代码，是串行的；GPU部分称为device代码（注：NVIDIA把device代码称为kernel），是并行的。两者都可用C语言实现，前者在CPU上串行执行，后者在GPU上并行执行。图3是一个device代码的示意图。
图3：Device代码示意图（Grid ->Block -> Thread三层）
最后，阳教授介绍了CUDA并行程序的优化策略，主要有三个方面：任务划分优化，存储器访问优化和指令流优化。
1.任务划分优化
主要思想有：逻辑复杂任务由CPU完成，大规模数据并行和计算密集型任务由GPU完成；一次GPU执行的任务应该有充足的计算量，尽量减少通信开销；计算任务和数据规模要协调，避免硬件资源占有率低的情况。
具体Grid、Block和Thread的设计策略：Grid采用二维和三维结构能够较好地使用共享存储器；Block中的线程数最好是32的倍数，以减少计算访存地址的运算；每个SM至少要有6个active warp才能有效隐藏访存时间，同时每个SM至少要有2个active block才能使得计算资源不闲置。
2.存储器访问优化
存储器访问优化主要有四个方面，分别是：主机-设备通信优化；全局储存器访问优化；共享存储器访问优化；使用纹理储存器加速。
其中如何降低bank冲突是比较重要的一点。Bank冲突：共享存储器（shared memory）由多个等大小的内存模块组成，这些模块被称为bank。这些bank可以被同时访问。当多个线程同时访问一个bank时，这些访存指令将会串行执行，这种情况被称为bank冲突。避免bank冲突的主要思路有：改变数据存储方式，尽可能使不同的线程访问不同的bank；利用多播（multicast）的方式，使得访问同一个bank的多个线程的访存指令只执行一次，以避免bank冲突。
3.指令流优化
主要优化策略可概括为：针对算术指令，可将核函数中的单精度计算函数替换为CUDA内部实现的高速版本（但会影响计算精度），同时尽量避免使用整数除和模运算，可用位运算代替；针对控制流指令，确保warp内所有线程执行相同的指令，尽量减少在程序中使用分支语句，同时在循环中尽可能不使用__syncthreads()；针对访存指令，可减少单个数据的赋值操作，对于大数据的访问，尽量拆成最大访问宽度的长度。
撰稿人：李宇明
排版：吴双
