第六章 性能优化

《大规模并行处理器编程实战》学习,其他章节关注专栏 CUDA C

CUDA C 编程友情链接:

6.1 WARP 和线程执行

由于 SM(Streaming Multiprocessor) 中实际由多个流处理单元(Streaming Processor, SP)进行单指令、多数据(SIMD)模式的执行,因此将线程块划分为 warp,方便调度给不同的 SP 执行。一般 warp 是 32 个线程。
在将线程划分时,按照多维数组变一维数组的映射进行32个线程的转化划分。
第六章 CUDA性能优化_核函数

如果是三维,则先对 threadIdx.z=0 的二维进行线性排列,再对1,2…排列。
当一个 warp 中存在 if-else,循环次数不定等情况时,warp 就要进行分支(diverge),判断情况并执行。以加法归约算法说明分支的存在:
第六章 CUDA性能优化_迭代_02

这是一个求和的并行加速算法,其本质是

第六章 CUDA性能优化_核函数_03

可以发现,在第2,3,…次迭代中,越来越多的线程被闲置,处于分支状态。其算法复杂度如下:
第六章 CUDA性能优化_核函数_04

可以通过对归约算法进行改进,从而提高硬件利用率,改进后的算法能让被使用的线程块越来越少,同时使用中的线程kernel不存在分支,多余的线程块block在逐次迭代中被释放出来,:
第六章 CUDA性能优化_卷积_05

第六章 CUDA性能优化_数据_06

6.2 全局存储器的带宽

为了方便整块数据的存取,矩阵乘法中(b)的访问形式能够合并数据,即其threadIdx.y实变动的,threadIdx.x和x+1,x+2…能被一起取出,从而加速。而(a)中循环主体是threadIdx.x,不能连续取出x+1, x+2,无法进行存储加速。
第六章 CUDA性能优化_数据_07

本章学习时尚未记录,如要学习建议查看原书(百度云链接,提取码:cuda,有我做的笔记,望谅解),顺便送一本CUDA指南。