大规模并行处理器编程实战笔记4

1：更多关于线程执行的问题

warp的一些限制：

对于那些大小不是32的整数倍的块，最后一个warp会用别的块中的线程来填充，直到满足32个线程为止

warp divergence的问题：

如果if-then-else结构中同一个warp中的线程不满足同一个条件，那么，一个warp中的线程有一部分通过then路径，一部分通过else路径，warp按照顺序来通过这些路径，导致执行时间的增加。

reduction中warp divergence的例子：

eg.

__shared__ float partialSum[]

unsigned int t = threadIdx.x;

for(unsigned int stride = 1; stride < blockDim.x; stride *= 2)

{

    __syncthreads();

    if(t % (2 * stride) == 0)

        partialSum[t] += partialSum[t+stride];

}

上述例子中每个warp中的线程在每一次计算中都要进行分支执行，减缓了执行的速度

eg.

__shared__ float partialSum[]

unsigned int t = threadIdx.x;

for(unsigned int stride = blockDim.x >> 1; stride > 0; stride >>= 1)

{

    __syncthreads();

    if(t < stride)

        partialSum[t] += partialSum[t+stride];

}

上述例子中不同的warp分别执行或者不执行加法，每个warp中线程的行为一致（当stride>32的时候），减少了分支的执行，速度加快

2：全局存储器带宽

DRAM的工作过程：

CUDA的global memory采用DRAM实现，而现代DRAM采用一个并行进程来提高数据访问的速度（由物理上的设计决定，如果每次只提取一个单元，会浪费很多资源），每次访问一个单元的时候，很多包括请求的单元在内的数据将以很快的速度传输到处理器中，所以，为了尽可能接近峰值，kernel函数必须安排数据的访问顺序，一边对于DRAM的每个请求都会访问DRAM中的大量连续单元。

warp对于DRAM的访问：

warp中所有线程执行同一条指令访问全局存储器中来连续的单元时，获得最有利的访问模式，硬件把所有这些访问结合或者合并成一个对DRAM连续单元的综合访问（比如：t1访问n，t2访问n+1，...）

合并访问：

行存储的矩阵需要按列进行访问，这样对于物理上连续的数据，将会被不同的线程连续合并访问，提高访问的效率，如果按照行进行访问，那么连续的线程将会对物理上的数据进行不连续的访问，无法合并，降低访问效率。所以，无论是直接访问global memory还是降global memory加载到shared memory再访问都要坚持合并访问。

3：SM资源的动态划分

eg.

GT200中，每个SM最多1024个线程，8个块，所以，可以分配4*256，也可以分配8*128，动态地分配资源，但是后者更加充分地利用了线程插槽和块插槽，效率更高

eg.

GT200中，每个SM拥有的寄存器只有8192个，如果每个线程占用的寄存器过多（声明的自动变量），那么一个SM上面能够运行的快就少了，如果每个线程10个寄存器，16*16的块，那么可以满足三个块：3*16*16*10 < 8192，但是如果每个线程11个寄存器，16*16的块，只能满足两个块：3*16*16*11 > 8192，这样可能会造成性能悬崖

eg.

增加寄存器虽然减少块的个数，但是也可能提高性能：

4条独立指令，每个指令4个时钟周期，global memory延迟为200个时钟周期，则隐藏延时需要的warp数目为：200/(4*4)=14

增加寄存器的使用，是独立指令数目上升到8，则隐藏延迟只需要warp数目为：200/(4*8)=7，则就算块数目从3减到2，warp数目从24减到16，也能够满足充分利用执行单元的要求，性能实际上还是提高了

4：数据预取

当所有的线程都在等待存储器访问结果的时候，访问时延将无法被容许，这个时候，可以采用方法：使用当前元素时预取下一个数据元素，增加在存储器访问和以访问的数据使用指令之间的独立指令数目

eg.

Code1：

Loop{

把当前块加载到shared memory中

__syncthreads();

计算当前块

__syncthreads();

}

Code2：

把global memory中的第一块加载到寄存器中

Loop{

把寄存器中的块存放在shared memory中

__syncthreads();

从global memory中把下一个块加载到寄存器中

计算当前块

__syncthreads();

}

Code2中的过程将Code1中的第一部分拆开，避免了长时间的等待，在数据从global memory加载到register，再从register加载到shared memory的这个原来是组合起来的过程拆开，并且在中间插入了计算指令，有效地隐藏了时延（但是增加了两个额外的自动变量，也就是寄存器）

5：指令混合

一个循环里面的指令常常混合了多种不同类型的指令：浮点运算，循环分支，地址运算，计数

通过将循环展开成一个长度的表达式可以消除这种限制，提高执行速度

eg.

混合指令：

for(int k = 0; k < BLOCK_SIZE; ++k){

    pvalue += Mx[ty][k] * Ns[k][tx];

}

消除混合的指令：

pvalue += Ms[ty][0]*Nx[0][tx] + ... + Ms[ty][15]*Ns[15][tx];

6：线程粒度

由于每个线程块计算的时候都要重复加载矩阵A中的块，如果让每个线程块计算C的水平方向上的相邻几块，就可以减少对A同一块的重复加载，这样可以增大线程的粒度，有时候可以提升性能，因为增加处理块的数目将会造成每个块使用更多的register和shared memory，也可能造成性能下降，所以，需要找一个平衡点

7：可度量的性能和小结

技术：分块，循环展开，数据预取，线程粒度

1：块大小在性能中其主要作用

2：块大小已经足够的情况下，循环展开和数据预取将更重要

3：对于线程粒度而言，如果使用得register超过了SM中寄存器总数，那么将会影响到其他优化的发挥

4：各种优化的调整技术互相影响，要不断结合找到最好的优化方法