写在最前

本篇是简单的CUDA编程相关的入门,比较浅显易懂,如需详细了解,请转向Programming Massively Parallel Processors这本书,或者相关文章

Why is GPU 为什么选择GPU

根据计算机体系结构的差异，GPU将大量的晶体管分配制作ALU(算术逻辑单元),而CPU则大量用于控制单元.实际体现出来则是GPU对于计算更加强力.但也因此,由于缺少控制单元,GPU编程则需要将更多的精力放在逻辑控制处理上.

GPU是为了提高吞吐量而诞生的,其能够处理更多的数据量,然而代价则是其计算延迟比较大(实际上,GPU一次计算大约是4个时钟周期,而CPU仅需1个时钟周期).因此,GPU选择同时处理多个任务,通过在一个计算的延迟时间去计算其他任务来隐藏实际的延迟,在整体上体现出来GPU没有那么大的延迟.

GPU的主存通常是HBM(High Bandwidth Memory),一条HBM3E的带宽是DDR5的20倍.为了达到HBM的标称访问速度,必须做地址连续的大块访问(coalesced access),以免浪费带宽.

What is CUDA 什么是CUDA

详情请参考Programming Massively Parallel Processors,会有更详细讲解

CUDA,即Compute Unified Device Architecture,是2006年NVIDIA推出的支持GPU编程的一系列图形化API工具,使对GPU编程运算更加友好便捷.

在CUDA中,程序分配规模从小到大依次是thread,warp,block,SM,grid,其中thread是程序运行的基本单元,一个warp含有32个threads,执行相同的命令.block又叫做CTA(Cooperative Thread Array),是线程间的合作范围.SM是执行block/warp的实际硬件结构,如下图所示.

grid则是block的集合,而block的数量可以远多于SM的数量,其调度顺序是未定义的,根据实际计算情况决定.block之间都是可以独立执行的程序,没有数据依赖,是并行关系.

而对于CUDA的编译,CUDA扩展于C/C++,由nvcc进行编译.CUDA程序需要在CPU上准备数据并拷贝到GPU上,根据配置的kernal函数进行GPU上的启动计算,执行结果需要拷贝回CPU从而结束整个计算.

CUDA函数的API,详情请参考CUDA API手册

CUDA采用的是SIMT&SIMD的方式.从CC 7.0开始,每一个thread均有自己的program counter,也就是可以执行不同的指令,而SIME只有一个program counter.然而,一个warp内每次激活的线程应该具有相同的program counter.如果分支太多,将导致program counter不同的取值太多,那么SIMT也是低效的,因此最有利于SIMT发挥效率的编程方式仍然是SIMD

Memory Hierarchy 内存层次结构

Warp-Level Primitives

Warp内可以高效进行数据交换和同步,以实现reduce操作,代码如下:

#define FULL_MASK 0xffffffff
for (int offset = 16; offset > 0; offset /= 2)
    val += __shfl_down_sync(FULL_MASK, val, offset);
/*
__shfl_down_sync(unsigned mask, T var, unsigned int delta, int width = warpSize)参数解析如下:
mask 一个三十二位整数,每一位代表warp中一个线程Lane,位1代表该线程参与操作,示例0xffffffff即32位全为1,代表warp内所有线程参与此次洗牌操作
var 当前线程想要发送出去的寄存器变量(按值传递)
delta 偏移量/步长.数据移动的通道数,核心规则为当前线程的ID为lane_id,调用函数后它会获得同一个warp内lane_id+delta那个线程的var值
width 组宽,可选,默认为32,用于将一个warp划分为多个独立的子组(如16,8,4,2)进行组内的洗牌
*/