PS：最近不務正業(yè)對GPU產生了濃厚的興趣，可能是AI需要軟硬結合的原因？對這本書（2017年出版不算太舊）進行整理，以備不時之需。

# 前言

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

對于圖形應用程序對計算和帶寬的巨大需求已經(jīng)導致GPU成為主要的大規(guī)模并行架構。如今，具有浮點運算能力的通用GPU（GPGPU）在具有混合CPU-GPU架構的大多數(shù)頂級機器上廣泛用于通用計算。這樣的GPGPU新產品NVIDIA Pascal擁有3840個核心。

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

本書聚焦于基于GPU的系統(tǒng)的研究和實踐，并試圖解決它們的重要問題。涵蓋的主題涵蓋從硬件和架構問題到直接涉及應用程序或系統(tǒng)用戶的高級問題，包括并行編程工具和中間件支持的等等這種計算系統(tǒng)。本書分為四個部分，每個部分都包含由知名研究人員撰寫的章節(jié)，他們在該部分的主題下分享了他們的最新研究成果。

第1部分涉及不同的編程問題和工具。該部分由5個章節(jié)組成：1-5。第1章討論了如何以可靠的方式編寫GPU程序。盡管GPU程序在不同領域提供高計算吞吐量，但正確優(yōu)化代碼有些困難，這是由于GPU并發(fā)性的微妙性。因此，本章的重點是討論并解決一些事情，以使GPU編程更容易。該章節(jié)還提供了一些關于GPU軟件形式分析的最新進展。后者是一個重要問題，隨著GPU編程的進步。因此，新的GPU程序員應該學習驗證方法。

第2章介紹了異構集群的統(tǒng)一開放計算語言（OpenCL）框架（SnuCL）。它是一個免費的，開源的OpenCL框架，用于異構集群?；旧?，OpenCL是一種針對異構并行計算系統(tǒng)的編程模型，它定義了傳統(tǒng)處理器和加速器之間的抽象層。使用OpenCL的優(yōu)點是程序員編寫一個OpenCL應用程序，然后在任何支持OpenCL的系統(tǒng)上運行它。然而，它缺乏針對異構集群的目標能力。

SnuCL為程序員提供了一個單一、統(tǒng)一的OpenCL平臺映像，用于集群。借助SnuCL，OpenCL應用程序能夠利用計算節(jié)點中的計算設備，就像它們在主機節(jié)點中一樣。此外，使用SnuCL，可以將來自不同供應商的多個OpenCL平臺集成到單個平臺中。因此，OpenCL對象在不同供應商的計算設備之間共享，能夠實現(xiàn)異構系統(tǒng)的高性能和易編程性。

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

第4章著重討論GPU上的軟件級任務調度。為了充分利用眾核處理器的潛力，任務調度是一個關鍵問題。與CPU相反，GPU缺乏程序員或編譯器控制調度的必要API。因此，使用現(xiàn)代GPU上的硬件調度器以靈活的方式是困難的。本章介紹了一個編譯器和運行時框架，以自動轉換和優(yōu)化GPU程序，以實現(xiàn)對流式多處理器（SM）的可控任務調度?？蚣艿闹行氖且許M為中心的轉換，它解決了硬件調度器的復雜性并提供了調度能力。該框架為新的優(yōu)化提供了許多機會，其中本章介紹了三個用于優(yōu)化并行性、局部性和處理器分區(qū)的優(yōu)化。廣泛的實驗表明，這些優(yōu)化可以在多種場景下顯著提高一組GPU程序的性能。

第5章研究了GPU上數(shù)據(jù)放置的復雜性，并介紹了PORPLE，一個軟件框架，展示了如何自動解決一個給定的GPU應用程序的復雜性。數(shù)據(jù)放置的概念是一個重要的問題，因為現(xiàn)代GPU內存系統(tǒng)由許多具有不同特性的組件組成，問題是如何將數(shù)據(jù)放置在各種內存組件上。

第2部分介紹了一些針對GPU的有用算法和應用。它包括第6-14章。第6章關注生物序列分析。高通量的DNA測序技術導致生物數(shù)據(jù)庫的指數(shù)增長。必須分析和解釋這些生物序列。

為了確定生物數(shù)據(jù)的功能和結構，需要進行生物序列分析。然而，生物數(shù)據(jù)庫的增長速度遠遠超過了單核處理器的性能。隨著眾核處理器的出現(xiàn)，可以利用生物序列分析工具。本章討論了GPU在兩個主要序列比較問題中的最新進展：序列對比和序列-配置文件比較。

第13章討論了GPU上的圖算法。首先介紹和比較了表示和分析圖的主要數(shù)據(jù)結構和技術。接著討論了GPU的高效圖算法的理論和最新研究。本章所關注的算法主要是遍歷算法（廣度優(yōu)先搜索）、單源最短路徑（Dijkstra、Bellman-Ford、delta stepping、混合算法）和全源最短路徑（Floyd-Warshall）等。隨后，本章討論負載平衡和存儲器訪問技術，概述其主要問題和管理技術。

第8章考慮了使用GPU對序列進行對齊，探討了使用GPU對兩個和三個序列進行最優(yōu)對齊的方法。對兩個序列進行對齊的問題通常稱為序列對齊。該章節(jié)還介紹了在NVIDIA Tesla C2050上進行的實驗結果。

第9章介紹了GPU上解決三對角系統(tǒng)的增強塊Cimmino分布式（ABCD）算法。三對角系統(tǒng)的特殊結構在科學和工程問題中經(jīng)常出現(xiàn)，例如交替方向隱式方法、流體模擬和泊松方程。本章介紹了在GPU上解決三對角系統(tǒng)的ABCD方法的并行化。在各種方面中，本章探討了邊界填充技術以消除GPU上的執(zhí)行分支等。還采用了各種性能優(yōu)化技術，如內存合并，以進一步提高性能。性能評估顯示，GPU實現(xiàn)比傳統(tǒng)CPU版本提高了超過24倍的速度。

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

第11章考慮了平面圖最短路徑計算的加速實現(xiàn)。針對兩種最短路徑問題，描述了三種算法及其相關的GPU實現(xiàn)。第一個算法解決了全源最短路徑問題，而第二個算法則為更好的并行擴展屬性和支持存儲器訪問而容易。第三個算法解決了單源最短路徑查詢問題。實現(xiàn)結果顯示，同時利用256個GPU的計算能力。因此，改進間隔比現(xiàn)有的并行方法高出一個數(shù)量級。

第12章討論了GPU上的排序算法。簡要介紹了CUDA編程、記憶和計算模型以及通用GPU。

第三部分集中討論架構和性能問題，其中包括15-18章。在第15章中，引入了一個框架來加速GPU應用程序中的瓶頸。困難的是，由于異構應用程序的要求不同，在GPU中資源的利用存在不平衡，因此在執(zhí)行過程中出現(xiàn)了不同的瓶頸。在本章中，介紹了一個核心助攻瓶頸加速（CABA）框架，它利用空閑的片上資源來消除GPU執(zhí)行中的不同瓶頸。 CABA提供靈活的機制來自動生成在GPU核心上執(zhí)行特定任務的“輔助卷積”。換句話說，它使用空閑的計算單元和管道來緩解內存帶寬瓶頸。因此，它提高了GPU的性能和效率。 CABA架構然后被深入地討論和評估，以在GPU存儲器子系統(tǒng)層次結構中有效和靈活地執(zhí)行數(shù)據(jù)壓縮。結果表明，使用CABA進行數(shù)據(jù)壓縮，平均可提高41.7％（最大2.6倍）的性能，跨越寬范圍的內存帶寬敏感GPGPU應用程序。

第16章考慮通過神經(jīng)算法變換來加速GPU。基本上，通過利用可編程網(wǎng)格，可擴展的設計，具有靈活的內存模型和具有高度優(yōu)化的編譯器，實現(xiàn)GPU的高性能增益。因此，算法變換技術可以有效地調整程序以更好地利用GPU資源，從而提高性能和效率。本章重點介紹了Algorithm Transformation Toolkit（ATT），它是一個用于加速GPU應用程序的工具包，該工具包實現(xiàn)了一組算法變換，可以幫助程序員更有效地優(yōu)化GPU應用程序。 ATT的評估表明，它可以提供高達3.00倍的性能改進，平均改進1.51倍。

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

作者在第17章介紹了一種帶有動態(tài)帶寬分配的異構光子網(wǎng)絡芯片，用于GPU。未來的多核芯片預計將擁有數(shù)百個異構組件，包括處理節(jié)點、分布式內存、定制邏輯、GPU單元和可編程織物。由于未來芯片預計會同時運行多個不同的并行工作負載，因此不同的通信核心將需要不同的帶寬。因此，異構網(wǎng)絡芯片（NoC）架構的存在是必要的。最近的研究表明，光子互連能夠實現(xiàn)高帶寬和能效性能較好的芯片內數(shù)據(jù)傳輸。本章討論了一種動態(tài)異構光子NoC（d-HetPNOC）架構，采用動態(tài)帶寬分配，以實現(xiàn)與同質光子NoC架構相比更好的性能和能效。

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

一、GPUs IN SUPPORT OF PARALLEL COMPUTING

圖形處理單元（GPU）是當前應用對強制功耗和有限電線延遲硬件的計算需求的自然產物 [2]。GPU通過使用較簡單的核心來實現(xiàn)比CPU更高的計算效率，并通過切換停滯的線程來隱藏內存延遲。

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

GPU的演進速度是引人注目的：從1997年的3M晶體管Nvidia NV3開始，到2012年的7B晶體管Nvidia GK110（Kepler）[3]。Nvidia的CUDA編程模型于2007年推出，提供了比編寫像素和頂點著色器的巴洛克符號更高的一步，而最近的CUDA 7.5 [4]提供了多功能的并發(fā)原語。OpenCL是一種業(yè)界標準的編程模型 [5]，得到了包括Nvidia、AMD、ARM和Imagination Technologies在內的所有主要設備廠商的支持，對數(shù)萬個核心的計算意圖提供了簡單、可移植的映射。

GPU現(xiàn)在遠遠超出了圖形應用程序，成為我們在游戲、網(wǎng)絡搜索、基因測序和高性能超級計算等各個領域追求并行性的重要組成部分。

1、并行計算和GPU代碼中的錯誤和不足

編寫正確的程序一直是計算機科學的一個基本挑戰(zhàn)，即使在圖靈 [6] 的時代也是如此。使用CPU線程或消息傳遞（如MPI）編寫的并行程序比順序程序更容易出錯，因為程序員必須編碼線程之間的同步和通信邏輯，并安排共享資源（主要是內存）。這種情況導致錯誤很難定位和修復。

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

如果不加檢查，GPU錯誤可能會成為程序無法運行的阻礙，使由價值數(shù)百萬美元的科學項目產生的模擬結果毫無用處。突然的不可解釋的崩潰、不可重復的執(zhí)行以及不可重復的科學結果，通常被忽視在exascale計算的嘈雜聲中。

然而這些錯誤確實會發(fā)生，并且嚴重地使專家們感到擔憂，他們常常艱難地啟動新投入的機器，或者需要停止他們正在進行的有用科學研究而進行修復。

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

在高層次介紹GPU后（第2節(jié)），我們以能夠使程序分析和驗證社區(qū)受益為目的，以一種回顧某些關鍵GPU正確性問題的方式進行調查（第3節(jié)）。問題是：什么是正確性挑戰(zhàn)，我們如何從共同努力中受益以解決可擴展性和可靠性問題？然后回答了一個問題：我們如何建立可以處理今天和即將到來的異構并發(fā)形式的嚴格的正確性檢查工具？為此，我們討論了已有的有助于建立正確性的工具，提供了它們如何運作的高層次描述，并總結了它們的限制（第4節(jié)）。我們通過呼吁行動結束本章，展現(xiàn)了我們進一步開展這項工作的觀點，這需要通過（a）研究驅動的GPU加速軟件正確性檢查方法的進展來解決開放性問題，(b) 通過傳播和技術轉移活動來增加該工業(yè)分析工具的使用率（第5節(jié)）。

2、快速介紹GPU

GPU通常被用作并行協(xié)處理器，在一個異構系統(tǒng)中由主機CPU控制。在這種設置下，具有豐富并行性的任務可以作為核函數(shù)(offloaded to the GPU)放入GPU內：一個指定任意線程行為的模板。圖1展示了一個CUDA核函數(shù)，用于執(zhí)行兩個向量的并行點乘，取自CUDA 5.0 SDK [7]，我們將其用作一個示例說明。我們對GPU編程模型進行簡要概述。我們介紹每個概念和組件時，給出該概念或組件的CUDA術語，并在括號中用OpenCL術語（如果有不同），之后使用CUDA術語。

2.1 線程的組織

內核在GPU上通過許多輕量級線程(工作項)以一種分層的方式作為一個網(wǎng)格(NDRange)的線程塊(work-groups)組織起來執(zhí)行，如圖所示。

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

1.計算部分和：在每個線程塊內，每個線程都會計算一段連續(xù)的部分和，并保存在共享內存的accumResult數(shù)組中。具體地，每個線程從自己的tid開始，每次加上一個塊內線程數(shù)bdim，直到遍歷所有ACCUM_N個元素。對于每個元素，內部執(zhí)行一個循環(huán)，將d_A和d_B對應下標的元素相乘，并累加到sum中，最終將sum存儲在accumResult對應的位置。

2.歸約匯總：在所有線程都計算完部分和后，需要使用歸約操作得到整個向量的點積。具體地，每個線程對相鄰的兩個部分和執(zhí)行累加操作，最終得到整個向量的點積。歸約過程中需要保證線程同步，使用__syncthreads()函數(shù)實現(xiàn)。最后，主線程將整個向量的點積存儲在d_c中，并返回。

1. 定義一個名為 ACCUM_N 的宏，其值為 1024。

2. 定義一個名為 dotProduct 的全局函數(shù)，使用 GPU 加速計算矩陣 A 和矩陣 B 的向量點積，結果存儲在 d_c 指針指向的位置。

3. 使用 __shared__ 關鍵字定義一個名為 accumResult 的共享數(shù)組，長度為 ACCUM_N，用來存儲每個線程計算出的部分和。

4. 獲取當前線程在塊內的線程編號。

5. 獲取塊內的線程數(shù)。

6. 計算出當前線程需要計算的元素下標的增量 ACCUM_N。

7. 遍歷每個線程需要計算的部分和。

8. 對每個元素進行計算，將結果存儲在名為 sum 的變量中。

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

10. 完成對每個線程的部分和的計算。

11. 進行規(guī)約操作，將所有線程計算出的部分和相加，得到最終的結果。

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

13. 對于每個 stride，每個線程計算出當前位置的元素和距離當前位置 stride 的位置的元素的和。

14. 將結果存儲在 accumResult 中。

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

16. 如果當前線程是編號為 0 的線程，則將最終計算結果存儲在 d_c 指針指向的位置中。

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

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

2.2 內存空間

GPU上的線程可以訪問多個內存空間中的數(shù)據(jù)，這些空間按照反映線程組織的層次結構排列，如圖2所示。按下降順序排列，它們是大小、范圍-可見性和延遲，它們是：

。在CUDA中，一個網(wǎng)格中的所有線程都能看到的大型全局內存。在圖1中，CUDA指針內核參數(shù)指向全局內存數(shù)組，因此d_a、d_b和d_c是全局數(shù)組。內核計算d_A和d_B數(shù)組的點積并將結果存儲在d_c中（一個單元數(shù)組）。

。每個線程塊之間可見的線程塊共享（本地）內存。圖1中的__shared__注釋指定accumResult數(shù)組駐留在共享存儲器中。每個線程塊都有此數(shù)組的不同副本，用于累加部分點積值。

。一個小的每線程私有內存。圖1中的循環(huán)變量i、j和stride駐留在私有內存中。每個線程都有這些變量的單獨副本。程序員有責任在全局和共享內存空間之間協(xié)調數(shù)據(jù)移動。內存合并是出于性能原因的重要屬性。當相鄰的線程訪問連續(xù)的內存位置時，硬件可以將這些訪問合成更少的內存事務，從而增加帶寬。

雖然線程共享某些內存空間，但內存寫入不會立即對所有線程可見（因為這會嚴重影響性能）。在計算機體系結構中，內存一致性模型的概念用于明確解釋線程何時可以觀察到其他線程的寫入。GPU編程模型指定了一個弱內存一致性模型[4,5]。也就是說，給定線程的更新和訪問順序不能保證由其他線程觀察到。

2.3 屏障同步（跟flink的streaming水位機制類似）

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

屏障操作會導致線程在所有線程塊的所有線程都到達同一屏障之前停頓。實際上，所有線程必須在相同的控制流下達到屏障，才能避免屏障分歧問題，這會導致未定義的行為。點積核使用屏障同步（indicated by __syncthreads()）來確保更新共享數(shù)組accumResult的正確順序。

屏障不能用于不同線程塊之間的線程同步。相反，這要求程序員將工作負載分成多個在序列中調用的內核。不同線程塊中的線程可以使用原子操作通過全局內存進行通信。

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

在Nvidia GPU上，硬件通過動態(tài)分區(qū)將線程塊分成一組線程束；AMD GPU有類似波前的概念。目前，Nvidia指定一個線程束是線程塊的32個相鄰線程的集合。同一線程束中的線程以鎖步模式執(zhí)行，因此隱式地進行同步；我們在第3節(jié)中對此現(xiàn)象的機會和風險進行評論。

2.5 點積的例子

現(xiàn)在我們已經(jīng)有了適當?shù)难b備來討論圖1中的內核?？紤]使用128個線程塊（即blockDim.x = 128）調用內核時，輸入數(shù)組的長度n等于4096。

內核有兩個并行階段。在第一階段，部分積累到共享數(shù)組accumResult中。外部循環(huán)為accumResult數(shù)組的每個元素分配一個不同的線程。由于元素(ACCUM_N)比線程多，外循環(huán)為每個線程分配了ACCUM_N/blockDim.x = 8個元素。例如，線程0分配元素0、128、256、…、896。輸出每個元素i的結果在線程私有變量sum中累加，并由內循環(huán)累加。內循環(huán)在ACCUM_N間隔內執(zhí)行n /ACCUM_N = 4個部分積。例如，當i = 0時，線程0的內循環(huán)將計算{j?0、1024、2048、3072}中的∑aj bj 的積。存取的步長stride確保內存在全局存儲器中訪問是合并的。

在第二個階段，部分積被縮減為最終期望結果。內核使用并行樹縮減，而不是逐個序列求和ACCUM_N元素?？s減是使用邏輯樹執(zhí)行的。給出了8個元素的簡化縮減樹，如圖3所示。循環(huán)的每次迭代——使用降冪二步長值——對應于樹的不同級別。屏障確保給定級別的更新在下一級的任何訪問之前有序，并因此成為線程間通信的一種形式。

3、GPU編程中的正確性問題

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

3.1 數(shù)據(jù)競爭

不足或錯誤的障礙物barrier可能導致數(shù)據(jù)競爭。如果兩個線程訪問相同的內存位置（在全局或共享內存中），其中至少一個訪問是非原子訪問，至少一個訪問修改該位置，并且沒有涉及兩個線程的干預屏障同步，那么就會發(fā)生數(shù)據(jù)競爭（在OpenCL 2.0中，某些原子操作之間的同步也可以用于避免競速;我們在這里不深入討論細節(jié)）。例如，考慮圖1中點積內核的原地樹規(guī)約。第20行的屏障確保給定迭代中（當stride = k時，例如）的所有訪問在下一次迭代（當stride = k / 2時）的任何訪問之前排序。如果省略該屏障，則可能發(fā)生數(shù)據(jù)競爭。例如，帶有stride 2的線程0和帶有stride 4的線程2將在累積結果[8]上進行競爭。

在大多數(shù)并發(fā)程序中，競爭是某些代碼嚴重問題的明顯指標，包括由于不充分的原子性而可能導致不確定結果或語義上不連貫的更新。 GPU編程模型（例如，OpenCL）要求程序員編寫沒有數(shù)據(jù)競爭的代碼。包含數(shù)據(jù)競爭的程序具有未定義的語義。因此，許多編譯器優(yōu)化假定無競爭性;在存在競爭的情況下，它們可能會產生意外的結果。

可以通過保守的屏障放置來防止數(shù)據(jù)競爭，但屏障的過度使用可能有兩個問題。首先，屏障的執(zhí)行成本很高，因此不必要的同步引入不必要的性能開銷。第二，將屏障放置在條件代碼中可能是危險的，因為它涉及到第2節(jié)中討論的屏障分歧問題。

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

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];

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

盡管如此，一些隨CUDA 5.0 SDK一起提供的示例依賴于隱式warp內同步。更令人驚訝的是，開源基準套件Parboil [10]和SHOC [11]中的OpenCL內核刪除了前面提到的屏障。這顯然是錯誤的，因為鎖步warp的概念并不是OpenCL規(guī)范的一部分。

最近的研究表明，Nvidia和AMD GPU表現(xiàn)出弱內存行為（如第2節(jié)所討論的），可以觀察到非順序一致的執(zhí)行方式，這是引起微妙軟件缺陷的來源[12]。因為弱內存效應而產生的錯誤將越來越相關，因為GPU應用程序將向利用細粒度并發(fā)代替屏障同步移動。

3.2 缺乏前向保證

GPU架構中線程調度不公平是另一個缺陷的來源。圖4展示了在CUDA中實現(xiàn)inter-block barrier的嘗試。這個眾所周知的策略背后的思想如下。每個線程在塊內同步（第3行），之后每個線程塊的領袖（threadIdx.x = 0的線程）原子性地減少一個計數(shù)器（第5行）。假設計數(shù)器初始值為總線程塊數(shù)（在第1行進行注釋），那么旋轉直到計數(shù)器達到零（第6行）可能會足以確保每個領袖已通過減量（第5行）。 （請注意，計數(shù)器值是通過向計數(shù)器原子添加零來檢索的； atomicAdd返回操作的位置的舊值。使用原子操作而不是普通的load操作可以避免訪問計數(shù)器時的數(shù)據(jù)競爭。）如果旋轉直到計數(shù)器達到零確實可以確保每個領袖已通過減量，則已實現(xiàn)全局同步，領袖可以重新與其塊中的其余部分同步（第8行），從而允許執(zhí)行繼續(xù)。柵欄（第2和第9行）的目的是確保全局同步點之前的內存訪問操作在全局同步點之后的內存訪問操作之前生效。關于此策略及其問題的社區(qū)討論，請參見例如Ref. [13]。

這個屏障的問題在于它假定了線程塊之間的前進。然而，如果請求足夠數(shù)量的CUDA塊，則它們將被分批調度：一定數(shù)量的塊將被調度。

并且必須在計算單元釋放用于安排更多塊之前完全運行。這種情況會導致死鎖：在第一輪調度的塊的領導者會在第六行處自旋，等待計數(shù)器由附加塊的領導者遞減；這些塊反過來將不能被調度，直到初始塊完成執(zhí)行。嘗試在OpenCL中實現(xiàn)全局同步的類似方法也無法工作，原因是相同的。

這個例子說明了線程塊之間的不公平調度。由于線程塊被分成了warp，使得相同warp中的線程以鎖步方式執(zhí)行，執(zhí)行相同的指令序列，因此在CUDA線程塊內的線程之間也會出現(xiàn)缺乏進展的問題。如果程序員試圖強制一個線程在同一warp中的另一個線程上忙等待，則會導致死鎖；進展需要線程在運行時真正執(zhí)行不同的指令。這破壞了基于忙等待的塊內關鍵部分的天真實現(xiàn)。有關此問題的示例和社區(qū)討論，請參見參考文獻[14]。

3.3 浮點精度

對于在浮點數(shù)據(jù)上計算的內核而言，實現(xiàn)與參考實現(xiàn)的等價性可能會更具挑戰(zhàn)性，這在高性能計算中很常見。

在設計并行算法時，假設浮點算符具有實數(shù)的代數(shù)性質可能很方便。例如，求和歸約操作，其中將數(shù)組的所有元素相加，可以通過計算樹并行化，如圖5所示。如果并行處理元素的數(shù)量超過數(shù)組的大小，基于樹的方法允許在對數(shù)步驟中進行歸約。該圖顯示了一個8個元素數(shù)組的兩個可能的計算樹。如果加法滿足結合律，即對于所有浮點值x、y和z，滿足法律（x + y）+ z = x +（y + z），其中+表示浮點加法，則計算樹將產生相同的結果。

然而，正向順序實現(xiàn)得到的左關聯(lián)和的結果與并行算法的結果相同。然而，眾所周知，浮點加法不可結合，因此結果預計會因并行算法的結構而有所不同。

雖然這些問題也適用于中央處理器，但對于圖形處理器可能更為重要的原因在于它們以吞吐量為導向，并且緩存較小。因此，在GPU上使用雙精度算術代替幾乎所有的單精度算術會增加懲罰，特別是由于GPU上雙精度單元的數(shù)量通常較少。與GPU相關的浮點精度語言特定問題也存在。例如，在OpenCL中，有關是否準確表示非正規(guī)化數(shù)的問題是實現(xiàn)定義的，并且OpenCL提供了一種半精度數(shù)據(jù)類型，僅為相關運算符指定了最?。ǘ皇蔷_）精度要求。實現(xiàn)差異使編寫在多個GPU平臺上行為精度可接受的高性能浮點代碼變得困難。

（待補充中）