※ 引述《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/cn.aspx?n=bbs/C_and_CPP/M.1685340827.A.DAC.html