《大规模并行处理器编程实战》学习,其他章节关注专栏 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指南。