PS：最近不務(wù)正業(yè)對(duì)GPU產(chǎn)生了濃厚的興趣缩膝，可能是AI需要軟硬結(jié)合的原因诫舅？對(duì)這本書(shū)（2017年出版不算太舊）進(jìn)行整理卦羡，以備不時(shí)之需鲁森。

# 前言

眾所周知，由于架構(gòu)和能源的限制振惰，將多核心中央處理器（CPU）的核心數(shù)量增加到數(shù)千個(gè)是不可能的歌溉，而圖形處理器（GPU）可以輕松地利用數(shù)千個(gè)簡(jiǎn)單有效的核心，每個(gè)核心執(zhí)行單個(gè)指令線程骑晶。

對(duì)于圖形應(yīng)用程序?qū)τ?jì)算和帶寬的巨大需求已經(jīng)導(dǎo)致GPU成為主要的大規(guī)模并行架構(gòu)痛垛。如今，具有浮點(diǎn)運(yùn)算能力的通用GPU（GPGPU）在具有混合CPU-GPU架構(gòu)的大多數(shù)頂級(jí)機(jī)器上廣泛用于通用計(jì)算桶蛔。這樣的GPGPU新產(chǎn)品NVIDIA Pascal擁有3840個(gè)核心匙头。

這樣強(qiáng)大的計(jì)算芯片的主要問(wèn)題是管理如此龐大的核心數(shù)量。為了滿足未來(lái)的需求仔雷，需要新的架構(gòu)技術(shù)來(lái)進(jìn)一步增加核心數(shù)量蹂析，同時(shí)考慮影響其可靠性，性能和能耗特性的技術(shù)限制碟婆。而缺乏有效的編程模型和系統(tǒng)軟件以利用這樣大量的核心在實(shí)際應(yīng)用中的全部性能电抚，是最具挑戰(zhàn)性的問(wèn)題。

本書(shū)聚焦于基于GPU的系統(tǒng)的研究和實(shí)踐竖共，并試圖解決它們的重要問(wèn)題蝙叛。涵蓋的主題涵蓋從硬件和架構(gòu)問(wèn)題到直接涉及應(yīng)用程序或系統(tǒng)用戶(hù)的高級(jí)問(wèn)題，包括并行編程工具和中間件支持的等等這種計(jì)算系統(tǒng)公给。本書(shū)分為四個(gè)部分借帘，每個(gè)部分都包含由知名研究人員撰寫(xiě)的章節(jié)蜘渣，他們?cè)谠摬糠值闹黝}下分享了他們的最新研究成果。

第1部分涉及不同的編程問(wèn)題和工具肺然。該部分由5個(gè)章節(jié)組成：1-5蔫缸。第1章討論了如何以可靠的方式編寫(xiě)GPU程序。盡管GPU程序在不同領(lǐng)域提供高計(jì)算吞吐量狰挡，但正確優(yōu)化代碼有些困難捂龄，這是由于GPU并發(fā)性的微妙性。因此加叁，本章的重點(diǎn)是討論并解決一些事情倦沧，以使GPU編程更容易。該章節(jié)還提供了一些關(guān)于GPU軟件形式分析的最新進(jìn)展它匕。后者是一個(gè)重要問(wèn)題展融，隨著GPU編程的進(jìn)步。因此豫柬，新的GPU程序員應(yīng)該學(xué)習(xí)驗(yàn)證方法告希。

第2章介紹了異構(gòu)集群的統(tǒng)一開(kāi)放計(jì)算語(yǔ)言（OpenCL）框架（SnuCL）。它是一個(gè)免費(fèi)的烧给，開(kāi)源的OpenCL框架燕偶，用于異構(gòu)集群〈〉眨基本上指么，OpenCL是一種針對(duì)異構(gòu)并行計(jì)算系統(tǒng)的編程模型，它定義了傳統(tǒng)處理器和加速器之間的抽象層榴鼎。使用OpenCL的優(yōu)點(diǎn)是程序員編寫(xiě)一個(gè)OpenCL應(yīng)用程序伯诬，然后在任何支持OpenCL的系統(tǒng)上運(yùn)行它。然而巫财，它缺乏針對(duì)異構(gòu)集群的目標(biāo)能力盗似。

SnuCL為程序員提供了一個(gè)單一、統(tǒng)一的OpenCL平臺(tái)映像平项，用于集群赫舒。借助SnuCL，OpenCL應(yīng)用程序能夠利用計(jì)算節(jié)點(diǎn)中的計(jì)算設(shè)備闽瓢，就像它們?cè)谥鳈C(jī)節(jié)點(diǎn)中一樣号阿。此外，使用SnuCL鸳粉，可以將來(lái)自不同供應(yīng)商的多個(gè)OpenCL平臺(tái)集成到單個(gè)平臺(tái)中扔涧。因此，OpenCL對(duì)象在不同供應(yīng)商的計(jì)算設(shè)備之間共享，能夠?qū)崿F(xiàn)異構(gòu)系統(tǒng)的高性能和易編程性枯夜。

第3章探討了在大規(guī)模并行GPU上的線程通信和同步弯汰。在GPU上，數(shù)千個(gè)線程同時(shí)運(yùn)行湖雹，處理器的性能很大程度上取決于線程間通信和同步的效率咏闪。了解現(xiàn)代GPU支持哪些機(jī)制以及它們對(duì)算法設(shè)計(jì)的影響是編寫(xiě)有效GPU代碼的關(guān)鍵問(wèn)題。由于傳統(tǒng)的GPGPU工作量是大規(guī)模并行的摔吏，線程之間幾乎沒(méi)有協(xié)作鸽嫂，早期的GPU只支持粗粒度的線程通信和同步。如今征讲，當(dāng)前的趨勢(shì)是加速更多的多樣化工作負(fù)載据某，這意味著粗粒度機(jī)制是利用并行性的主要限制因素。最新的工業(yè)標(biāo)準(zhǔn)編程框架OpenCL 2.0引入了細(xì)粒度線程通信和同步支持诗箍，以解決這個(gè)問(wèn)題癣籽。本章討論了現(xiàn)代GPU上可用的粗粒度和細(xì)粒度線程同步和通信機(jī)制。

第4章著重討論GPU上的軟件級(jí)任務(wù)調(diào)度滤祖。為了充分利用眾核處理器的潛力筷狼，任務(wù)調(diào)度是一個(gè)關(guān)鍵問(wèn)題。與CPU相反匠童，GPU缺乏程序員或編譯器控制調(diào)度的必要API埂材。因此，使用現(xiàn)代GPU上的硬件調(diào)度器以靈活的方式是困難的汤求。本章介紹了一個(gè)編譯器和運(yùn)行時(shí)框架俏险，以自動(dòng)轉(zhuǎn)換和優(yōu)化GPU程序，以實(shí)現(xiàn)對(duì)流式多處理器（SM）的可控任務(wù)調(diào)度首昔」押龋框架的中心是以SM為中心的轉(zhuǎn)換糙俗，它解決了硬件調(diào)度器的復(fù)雜性并提供了調(diào)度能力勒奇。該框架為新的優(yōu)化提供了許多機(jī)會(huì)，其中本章介紹了三個(gè)用于優(yōu)化并行性巧骚、局部性和處理器分區(qū)的優(yōu)化赊颠。廣泛的實(shí)驗(yàn)表明，這些優(yōu)化可以在多種場(chǎng)景下顯著提高一組GPU程序的性能劈彪。

第5章研究了GPU上數(shù)據(jù)放置的復(fù)雜性竣蹦，并介紹了PORPLE，一個(gè)軟件框架沧奴，展示了如何自動(dòng)解決一個(gè)給定的GPU應(yīng)用程序的復(fù)雜性痘括。數(shù)據(jù)放置的概念是一個(gè)重要的問(wèn)題，因?yàn)楝F(xiàn)代GPU內(nèi)存系統(tǒng)由許多具有不同特性的組件組成，問(wèn)題是如何將數(shù)據(jù)放置在各種內(nèi)存組件上纲菌。

第2部分介紹了一些針對(duì)GPU的有用算法和應(yīng)用挠日。它包括第6-14章。第6章關(guān)注生物序列分析翰舌。高通量的DNA測(cè)序技術(shù)導(dǎo)致生物數(shù)據(jù)庫(kù)的指數(shù)增長(zhǎng)嚣潜。必須分析和解釋這些生物序列。

為了確定生物數(shù)據(jù)的功能和結(jié)構(gòu)椅贱，需要進(jìn)行生物序列分析懂算。然而，生物數(shù)據(jù)庫(kù)的增長(zhǎng)速度遠(yuǎn)遠(yuǎn)超過(guò)了單核處理器的性能庇麦。隨著眾核處理器的出現(xiàn)计技，可以利用生物序列分析工具。本章討論了GPU在兩個(gè)主要序列比較問(wèn)題中的最新進(jìn)展：序列對(duì)比和序列-配置文件比較女器。

第13章討論了GPU上的圖算法合蔽。首先介紹和比較了表示和分析圖的主要數(shù)據(jù)結(jié)構(gòu)和技術(shù)薄货。接著討論了GPU的高效圖算法的理論和最新研究。本章所關(guān)注的算法主要是遍歷算法（廣度優(yōu)先搜索）、單源最短路徑（Dijkstra污尉、Bellman-Ford、delta stepping沥曹、混合算法）和全源最短路徑（Floyd-Warshall）等髓帽。隨后，本章討論負(fù)載平衡和存儲(chǔ)器訪問(wèn)技術(shù)驳阎，概述其主要問(wèn)題和管理技術(shù)抗愁。

第8章考慮了使用GPU對(duì)序列進(jìn)行對(duì)齊，探討了使用GPU對(duì)兩個(gè)和三個(gè)序列進(jìn)行最優(yōu)對(duì)齊的方法呵晚。對(duì)兩個(gè)序列進(jìn)行對(duì)齊的問(wèn)題通常稱(chēng)為序列對(duì)齊蜘腌。該章節(jié)還介紹了在NVIDIA Tesla C2050上進(jìn)行的實(shí)驗(yàn)結(jié)果。

第9章介紹了GPU上解決三對(duì)角系統(tǒng)的增強(qiáng)塊Cimmino分布式（ABCD）算法饵隙。三對(duì)角系統(tǒng)的特殊結(jié)構(gòu)在科學(xué)和工程問(wèn)題中經(jīng)常出現(xiàn)撮珠，例如交替方向隱式方法、流體模擬和泊松方程金矛。本章介紹了在GPU上解決三對(duì)角系統(tǒng)的ABCD方法的并行化芯急。在各種方面中，本章探討了邊界填充技術(shù)以消除GPU上的執(zhí)行分支等驶俊。還采用了各種性能優(yōu)化技術(shù)娶耍，如內(nèi)存合并，以進(jìn)一步提高性能饼酿。性能評(píng)估顯示榕酒，GPU實(shí)現(xiàn)比傳統(tǒng)CPU版本提高了超過(guò)24倍的速度胚膊。

第10章討論了線性和混合整數(shù)規(guī)劃方法。它表明運(yùn)營(yíng)研究（OR）社區(qū)中的復(fù)雜問(wèn)題可以從GPU中獲益想鹰。作者還通過(guò)突出不同作者如何克服實(shí)現(xiàn)困難澜掩，介紹了應(yīng)用于線性和混合整數(shù)規(guī)劃的GPU計(jì)算的主要貢獻(xiàn)的調(diào)查結(jié)果。

第11章考慮了平面圖最短路徑計(jì)算的加速實(shí)現(xiàn)杖挣。針對(duì)兩種最短路徑問(wèn)題肩榕，描述了三種算法及其相關(guān)的GPU實(shí)現(xiàn)。第一個(gè)算法解決了全源最短路徑問(wèn)題惩妇，而第二個(gè)算法則為更好的并行擴(kuò)展屬性和支持存儲(chǔ)器訪問(wèn)而容易株汉。第三個(gè)算法解決了單源最短路徑查詢(xún)問(wèn)題。實(shí)現(xiàn)結(jié)果顯示歌殃，同時(shí)利用256個(gè)GPU的計(jì)算能力乔妈。因此，改進(jìn)間隔比現(xiàn)有的并行方法高出一個(gè)數(shù)量級(jí)氓皱。

第12章討論了GPU上的排序算法路召。簡(jiǎn)要介紹了CUDA編程、記憶和計(jì)算模型以及通用GPU波材。

第三部分集中討論架構(gòu)和性能問(wèn)題股淡，其中包括15-18章。在第15章中廷区，引入了一個(gè)框架來(lái)加速GPU應(yīng)用程序中的瓶頸唯灵。困難的是，由于異構(gòu)應(yīng)用程序的要求不同隙轻，在GPU中資源的利用存在不平衡埠帕，因此在執(zhí)行過(guò)程中出現(xiàn)了不同的瓶頸。在本章中玖绿，介紹了一個(gè)核心助攻瓶頸加速（CABA）框架敛瓷，它利用空閑的片上資源來(lái)消除GPU執(zhí)行中的不同瓶頸。 CABA提供靈活的機(jī)制來(lái)自動(dòng)生成在GPU核心上執(zhí)行特定任務(wù)的“輔助卷積”斑匪。換句話說(shuō)呐籽，它使用空閑的計(jì)算單元和管道來(lái)緩解內(nèi)存帶寬瓶頸。因此秤标，它提高了GPU的性能和效率绝淡。 CABA架構(gòu)然后被深入地討論和評(píng)估宙刘，以在GPU存儲(chǔ)器子系統(tǒng)層次結(jié)構(gòu)中有效和靈活地執(zhí)行數(shù)據(jù)壓縮苍姜。結(jié)果表明，使用CABA進(jìn)行數(shù)據(jù)壓縮悬包，平均可提高41.7％（最大2.6倍）的性能衙猪，跨越寬范圍的內(nèi)存帶寬敏感GPGPU應(yīng)用程序。

第16章考慮通過(guò)神經(jīng)算法變換來(lái)加速GPU〉媸停基本上丝格，通過(guò)利用可編程網(wǎng)格，可擴(kuò)展的設(shè)計(jì)棵譬，具有靈活的內(nèi)存模型和具有高度優(yōu)化的編譯器显蝌，實(shí)現(xiàn)GPU的高性能增益。因此订咸，算法變換技術(shù)可以有效地調(diào)整程序以更好地利用GPU資源曼尊，從而提高性能和效率。本章重點(diǎn)介紹了Algorithm Transformation Toolkit（ATT）脏嚷，它是一個(gè)用于加速GPU應(yīng)用程序的工具包骆撇，該工具包實(shí)現(xiàn)了一組算法變換，可以幫助程序員更有效地優(yōu)化GPU應(yīng)用程序父叙。 ATT的評(píng)估表明神郊，它可以提供高達(dá)3.00倍的性能改進(jìn)，平均改進(jìn)1.51倍趾唱。

實(shí)現(xiàn)大量數(shù)據(jù)級(jí)并行和采用單指令多線程（SIMT）執(zhí)行模型涌乳。許多應(yīng)用領(lǐng)域都受益于GPU，包括識(shí)別甜癞、游戲爷怀、數(shù)據(jù)分析、天氣預(yù)測(cè)和多媒體等带欢。大多數(shù)這些應(yīng)用都是適合進(jìn)行近似計(jì)算的優(yōu)秀候選者运授。因此，采用近似計(jì)算技術(shù)可以提高GPU的性能和能效乔煞。在各種近似技術(shù)中吁朦，神經(jīng)加速器能夠帶來(lái)顯著的性能和效率提升。本章介紹了一種神經(jīng)加速的GPU（NGPU）架構(gòu)渡贾，將神經(jīng)加速嵌入到GPU加速器中逗宜，同時(shí)不影響它們的SIMT執(zhí)行。

作者在第17章介紹了一種帶有動(dòng)態(tài)帶寬分配的異構(gòu)光子網(wǎng)絡(luò)芯片空骚，用于GPU纺讲。未來(lái)的多核芯片預(yù)計(jì)將擁有數(shù)百個(gè)異構(gòu)組件，包括處理節(jié)點(diǎn)囤屹、分布式內(nèi)存熬甚、定制邏輯、GPU單元和可編程織物肋坚。由于未來(lái)芯片預(yù)計(jì)會(huì)同時(shí)運(yùn)行多個(gè)不同的并行工作負(fù)載乡括，因此不同的通信核心將需要不同的帶寬肃廓。因此，異構(gòu)網(wǎng)絡(luò)芯片（NoC）架構(gòu)的存在是必要的诲泌。最近的研究表明盲赊，光子互連能夠?qū)崿F(xiàn)高帶寬和能效性能較好的芯片內(nèi)數(shù)據(jù)傳輸。本章討論了一種動(dòng)態(tài)異構(gòu)光子NoC（d-HetPNOC）架構(gòu)敷扫，采用動(dòng)態(tài)帶寬分配哀蘑，以實(shí)現(xiàn)與同質(zhì)光子NoC架構(gòu)相比更好的性能和能效。

第一章 可靠GPU編程的形式化分析技術(shù)：現(xiàn)有解決方案和后續(xù)的改進(jìn)方向

一葵第、GPUs IN SUPPORT OF PARALLEL COMPUTING

圖形處理單元（GPU）是當(dāng)前應(yīng)用對(duì)強(qiáng)制功耗和有限電線延遲硬件的計(jì)算需求的自然產(chǎn)物 [2]递礼。GPU通過(guò)使用較簡(jiǎn)單的核心來(lái)實(shí)現(xiàn)比CPU更高的計(jì)算效率，并通過(guò)切換停滯的線程來(lái)隱藏內(nèi)存延遲羹幸。

總體而言脊髓，GPU面向吞吐量的執(zhí)行模型非常適合許多數(shù)據(jù)并行應(yīng)用程序。

GPU的演進(jìn)速度是引人注目的：從1997年的3M晶體管Nvidia NV3開(kāi)始栅受，到2012年的7B晶體管Nvidia GK110（Kepler）[3]将硝。Nvidia的CUDA編程模型于2007年推出，提供了比編寫(xiě)像素和頂點(diǎn)著色器的巴洛克符號(hào)更高的一步屏镊，而最近的CUDA 7.5 [4]提供了多功能的并發(fā)原語(yǔ)依疼。OpenCL是一種業(yè)界標(biāo)準(zhǔn)的編程模型 [5]，得到了包括Nvidia而芥、AMD律罢、ARM和Imagination Technologies在內(nèi)的所有主要設(shè)備廠商的支持，對(duì)數(shù)萬(wàn)個(gè)核心的計(jì)算意圖提供了簡(jiǎn)單棍丐、可移植的映射误辑。

GPU現(xiàn)在遠(yuǎn)遠(yuǎn)超出了圖形應(yīng)用程序，成為我們?cè)谟螒蚋璺辍⒕W(wǎng)絡(luò)搜索巾钉、基因測(cè)序和高性能超級(jí)計(jì)算等各個(gè)領(lǐng)域追求并行性的重要組成部分。

1秘案、并行計(jì)算和GPU代碼中的錯(cuò)誤和不足

編寫(xiě)正確的程序一直是計(jì)算機(jī)科學(xué)的一個(gè)基本挑戰(zhàn)砰苍，即使在圖靈 [6] 的時(shí)代也是如此。使用CPU線程或消息傳遞（如MPI）編寫(xiě)的并行程序比順序程序更容易出錯(cuò)阱高，因?yàn)槌绦騿T必須編碼線程之間的同步和通信邏輯赚导，并安排共享資源（主要是內(nèi)存）。這種情況導(dǎo)致錯(cuò)誤很難定位和修復(fù)赤惊。

與通用并發(fā)程序相比吼旧，GPU程序是“尷尬的并行”的，線程受到受控規(guī)則的控制和同步荐捻。然而黍少，GPU程序提出了某些獨(dú)特的調(diào)試難題，這些挑戰(zhàn)尚未得到足夠的關(guān)注处面，正如第2和第3節(jié)中所詳細(xì)討論GPU錯(cuò)誤的情況所示厂置。

如果不加檢查，GPU錯(cuò)誤可能會(huì)成為程序無(wú)法運(yùn)行的阻礙魂角，使由價(jià)值數(shù)百萬(wàn)美元的科學(xué)項(xiàng)目產(chǎn)生的模擬結(jié)果毫無(wú)用處昵济。突然的不可解釋的崩潰、不可重復(fù)的執(zhí)行以及不可重復(fù)的科學(xué)結(jié)果野揪，通常被忽視在exascale計(jì)算的嘈雜聲中访忿。

然而這些錯(cuò)誤確實(shí)會(huì)發(fā)生，并且嚴(yán)重地使專(zhuān)家們感到擔(dān)憂斯稳，他們常常艱難地啟動(dòng)新投入的機(jī)器海铆，或者需要停止他們正在進(jìn)行的有用科學(xué)研究而進(jìn)行修復(fù)。

本章的目的是描述這些挑戰(zhàn)的一些解決方案挣惰，提供對(duì)正在開(kāi)發(fā)的解決方案的理解卧斟，并描述仍需完成的工作內(nèi)容。

在高層次介紹GPU后（第2節(jié)）憎茂，我們以能夠使程序分析和驗(yàn)證社區(qū)受益為目的珍语，以一種回顧某些關(guān)鍵GPU正確性問(wèn)題的方式進(jìn)行調(diào)查（第3節(jié)）。問(wèn)題是：什么是正確性挑戰(zhàn)竖幔，我們?nèi)绾螐墓餐χ惺芤嬉越鉀Q可擴(kuò)展性和可靠性問(wèn)題板乙？然后回答了一個(gè)問(wèn)題：我們?nèi)绾谓⒖梢蕴幚斫裉旌图磳⒌絹?lái)的異構(gòu)并發(fā)形式的嚴(yán)格的正確性檢查工具？為此拳氢，我們討論了已有的有助于建立正確性的工具募逞，提供了它們?nèi)绾芜\(yùn)作的高層次描述，并總結(jié)了它們的限制（第4節(jié)）馋评。我們通過(guò)呼吁行動(dòng)結(jié)束本章凡辱，展現(xiàn)了我們進(jìn)一步開(kāi)展這項(xiàng)工作的觀點(diǎn)，這需要通過(guò)（a）研究驅(qū)動(dòng)的GPU加速軟件正確性檢查方法的進(jìn)展來(lái)解決開(kāi)放性問(wèn)題栗恩，(b) 通過(guò)傳播和技術(shù)轉(zhuǎn)移活動(dòng)來(lái)增加該工業(yè)分析工具的使用率（第5節(jié)）透乾。

2、快速介紹GPU

GPU通常被用作并行協(xié)處理器磕秤，在一個(gè)異構(gòu)系統(tǒng)中由主機(jī)CPU控制乳乌。在這種設(shè)置下，具有豐富并行性的任務(wù)可以作為核函數(shù)(offloaded to the GPU)放入GPU內(nèi)：一個(gè)指定任意線程行為的模板市咆。圖1展示了一個(gè)CUDA核函數(shù)汉操，用于執(zhí)行兩個(gè)向量的并行點(diǎn)乘，取自CUDA 5.0 SDK [7]蒙兰，我們將其用作一個(gè)示例說(shuō)明磷瘤。我們對(duì)GPU編程模型進(jìn)行簡(jiǎn)要概述芒篷。我們介紹每個(gè)概念和組件時(shí)，給出該概念或組件的CUDA術(shù)語(yǔ)采缚，并在括號(hào)中用OpenCL術(shù)語(yǔ)（如果有不同）针炉，之后使用CUDA術(shù)語(yǔ)。

2.1 線程的組織

內(nèi)核在GPU上通過(guò)許多輕量級(jí)線程(工作項(xiàng))以一種分層的方式作為一個(gè)網(wǎng)格(NDRange)的線程塊(work-groups)組織起來(lái)執(zhí)行扳抽，如圖所示篡帕。

這段代碼實(shí)現(xiàn)了一個(gè)向量點(diǎn)積運(yùn)算。其中ACCUM_N定義了每個(gè)線程塊中共享內(nèi)存的大小贸呢。主要分為兩個(gè)部分：

1.計(jì)算部分和：在每個(gè)線程塊內(nèi)镰烧，每個(gè)線程都會(huì)計(jì)算一段連續(xù)的部分和，并保存在共享內(nèi)存的accumResult數(shù)組中楞陷。具體地怔鳖，每個(gè)線程從自己的tid開(kāi)始，每次加上一個(gè)塊內(nèi)線程數(shù)bdim固蛾，直到遍歷所有ACCUM_N個(gè)元素败砂。對(duì)于每個(gè)元素，內(nèi)部執(zhí)行一個(gè)循環(huán)魏铅，將d_A和d_B對(duì)應(yīng)下標(biāo)的元素相乘昌犹，并累加到sum中，最終將sum存儲(chǔ)在accumResult對(duì)應(yīng)的位置览芳。

2.歸約匯總：在所有線程都計(jì)算完部分和后斜姥，需要使用歸約操作得到整個(gè)向量的點(diǎn)積。具體地沧竟，每個(gè)線程對(duì)相鄰的兩個(gè)部分和執(zhí)行累加操作铸敏，最終得到整個(gè)向量的點(diǎn)積。歸約過(guò)程中需要保證線程同步悟泵，使用__syncthreads()函數(shù)實(shí)現(xiàn)杈笔。最后，主線程將整個(gè)向量的點(diǎn)積存儲(chǔ)在d_c中糕非，并返回蒙具。

1. 定義一個(gè)名為 ACCUM_N 的宏，其值為 1024朽肥。

2. 定義一個(gè)名為 dotProduct 的全局函數(shù)禁筏，使用 GPU 加速計(jì)算矩陣 A 和矩陣 B 的向量點(diǎn)積，結(jié)果存儲(chǔ)在 d_c 指針指向的位置衡招。

3. 使用 __shared__ 關(guān)鍵字定義一個(gè)名為 accumResult 的共享數(shù)組篱昔，長(zhǎng)度為 ACCUM_N，用來(lái)存儲(chǔ)每個(gè)線程計(jì)算出的部分和。

4. 獲取當(dāng)前線程在塊內(nèi)的線程編號(hào)州刽。

5. 獲取塊內(nèi)的線程數(shù)空执。

6. 計(jì)算出當(dāng)前線程需要計(jì)算的元素下標(biāo)的增量 ACCUM_N。

7. 遍歷每個(gè)線程需要計(jì)算的部分和穗椅。

8. 對(duì)每個(gè)元素進(jìn)行計(jì)算辨绊，將結(jié)果存儲(chǔ)在名為 sum 的變量中。

9. 將計(jì)算得到的部分和存儲(chǔ)在共享數(shù)組 accumResult 中房待。

10. 完成對(duì)每個(gè)線程的部分和的計(jì)算邢羔。

11. 進(jìn)行規(guī)約操作驼抹，將所有線程計(jì)算出的部分和相加桑孩，得到最終的結(jié)果。

12. 使用 __syncthreads() 函數(shù)同步所有線程框冀。

13. 對(duì)于每個(gè) stride流椒，每個(gè)線程計(jì)算出當(dāng)前位置的元素和距離當(dāng)前位置 stride 的位置的元素的和。

14. 將結(jié)果存儲(chǔ)在 accumResult 中明也。

15. 完成規(guī)約操作宣虾。

16. 如果當(dāng)前線程是編號(hào)為 0 的線程，則將最終計(jì)算結(jié)果存儲(chǔ)在 d_c 指針指向的位置中温数。

例如绣硝，使用每個(gè)256個(gè)線程的4個(gè)線程塊的網(wǎng)格將產(chǎn)生1024個(gè)線程，每個(gè)線程都運(yùn)行核函數(shù)的副本撑刺。圖1中的__global__注釋表示dotProduct是一個(gè)核函數(shù)鹉胖。

在核函數(shù)內(nèi)部，一個(gè)線程可以使用內(nèi)置函數(shù)（如threadIdx（get_local_id）和blockDim（get_local_size））查詢(xún)其在網(wǎng)格層次結(jié)構(gòu)中的位置（以及網(wǎng)格和線程塊的尺寸）够傍。網(wǎng)格和塊可以是多維的甫菠，例如，blockDim.x（get_local_size（0））和threadIdx.x（get_local_id（0））分別提供線程塊的大小和第一維度內(nèi)的線程的ID冕屯。這使線程可以對(duì)不同的數(shù)據(jù)進(jìn)行操作寂诱，并通過(guò)核函數(shù)跟蹤不同的執(zhí)行路徑。

2.2 內(nèi)存空間

GPU上的線程可以訪問(wèn)多個(gè)內(nèi)存空間中的數(shù)據(jù)安聘，這些空間按照反映線程組織的層次結(jié)構(gòu)排列痰洒，如圖2所示。按下降順序排列浴韭，它們是大小带迟、范圍-可見(jiàn)性和延遲，它們是：

囱桨。在CUDA中仓犬，一個(gè)網(wǎng)格中的所有線程都能看到的大型全局內(nèi)存。在圖1中舍肠，CUDA指針內(nèi)核參數(shù)指向全局內(nèi)存數(shù)組搀继，因此d_a窘面、d_b和d_c是全局?jǐn)?shù)組。內(nèi)核計(jì)算d_A和d_B數(shù)組的點(diǎn)積并將結(jié)果存儲(chǔ)在d_c中（一個(gè)單元數(shù)組）叽躯。

财边。每個(gè)線程塊之間可見(jiàn)的線程塊共享（本地）內(nèi)存。圖1中的__shared__注釋指定accumResult數(shù)組駐留在共享存儲(chǔ)器中点骑。每個(gè)線程塊都有此數(shù)組的不同副本酣难，用于累加部分點(diǎn)積值。

黑滴。一個(gè)小的每線程私有內(nèi)存憨募。圖1中的循環(huán)變量i、j和stride駐留在私有內(nèi)存中袁辈。每個(gè)線程都有這些變量的單獨(dú)副本菜谣。程序員有責(zé)任在全局和共享內(nèi)存空間之間協(xié)調(diào)數(shù)據(jù)移動(dòng)。內(nèi)存合并是出于性能原因的重要屬性晚缩。當(dāng)相鄰的線程訪問(wèn)連續(xù)的內(nèi)存位置時(shí)尾膊，硬件可以將這些訪問(wèn)合成更少的內(nèi)存事務(wù)，從而增加帶寬荞彼。

雖然線程共享某些內(nèi)存空間冈敛，但內(nèi)存寫(xiě)入不會(huì)立即對(duì)所有線程可見(jiàn)（因?yàn)檫@會(huì)嚴(yán)重影響性能）。在計(jì)算機(jī)體系結(jié)構(gòu)中鸣皂，內(nèi)存一致性模型的概念用于明確解釋線程何時(shí)可以觀察到其他線程的寫(xiě)入抓谴。GPU編程模型指定了一個(gè)弱內(nèi)存一致性模型[4,5]。也就是說(shuō)签夭，給定線程的更新和訪問(wèn)順序不能保證由其他線程觀察到齐邦。

2.3 屏障同步（跟flink的streaming水位機(jī)制類(lèi)似）

屏障可以安全地在同一線程塊的線程之間進(jìn)行通信。

屏障操作會(huì)導(dǎo)致線程在所有線程塊的所有線程都到達(dá)同一屏障之前停頓第租。實(shí)際上措拇，所有線程必須在相同的控制流下達(dá)到屏障，才能避免屏障分歧問(wèn)題慎宾，這會(huì)導(dǎo)致未定義的行為丐吓。點(diǎn)積核使用屏障同步（indicated by __syncthreads()）來(lái)確保更新共享數(shù)組accumResult的正確順序。

屏障不能用于不同線程塊之間的線程同步趟据。相反券犁，這要求程序員將工作負(fù)載分成多個(gè)在序列中調(diào)用的內(nèi)核。不同線程塊中的線程可以使用原子操作通過(guò)全局內(nèi)存進(jìn)行通信汹碱。

2.4 線程束和鎖步執(zhí)行

在Nvidia GPU上粘衬，硬件通過(guò)動(dòng)態(tài)分區(qū)將線程塊分成一組線程束；AMD GPU有類(lèi)似波前的概念。目前稚新，Nvidia指定一個(gè)線程束是線程塊的32個(gè)相鄰線程的集合勘伺。同一線程束中的線程以鎖步模式執(zhí)行，因此隱式地進(jìn)行同步褂删；我們?cè)诘?節(jié)中對(duì)此現(xiàn)象的機(jī)會(huì)和風(fēng)險(xiǎn)進(jìn)行評(píng)論飞醉。

2.5 點(diǎn)積的例子

現(xiàn)在我們已經(jīng)有了適當(dāng)?shù)难b備來(lái)討論圖1中的內(nèi)核⊥头В考慮使用128個(gè)線程塊（即blockDim.x = 128）調(diào)用內(nèi)核時(shí)缅帘，輸入數(shù)組的長(zhǎng)度n等于4096。

內(nèi)核有兩個(gè)并行階段难衰。在第一階段钦无，部分積累到共享數(shù)組accumResult中。外部循環(huán)為accumResult數(shù)組的每個(gè)元素分配一個(gè)不同的線程召衔。由于元素(ACCUM_N)比線程多常挚，外循環(huán)為每個(gè)線程分配了ACCUM_N/blockDim.x = 8個(gè)元素寸爆。例如，線程0分配元素0粮彤、128兵志、256醇蝴、…、896想罕。輸出每個(gè)元素i的結(jié)果在線程私有變量sum中累加悠栓，并由內(nèi)循環(huán)累加。內(nèi)循環(huán)在ACCUM_N間隔內(nèi)執(zhí)行n /ACCUM_N = 4個(gè)部分積按价。例如惭适，當(dāng)i = 0時(shí)，線程0的內(nèi)循環(huán)將計(jì)算{j?0楼镐、1024癞志、2048、3072}中的∑aj bj 的積框产。存取的步長(zhǎng)stride確保內(nèi)存在全局存儲(chǔ)器中訪問(wèn)是合并的凄杯。

在第二個(gè)階段，部分積被縮減為最終期望結(jié)果秉宿。內(nèi)核使用并行樹(shù)縮減戒突，而不是逐個(gè)序列求和ACCUM_N元素∶枘溃縮減是使用邏輯樹(shù)執(zhí)行的膊存。給出了8個(gè)元素的簡(jiǎn)化縮減樹(shù)，如圖3所示。循環(huán)的每次迭代——使用降冪二步長(zhǎng)值——對(duì)應(yīng)于樹(shù)的不同級(jí)別隔崎。屏障確保給定級(jí)別的更新在下一級(jí)的任何訪問(wèn)之前有序嗡载，并因此成為線程間通信的一種形式。

3仍稀、GPU編程中的正確性問(wèn)題

現(xiàn)在我們概述四個(gè)關(guān)鍵的正確性問(wèn)題洼滚，這些問(wèn)題可能使從CPU到GPU的轉(zhuǎn)換困難。

3.1 數(shù)據(jù)競(jìng)爭(zhēng)

不足或錯(cuò)誤的障礙物barrier可能導(dǎo)致數(shù)據(jù)競(jìng)爭(zhēng)技潘。如果兩個(gè)線程訪問(wèn)相同的內(nèi)存位置（在全局或共享內(nèi)存中）遥巴，其中至少一個(gè)訪問(wèn)是非原子訪問(wèn)，至少一個(gè)訪問(wèn)修改該位置享幽，并且沒(méi)有涉及兩個(gè)線程的干預(yù)屏障同步铲掐，那么就會(huì)發(fā)生數(shù)據(jù)競(jìng)爭(zhēng)（在OpenCL 2.0中，某些原子操作之間的同步也可以用于避免競(jìng)速;我們?cè)谶@里不深入討論細(xì)節(jié)）值桩。例如摆霉，考慮圖1中點(diǎn)積內(nèi)核的原地樹(shù)規(guī)約。第20行的屏障確保給定迭代中（當(dāng)stride = k時(shí)奔坟，例如）的所有訪問(wèn)在下一次迭代（當(dāng)stride = k / 2時(shí)）的任何訪問(wèn)之前排序携栋。如果省略該屏障，則可能發(fā)生數(shù)據(jù)競(jìng)爭(zhēng)咳秉。例如婉支，帶有stride 2的線程0和帶有stride 4的線程2將在累積結(jié)果[8]上進(jìn)行競(jìng)爭(zhēng)。

在大多數(shù)并發(fā)程序中澜建，競(jìng)爭(zhēng)是某些代碼嚴(yán)重問(wèn)題的明顯指標(biāo)向挖，包括由于不充分的原子性而可能導(dǎo)致不確定結(jié)果或語(yǔ)義上不連貫的更新。 GPU編程模型（例如炕舵，OpenCL）要求程序員編寫(xiě)沒(méi)有數(shù)據(jù)競(jìng)爭(zhēng)的代碼何之。包含數(shù)據(jù)競(jìng)爭(zhēng)的程序具有未定義的語(yǔ)義。因此咽筋，許多編譯器優(yōu)化假定無(wú)競(jìng)爭(zhēng)性;在存在競(jìng)爭(zhēng)的情況下溶推，它們可能會(huì)產(chǎn)生意外的結(jié)果。

可以通過(guò)保守的屏障放置來(lái)防止數(shù)據(jù)競(jìng)爭(zhēng)晤硕，但屏障的過(guò)度使用可能有兩個(gè)問(wèn)題悼潭。首先，屏障的執(zhí)行成本很高舞箍，因此不必要的同步引入不必要的性能開(kāi)銷(xiāo)舰褪。第二，將屏障放置在條件代碼中可能是危險(xiǎn)的疏橄，因?yàn)樗婕暗降?節(jié)中討論的屏障分歧問(wèn)題占拍。

在Nvidia GPU上略就，許多CUDA程序員選擇在僅需要同步同一warp中的線程之間時(shí)省略屏障。例如晃酒，點(diǎn)積內(nèi)核中第二階段的樹(shù)規(guī)約循環(huán)的最后六次迭代（請(qǐng)參見(jiàn)圖1）可能被替換為以下語(yǔ)句序列表牢，以避免當(dāng)stride小于或等于32時(shí)的顯式同步：

if(tid < 32) accumResult[tid] += accumResult[tid + 32];

if(tid < 16) accumResult[tid] += accumResult[tid + 16];

if(tid < 8) accumResult[tid] += accumResult[tid + 8];

if(tid < 4) accumResult[tid] += accumResult[tid + 4];

if(tid < 2) accumResult[tid] += accumResult[tid + 2];

if(tid < 1) accumResult[tid] += accumResult[tid + 1];

此做法依賴(lài)于編譯器在優(yōu)化過(guò)程中保留隱式的warp內(nèi)同步。目前尚不清楚是否存在這種情況贝次。在CUDA編程指南5.0版中崔兴，有關(guān)warp內(nèi)同步的建議已被刪除，并且從業(yè)者對(duì)當(dāng)前平臺(tái)在這方面提供的保證存在分歧（例如蛔翅，有關(guān)該主題的NVIDIA論壇討論[9]）敲茄。

盡管如此，一些隨CUDA 5.0 SDK一起提供的示例依賴(lài)于隱式warp內(nèi)同步山析。更令人驚訝的是堰燎，開(kāi)源基準(zhǔn)套件Parboil [10]和SHOC [11]中的OpenCL內(nèi)核刪除了前面提到的屏障。這顯然是錯(cuò)誤的笋轨，因?yàn)殒i步warp的概念并不是OpenCL規(guī)范的一部分秆剪。

最近的研究表明，Nvidia和AMD GPU表現(xiàn)出弱內(nèi)存行為（如第2節(jié)所討論的）爵政，可以觀察到非順序一致的執(zhí)行方式仅讽，這是引起微妙軟件缺陷的來(lái)源[12]。因?yàn)槿鮾?nèi)存效應(yīng)而產(chǎn)生的錯(cuò)誤將越來(lái)越相關(guān)茂卦，因?yàn)镚PU應(yīng)用程序?qū)⑾蚶眉?xì)粒度并發(fā)代替屏障同步移動(dòng)何什。

3.2 缺乏前向保證

GPU架構(gòu)中線程調(diào)度不公平是另一個(gè)缺陷的來(lái)源组哩。圖4展示了在CUDA中實(shí)現(xiàn)inter-block barrier的嘗試等龙。這個(gè)眾所周知的策略背后的思想如下。每個(gè)線程在塊內(nèi)同步（第3行）伶贰，之后每個(gè)線程塊的領(lǐng)袖（threadIdx.x = 0的線程）原子性地減少一個(gè)計(jì)數(shù)器（第5行）蛛砰。假設(shè)計(jì)數(shù)器初始值為總線程塊數(shù)（在第1行進(jìn)行注釋?zhuān)敲葱D(zhuǎn)直到計(jì)數(shù)器達(dá)到零（第6行）可能會(huì)足以確保每個(gè)領(lǐng)袖已通過(guò)減量（第5行）黍衙。 （請(qǐng)注意泥畅，計(jì)數(shù)器值是通過(guò)向計(jì)數(shù)器原子添加零來(lái)檢索的； atomicAdd返回操作的位置的舊值琅翻。使用原子操作而不是普通的load操作可以避免訪問(wèn)計(jì)數(shù)器時(shí)的數(shù)據(jù)競(jìng)爭(zhēng)位仁。）如果旋轉(zhuǎn)直到計(jì)數(shù)器達(dá)到零確實(shí)可以確保每個(gè)領(lǐng)袖已通過(guò)減量，則已實(shí)現(xiàn)全局同步方椎，領(lǐng)袖可以重新與其塊中的其余部分同步（第8行）聂抢，從而允許執(zhí)行繼續(xù)。柵欄（第2和第9行）的目的是確保全局同步點(diǎn)之前的內(nèi)存訪問(wèn)操作在全局同步點(diǎn)之后的內(nèi)存訪問(wèn)操作之前生效棠众。關(guān)于此策略及其問(wèn)題的社區(qū)討論琳疏，請(qǐng)參見(jiàn)例如Ref. [13]有决。

這個(gè)屏障的問(wèn)題在于它假定了線程塊之間的前進(jìn)。然而空盼，如果請(qǐng)求足夠數(shù)量的CUDA塊书幕，則它們將被分批調(diào)度：一定數(shù)量的塊將被調(diào)度。

并且必須在計(jì)算單元釋放用于安排更多塊之前完全運(yùn)行揽趾。這種情況會(huì)導(dǎo)致死鎖：在第一輪調(diào)度的塊的領(lǐng)導(dǎo)者會(huì)在第六行處自旋台汇，等待計(jì)數(shù)器由附加塊的領(lǐng)導(dǎo)者遞減；這些塊反過(guò)來(lái)將不能被調(diào)度篱瞎，直到初始?jí)K完成執(zhí)行励七。嘗試在OpenCL中實(shí)現(xiàn)全局同步的類(lèi)似方法也無(wú)法工作，原因是相同的奔缠。

這個(gè)例子說(shuō)明了線程塊之間的不公平調(diào)度掠抬。由于線程塊被分成了warp，使得相同warp中的線程以鎖步方式執(zhí)行校哎，執(zhí)行相同的指令序列两波，因此在CUDA線程塊內(nèi)的線程之間也會(huì)出現(xiàn)缺乏進(jìn)展的問(wèn)題。如果程序員試圖強(qiáng)制一個(gè)線程在同一warp中的另一個(gè)線程上忙等待闷哆，則會(huì)導(dǎo)致死鎖腰奋；進(jìn)展需要線程在運(yùn)行時(shí)真正執(zhí)行不同的指令。這破壞了基于忙等待的塊內(nèi)關(guān)鍵部分的天真實(shí)現(xiàn)抱怔。有關(guān)此問(wèn)題的示例和社區(qū)討論劣坊，請(qǐng)參見(jiàn)參考文獻(xiàn)[14]。

3.3 浮點(diǎn)精度

對(duì)于在浮點(diǎn)數(shù)據(jù)上計(jì)算的內(nèi)核而言屈留，實(shí)現(xiàn)與參考實(shí)現(xiàn)的等價(jià)性可能會(huì)更具挑戰(zhàn)性局冰，這在高性能計(jì)算中很常見(jiàn)。

在設(shè)計(jì)并行算法時(shí)灌危，假設(shè)浮點(diǎn)算符具有實(shí)數(shù)的代數(shù)性質(zhì)可能很方便康二。例如，求和歸約操作勇蝙，其中將數(shù)組的所有元素相加沫勿，可以通過(guò)計(jì)算樹(shù)并行化，如圖5所示味混。如果并行處理元素的數(shù)量超過(guò)數(shù)組的大小产雹，基于樹(shù)的方法允許在對(duì)數(shù)步驟中進(jìn)行歸約。該圖顯示了一個(gè)8個(gè)元素?cái)?shù)組的兩個(gè)可能的計(jì)算樹(shù)翁锡。如果加法滿足結(jié)合律蔓挖，即對(duì)于所有浮點(diǎn)值x、y和z盗誊，滿足法律（x + y）+ z = x +（y + z）时甚，其中+表示浮點(diǎn)加法隘弊，則計(jì)算樹(shù)將產(chǎn)生相同的結(jié)果。

然而荒适，正向順序?qū)崿F(xiàn)得到的左關(guān)聯(lián)和的結(jié)果與并行算法的結(jié)果相同梨熙。然而，眾所周知刀诬，浮點(diǎn)加法不可結(jié)合咽扇，因此結(jié)果預(yù)計(jì)會(huì)因并行算法的結(jié)構(gòu)而有所不同。

雖然這些問(wèn)題也適用于中央處理器陕壹，但對(duì)于圖形處理器可能更為重要的原因在于它們以吞吐量為導(dǎo)向质欲，并且緩存較小。因此糠馆，在GPU上使用雙精度算術(shù)代替幾乎所有的單精度算術(shù)會(huì)增加懲罰嘶伟，特別是由于GPU上雙精度單元的數(shù)量通常較少。與GPU相關(guān)的浮點(diǎn)精度語(yǔ)言特定問(wèn)題也存在又碌。例如九昧，在OpenCL中，有關(guān)是否準(zhǔn)確表示非正規(guī)化數(shù)的問(wèn)題是實(shí)現(xiàn)定義的毕匀，并且OpenCL提供了一種半精度數(shù)據(jù)類(lèi)型铸鹰，僅為相關(guān)運(yùn)算符指定了最小（而不是精確）精度要求皂岔。實(shí)現(xiàn)差異使編寫(xiě)在多個(gè)GPU平臺(tái)上行為精度可接受的高性能浮點(diǎn)代碼變得困難蹋笼。

（待補(bǔ)充中）