簡單加法示例：共享內存版本的縮減執行速度低於全局內存

Question

我已經實現了兩個版本的add。兩者的加法概念完全相同。唯一的區別是在一個代碼中（下面的第一個代碼）我使用全局內存，而第二個代碼使用共享內存。正如在幾個地方提到的那樣，共享內存版本應該更快，但就我而言，全局內存版本更快。 請告訴我哪裏出錯了。注意：我有一個cc 2.1的GPU。因此，對於共享內存，我有32家銀行。由於我在示例中僅使用了16個整數，所以我的代碼不應該有銀行衝突。 請讓我知道這是否正確。

全球版本

#include<stdio.h> 
__global__ void reductionGlobal(int* in, int sizeArray, int offset){ 

    int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    if(tid < sizeArray){ 
     if(tid % (offset * 2) == 0){ 
      in[tid] += in[tid+offset]; 
     } 

    } 

} 
int main(){ 
    int size = 16; // size of present input array. Changes after every loop iteration 
    int cidata[] = {1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16}; 

    int* gidata; 
    cudaMalloc((void**)&gidata, size* sizeof(int)); 
    cudaMemcpy(gidata,cidata, size * sizeof(int), cudaMemcpyHostToDevice); 
    int offset = 1; 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start, 0); 
    while(offset < size){ 
     //use kernel launches to synchronize between different block. syncthreads() will not work 
     reductionGlobal<<<4,4>>>(gidata,size,offset); 
     offset *=2; 

    } 
    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime , start, stop); 
    printf("time is %f ms", elapsedTime); 
    int* output = (int*)malloc(size * sizeof(int)); 
    cudaMemcpy(output, gidata, size * sizeof(int), cudaMemcpyDeviceToHost); 
    printf("The sum of the array using only global memory is %d\n",output[0]); 
    getchar(); 
    return 0; 
} 

共享內存版本：

#include<stdio.h> 

__global__ void computeAddShared(int *in , int *out, int sizeInput){ 
    extern __shared__ float temp[]; 

    int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    int ltid = threadIdx.x; 
    temp[ltid] = 0; 
    while(tid < sizeInput){ 
     temp[ltid] += in[tid]; 
     tid+=gridDim.x * blockDim.x; // to handle array of any size 
    } 
    __syncthreads(); 
    int offset = 1; 
    while(offset < blockDim.x){ 
     if(ltid % (offset * 2) == 0){ 
      temp[ltid] = temp[ltid] + temp[ltid + offset]; 
     } 
     __syncthreads(); 
     offset*=2; 
    } 
    if(ltid == 0){ 
     out[blockIdx.x] = temp[0]; 
    } 

} 

int main(){ 

    int size = 16; // size of present input array. Changes after every loop iteration 
    int cidata[] = {1,2,3,4,5,6,7,8,9,10,11,12,13,14,15,16}; 

    int* gidata; 
    int* godata; 
    cudaMalloc((void**)&gidata, size* sizeof(int)); 
    cudaMemcpy(gidata,cidata, size * sizeof(int), cudaMemcpyHostToDevice); 
    int TPB = 4; 
    int blocks = 10; //to get things kicked off 
    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 
    cudaEventRecord(start, 0); 
    while(blocks != 1){ 
     if(size < TPB){ 
      TPB = size; // size is 2^sth 
     } 
     blocks = (size+ TPB -1)/TPB; 
     cudaMalloc((void**)&godata, blocks * sizeof(int)); 
     computeAddShared<<<blocks, TPB,TPB>>>(gidata, godata,size); 
     cudaFree(gidata); 
     gidata = godata; 
     size = blocks; 
    } 

    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime , start, stop); 
    printf("time is %f ms", elapsedTime); 
    int *output = (int*)malloc(sizeof(int)); 
    cudaMemcpy(output, gidata, sizeof(int), cudaMemcpyDeviceToHost); 
    //Cant free either earlier as both point to same location 
    cudaFree(godata); 
    cudaFree(gidata); 
    printf("The sum of the array is %d\n", output[0]); 
    getchar(); 
    return 0; 
} 

Answer 1

有很多錯在這裏。首先，一些一般性評論：

• 您正在對16個數字進行縮減，這是一個可笑的 小輸入大小。 CUDA在主機 和設備端都有很多固定開銷。您爲設備執行的工作量太小，您所測量的只是那些開銷，而不是GPU的執行時間。您看到 的兩個代碼之間的差異可能僅僅是由於在共享內存版本爲 的情況下增加了設置開銷。當然與代碼本身無關。如果您想測量代碼的實際性能，那麼您爲該代碼提供的工作量必須足夠大，以確保執行時間遠大於設置時間。請放心，即使在一個小型GPU上，對於這個問題你也有5個數量級的工作量。
• 你已經提到銀行衝突，但這是你正在使用的架構上的稻草人。與舊的硬件相比，費米擁有完全不同的共享內存佈局，並且在銀行衝突方面只有相對較小的問題。在這種情況下肯定沒什麼可擔心的。

至於實際減少代碼本身：

• 如果你不能想出辦法對於要在單個內核啓動 降至輸入數組每個線程一個部分和，那麼你真的還沒有想到 足夠的問題。您在「全球」 和「共享」版本中的當前方法極其低效。並行壓縮是一個解決的問題，CUDA SDK附帶一份關於優化和降低GPU性能的優秀白皮書。你應該在做任何事之前閱讀它。
• 一旦您達到每個線程有一個部分總和的點，您想要執行共享內存減少每塊以減少每個塊發出一個部分總和。這將只需要兩個內核啓動 來計算完整的減少。
• 您的「共享」版本有一個緩衝區溢出，應該導致運行時錯誤 。在啓動 時間指定的動態共享內存大小是字節，而不是單詞。如果你的代碼有錯誤檢查，你 已經找到了這個。費米擁有出色的共享內存保護，如果您嘗試在靜態或動態分配的內容之外編寫 ，則會產生運行時錯誤。