Re: [問題] CUDA 程式
※ 引述《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://www.ptt.cc/bbs/C_and_CPP/M.1685340827.A.DAC.html
推
05/30 22:16,
11月前
, 1F
05/30 22:16, 1F
→
05/30 22:18,
11月前
, 2F
05/30 22:18, 2F
→
05/30 22:21,
11月前
, 3F
05/30 22:21, 3F
推
06/02 23:16,
11月前
, 4F
06/02 23:16, 4F
→
06/02 23:19,
11月前
, 5F
06/02 23:19, 5F
討論串 (同標題文章)