- 通过配置文件驱动的方法优化内核
- 理解线程束执行的本质
- 增大GPU的并行性
- 掌握网格和线程块的启发式配置
- 学习多种CUDA的性能指标和事件
- 了解动态并行与嵌套执行
一、CUDA执行模型概述
CUDA编程模型中两个主要的抽象概念:内存层次结构和线程层次结构。
CUDA执行模型有助于在指令吞吐量和内存访问方面提高代码效率。
1. GPU架构概述
GPU架构围绕一个流式多处理器(SM)的可扩展阵列搭建的,可以通过复制这种架构的构建块来实现GPU的硬件并行。
Fermi SM的关键组件:
(1)CUDA核心
(2)共享内存/一级缓存
(3)寄存器文件
(4)加载/存储单元
(5)特殊功能单元
(6)线程束调度器
2. Fermi架构
(1)CUDA核心:多达512个加速器核心(CUDA核心),每个CUDA核心都有一个全流水线的整数算术逻辑单元(ALU)和一个浮点运算单元(FPU)。
(2)SM:CUDA核心被组织到16个SM中,每一个SM中含有32个CUDA核心。
(3)DRAM:6个384位的GDDR5 DRAM存储器接口,支持多达6GB的全局机载内存。
(4)Host Interface:主机接口通过PCIe总线将GPU与CPU相连。
(5)GigaThread:全局调度器,用来分配线程块到SM线程束调度器上。
Fermi架构包含一个耦合的768KB的二级缓存,被16个SM所共享,一个垂直矩形条表示一个SM。
3. Kepler架构
4. 配置文件驱动优化
性能分析是通过检测来分析程序性能的行为:(1)应用程序代码的空间(内存)或时间复杂度;(2)特殊指令的使用;(3)函数调用的频率和持续时间
开发一个HPC应用程序通常包括两个主要步骤:
1. 提高代码的正确性
2. 提高代码的性能
CUDA提供了两个主要的性能分析工具:(1)nvvp:独立的可视化分析器;(2)nvprof:命令行分析器。可以获得CPU与GPU上CUDA关联活动的时间表。
二、理解线程束执行的本质
1. 线程束和线程块
线程束是SM中基本的执行单元。
(1)当一个内核网格被启动后,网格中的线程块被分布在SM中。
(2)一旦线程块被调度到SM上,线程块中的线程会被进一步划分为线程束。
(3)一个线程束由32个连续的线程组成。在一个线程束中,所有的线程按照单指令多线程方式执行。所有的线程都执行相同的指令,每个线程在私有数据上进行操作。
硬件总是给一个线程块分配一定数量的线程束。线程束不会在不同的线程块之间分离。如果线程块的大小不是线程束大小的偶数倍,那么在最后的线程束里有些线程就不会活跃。
从逻辑角度来看,线程块是线程的集合,可以被组织为一维、二维或三维布局。
从硬件角度来看,线程块是一维线程束的集合。在线程块中线程被组织成一维布局,每32个连续线程组成一个线程束。
2. 线程束分化
控制流是高级编程语言的基本构造的一种,GPU支持传统的、C风格的、显式的控制流结构。
如果一个线程束中的线程产生分化,线程束将连续执行每一个分支路径,而禁用不执行这一路径的线程。
条件分支越多,并行性削弱越严重。为了获得最佳的性能,应该避免在同一线程束中有不同的执行路线。
- GPU中设备代码支持与C语言类似的分支控制语句
- 在一个GPU时钟周期内,线程束(warps)中所有的线程必须执行相同的指令
- 线程束中的线程执行不同的分支指令称为线程束分支
- 线程束分支会降低GPU实际的计算能力
- 线程束分支对程序性能影响通过分支效率(branch efficiency)衡量
__global__ void mathKernel(double *c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float ia, ib, ic;
ia = ib = ic = 0.0f;
if(tid % 3 == 0)
{
ia = 100.0f;
}
else if(tid % 3 == 1)
{
ib = 200.0f;
}
else
{
ic = 300.0f;
}
c[tid] = ia + ib + ic;
}
sudo nvprof --metrics branch_efficiency ./main
nvcc对编译器进行了优化,去除nvcc对于主机和设备代码的优化:
nvcc -g -G -arch=sm_60 main.cu -o main
sudo nvprof --metrics branch_efficiency ./main
NVCC对于线程束分支的优化原理:
- 将分支指令替换为预测指令
- 根据程序运行状态将预测变量值设置为1或0
- 优化后的所有分支代码都会执行,但是只有预测变量值为1的分支下的代码才会被线程执行
- 编译器对线程分支的优化能力有限,只有当分支下的代码量很少时优化才会起作用
__global__ void mathKernel2(double *c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float ia, ib, ic;
ia = ib = ic = 0.0f;
bool ipred_ia = (tid % 3 == 0);
bool ipred_ib = (tid % 3 == 1);
bool ipred_ic = (!ipred_ia) & (!ipred_ib);
if(ipred_ia)
{
ia = 100.0f;
}
if(ipred_ib)
{
ib = 200.0f;
}
if(ipred_ic)
{
ic = 300.0f;
}
c[tid] = ia + ib + ic;
}
分支特点:
- 分支只会发生在同一个线程束内
- 不同线程束中的条件判断值,不会造成线程束分支
- 如果代码中条件的判断值与线程id关联,则以线程束为基本单元访问数据
__global__ void mathKernel3(double *c)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float a, b;
a = b = 0.0f;
if((tid / warpSize) % 2 == 0) // 以线程束作为一个整体
{
a = 100.0f;
}
else
{
b = 200.0f;
}
c[tid] = a + b;
}
线程束资源:
1. 线程束计算资源包括:
(1)程序计数器;(2)寄存器;(3)共享内存
2. 线程束所需的计算资源属于片上(on-chip)资源,因此GPU在不同线程束间切换的成本可以忽略
3. 流处理器具有多个32位的寄存器和一定量的共享内存
- 寄存器在线程间分配
- 共享内存在线程块间分配
- 如果SM中的资源无法满足至少一个线程块的需求,内核就无法运行
- 已经分配资源的线程块称为活动块(active block),所包含的线程束称为活动线程束(active warps)。
- 活动线程束根据状态:(1)被选择(selected)线程束;(2)闲置(stalled)线程束;(3)就绪线程束。
四、避免分支分化
1. 并行归约问题
在向量中执行满足交换律和结合律的运算,被称为归约问题。并行归约问题是这种运算的并行执行。
例如:向量求和:
并行加法的一个常用方法是使用迭代成对实现。一个数据块只包含一对元素,并且一个线程对这两个元素求和产生一个局部结果。然后,这些局部结果在最初的输入向量中就地保存。这些新值被作为下一次迭代求和的输入值。因为输入值的数量在每一次迭代后会减半,当输出向量的长度达到1时,最终的求和就已经被计算出来了。
2. 并行归约中的分化
每个线程块在数组的一部分上独立地执行操作。
循环中迭代一次执行一个归约步骤。
归约是在就地完成的,每一步全局内存中的值都被部分和替代。
__syncthreads语句可以保证,线程块中的任一线程在进入下一次迭代之前,在当前迭代里的每个线程的所有部分和都被保存在全局内存中。
进入下一次迭代的所有线程都使用上一步产生的数值。
在最后一个循环以后,整个线程块的和都被保存在全局内存中。