CUDA程序并發(fā)錯(cuò)誤的自動(dòng)檢測(cè)
發(fā)布時(shí)間:2020-12-30 15:29
在當(dāng)前的工業(yè)界,GPU程序的開發(fā)和維護(hù)已經(jīng)成為了重要的新平臺(tái)。機(jī)器學(xué)習(xí)的任務(wù)在GPU平臺(tái)上能夠獲得更加優(yōu)秀的效果。CUDA(Computer Unified Device Architecture)是n Vidia開發(fā)的一個(gè)GPU計(jì)算平臺(tái)。雖然CUDA已成為通用GPU計(jì)算的主要并行計(jì)算平臺(tái)和編程模型,但CUDA引發(fā)的錯(cuò)誤模式尚未有成熟的自動(dòng)化解決方案。在此篇論文中,我們將提出一個(gè)輕量級(jí)的CUDA并發(fā)錯(cuò)誤檢測(cè)工具Simulee,通過自動(dòng)生成觸發(fā)內(nèi)存訪問的輸入,追蹤并分析實(shí)際被執(zhí)行的CUDA程序語句序列,并基于該序列,收集不同線程的內(nèi)存訪問信息并做出是否存在并發(fā)錯(cuò)誤的判斷。為了從效能和效率兩個(gè)維度評(píng)價(jià)Simulee的性能,本文還從Github上選擇了7個(gè)流行的CUDA相關(guān)項(xiàng)目,來作為測(cè)試的對(duì)象。實(shí)驗(yàn)結(jié)果證明,Simulee可以從24個(gè)手動(dòng)分析的并發(fā)bug中成功檢測(cè)21個(gè)出來,同時(shí)還檢測(cè)出了24個(gè)尚未發(fā)現(xiàn)的并發(fā)錯(cuò)誤bug,其中的10個(gè)已經(jīng)被開發(fā)者確認(rèn)。與此同時(shí),Simulee在檢測(cè)中體現(xiàn)的性能都顯著強(qiáng)于已知的CUDA并發(fā)錯(cuò)誤檢測(cè)的方案。
【文章來源】:哈爾濱工業(yè)大學(xué)黑龍江省 211工程院校 985工程院校
【文章頁數(shù)】:56 頁
【學(xué)位級(jí)別】:碩士
【部分圖文】:
CUDA核函數(shù)并行計(jì)算機(jī)
中看出,“if”圖1-2寫寫沖突
哈爾濱工業(yè)大學(xué)工程碩士學(xué)位論文-4-能。圖1-3展示了一個(gè)來自Github項(xiàng)目arrayfire中的修復(fù)補(bǔ)叮可以觀察到,在這個(gè)程序里,Block被定義為一維,因此,“tid”的值僅由“threadIdx.x”指定,又因?yàn)锽lock為一維,所以這表示“tid”在同一個(gè)塊的不同線程之間是相同的。因此,只能通過一個(gè)線程訪問“s_median[tid]”和“s_idx[tid]”,從而導(dǎo)致第4行中出現(xiàn)冗余的屏障函數(shù),因?yàn)閯h除該屏障函數(shù)后,“s_median”或“s_idx”中沒有數(shù)據(jù)競(jìng)爭(zhēng),也就不會(huì)發(fā)生任何并發(fā)錯(cuò)誤。當(dāng)同一個(gè)Block內(nèi)的線程,一個(gè)屏障函數(shù)對(duì)于一部分線程可達(dá)而對(duì)于另一部分線程不可達(dá),或者一部分線程已經(jīng)結(jié)束而另一部分線程有執(zhí)行道這個(gè)屏障函數(shù),則稱之為屏障分化。圖1-4展示了一個(gè)同樣來自Github的CUDA項(xiàng)目arrayfire,修復(fù)一個(gè)屏障分化的代碼。從圖1-4中可以看出,開發(fā)人員通過將“__syncthreads()”語句移到給定分支之外,確保同一Block中的所有線程在內(nèi)核函數(shù)的每個(gè)執(zhí)行過程中都達(dá)到相同的屏障函數(shù)。否則,將會(huì)出現(xiàn)未定義行為。圖1-3冗余的屏障函數(shù)圖1-4屏障分化
【參考文獻(xiàn)】:
期刊論文
[1]多核多線程技術(shù)綜述[J]. 眭俊華,劉慧娜,王建鑫,秦慶旺. 計(jì)算機(jī)應(yīng)用. 2013(S1)
[2]CPU/GPU協(xié)同并行計(jì)算研究綜述[J]. 盧風(fēng)順,宋君強(qiáng),銀福康,張理論. 計(jì)算機(jī)科學(xué). 2011(03)
[3]多線程技術(shù)及其應(yīng)用的研究[J]. 伍光勝,宋信忠,鄭明輝. 計(jì)算機(jī)應(yīng)用研究. 2001(01)
本文編號(hào):2947877
【文章來源】:哈爾濱工業(yè)大學(xué)黑龍江省 211工程院校 985工程院校
【文章頁數(shù)】:56 頁
【學(xué)位級(jí)別】:碩士
【部分圖文】:
CUDA核函數(shù)并行計(jì)算機(jī)
中看出,“if”圖1-2寫寫沖突
哈爾濱工業(yè)大學(xué)工程碩士學(xué)位論文-4-能。圖1-3展示了一個(gè)來自Github項(xiàng)目arrayfire中的修復(fù)補(bǔ)叮可以觀察到,在這個(gè)程序里,Block被定義為一維,因此,“tid”的值僅由“threadIdx.x”指定,又因?yàn)锽lock為一維,所以這表示“tid”在同一個(gè)塊的不同線程之間是相同的。因此,只能通過一個(gè)線程訪問“s_median[tid]”和“s_idx[tid]”,從而導(dǎo)致第4行中出現(xiàn)冗余的屏障函數(shù),因?yàn)閯h除該屏障函數(shù)后,“s_median”或“s_idx”中沒有數(shù)據(jù)競(jìng)爭(zhēng),也就不會(huì)發(fā)生任何并發(fā)錯(cuò)誤。當(dāng)同一個(gè)Block內(nèi)的線程,一個(gè)屏障函數(shù)對(duì)于一部分線程可達(dá)而對(duì)于另一部分線程不可達(dá),或者一部分線程已經(jīng)結(jié)束而另一部分線程有執(zhí)行道這個(gè)屏障函數(shù),則稱之為屏障分化。圖1-4展示了一個(gè)同樣來自Github的CUDA項(xiàng)目arrayfire,修復(fù)一個(gè)屏障分化的代碼。從圖1-4中可以看出,開發(fā)人員通過將“__syncthreads()”語句移到給定分支之外,確保同一Block中的所有線程在內(nèi)核函數(shù)的每個(gè)執(zhí)行過程中都達(dá)到相同的屏障函數(shù)。否則,將會(huì)出現(xiàn)未定義行為。圖1-3冗余的屏障函數(shù)圖1-4屏障分化
【參考文獻(xiàn)】:
期刊論文
[1]多核多線程技術(shù)綜述[J]. 眭俊華,劉慧娜,王建鑫,秦慶旺. 計(jì)算機(jī)應(yīng)用. 2013(S1)
[2]CPU/GPU協(xié)同并行計(jì)算研究綜述[J]. 盧風(fēng)順,宋君強(qiáng),銀福康,張理論. 計(jì)算機(jī)科學(xué). 2011(03)
[3]多線程技術(shù)及其應(yīng)用的研究[J]. 伍光勝,宋信忠,鄭明輝. 計(jì)算機(jī)應(yīng)用研究. 2001(01)
本文編號(hào):2947877
本文鏈接:http://sikaile.net/kejilunwen/jisuanjikexuelunwen/2947877.html
最近更新
教材專著