Professional CUDA C Programing

代码下载：http:www.wrox.com/go/procudac
本章的主要内容：
➤了解warp执行的本质
➤将更多的并行性暴露给GPU
➤掌握网格和块配置的设置方法
➤学习各种CUDA性能指标和事件
➤探测动态并行和嵌套执行

GPU的硬件结构

GPU是由Streaming Multiprocessors (SM)组成的，每个SM如下：
➤ CUDA Cores
➤ Shared Memory/L1 Cache
➤ Register File
➤ Load/Store Units
➤ Special Function Units
➤ Warp Scheduler

Fermi SM
GPU中的每个SM都支持数百个线程的并发执行，通常是每个GPU有多个SM，所以有可能有数千个线程并发执行。
CUDA采用了SIMT单指令多线程执行，一个指令32个线程执行，32个线程组织成warp。一个warp中的线程同一时刻执行同一个指令。每个线程有自己的指令技术计数器和寄存器，在自己的数据上执行指令。
SIMT 和 SIMD最大的差异：
➤ 每个线程有自己独立的指令寄存器
➤ 每个线程有自己独立的寄存器状态
➤ 每个线程有独立的执行路径
一个线程块只能分配到一个SM上执行，一个SM可以同时允许多个线程块。
logical view and hardware view
共享存储器和寄存器都是SM上珍贵的资源，共享存储器按线程块进行划分，同一个线程块中的线程可以通过共享内存互相通信，在逻辑上同一个线程块中的所有线程同时执行，但是在物理上，同一个线程块中的所有线程并不是同时执行的，所以同一个线程块中的线程并不是同时执行结束的。While all threads in a thread block run logically in parallel, not all threads can execute physically at the same time. As a result, different threads in a thread block may make progressat a different pace.
共享内存可能会导致线程之间的竞争：多个线程同时访问某个数据。CUDA提供了线程块内的同步，保证同一个线程块中的线程在下一步执行前都完成了上一步的执行。但是线程块之间无法同步。
在SM1中warp1正在执行，但是warp1需要从device中读取数据，此时SM1将调用warp2继续执行，warp1和warp2之间的转换开销不大（SM的资源为所有线程共享），由于warp间并发的执行提高了SM的利用率。（一个SM中真正执行的warp数目和GPU的资源有关）
Fermi Architecture
Fermi Architecture
Fermi有16个SM，每个SM有32个CUDA core（一个warp32个线程），每个CUDA core有ALU和FPU。当一个线程块分配到一个SM上时，线程块被组织成warps，SM上的warp调度器选择合适的warp执行。
Screenshot from 2017-04-26 12:25:10.png
对于计算能力2.0以上的Fermi结构，一个SM最多同时处理48个warps。
Fermi的两个关键点：
➤ 可以通过CUDA runtime API 设置共享内存和L1cache
➤ 支持并发的内核执行：多个小的kernel可以并发执行，最多16个kernels同时在设备上运行。

Kepler Architecture
➤ 15个SM
➤ 每个SM：192 单精度CUDA core，64个双精度计算单元，32个特殊功能计算单元，32个load／store计算单元。4个warp调度器，8个指令分配器。
➤ 计算能力3.5每个SM一次可以调度64个warps驻留在SM上。
➤ 动态并行性。一个kernel可以创建其它的kernel

➤ Hyper-Q。Hyper-Q在CPU和GPU之间增加了更多同步的硬件连接，从而实现了CPU核心同时在GPU上运行更多任务。 因此，可以增加GPU 使用率。 费米GPU依靠单一硬件工作队列将任务从CPU传递到GPU，这可能导致单个任务阻止所有其他任务落后于队伍中取得进展。 开普勒Hyper-Q消除了这个限制。Kepler GPU在主机和主机之间提供32个硬件工作队列GPU。 Hyper-Q可以在GPU上实现更多的并发性，最大限度地提高GPU的利用率。。

性能优化

➤ 时间复杂度、空间复杂度
➤ 特殊指令的使用
➤ 调用函数的频率

性能优化的必要性：

➤简单的内核实现通常不会产生最佳性能。 性能调优工具可以帮助您查找代码中的关键区域，这些区域是性能瓶颈。
➤CUDA中的SM资源在多个驻留线程块中分分配。此分配可能会导致一些资源成为性能限制。 Profiling工具可以帮助您深入了解如何利用计算资源。
➤CUDA提供了硬件架构的抽象，使您能够控制线程并发性 。Profiling工具可以帮助您测量，可视化和指导您的优化。
nvvp：可视化性能分析工具
nvprof：命令行性能那分析工具
**注意：＊＊
1.很多性能指标都是针对的每个SM并不是整个GPU。
2.运行一次可能只会得到某些参数，多次运行可以收集完整。
3.多次运行的结果可能会不同。
考虑的因素：
1.存储器带宽
2.计算资源
3.指令和存储的时延

Warp的执行方式

当创建了一个kernel时，从逻辑上理解为kernel中的所有线程都在并行，但是从硬件物理条件上看同一时刻并不是所有的线程都在执行。。
Warp和线程块

Screenshot from 2017-04-26 19:18:01.png
warp是SM上的基本执行单元。warp一定是同一个block中的，如果一个block中的threads不足32个，则补足成为32个构成一个warp。
Screenshot from 2017-04-26 19:26:33.png
如图所示，本来只需要80个线程，但是实际上仍然需要32＊3=96个threads，尽管最后一个warp的16个线程没有使用，但是仍然会消耗SM上的资源，比如共享存储器、寄存器。
Warp分支
定义：一个warp中的线程执行不同的指令，叫做warp分支。
如果warp发生分支，则需要顺序执行每个分支路径。
Screenshot from 2017-04-26 19:36:30.png
在一个warp中所有线程都必须具有两个分支if...else....一个warp中如果有线程的条件为true，则执行if子句，其它为false的线程将等待if执行完成。然后执行else语句，当条件为true的线程则等待else执行完成。
为了获得更高的性能，尽量避免warp分支，warp是32个连续的线程，在算法允许的情况下，可以将数据分割，使同一个warp避免分支。
Example
实现偶数的线程计算结果为100，奇数线程的计算结果为200.

 // set up data size
    int size = 64;
    int blocksize = 64;

//线程分支
__global__ void mathKernel1(float *c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float ia, ib;
    ia = ib = 0.0f;
    if (tid % 2 == 0)
    {
        ia = 100.0f;
    }
    else
    {
        ib = 200.0f;
    }
    c[tid] = ia + ib;
}

//没有warp分支，设备利用率更高，计算结果相同，但是顺序不同。
__global__ void mathKernel2(float *c)
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    float ia, ib;
    ia = ib = 0.0f;
    if ((tid / warpSize) % 2 == 0)
    {
        ia = 100.0f;
    }
    else
    {
        ib = 200.0f;
    }
    c[tid] = ia + ib;
}

分支效率：

???不知道为什么，我的电脑运行结果很奇怪Tesla K80，反而是kernel1运行时间更短，kernel2运行时间更长。
warmingup:不分支
mathKernel1：分支
mathKernel2：不分支
mathKernel3：分支
mathKernel4：不分支
以前的nvprof计算warp分支的效率，但是我的CUDA8.0已经提示没有该metrics了和events。

$ nvprof --metrics branch_efficiency 
$ nvprof --events branch,divergent_branch 

Screenshot from 2017-04-26 20:44:59.png
Screenshot from 2017-04-26 20:47:37.png