《CUDA并行程序设计-GPU编程指南》读书笔记--(2)CUDA内存处理

CUDA内存处理

高速缓存

---

不仅要思考如何高效地访问全局内存，也要时刻想办法减少对全局内存的访问次数，尤其在数据会被重复利用的时候。

---

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

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

---

如果一个内核函数中的每个线程需要的寄存器过多，在每个SM中GPU能够调度的线程块的数量就会受到限制，因此总的可以执行的线程数量也会受到限制。

---

for(i=0;i<31;<i++){    packed_result |= (pack_array[i]<<i);}

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

寄存器版本：

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

共享内存

GPU执行的是一种内存的加载/存储模型(load-store model)，即所有的操作都要 在指令载入寄存器之后才能执行。因此，加载数据到共享内存与加载数据到寄存器中不同，只有当数据重复利用、全局内存合并，或线程之间有共享数据时使用共享内存才更合适。

__device__ void merge_array(){    __shared__ u32 list_indexes[MAX_NUM_LISTS];    // do something}

此处GPU内核是以一个设备函数的形式编写的。设备函数即只能被GPU内核调用的函数。它相当于C语言函数声明之前添加一个”static”，或C++中的”private”。

常量内存
常量内存其实只是全局内存的一种虚拟地址形式，并没有特殊保留的常量内存块。常量内存有两个特性，一个是高速缓存，另一个是它支持将单个值广播到线程束中的每个线程。

常量内存是只读内存。这种类型的内存要么是在编译时声明为只读内存，要么是在运行时通过主机端定义为只读内存。常量只是从GPU内存的角度而言。常量内存的大小被限制为64K。

__constant__ float my_array[1024] = {0.0F,1.0F,2.0F,...};

如果要在运行时改变常量内存区中的内容，只需在调用GPU内核之前简单地调用cudaCopyToSymbol函数。

如果一个常量只是字面值，那么最好用#define对字面值进行定义，因为这样可以减少常量内存的使用。所以，尽量不要把诸如PI这样的字面值放到常量内存中，而是用#define进行宏定义。事实上，这两种方案只是使用了不同的内存，但在速度上会产生一定影响。

注意cudaMemcpyToSymbol函数的工作原理。该函数可以将数据复制到GPU上任何以全局符号命名的内存区域，无论该符号是全局内存还是常量内存。因此，我们可以将一块64K大小的数据块复制一个64K大小的内存区上，从而通过常量内存缓存进行访问。当所有线程访问同一数据元素时，这种访问方式非常有用，因为我们可以借助缓存技术从常量内存区获取数据然后再广播到每个线程中。

全局内存
CPU主机端处理器可以通过以下三种方式对GPU上的内存进行访问：

• 显式地阻塞传输；
• 显式地非阻塞传输；
• 隐式地使用零内存复制。

---

通常的执行模型是CPU将一个数据块传输到GPU，GPU内核对其进行处理，然后再由CPU将数据块传输回主机端内存中。比较高级的模型是使用流(稍后将进行介绍)将数据传输和内核执行部分重叠，以保证GPU一直在工作。

---

所谓合并访问就是所有线程访问连续的对齐的内存块。假定以基于字节的方式对数据进行访问，此处显示的Addr即从基地址算起的逻辑偏移地址。TID表示线程标号。如果我们对内存进行一对一连续对齐访问，则每个线程的访问地址可合并起来，只需一次存储事务即可解决问题。假设我们访问一个4字节的内存块。内存会基于线程束的方式进行合并(在老式的G80硬件上使用半个线程束)，也就是说访问一次内存将得到32x4=128个字节的数据。

合并大小支持32字节、64字节以及128字节，分别表示线程束中每个线程以一个字节、16位以及32位为单位读取数据，但前提是访问必须连续，并且是以32字节为基准对齐的。

---

将标准的cudaMalloc替换成cudaMallocPitch，使用这种特殊的分配内存指令可以得到对齐的内存块。

extern __host__ cudaError_t CUDARTAPI cudaMallocPitch(void **devPtr,            size_t *pitch, size_t width, size_t heigth);

cudaMallocPitch的第一个参数表示指向设备内存指针的指针，第二个参数表示指向对齐之后每行真实字节数的指针，第三个参数为需要开辟的数组的宽度，单位为字节，最后一个参数为数组的高度。

---

对齐的访问将导致多次内存获取。当等待内存获取时，线程束中的所有线程将阻塞直到所有的内存获取都从硬件中返回。因此，为了获得最好的吞吐量，我们需通过对齐并且连续合并的访问方式，将大量的内存获取请求合并，达到减少内存获取次数的效果。

typedef struct{    int a;    int b;    int c;    int d;} MY_TYPE_T;typedef INTERLEAVED_T MY_TYPE_T[1024]; typedef int ARRAY_T[1024];typedef struct{    ARRAY_T a;    ARRAY_T b;    ARRAY_T c;    ARRAY_T d;}NON_INTERLEAVED_T;

首先是INTERLEAVED_T，它表示一个结构体数组，结构体中的成员为a、b、c和d。然后我们声明了NON_INTERLEAVED_T作为一个结构体，结构体中的每个元素是一个数组，分别包含了4个数组，a、b、c和d。如类型名所示，第一种类型的数据在内存中以交错的方式分布，而第二种类型的数据则是以连续的方式在内存中分布。首先，我们将对结构体数组的元素进行加和运算，然后对结构体中的数组进行加和运算。

通过Parallel Nsight这样的分析工具我们可以更深入地理解为什么交错的访问方式如此慢而非交错的访问方式却要快得多。在Parallel Nsight中我们可以看到，非交错版本中存储事务(CUDA内存统计实验)的数量大约为交错版本的1/4，这意味着非交错版本中内存数据读/写操作次数只为交错版本的1/4。

另一个比较有趣的现象是CPU版本的执行时间恰好相反。这看起来很奇怪，但如果了解了访问方式以及缓存重用，你就不会这么觉得了。在交错访问的例子中，CPU访问元素a的同时会将结构体中元素b、c以及d读入缓存中，使它们在相同的缓存行中。然而，非交错版本则需要对4个独立的物理内存区进行访问。这意味着存储事务的数目为交错版本的4倍，并且CPU使用的任何预读策略都不会起作用。