CUDA 内存处理

第六章 CUDA内存处理

6.1 高速缓存

G80 与 GT200 系列没有与CPU中高速缓存等价的存储器。在Fermi架构的GPU实现中，第一次引入了不基于程序员托管的数据缓存这个概念。这个架构的GPU中每个SM有一个一级缓存，这个一级缓存既是基于程序员托管的又是基于硬件托管的。在所有的SM之间有一个共享的二级缓存。

缓存在处理器的核或SM之间共享有什么意义？

• 主要是为了让设备之间能够通过相同的共享缓存进行通信。共享缓存允许处理器之间不需要每次都通过全局内存进行通信。这在进行原子操作的时候特别有用，由于二级缓存是统一的，所有的SM在给出的内存地址处获取一致版本的数据，处理器无须将数据写回缓慢的全局内存中，然后再读出来，它只需要保证处理器核之间的数据一致性。

GPU提供了不同层次的若干区域供程序员存放数据

• 大多数CUDA程序都是逐渐成熟的。一开始使用全局内存初始化，初始化完毕之后再考虑使用其他类型的内存，例如，零复制内存、共享内存、常量内存、最终寄存器也被考虑进来。
• 为了优化一个程序，我们需要在开发过程中思考这些问题。在程序之初就要考虑使用速度较快的存储器，并且准确知道在何处以及如何提高程序性能，而不是在程序写完之后才想到用哪些快速的存储器对程序进行优化。
• 另外，不仅要思考如何高效地访问全局内存，也要时刻想办法减少对全局内存的访问次数，尤其在数据会被重复利用的时候。（可见上一篇文章的最后直方图优化的过程）

6.2 寄存器的用法

一般地，CPU每个核会支持一到两个硬件线程。相比之下，GPU的每个SM可能有8~192个SP，这意味着每个SM在任何时刻都能同时运行这些数目的硬件线程。
CPU 与 GPU 架构的一个主要区别就是CPU 与 GPU 映射寄存器的方式。

CPU

• CPU 通过使用寄存器重命名和栈来执行多线程。为了运行一个新任务，CPU需要进行上下文切换，将当前所有寄存器的状态保存到栈（系统内存）上，然后从栈中恢复当前需要执行的新线程上次的执行状态。这些操作通常需要花费上百个CPU时钟周期。 如果在CPU上开启过多的线程，时间几乎都花费在上下文切换过程中寄存器内容的换进换出上。因此，如果在CPU上开启过多的线程，有效工作的吞吐量将会快速降低。

寄存器重命名是计算机CPU的微体系结构（Microarchitecture）中的一种技术，避免了机器指令或者微操作不必要的顺序化执行，从而提高了处理器的指令级并行的能力。

GPU

• 然而，GPU却恰恰相反。GPU利用多线程隐藏了内存获取与指令执行带来的延迟。因此，在GPU上开启过少的线程反而会因为等待内存事务使GPU处于闲置状态（因为在GPU上线程切换的代价很小，所以当线程束被挂起时，立马调度下一个线程束）。
• 此外，GPU也不使用寄存器重命名的机制，而是致力于为每一个线程都分配真实的寄存器(因为其寄存器数量很多)。因此，当需要上下文切换的时候，所需要的操作就是将指向当前寄存器组的选择器（或指针）更新，以指向下一个执行的线程束的寄存器组，因此几乎是 零开销 。

在SM层，线程块即若干个独立线程束的逻辑组。编译时会计算出每个内核线程需要的寄存器数目(寄存器总数➗线程块数量➗每个线程块的线程数量)。所有的线程块都具有相同的大小，并拥有已知数目的线程，每个线程块需要的寄存器数目也就是已知和固定的。因此，GPU就能为在硬件上调度的线程块分配固定数目的寄存器组。

• 如果一个内核函数中的每个线程需要的寄存器过多，则每个SM中GPU能够调度的线程块的数量就会收到限制，因此总的可以执行的线程数量也就会受到限制。
• 开启的线程数量过少会造成硬件无法被充分利用，性能急剧下降，但开启过多又意味着资源可能短缺，调度到SM上的线程块数量就会减少。

牢记，每个线程中每个变量会占用一个寄存器。因此，C语言中的一个浮点型变量就会占用N个寄存器，其中N代表调度的线程数量。

如果能够最大化地利用寄存器，例如，使用寄存器对一个数组的某一块进行计算，会非常高效。由于这一组值通常是数据集中的N个元素，元素之间是相互独立的，因此可以在单个线程中实现指令级的并行(ILP)。这是由硬件将许多独立的指令流水线化实现的。

一个例子

一个循环根据一些布尔变量依次设置某个值的每一位。高效的方法是将32个bool封装到一个32位的字中(此例子为pack_array？ 此处有点不解pack_array数组的元素类型是什么？每一个都是int，都是32个bool的封装？还是该数组元素的类型为bool？)，然后解封装。可以这样写一个循环，每次根据新的bool修改内存中的内容，做移位操作，移至字中正确的位置。

// 从数组中读出第i个元素，然后将其封装到一个整型数 packed_result 中
int packed_result;
for (int i = 0; i < 31; i++) {
  packed_result |= (pack_array[i] << i);
}

• 如果变量packed_result存于内存中，则需要做32次读/写内存的操作。但如果将变量packed_result设置为局部变量，编译器会将其放入寄存器中，在寄存器中而不是在主内存中做操作，最后再将结果写回主内存中，因此可节省31次内存读/写的操作。

在sum，min，max等普通的归约操作中也会看到类似的关系。所谓归约操作，即利用函数将某个较大的数据集减少为较小的集合，通常减少到一个单项。将结果累计在寄存器中可省去大量的内存读写操作。

• 寄存器版本

__global__ void test_gpu_register(u32 * const data, const u32 num_elements){
    const u32 tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(tid < num_elements){
        u32 d_tmp = 0;
        for(int i=0;i<KERNEL_LOOP;i++){
            d_tmp |= (packed_array[i]<<i);
        }
        data[tid] = d_tmp;
    }
}

• 全局内存版本

__devicd__ static u32 d_tmp[NUM_ELEM];
__global__ void test_gpu_register(u32 * const data, const u32 num_elements){
    const u32 tid = (blockIdx.x * blockDim.x) + threadIdx.x;
    if(tid < num_elements){
        for(int i=0;i<KERNEL_LOOP;i++){
            d_tmp[tid] |= (packed_array[i]<<i);
        }
        data[tid] = d_tmp[tid];
    }
}

显卡	寄存器版本(ms)	GMEM版本(ms)	加速比
GTX470	0.26	0.51	2
9800GT	9.27	10.31	1.1
GTX260	0.62	1.1	1.8
GTX460	0.34	0.62	1.8
平均			1.7