Re: [問題] 關於CUDA的bank conflict

看板C_and_CPP (C/C++)作者lgd (嚕嚕嚕)時間13年前 (2013/03/26 16:39)推噓0(0推 0噓 0→)

留言0則, 0人參與討論串2/4 (看更多)

※ 引述《lukelan (藍丁丁)》之銘言： : 開發平台(Platform): (Ex: VC++, GCC, Linux, ...) : CUDA : 額外使用到的函數庫(Library Used): (Ex: OpenGL, ...) : 問題(Question)： : Bank Conflict : 餵入的資料(Input)： : sum[] array : 預期的正確結果(Expected Output)： : sum array 的每個位置都能夠正確的被累加 : 錯誤結果(Wrong Output)： : 由於多個thread同時access到sum，會導致sum的值無法正確地被累加 : 程式碼(Code)：(請善用置底文網頁, 記得排版) : __global__ kernel_function() : { : int n = blockIdx.x * blockDim.x + threadIdx.x; : int m= gridDim.x*blockDim.x; : for( int i = threadIdx.x ; i < maxLine -1 ; i += blockDim.x ) : { : if( flag[i] == 1 ) : { : device_function_1( sum ); : } : else if( flag[i] == 2 ) : { : device_function_2( sum ); : } : } : __device__ device_function_1( int sum[] ) : { : if(達成某些條件) : sum[ 部分位置 ]++; : } : __device__ device_function_2( int sum[] ) : { : if(達成某些條件) : sum[ 部分位置 ]++; : } : 補充說明(Supplement)： : 我的sum這個共用變數，是希望當成一個hash table的方式，讓所有資料執行的結果，存 : 在這個array中，所以會在device function都傳入這個變數，來把結果儲存起來。 : 此時會發生問題就是，所有的thread在進行運算時，都會對他做存取，導致結果的值會少 : 非常多。 : 以上，還請大家多多指教這應該不算bank conflict 這應該是大部分人開始寫CUDA的時候容易幹的蠢事（我也這麼蠢過）因為我們想法從sequencial 轉到 parallel很容易沒想好 == 簡單的解法的概念是宣告: sum[部份位置][thread number] 每一個thread有一個variable，等所有thread 算完之後再加總一次就可以了 == 雖然這根據你的用途還是會產生其他的問題，不過應該會比你本來的狀況好其他可能問題: * sum 要放在哪邊 global, shared or local memory。要怎麼放怎麼搬？ * if(a) else if(b) <= 這種寫法適不適合你的用途? 這樣可能會有一半的thread是沒動作的 * task的分法適合你的用途嗎？block,thread的分法適不適合你的用途? 因為我看到你宣告n, m 卻沒用所以我不確定你這邊的演算法不過這都得看你實際的case來決定，譬如memory使用量，甚至register的使用量都有差 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 220.135.58.34