笔记03：CUDA执行模型

• 通过配置文件驱动的方法优化内核
• 理解线程束执行的本质
• 增大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位的寄存器和一定量的共享内存

• 寄存器在线程间分配
• 共享内存在线程块间分配

1. 如果SM中的资源无法满足至少一个线程块的需求，内核就无法运行
2. 已经分配资源的线程块称为活动块（active block），所包含的线程束称为活动线程束（active warps）。
3. 活动线程束根据状态：（1）被选择（selected）线程束；（2）闲置（stalled）线程束；（3）就绪线程束。

四、避免分支分化

1. 并行归约问题

在向量中执行满足交换律和结合律的运算，被称为归约问题。并行归约问题是这种运算的并行执行。

例如：向量求和：

 并行加法的一个常用方法是使用迭代成对实现。一个数据块只包含一对元素，并且一个线程对这两个元素求和产生一个局部结果。然后，这些局部结果在最初的输入向量中就地保存。这些新值被作为下一次迭代求和的输入值。因为输入值的数量在每一次迭代后会减半，当输出向量的长度达到1时，最终的求和就已经被计算出来了。

2. 并行归约中的分化

每个线程块在数组的一部分上独立地执行操作。

循环中迭代一次执行一个归约步骤。 

归约是在就地完成的，每一步全局内存中的值都被部分和替代。

__syncthreads语句可以保证，线程块中的任一线程在进入下一次迭代之前，在当前迭代里的每个线程的所有部分和都被保存在全局内存中。

进入下一次迭代的所有线程都使用上一步产生的数值。

在最后一个循环以后，整个线程块的和都被保存在全局内存中。

五、展开循环