待解决的问题：代码3-1中运行时间为0。

第三章 CUDA执行模型

3.1 CUDA执行模型概述

3.1.1 GPU架构概述

核心：流式多处理器SM
每个GPU有多个SM，每个SM支持数百个线程并行执行
其中，共享内存和寄存器是非常重要的资源，共享内存分配在常驻block中，寄存器在thread中被分配，活跃的warp数量对应SM上的并行量
当启动一个gird时，block被分配在SM上执行，多个block可以被分配到同一个SM，一个SM可以容纳多个block

线程束warp+单指令多线程SIMT：每32个线程为一组，所有线程执行相同的指令，每个线程有自己的指令计数器和寄存器状态，利用自身的数据执行当前指令，每个SM将划分给它的block划分到warp里，然后在硬件上调度执行

在并发的warp之间切换没有开销，因为硬件资源分配到了所有block和thread中，最新被调度的warp的状态以及存储在SM上

SIMT（单指令多线程）与SIMD（单指令多数据）的一异同：
同：多个单元执行相同的指令
异：SIMD要求同一个向量的所有数据在统一的同步组里面一起执行
SIMT允许统一warp的多个线程独立执行，因此单独的线程可能有不同的行为，实现线程级并行

软硬件的对应：
线程——cuda核
block——SM
grid——设备

3.1.2 Fermi架构

特点：

• 16个SM*每个上32个core核=512个加速核心，每个核都有整数ALU和浮点数运算单元FPU，每个时钟周期执行一个整数/浮点数指令
• 6GB全局内存，全局调度器GigaThread引擎分配线程块到SM线程束调度器上
• 64KB的片内可配置存储器：在共享内存和一级缓存之间

每个SM包含：

• 32个cuda核
• 调度warp的调度器+调度单元
• 共享内存+寄存器+一级缓存
• 16个加载/存储单元——>每个时钟周期可以有16个线程（warp的一半）计算源地址和目的地址
• 4个特殊功能单元SFU：执行固有指令
• 2个线程束调度器+2个指令调度单元：选择2个warp，把1个指令从warp中发射到一个组上
• 可同时处理48个warp

3.1.3 Kepler架构

创新点：

• 强化SM
• 动态并行
• Hyper-Q技术

15个SM+6个64位内存控制器

1、每个强化SM包含：

• 192个cuda核+64个双精度单元+32个特殊功能单元SFU+32个加载/存储单元LD/ST
• 4个线程束调度器+8个指令调度单元：一个SM上可同时发送执行4个warp

2、动态并行
运行GPU动态启动新的grid，每个内核都能启动新的内核，便于递归

3、Hyper-Q技术
原来的Fermi架构只有一个硬件工作队列在CPU和GPU之间传送任务，如果有个任务阻塞，该任务后的其他任务也被阻塞。
Hyper-Q技术提供了32个队列，可以在GPU上有更多的并发执行。

3.1.4 配置文件驱动优化

性能分析工具：nvvp独立的可视化分析器+nvprof命令行分析器
事件：可计算的活动，对应内核执行期间被收集的硬件计数器，通过SM报告
指标：内核的特征，由一个或多个事件计算得到

常见的限制内核性能的因素：①存储带宽②计算资源③指令和内存延迟（本章主要）

3.2 理解线程束执行的本质

3.2.1 线程束和线程块

warp是SM中的基本执行单元，一个warp中的所有线程执行相同的指令，在线程对应的私有数据上操作

连续值的线程被分配到warp中
一维线程块：线程id=threadIdx.x;
二维线程块：线程id=threadIdx.y * threadDim.x + threadIdx.x;
三维线程块：线程id=(threadIdx.z * threadDim.y * threadDim.x)+ (threadIdx.y * threadDim.x) + threadIdx.x;

一个block中warp的数量：block中的线程数/warp大小
如果block大小不是warp大小的整数倍，最后的warp部分线程不活跃，虽然不活跃，仍然消耗SM的资源，eg：寄存器

逻辑角度：block是thread的集合，可以是一维、二维、三维
硬件角度：block是一维warp的集合

3.2.2 线程束分化

GPU不像CPU有分支预测机制，一个warp中的所有thread在同一周期执行相同的指令。
线程束分化：同一线程束中的线程执行不同的指令。eg：warp中的线程碰到if else语句时，一半执行if语句块的指令，另一半执行else语句块的指令。
线程束分化会导致性能明显地下降，上例中只有16个线程同时活跃，条件分支越多，并行性削弱越严重。
线程束分化只发生在同一个线程束内。
为了避免线程束分化，应该避免同一线程束有不同的执行路线。例如：

if(tid % 2 == 0){
	代码块
}else{
	代码块
}

这段代码会导致线程束分化，一个线程束中，奇数号线程执行else内的代码，偶数号线程执行if内的代码，下面的代码对这种情况进行了优化。

if((tid / warpSize) % 2 == 0){
	代码块
}else{
	代码块
}

比如(0/32)%2=0，(1/32)%2=0，直到(31/32)%2=0，这是第一个warp块，前32个线程，也就是0号warp块都执行if内的代码。
从第32个线程开始，(32/32)%2=1，(33/32)%2=1，直到(63/32)%2=1，这是第二个warp块，第二个32个线程，也就是1号warp块都执行else内的代码。
这样一个warp块内都执行相同的代码，就不会出现线程束分化，使每个线程束的利用率都达到100%；奇数编号的warp块都执行else代码，偶数编号的warp块都执行if代码。

正确代码3-1 simpleDivergence.cu【有问题】

代码3-1遇到的问题

报错1、执行nvcc编译后的文件时，耗费时间为0；复制了官方的代码之后，运行时间还是0

$ ./3-1simpleDivergence
./3-1simpleDivergence using Device 0: NVIDIA GeForce RTX 3090
Data size 64Execution Configure (block 64 grid 1)
warmup           <<<   1   64>>> elapsed 0 sec 
mathKernel1 <<<   1   64>>> elapsed 0 sec 
mathKernel2 <<<   1   64>>> elapsed 0 sec 
mathKernel3 <<<   1   64>>> elapsed 0 sec 
mathKernel4 <<<   1   64>>> elapsed 0 sec 

用以下命令查看分化情况：

nvprof --metrics branch_efficiency ./执行的程序

分支效率：未分化的分支/全部分支
一般来说，分支效率都是100%，也就是说，没有分支分化，这是因为CUDA编译器的优化，把短的、有条件的代码段的断定指令取代了分支指令。

分支预测：把断定变量设置成1或0，为1的执行，为0的不执行，但相应的进程也不会停止。与实际的分支指令有区别。
很长的代码路径肯定会导致线程束分化，因为当条件语句的指令数小于某个阈值时，编译器才用断定指令替换分支指令。

3.2.3 资源分配

SM负责处理warp本地执行上下文（PC+寄存器+shared memory），整个生命周期保存在芯片上，因此切换没有损失。
每个SM有32为寄存器组，在寄存器文件中，给thread进行分配；shared memory在block中进行分配。
对于一个kernel来说，决定在一个SM中可用的block和warp的数量取决于：可用的 寄存器+shared memory 的数量。

block+分配资源（限制条件）=活跃的block—包含—>活跃的warp：
①选定的warp：在执行单元执行
②阻塞的warp：没有准备好
③符合条件的warp：准备好但还未执行

warp可执行的条件：
①32个kernel可用于执行
②所有参数就绪

3.2.4 延迟隐藏

1、指令延迟
利用率与常驻warp的数量有关。
指令延迟：从发出到完成的时钟周期。
如果每个时钟周期所有的线程调度器都有符合条件的warp时，计算资源完全利用。在其他常驻warp中发布其它指令，隐藏每个指令的延迟（类似指令流水线）。
估算隐藏延迟所需的活跃warp的数量：延迟*吞吐量
带宽和吞吐量的区别：带宽是理论上的最大值；吞吐量是实际的值，不一定会达到最大。

指令类型：

• 算数指令：从开始到输出的时间，10~20个周期。
• 内存指令：从发出指令到数据到目的地的时间，400~800个周期。

吞吐量由SM中每个周期的操作数量确定，执行一条指令的一个warp有32个线程操作。
提高并行的方法：

• 指令级并行ILP：一个thread有多条独立指令
• 线程级并行TLP：多个符合条件的thread
  延迟隐藏取决于每个SM活跃的warp数，由执行配置和资源约束 隐式决定。

2、内存延迟
对内存来说，所需的并行为每个周期内隐藏内存延迟所需的字节数，这是对于整个设备而言，而不是对于一个SM而言。
与指令延迟类似，内存提高并行的方法：①一个thread有多个独立的内存操作 ②创建更多并发的thread/warp。

3.2.5 占用率

占用率=活跃warp数/最大warp数
检测设备中每个SM最大warp数的函数：

cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop, int device);

设备中各种数据在cudaDeviceProp结构中返回，每个SM中thread的最大值在maxThreadPerMultiProcessor变量中返回，该值/32即为最大warp数。

正确代码3-2 simpleDeviceQuery.cu

#include <stdio.h>
#include <cuda_runtime.h>

int main(int argc, char *argv[]){
    int iDev = 0;
    cudaDeviceProp iProp;
    cudaGetDeviceProperties(&iProp, iDev);

    printf("Device %d: %s\n",iDev,iProp.name);
    printf("Number of multiprocessors:%d\n",iProp.multiProcessorCount);
    printf("Total amount of constant memory: %4.2f KB\n",iProp.totalConstMem/1024.0);
    printf("Total amount of shared memory: %4.2f KB\n",iProp.sharedMemPerBlock/1024.0);
    printf("Total number of registers available per block: %d\n",iProp.regsPerBlock);
    // printf("Warp size%d\n",deviceProp.warpSize);//按照书上敲的，报错deviceProp未定义
    printf("Warp size:%d\n",iProp.warpSize);
    printf("Maximum number of threads per block: %d\n",iProp.maxThreadsPerBlock);
    printf("Maximum number of threads per multiprocessor: %d\n",iProp.maxThreadsPerMultiProcessor);
    printf("Maximum number of warps per multiprocessor: %d\n",iProp.maxThreadsPerMultiProcessor/32);
    return EXIT_SUCCESS;
}

CUDA占用率计算器
查看每个线程的寄存器和每个块的共享内存情况指令：–ptxas-options=-v（链接: 使用参考）
手动设置告诉寄存器每个线程使用的寄存器上限NUM：-maxrregcount=NUM

常用准则：
每个block中thread是32（warp中的thread）的倍数
避免块太小，每个块至少128或256个thread
block数 >> SM数，保证有足够的并行
通过实验得到最佳配置和资源使用情况

占用率注重的是每个SM中并发thread和warp的数量，充分的占用率不是唯一目标，一旦达到一定级别，占用率再增加也不会改进性能，这时候应该调整其他因素。

3.2.6 同步

栅栏同步原语可以在两个级别执行：

• 系统级：等主机和设备完成所有的工作
• 块级：执行中等待block中的所有thread到达同一点

对主机来说，许多CUDA API调用和主机启动不是同步的，用cudaDeviceSynchronize函数阻塞主机应用程序，等待CUDA操作完成。

一个block的warp以未定义的顺序被执行，可以使用及局部栅栏同步它们的执行，在kernel中标记同步点。

__device__ void __syncthreads(void);

__syncthreads被调用时，同一个block中的每个thread都需要等其他线程执行到这个同步点。
栅栏之后的block中所有thread可以看见栅栏前所有thread的全局内存和共享内存，可以用来协调通信（同一个block中的所有thread）。
同一个block中的thread可以通过共享内存和寄存器共享数据；不同block间的thread不能同步。
不同block之间没有线程同步。块间同步的唯一方法是在每个kernel结束端使用全局同步点，终止当前kernel，开启新kernel。

3.2.7 可扩展性

可扩展性：为并行程序提供的额外的硬件资源，从而加速
透明可扩展性：在可变数量的计算kernel上执行相同的应用程序代码