亲爱的朋友们, 新年快乐!! 感谢过去一年热情支持与协助的朋友们, 提供国网光碟 DVD 压缩、FTP、BT 下载, 以及许许多多的资讯, 并且校正我在硬体认识上的谬误, 使小弟获益良多. 在这新的一年, 我们迈入 CUDA 进阶讨论, 其它基本 CUDA 程设, 会再慢慢补足, 比较急的话不妨先看 CUDA ToolKit 的 programming guide, 也可到 google 上 搜寻「CUDA 教学」, 有很多很赞的部落格, 他们整理得比我好, 凌乱的生活态度, 想不到竟反应在我写的文章上, 真是伤脑筋...XD 希望在这新的一年, 能扩大对 GPGPU 这方面计算感兴趣的朋友阵容, 能空出时间 在 ATI GPGPU 尝试 programming 与 Apple OpenCL, 不论是进行研究 or 切磋, 期待中文化文件的发展, 以及在 PTT 上成立正式 GPGPU 版, 让大家能有更好的 交流空间. Yours Truly, 咖啡里的海洋蓝 2009 ◆ 何谓 warps ============================================================================= 现阶段 warp 指 32 个执行绪, 在硬体上为实际的 SIMT 群组, SM (串流多处理器, Stream Multiprocessor) 的 8 个执行单元, 轮流执行这 32 个执行绪, 达成在 4 个 clocks 之下完成整个 warp 操作 (一般的运算指令, 如果是 global memory 的存取, 当然会需要更多的时间, 但 warp 会在 issue 完指令後, 把工作交给记忆体管线处理, 然後 free to work on 其它的 warp). ◆ Grid/Block/Warp/Thread 架构 ============================================================================= 我们提过 __syncthreads() 对效能很负面, 因为一旦使用这个指令, multi-threads 就无法隐藏 memory access latency, 而这却是 SIMT 比起传统设计最大的优势之一, 解决的办法之一就是让 SM 可以同时执行数个 blocks, 使得部份 block 里的 threads hang 住时, SM 仍可执行其它的 block 来隐藏 latency. 一般人通常认为一个 SM 只能执行一个 block, 这是错误的观念, 事实上一个 SM 具有 同时执行数个 blocks 的能力, 但设定上也要做更多调整, 毕竟天下没有白吃的午餐, 为了简化 software 设计起见, 初学者并不建议在这方面涉入太深, 但进阶最佳化时, 这却是基本的(programming guide 都有写), 不过仍属於控制 blocks 阶段. 然而, 绝少文献提及, SIMT 的群组模型必需要视为 grid/block/warp/thread 架构, 大部份都只提 grid/block/thread (包括 programming guide), 甚至 CUDA 实作时, warps 也被排除在正规软体架构之外, 只是视为是 hardware 的一部份, 理解上也 被当作是 optional 的. (早期的 CUDA 版本中, warp 并不以明显的方式呈现, 直到 1.2 版後, 才开始有 warp vote 等指令出现) ◆ 对 warp 进行精确控制的好处 ============================================================================= 可曾想过我们可以不用 __syncthreads(), 就可对 shared memory 做存取吗? Well... 这的确是可以做到的, 因为 warp 是 hardware locking 的同步执行群组, 本身就 synchronized 了. 也就是透过对 warps 的精确控制, 我们可以在标准的 grid/block/thread 架构下, 再细切一层 warp, 在这层中不需要 __syncthreads(). 例如在 SDK 的 scan 这支程式中, 实做 tree reduction 时必需 __syncthreads(), 但我们可以在 warp 层先做不用 __syncthreads() 的 reduction, 而且因为 warp size 是固定, 使用 32 个 threads 来 tree reduce 64 笔 data, 可以直接 写入 6 个指令, 然後再把 global 的那套搬到 block 层再搞一次, 就可以有效的 去除许多 __syncthreads(), 而且各 warps 也可以非同步执行. 另一个好处是, 我们在 global memory access 之後, 到 __syncthreads() 之间, 利用 warp 的这种特性可插入许多指令, 更能有效的缓冲 latency hidden 的压力, 使得 SIMT 的效果呈现得更好. 再来就是, 我们以前时常用到的老招数, 应用在某些地方只用到少量 threads 时, 例如使用单一 thread 载入数个边界点, 可以分散在数个 warp 中, 避免只使用一个 thread 造成的 I/O serial 现象, 有效的运用 SM 中的多个记忆体单元. 还有很多应用不胜枚举, 小弟也还在探索中... ◆ warp 的执行绪定址 ============================================================================= 如果指定 blockDim = 512, 则它包含了 16 个 warps, 执行绪定址如下 warp 0: tid=0~31 warp 1: tid=32~63 warp 2: tid=64~95 ... warp 15: tid=480~511 当指定的 blockDim 不是 32 的整数倍时, 最後的 warp 不完整, 硬体上仍占一个 warp, 只是多余的单元被 disable 掉而己, 例如 blockDim = 100 包含了 4 个 warps warp 0: tid=0~31 warp 1: tid=32~63 warp 2: tid=64~95 warp 3: tid=96~100 (incomplete, 101~127 disabled) 为了方便操作, 我通常会把 blockDim 设成 2D/3D 的型式, CUDA 的 lowest dimention 是由 x 算起, 所以 threadIdx.(y,z) 会形成 warp 的 2D index, 例如 blockDim=Dim3(32, 5, 1) -- 2D 的方式, 配置 5 个 warp 的阵列 warp 0: threadIdx.y=0 warp 1: threadIdx.y=1 ... warp 5: threadIdx.y=5 blockDim=Dim3(32, 3, 4) -- 3D 的方式, 配置 3x4 的 warp 阵列 warp (0,0): threadIdx.y=0 threadIdx.z=0 warp (0,1): threadIdx.y=0 threadIdx.z=1 warp (0,2): threadIdx.y=0 threadIdx.z=2 warp (0,3): threadIdx.y=0 threadIdx.z=3 ... warp (2,0): threadIdx.y=2 threadIdx.z=0 warp (2,1): threadIdx.y=2 threadIdx.z=1 warp (2,2): threadIdx.y=2 threadIdx.z=2 warp (2,3): threadIdx.y=2 threadIdx.z=3 ◆ 硬体执行配置 1 (compute 规格) ============================================================================= 一个 SM 可同时执行数个 blocks, 其限制在於 register 和 shared memory 的大小, 以及一些规格上的限制 (see compute 版本的规格), 例如在 1.0 版中规定 max # of active blocks per SM = 8 max # of active warps per SM = 24 max # of active threads per SM = 768 这里的 active 指的是 SM 当下正在执行它, 也就是已经 load 到 SM 中执行了, 其中第 3 条 rule 比较没有义意, 因为由第 2 条已决定 (24*32=768), 且当存在 incomplete warps 时, 通常是 meet 不到的, 但在纯 software 的观点上 (不理会 warp 时), 有时候会觉得它反而比较容易理解. Anyway, 这些 rules 是用来避免 GPU 配置过多执行绪给 SM 的硬性限制, 否则当 shared memory = 0 时, 所有的 blocks 全都发给同一个 SM 执行, 那还得了. ps. 在支援 compute 1.2 的硬体上, 对这些规格有些异动 max # of active blocks per SM = 8 max # of active warps per SM = 32 max # of active threads per SM = 1024 ◆ 硬体执行配置 2 (shared memory 的影响) ============================================================================= 再来就是 shared memory 限制, 它是实作在每个 SM 上的, 现阶段总量只有 16KB, 未使用或 8 blocks 的使用量不满 16KB 时, SM 就会按照 default 规格来发配 warps, 否则 GPU 会切割出最大的 block 数来配置, 例如 shared = 3KB 时 16KB / 3KB = 5 所以 1 SM 会配置 5 blocks, 这些 blocks 虽然在同一个 SM 中执行, 但为了维持 Grid/Block 架构的弹性, 它们仍是无法再透过 shared memory 彼此沟通的. 当 shared > 8KB (超过 SM 总量的一半) 时, 一个 SM 就只能执行一个 block. 另外通常系统会使用到一些 shared memory, 所以要用 --ptxasoptions=-v 选项 来看实际的 shared memory 使用量. ◆ 硬体执行配置 3 (register 的影响) ============================================================================= 在 registers 的用量方面, 影响的不只 compile 过不过, 也影响 SM 对 block 的 配置, 考虑如下问题: G200 有 64KB register space, 若 1 个 block 使用 registers 大於 SM 总量之半 (例如 50KB), 假设不使用任何的 shared memory, 且每个 block 使用的 threads 和 warps 很少 (使它 bypass 前面的限制, 例如 blockDim = 20), 难道 GPU 要配给 SM 8 blocks 吗? 这样一来, 只有其中一个 block 的 register 能够 in core, 其它的 7 blocks 呢? 所以在硬体设计上, 让 register space 对 SM 分配的 block 进行限制是必要的. ◆ 如何 fulfill max # of active threads per SM ============================================================================= 也许你会说： God damn... @#$^%, 这麽多限制条件, 有什麽办法让 SM 能开启全部 threads? 管它是 1.0 版的 768 还是 1.2 版的 1024 threads. Well...这是可行的: (1) 首先在 block 配置时必需要是 complete warp (blockDim = 32 的整数倍), 不然 1024 个 threads 一定有些会被 disabled. (2) 再来要理解 block 数不必完全填满, 不然在 shared memory 控制上压力会很大, 例如在 1.2 版中, 宁愿配 2 个 blockDim=512, 也不愿 8 个 blockDim=128, 起码前者可用的 shared 有到 8 KB, 後者只有 2KB, 当然这要看应用而定. (3) 限制 register 和 shared memory 的使用量. (4) 在 launch 时使用 device query 的 API 来决定 blockDim 大小, 让程式对 後续版本具有相容性. (optional) 例如向量加法就很容易这样做, 因为它使用的 register 很少且 shared memory = 0 compute 1.0/1.1 版: blockDim=384 -> 2 blocks, full 768 threads compute 1.2/1.3 版: blockDim=512 -> 2 blocks, full 1024 threads 也可以配置 blockDim=256 (768,1024 的最大公因数, 8 warps) compute 1.0/1.1 版: blockDim=256 -> 3 blocks, full 768 threads compute 1.2/1.3 版: blockDim=256 -> 4 blocks, full 1024 threads note: 配置多点的 block 在 SM 上较可隐藏 __syncthreads() 的负面效果, 所以要对 shared memory 使用量进行妥协. --

※ 发信站: 批踢踢实业坊(ptt.cc) ◆ From: 114.45.219.173 ※ a5000ml:转录至看板 C_and_CPP 01/04 10:42