CUDA程序并发错误的自动检测

发布时间：2020-12-30 15:29

　　在当前的工业界,GPU程序的开发和维护已经成为了重要的新平台。机器学习的任务在GPU平台上能够获得更加优秀的效果。CUDA(Computer Unified Device Architecture)是n Vidia开发的一个GPU计算平台。虽然CUDA已成为通用GPU计算的主要并行计算平台和编程模型,但CUDA引发的错误模式尚未有成熟的自动化解决方案。在此篇论文中,我们将提出一个轻量级的CUDA并发错误检测工具Simulee,通过自动生成触发内存访问的输入,追踪并分析实际被执行的CUDA程序语句序列,并基于该序列,收集不同线程的内存访问信息并做出是否存在并发错误的判断。为了从效能和效率两个维度评价Simulee的性能,本文还从Github上选择了7个流行的CUDA相关项目,来作为测试的对象。实验结果证明,Simulee可以从24个手动分析的并发bug中成功检测21个出来,同时还检测出了24个尚未发现的并发错误bug,其中的10个已经被开发者确认。与此同时,Simulee在检测中体现的性能都显著强于已知的CUDA并发错误检测的方案。 

【文章来源】：哈尔滨工业大学黑龙江省 211工程院校 985工程院校

【文章页数】：56 页

【学位级别】：硕士

【部分图文】：

CUDA核函数并行计算机

中看出，“if”图1-2写写冲突

哈尔滨工业大学工程硕士学位论文-4-能。图1-3展示了一个来自Github项目arrayfire中的修复补叮可以观察到，在这个程序里，Block被定义为一维，因此，“tid”的值仅由“threadIdx.x”指定，又因为Block为一维，所以这表示“tid”在同一个块的不同线程之间是相同的。因此，只能通过一个线程访问“s_median[tid]”和“s_idx[tid]”，从而导致第4行中出现冗余的屏障函数，因为删除该屏障函数后，“s_median”或“s_idx”中没有数据竞争，也就不会发生任何并发错误。当同一个Block内的线程，一个屏障函数对于一部分线程可达而对于另一部分线程不可达，或者一部分线程已经结束而另一部分线程有执行道这个屏障函数，则称之为屏障分化。图1-4展示了一个同样来自Github的CUDA项目arrayfire，修复一个屏障分化的代码。从图1-4中可以看出，开发人员通过将“__syncthreads()”语句移到给定分支之外，确保同一Block中的所有线程在内核函数的每个执行过程中都达到相同的屏障函数。否则，将会出现未定义行为。图1-3冗余的屏障函数图1-4屏障分化

【参考文献】：
期刊论文
[1]多核多线程技术综述[J]. 眭俊华,刘慧娜,王建鑫,秦庆旺.  计算机应用. 2013(S1)
[2]CPU/GPU协同并行计算研究综述[J]. 卢风顺,宋君强,银福康,张理论.  计算机科学. 2011(03)
[3]多线程技术及其应用的研究[J]. 伍光胜,宋信忠,郑明辉.  计算机应用研究. 2001(01)



本文编号：2947877