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\n", 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\n", 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%.
