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

1：存储器访问效率的重要性

提高CGMA（Compute to Global Memory Access）比值：

示例：

__global__ void MatrixMultiplication_kernel(flaot * Md, float * Nd, flaot *   Pd, int width)

{

//compute row index of Pd and Md

int Row = blockIdx.y * TILE_WIDTH + threadIdx.y;

//compute column index of Pd and Nd

int Col = blockIde.x * ITLE_WIDTH + threadIdx.x;

float pValue = 0;

//every thread compute one element of subMatrix

for(int k = 0; k < Width; k++)

pValue += Md[Row * Width + k] * Nd[k * Width + Col];

pd[Row * Width + Col] = pValue;

}

上述kernel函数中最重要的部分是执行点积运算的for循环的执行时间。

循环的每一次迭代都要访问两次全局存储器，一次执行浮点乘法，一次执行浮点加法，从而浮点运算与全局存储器访问操作之比为1:1（或者说比值为1.0）。

CGMA比值能够比较准确地体现CUDA中kernel函数的性能。

重要性：

例如：G80的全局存储器访问带宽为86.4GB/s，每个单精度浮点数占4B，因此对单精度数据的加载速度不会超过21.6GB/s（86.4/4），那么，如果CGMA比值为1.0，则矩阵乘法中的kernel函数每秒可以执行的浮点操作不会超过21.6Gflops（因为每个浮点操作用到的单精度数据都要访问全局存储器一次才能得到），然而，相对于G80的峰值性能367Gflops来说，这只是很小的一部分，所以，应该尽量提高CGMA比值。

2：CUDA设备存储器的类型

声明CUDA变量时，通过对CUDA存储器类型的选择，可以指定对应变量的可见性和访问速度：

设备代码可以：

——读写每个线程的寄存器

——读写每个线程的局部存储器

——读写每个block的共享存储器

——读写每个grid的全局存储器

——只读每个grid的常数存储器

主机代码可以：

——在每个grid的全局存储器和常数存储器中来回传输数据

CUDA变量类型限定符

变量声明                           存储器             作用域         生命周期

数组以外的自动变量                 寄存器           线程           kernel函数

自动数组变量                   局部存储器       线程         kernel函数

__device__，__shared__，int SV；共享存储器       block         kernel函数

__device__，int GV；             全局存储器       grid         应用程序

__device__，__constant__，int CV；常数存储器     grid         应用程序

前面的tx，ty，pValue都属于自动变量，当线程终止时，它们不再存在，而自动变量数组很少使用。

共享变量声明前有关键字__shared__，也可以加上可选的__device__，通常驻留kernel或设备函数中，为一个block中所有线程共享。

常数变量声明前有关键字__constant__，也可以加上可选的__device__，作用域是所有的grid，声明周期是整个应用程序，存放在全局存储器中，单采用缓存提高访问效率。

如果一个变量声明前只有__device__，那么它是一个全局变量且存放在全局存储器中，访问速度慢，对所有kernel函数的所有线程可见，也常用于调用一个kernel函数和另一个kernel函数时传递信息。

3：减少全局存储器流量的一种策略

通过分块来实现减少对全局存储器的访问次数，取而代之的是共享存储器，从而提高访问效率，直接上代码，更直观：

__global__ void MatrixMultiplication_kernel(float * Md, float * Nd, float * Pd, int Width)

{

__shared__ flaot Mds[TILE_WIDTH][TILE_WIDTH];

__shared__ float Nds[TILE_WIDTH][TILE_WIDTH];

int bx = blockIdx.x;

int by = blockIdx.y;

int tx = threadIdx.x;

int ty = threadIdx.y;

int Row = by * TILE_WIDTH + ty;

int Col = bx * TILE_WIDTH + tx;

float pValue = 0;

for(int m = 0; m < Width/TILE_WIDTH; m++)

{

Mds[ty][tx] = Md[Row * Width + (m * TILE_WIDTH) + tx];

Nds[ty][tx] = Nd[(m * TILE_WIDTH + ty) * Width + Col];

__syncthreads();

for(int k = 0; k < TILE_WIDTH; k++)

pValue += Mds[ty][k] * Nds[k][tx];

__syncthreads();

}

Pd[Row * Width + Col] = pValue;

}

以上对于N*N大小的block可以加速N倍。

解析：每次复制Md和Nd中对应相乘的那个block的元素到共享存储器，然后通过访问共享存储器算出对应block乘积的中间过程，遍历次过程可以积累求出Pd对应block中的每一个元素的值，但是仅仅访问全局存储器TILE_WIDTH（block维度）次，比起直接访问全局存储器需要TILE_WIDTH*TILE_WIDTH次速度提高了许多。

4：存储器是一个限制并行的因素

寄存器，共享存储器，常数寄存器都可以有效地减少对全局存储器的访问，但是，如果每个线程存放了过多的单元在这些地方，将会导致每个SM上可以驻留的线程数很少（因为SM必须保证每个线程有足够的执行资源），整个处理器中驻留的线程数就很少，不利于并行。

5：小结

CUDA定义的寄存器，共享存储器和常数存储器与全局存储器相比，可以提供更快且并行度更高的访问方式。

通过重新设计算法可以有效地利用这些存储器（例如矩阵乘法分块，使用共享存储器提速）。

另一方面，又不能过多地使用这些存储器，由于每个线程的资源过多将会造成可同时执行的线程数量太少，使其变成了制约因素（但是矩阵乘法分块的这种思想是在所有并行系统中都很有效的，因为其很好地利用了局部性的优点）。