CUDA共享內存速度

這是一個與性能相關的問題。我已經寫了基於「CUDA按示例」示例代碼以下簡單的CUDA內核：CUDA共享內存速度

#define N 37426 /* the (arbitrary) number of hashes we want to calculate */ 
#define THREAD_COUNT 128 

__device__ const unsigned char *m = "Goodbye, cruel world!"; 

__global__ void kernel_sha1(unsigned char *hval) { 
    sha1_ctx ctx[1]; 
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    while(tid < N) { 
    sha1_begin(ctx); 
    sha1_hash(m, 21UL, ctx); 
    sha1_end(hval+tid*SHA1_DIGEST_SIZE, ctx); 
    tid += blockDim.x * gridDim.x; 
    } 
}

的代碼在我看來是正確的，確實吐出相同的散列的37426份（如基於預期在我閱讀第5章第5.3節時，我認爲寫入全局內存的每個線程都以「hval」的形式傳入，效率極低。

然後，我實現了我認爲會使用共享的性能提升緩存代碼修改如下：

#define N 37426 /* the (arbitrary) number of hashes we want to calculate */ 
#define THREAD_COUNT 128 

__device__ const unsigned char *m = "Goodbye, cruel world!"; 

__global__ void kernel_sha1(unsigned char *hval) { 
    sha1_ctx ctx[1]; 
    unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x; 
    __shared__ unsigned char cache[THREAD_COUNT*SHA1_DIGEST_SIZE]; 

    while(tid < N) { 
    sha1_begin(ctx); 
    sha1_hash(m, 21UL, ctx); 
    sha1_end(cache+threadIdx.x*SHA1_DIGEST_SIZE, ctx); 

    __syncthreads(); 
    if(threadIdx.x == 0) { 
     memcpy(hval+tid*SHA1_DIGEST_SIZE, cache, sizeof(cache)); 
    } 
    __syncthreads(); 
    tid += blockDim.x * gridDim.x; 
    } 
}

第二個版本似乎也能正常運行，但比初始版本慢幾倍。後者代碼在大約8.95毫秒內完成，而前者在大約1.64毫秒內執行。我對Stack Overflow社區的問題很簡單：爲什麼？

來源

2012-11-18 Shiba

我通過示例查看了CUDA，並找不到與此類似的任何內容。是的，在附錄中討論了一些GPU散列表，但它看起來不像這樣。所以我真的不知道你的函數做什麼，特別是sha1_end。如果此代碼與該書中的內容類似，請指出，我錯過了。

但是，如果sha1_end一次（每個線程）寫入全局內存並以一種合併的方式寫入全局內存，那麼沒有理由說它不是很有效。據推測，每條線索都會寫入不同的位置，所以如果它們相鄰或多或少，肯定有合併的機會。沒有進入合併的細節，只要說它允許多個線程在單個交易中將數據寫入存儲器即可。如果你要將你的數據寫入全球記憶體，你將不得不在至少一次的地方支付這個罰款。

對於您的修改，您已經徹底殺死了這個概念。您現在已經從單個線程執行所有數據複製，並且memcpy意味着後續數據寫入（整數或字符，無論）是否發生在單獨的事務中。是的，有一個緩存可以幫助解決這個問題，但是在GPU上完成它是完全錯誤的方法。讓每個線程更新全局內存，並利用機會並行執行它。但是當你強制單個線程上的所有更新時，那個線程必須順序複製數據。這可能是時差的最大單一成本因素。

__syncthreads（）的使用也會產生額外的成本。

CUDA by Example書的12.2.7節提到了可視化分析器（並提到它可以收集關於合併訪問的信息）。視覺分析器是一個很好的工具，可以幫助你回答這個問題。

如果你想了解更多關於高效內存技術和合並的知識，我推薦NVIDIA GPU計算webinar題爲「使用CUDA C - 高級1（2010）的GPU計算」。直接鏈接到here與slides。

來源

2012-11-18 22:20:43

非常感謝您經過深思熟慮的回覆。代碼鬆散地基於一些示例，但是我從頭開始編寫它。 sha1_ *函數是我沒有在我的文章中包含的庫的一部分。他們從一個輸入（在這種情況下是一個字符串）計算SHA1哈希值。它們不是CUDA運行時的一部分。 – Shiba

在百靈之際，我評論了__syncthreads（）調用來衡量其性能影響。差異可以忽略不計：其他因素導致5倍的性能損失。全局內存寫入每塊相鄰。感謝您參考合併：如果本書提到它，我錯過了它。我會看看你鏈接到的網絡研討會。 – Shiba

是的，sha1_end每個線程寫入全局內存一次。它採用ctx中的狀態並將160位SHA1哈希寫入hval。 – Shiba

CUDA共享內存速度

回答

相關問題