最近在学习GPU并行计算，对针对全局内存的并行归约和共享内存的并行归约的理解做一个小总结。以下代码出自<<CUDA C编程权威指南>>。
归约与线程分支分化
线程束分化
- 一个线程束（32个线程）在同一周期中必须执行相同的指令；
- 如果一个线程束分化（如tid为偶数时执行一个分支，tid为奇数时执行另一个分支），那么则整个线程束将会连续执行每一个分支路径，在执行时不满足这个分支的线程则会被禁用（这就浪费了线程)。
归约中的线程束分化
- 如上图所示，在进行相邻配对中，执行一步后，则只有tid=0，2，4，6...等偶数线程保存了下一步的值，并需要执行下一步归约，下一步则只有tid=0，4，8等间隔更大的线程执行归约。这就造成了每进行一次归约，线程束分化中有效的线程数越少
- 另外，这样也会影响到全绝内存的加载（见第4章：全局内存），由于内存加载时是32个相邻内存单元（对应于32个线程）同时加载（算一个内存事务）。如果不相邻则需要多个内存加载事务，上面的线程束分化会造成内存加载的多个事务。
线程束分化的归约示例
相邻配对的归约
造成线程束分化的相邻配对归约
// Neighbored Pair Implementation with divergence
__global__ void reduceNeighbored(int* g_idata, int* g_odata, unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
// convert global data pointer to the local pointer of this block
int* idata = g_idata + blockIdx.x * blockDim.x;
// boundary check
if (idx >= n) return;
// in-place reduction in global memory
for (int stride = 1; stride < blockDim.x; stride *= 2)
{
if ((tid % (2 * stride)) == 0)//造成了线程束的分化
{
idata[tid] += idata[tid + stride];
}
// synchronize within threadblock
__syncthreads();
}
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = idata[0];
}
上面的 if ((tid % (2 * stride)) == 0)语句会造成线程束的分化，stride越大，分化越严重。
消除线程束分化的相邻配对的规约
__global__ void reduceNeighboredLess(int* g_idata, int* g_odata,
unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
// convert global data pointer to the local pointer of this block
int* idata = g_idata + blockIdx.x * blockDim.x;
// boundary check
if (idx >= n) return;
// in-place reduction in global memory
for (int stride = 1; stride < blockDim.x; stride *= 2)
{
// convert tid into local array index
int index = 2 * stride * tid;
if (index < blockDim.x)
{
idata[index] += idata[index + stride];//而这里stride为定值，index为非连续值，内存加载会出现多次加载，效率低
}
// synchronize within threadblock
__syncthreads();
}
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = idata[0];
}
注意如上面代码中中文注释部分，由于不需要进行if判断，故未造成线程束分化。通过测试每个线程束执行的指令束平均值可知，未线程束分化的函数拥有更少的平均指令数。
但是相邻配对本身仍然存在缺陷由于线程id不连续，会需要多个内存加载事务，因而会使得速度变慢。
这里比如stride=1时，需要加载的内存的线程id为0，2，4，6，8，10... 等偶数，而内存加载则是加载0-31号线程的所有内存，造成了内存加载的浪费，越往后归约浪费的越多。交错配对的归约则解决了这个问题。代码如下：
交错配对的归约
__global__ void reduceInterleaved(int* g_idata, int* g_odata, unsigned int n)
{
// set thread ID
unsigned int tid = threadIdx.x;
unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
// convert global data pointer to the local pointer of this block
int* idata = g_idata + blockIdx.x * blockDim.x;
// boundary check
if (idx >= n) return;
// in-place reduction in global memory
for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
{
if (tid < stride)
{
idata[tid] += idata[tid + stride];
//这里stride为定值，
//tid连续，所以32个内存事务是连续的，内存加载连续，所以即使吞吐量比上一个小，其效率 高，无冲突，因此快
}
__syncthreads();
}
// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = idata[0];
}
而交错配对的归约则不然，内存加载事务是连续的，提高了指令的效率。从下图4中可以看出交错配对（reduceInterleaved)的内存吞吐量相较于前两个函数是降低的，但是其内存加载的都是需要用的有效内存，因而效率高，从执行时间来看其执行速度也更快。
循环展开
由以上知道，我们通过改善：
- 减少线程束分化；
- 减少内存事务的加载(增加内存的有效加载）
获得了运行速度的增加。但是还有没有更近一步的提高速度的手段呢？有，那就是减少指令消耗（展开循环）增加更多的独立指令调度(独立的内存加载事务等）来提高性能。这样能提高速度的原因是能够通过以上此举使得更多的线程达到可以符合执行条件的状态（指令就位或者内存数据就位），从而帮助隐藏指令延迟或者内存延迟。
在优化开始前，需要了解下延迟隐藏。
延迟隐藏
延迟隐藏包括算术指令隐藏和内存指令隐藏
延迟的定义
- 算术指令延迟是一个算术操作从开始到它产生输出之间的时间;
- 内存指令延迟是送出的加载或存储操作和数据到达目的地之间的时间。
所需要的线程数量有个计算公式：
所需线程数量=延迟*吞吐量
算术运算所需的并行
如对于float算术运算(a+bc)，Fermi gpu模型需要20个指令延迟，而一条指令是对于一个线程束而言的，即32个操作，因此在每个SM中需要的线程束数量为：2032/32=20个线程束。
对于Turing架构来说，每个sm上有64个cuda核，则需要20个线程束。
因此提高指令级并行，有如下两个方法：
- 指令级别并行：一个线程中有很多独立的指令；
- 线程级并行：很多并发的符合条件的线程。
循环展开减少了指令的执行，我的理解是循环展开一次指令操作了更多的数据块，从而减少了总的指令执行次数，从而加快了速度。
内存所需要的并行
内存所需的并行可以表示为在每个周期内隐藏内存延迟所需的字节数。
以答主的RTX 2070显卡为例，其吞吐量448GB/s,指令延迟800（假设），内存频率：7001MHz。
则把带宽转换为字节每周期为：
448GB/s/(7001MHz/1024)=66字节/周期
那么，用字节/周期(吞吐量）乘以延迟则得到需要的内存并行，即800*66/1024=51KB。
假如一个线程操作的数据为4字节大小，那么共需要的线程数量为：
51KB/4字节/线程=13105线程=13105/32=410个线程束，对于RTX2070来说，共有36个SM，则每个SM上分配的线程束为410/36=11.4线程束/SM。
而对于循环展开来说，展开的循环进行了多个独立的内存加载操作，我的理解是此时能够尽可能的利用带宽，同时传输更多的数据。书上（p152页）说展开并不影响内存操作的数量，但是会增大并发执行的数量，从而加速了运行。
对于以上计算，可以看出对于2070显卡来说，算术延迟每个sm要20个线程束，内存延迟要11.4个线程束，所以在以上假设情况下取其最大值，即20个线程束，即每个块的大小需要大于640个线程才能隐藏算术延迟（20个指令延迟）。
基于全局内存的展开
//todo
基于共享内存的展开
//todo
--2020.11.29
