※ 引述《goodzey (--)》之銘言： : 不知道有沒有高手可以解答以下問題? : 資料形式: 600列800行的隨機數 : 目的: 把每一行的數據加起來 : 初始化: : sum[600]={0.0} : data[600x800]= 上述資料 : CUDA程式1: 成功 : // dim3 gridsize(1, 1, 1); : // dim3 blocksize(600, 1, 1); : for (int j = 0; j < 800; j+= 1){ : sum[(blockDim.x*bdx + tdx)] = sum[(blockDim.x*bdx + tdx)] : + data[600*j + (blockDim.x*bdx +tdx)]; : } 這邊有一個簡單的最佳化，先把 tdx thread 負責的 row 之和放在 register 裡面，這樣可以減少一些不必要的 global memory write。 : CUDA程式2: 失敗 : // dim3 gridsize(40, 1, 1); : // dim3 blocksize(600, 1, 1); : for (int j = 0; j < 800; j+= 40){ : sum[0*(j + bdx) +tdx] = sum[0*(j + bdx) +tdx] : + data[600*(j + bdx) +tdx]; : } : 請問程式2失敗的原因是?可以怎麼寫呢? : 我自己猜測是: 例如, sum[1]無法同時處理40筆資料 : 請教大家, 謝謝 2 的話，每個 block 的 tdx thread 都會往 sum[tdx] 做加總，而 blocks 並沒有保證 結束的時間點，所以會需要用 atomicAdd 避免 race condition。但因為總是往 sum 做加 總，實際上在 kernel launch 前，還得把 sum 清零，因此在量測效能上，是需要計算 清零 + kernel 運行的時間。這邊提供修改過的 kernel 給你參考: __global__ void multipleBlockSum(float *sum, float *data, size_t m, size_t n) { const auto numBlocks = gridDim.x; const auto bdx = blockIdx.x; const auto tdx = threadIdx.x; float s{}; for (int j = 0; j < n; j += numBlocks) { s += data[m * (j + bdx) + tdx]; } float *dst = &sum[tdx]; atomicAdd(dst, s); } 不過這邊想拋磚引玉提供一點關於這種 reduction 問題 kernel 的做法: 1. 每個 block 劃分一塊區域(tiling)去做 reduction，以這個問題就是 row-wise sum 2. 先把 tile 讀進 shared memory 後，在 shared memory 做 reduction，如果 tile 無法覆蓋所有 columns，則用 tile 大小 loop 過所有 column。 reduction 結果要放 shared memory or register 都可以。 3. 寫出 reduction 結果。 kernel 大概會長這樣: 1 template<size_t TileM, size_t TileN> 2 __global__ void reductionSum(float* s, float* a, size_t m, size_t n) { 3 const auto blockReadOffset = blockIdx.x * TileM; 4 const auto row = threadIdx.x / TileN; 5 const auto col = threadIdx.x % TileN; 6 const auto blockWriteOffset = blockIdx.x * TileM + row; 7 const auto localWriteOffset = row * TileN + col; 8 const auto localReadOffset = row + col * m; 9 __shared__ float buf[TileM * TileN]; 10 __shared__ float sum[TileM]; 11 memset(sum, 0, sizeof(float) * TileM); 12 size_t nIter = 0; 13 14 while (nIter < n) { 15 buf[localWriteOffset] = a[blockReadOffset + nIter * m + 16 localReadOffset]; 17 __syncthreads(); 18 19 #pragma unroll 20 for (uint32_t s = (TileN >> 1); s >= 1; s >>= 1) { 21 if (nIter + col < n && ((nIter + col + s) < n) && col < s) { 22 buf[localWriteOffset] += buf[localWriteOffset + s]; 23 } 24 __syncthreads(); 25 } 26 27 if (col == 0) { 28 sum[row] += buf[localWriteOffset]; 29 } 30 __syncthreads(); 31 32 nIter += TileN; 33 } 34 35 if (col == 0) { 36 s[blockWriteOffset] = sum[row]; 37 } 38} 參數: - TileM, TileN, block 每次 loop 負責的區域，[TileM, TileN] - s: sum 結果，a: input matrix，m: # of rows，n: # of columns 程式碼的大致解說如下: L2~L8: global read/write，shared memory read/write 的位址計算。 L9~L11: shared memory 的配置與初始化，包含 reduction 與 sum 結果的 buffer。 L12~L13: 開始 N 方向的 iteration。 L15~L17: 讀取 global memory 的資料到 shared memory，用 __syncthreads() 來保證 block 所需要的資料都已讀進 shared memory。 L19~L32: shared memory 內的 reduction，只有 col == 0 的 thread 更新 sum buffer 的值。 reduction 的做法可以參考: https://developer.download.nvidia.com/assets/cuda/files/reduction.pdf 我這邊寫的 kernel 就簡單做而已，沒有最佳化到極致。 L35~L36: 將 sum 寫出至 global memory。 要 launch kernel 的話大概像這樣: block 數量是 m / TileM，不整除的話要 +1， e.g. m, n = 600, 800, TileM, TileN = 16, 16, # of blocks = 600 / 16 + 1 = 38 reductionSum<16, 16><<<38, 256>>>(...) --

※ 發信站: 批踢踢實業坊(ptt.cc), 來自: 1.162.155.177 (臺灣) ※ 文章網址: https://webptt.com/m.aspx?n=bbs/C_and_CPP/M.1685340827.A.DAC.html