缓存在处理器的核或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）中的一种技术，避免了机器指令或者微操作不必要的顺序化执行，从而提高了处理器的指令级并行的能力。

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

__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];
    }
}