第六章 CUDA性能优化

第六章 性能优化

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

CUDA C 编程友情链接：

• 第三章 CUDA 简介-CUDA C编程向量加法
• 第四章 CUDA数据并行执行模型
• 第五章 CUDA 存储器
• 第六章 CUDA性能优化（内附原书链接）
• 核函数：CUDA编程入门（一）-以图片运算看线程的组织和核函数的使用
• 拓展：CUDA卷积计算及其优化——以一维卷积为例

6.1 WARP 和线程执行

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

如果是三维，则先对 threadIdx.z=0 的二维进行线性排列，再对1,2…排列。
当一个 warp 中存在 if-else，循环次数不定等情况时，warp 就要进行分支（diverge），判断情况并执行。以加法归约算法说明分支的存在：

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

可以发现，在第2,3，…次迭代中，越来越多的线程被闲置，处于分支状态。其算法复杂度如下：

可以通过对归约算法进行改进，从而提高硬件利用率，改进后的算法能让被使用的线程块越来越少，同时使用中的线程kernel不存在分支，多余的线程块block在逐次迭代中被释放出来，：

6.2 全局存储器的带宽

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

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