6.1向量加法及其串行代码

• 向量加法是HPC领域最简单的运算,
• HPC领域的“ hello world!”

 

6.2 单block单thread向量加

• GPU是协处理器
• GPU程序开发与传统CPU开发最大的区别
  • 协处理器无法识别CPU端数据和代码
• 如何オ能让GPU跑起来
• 这是GPU并行程序开发的基础
• 向量加法为例,
• 如何让GPU向量加

 

• 就GPU单个线程而言,向量加与CPU类似,只是kernel函数需用
  _ global _限定符标识
• kernel函数调用时需<<<X,X>>>配置grid和block维度。
• 下面代码展示5.6节描述的简化的CUDA编程五步曲

6.3单 block多 thread向量加

• 6.2节已成功在GPU运行起来了,但依然是串行的。
• 5.1节可知, kernel函数启动时,1个grid可拥有多个 block,而1个 block又有多个 thread。
• 尝试用block内的多个thread并行运算向量加

 

• kernel函数用threadIdx.x获得线程索引号,
• <<<1, threadnum>>>指定1个block内threadnum个thread同时运算,
• 每个线程完成1次向量加法后索引tid根据线程总数( blockDim.x)跳步

6.4多 block多 thread向量加

• GPU中不只1个block,此时其他block是空闲的,要充分利用GPU资源就要利用多个 block进行计算

 

• 进行多block和多 thread改造,关键是对线程索引tid、跳步步长重新构造
• 线程索引不再是threadIdx.x获得的block内部索引,
  • 而需全局线程索引,
  • tid=blockIdx.x * blockDim.x+ threadIdx.x)
• 跳步也不再是block尺寸,
  • 而是grid内所有thread数量
  • gridDim.x * blockDim.x)

6.6实验结果分析与结论

6.6.1本书实验平台

• 2个8核Xeon CPU E5-2670
• 内存64G
• 2个NVIDIA Tesla K20C GPU。
• 软件环境ICC version13.0.0.079
• CUDA5.5

 

• 本书串行代码均最佳优化开关(-O2或-O3)且保证最内层循
  环向量化,即保证向量化

 

• 本书部分章节实验的编译环境是 ICC version1.1。
• 本书实验平台是大集群中的一个节点,有专人管理维护,笔者不知道ICC编译器是啥时候升级的,
• 代码编写的时间跨度较大,
• 笔者不确定具体哪个章节用的哪套环境

6.6.3结论

• (1)
  • GPU单核计算性能极差,计算时间为CPU的300多倍
• (2)
  • 满负载时的1个block性能跟1个向量运算的CPU核性能相近(稍强)
• (3)GPU满负载计算时,gpu_3和cublas性能接近
  • 当数据放大128倍时,
  • gpu_3计算时间是10.232ms,
  • cublas的计算时间是9.879ms
  • 说明当数据量大时, CUBLAS库性能更好
• (4)GPU并行计算时间明显比CPU串行计算时间短很多
  • 说明GPU并行计算性能高
• (5)CPU与GPU的通信时间比串行计算时间还长,
  • 说明向量加法的GPU移植瓶颈在数据传输,
  • 当前平台上若数据位于CPU端情况下不适合移植到GPU计算;
  • 本节代码通信时间还能进一步优化(用页锁定存储),但依然无法匹配串行计算时
  • Tesla P100GPU配置NVLink可令通信带宽高达160GB/s,
  • 用该产品可能向量加法中的通信瓶颈将大大削弱甚至不再存在