cuda Global Memory Access

GLobal memory的load/store都要经过L2缓存（在计算能力 < 3的卡还有L1缓存），所以目前的主流卡基本上都已经没有L1缓存了，所以后面就只以L2缓存为例.

对齐寻址和临近寻址：

　　Global memory在与L2做数据传输的最小单位为32bytes，称为L2 cahce line size.  例如有一段64bytes的global memory，如果warp中的一个线程要访问第4个字节的数据，那么L2缓存会把前32个字节的global memory都缓存过来，这种策略称为space locality （空间局部性：当前使用的数据临近的数据被使用的可能性更高），同理如果访问第33个bytes的内存，那么L2缓存会把后32个bytes的数据都缓存过来，相当于将global memory从起始地址按照32bytes分段，每次都缓存若干段的数据，所以就存在一个global memory efficiency的指标表示global memory的Load和store效率.

　　以一个warp的32个线程为例：假设有一段global memory，float* data，长度为160（5*32）.假设线程id为tid的线程去访问data[tid]位置的数据，那么总计需要访问data的前128字节的内容，按照L2以32字节为单位的缓存方式，正好缓存4段地址，这就属于对齐访问。

　　现在换一个访问方式：tid的线程去访问data[tid + offset]位置的内存，其中：0 < offset < 8.可以看到无论offset取多少，都将缓存全部5段内存，但是实际只用了4段，因此使用效率为80%，这种就属于没有对齐，所以对齐访问指的就是是否按照32字节访问global memory中的数据.

　　而临近访问指的就是让一个warp中的线程访问的内存尽量挨着，这样就能降低数据缓存的总量.访问的内存总量如果很高，就算efficiency很高，速度肯定也慢.

验证程序：

#include<stdio.h>
#include<time.h>
typedef float TYPE;
__global__ void readOffset(TYPE* A, TYPE* B, TYPE* C,
        int n, int offset) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int i = tid + offset;
    if ( i < n ) C[tid] = A[i] + B[i];
}

__global__ void warmup(TYPE* A, TYPE* B, TYPE* C,
    int n, int offset) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int i = tid + offset;
if ( i < n ) C[tid] = A[i] + B[i];
}

void InitValue(TYPE* p, int n) {
    for (int i = 0; i != n - 1; i++) {
        p[i] = i;
    }
}
int main(int argc, char** argv) {
    int n = 1 << 20;
    size_t nBytes = n*sizeof(TYPE);
    int blockSize = 512;
    int offset = 0;
    if (argc > 1) offset = atoi(argv[1]);
    int nBlocks = (n-1)/blockSize + 1;
    TYPE* a = (TYPE*)malloc(nBytes);    
    TYPE* b = (TYPE*)malloc(nBytes);
    InitValue(a, n);
    InitValue(b, n);
    TYPE* A, *B, *C;
    cudaMalloc(&A, nBytes);    
    cudaMalloc(&B, nBytes);    
    cudaMalloc(&C, nBytes);
    cudaMemcpy(A, a, nBytes, cudaMemcpyHostToDevice);    
    cudaMemcpy(B, b, nBytes, cudaMemcpyHostToDevice);    
    //warmup    
    clock_t start, end;
    start = clock();
    warmup<<<nBlocks, blockSize>>>(A, B, C, n, offset);
    cudaDeviceSynchronize();
    end = clock();    
    double dura = (double)(end - start) / CLOCKS_PER_SEC;
    printf("warmup  : %f
", dura);
    
    start = clock();
    readOffset<<<nBlocks, blockSize>>>(A, B, C, n, offset);
    cudaDeviceSynchronize();
    end = clock();    
    dura = (double)(end - start) / CLOCKS_PER_SEC;
    printf("offset %d: %f
", offset, dura);
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    free(a);
    free(b);

    cudaDeviceReset();
    return 0;    
}

编译：

nvcc -O3 aliged.cu -o align

分别取不同的offset运行程序：

./align
./align 4
./align 128

结果如下：（在垃圾游戏本上可能要多运行几次才能出现这么好的结果，在较好的比较稳定的服务器上几乎每次都能得到类似的结果）

warmup  : 0.000350
offset 0: 0.000137

warmup  : 0.000350
offset 4: 0.000154

warmup  : 0.000346
offset 128: 0.000137

使用nvprof进一步验证：

nvprof ./align

nvprof ./align 4

nvprof ./align 128

结果如下：

可以看到不管是warmup还是readOffset函数，当offset不对齐的时候执行时间都比对其的情况慢.

然后继续使用--metrics选项查看利用率，执行：

nvprof --metrics gld_efficiency,gst_efficiency ./align 4

结果如下：

可以看到efficiency为80%，正好符合预期，如果offset换为8的倍数，那么efficiency就都是100%.

另外简单修改代码就能实现gst_efficiency为80%，就是将readOffset中的C[tid]改为C[tid+offset]，修改之后重新编译，同样执行上面的命令，结果如下：

AoS VS SoA

什么意思：Array of struct && struct of Array

例如：

case AoS:
struct P {
  int x;
  int y;  
};
P* arr;
cudaMalloc(&arr, 1024*sizeof(P));

case SoA:
struct P {
  int x[1024];
  int y[1024];  
};
P* arr;
cudaMalloc(&arr, sizeof(P));

假设(x,y)代表点的坐标，假设实现两个kernel函数，分别给两个不同形式的arr赋相同的值，代码参考：http://www.wrox.com/WileyCDA/WroxTitle/Professional-CUDA-C-Programming.productCd-1118739329,descCd-DOWNLOAD.html->chapter04中的simpleMathAos.cu.

理论上两种表示都能实现这个功能，但是对内存的efficiency是天差地别的.

Aos与SoA的存储方式，如下图：

经过上面的例子可以简单分析一下：当以AoS格式请求arr[tid].x时，arr[tid].y也被隐式加载，因此利用率只有50%，而SoA则没有这个问题.

这里注意：gld_efficiency的计算公式为：

  其中Required为一个warp的线程一次请求需要的全部内存，而Requested为每次请求的使用显存大小，所以simpleMathAoS.cu的kernel每次都对x,y分别赋值，每次都只request了4字节的内存，因此最后结果是50%.