GLobal memory的load/store都要通過L2緩存（在計算能力 < 3的卡還有L1緩存），因此目前的主流卡基本上都已經沒有L1緩存了，因此後面就只以L2緩存爲例.html
對齊尋址和臨近尋址：緩存
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中的數據.ui
而臨近訪問指的就是讓一個warp中的線程訪問的內存儘可能挨着，這樣就能下降數據緩存的總量.訪問的內存總量若是很高，就算efficiency很高，速度確定也慢.spa
驗證程序：線程
#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; }
編譯：3d
nvcc -O3 aliged.cu -o align
分別取不一樣的offset運行程序：code
./align ./align 4 ./align 128
結果以下：（在垃圾遊戲本上可能要多運行幾回才能出現這麼好的結果，在較好的比較穩定的服務器上幾乎每次都能獲得相似的結果）htm
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%.