優化在異構性平台的OpenCL程式撰寫架構 -以光追蹤為例

全文

(1)國立臺灣師範大學資訊工程研究所碩士論文. 優化在異構性平台的 OpenCL 程式撰寫架構 -以光追蹤為例 Optimization of OpenCL Programming Structure On Heterogeneous Platforms – A Case Study of Ray Tracing. 研究生：指導教授：中華民國. 張潘立張鈞法 102. 撰博士年. 7月.

(2) 摘要現今使用 GPU(圖形運算單元)協助 CPU(中央運算單元)運算已經是常見的程式優化方式，GPU 藉由其核心數遠大於 CPU 並且可進行 SIMD(Single instruction， multiple data)平行運算之優勢可以大幅度的提升程式的效能。 OpenCL(Open Computing Language)則是一種異構性平台撰寫之框架，由於這是一個開放性的規範，因此各大廠牌之 GPU 與 CPU 皆有針對 OpenCL 框架下進行實作甚至已經有部分嵌入式系統的圖形處理器與中央處理器也逐漸開始支援，本篇研究即是以 OpenCL 平台下各個不同廠牌之 CPU&GPU 之運行為題，能夠找出優化各個硬體裝置於 OpenCL 平行化程式的一些方法與經驗。光追蹤是計算機圖學中一項非常泛用的技術，主要是以追蹤從眼睛開始灑出去無數的光線束藉以去嘗試碰撞空間中的立體物件，並經由一系列的物理模擬計算(如反射、折射、光影與其他……)，最後轉為每個點的顏色並產生呈現在眼睛上的影像。然而此種演算法的特色因為要灑出龐大的光線束量，並且幾乎都是獨立作業的方式分別運行計算，因此該演算法很符合使用 OpenCL 平行運算的特性，但是其中 OpenCL 還是有諸多使用上的條件與限制，這些會在本篇會有更詳細的說明。最後能藉由本篇論文能分享一些使用 OpenCL 平行化運算優化程式設計技巧以及設定相關參數的微調，以達到更好的效能。. 關鍵字：光追蹤(Ray tracing)、跨平台(Heterogeneous platforms)、開放計算語言 (OpenCL)、計算機圖學(Computer graphic)、優化(Optimization)、通用圖形處理器(GPGPU)、中央處理器(CPU)、平行計算(Parallel computing) II.

(3) 目錄附表目錄....................................................................................................................... V 附圖目錄......................................................................................................................VI 第一章緒論................................................................................................................ 1 1.1 研究背景 ......................................................................................................................... 1 1.2 研究目的 ......................................................................................................................... 2 1.3 論文架構 ......................................................................................................................... 2. 第二章文獻探討........................................................................................................ 3 2.1 光跡追蹤 Ray Tracing .................................................................................................. 3 2.2 基礎演算法分析 ............................................................................................................ 4. 第三章 OpenCL 優化架構 ......................................................................................... 7 3.1 開放式計算語言 OpenCL 介紹 .................................................................................... 7 3.2 OpenCL 使用條件與限制 ............................................................................................ 10 3.2.1 硬體限制與需求 .................................................................................................... 10 3.2.2 程式語言限制需求 ................................................................................................ 12 3.3 OpenCL 優化方針 ........................................................................................................ 13 3.3.1 通用性優化方針 ................................................................................................... 14 3.3.2. CPU 優化架構 .................................................................................................... 15. 3.3.3. GPU 優化架構 .................................................................................................... 16. 第四章相關技術實驗與結果.................................................................................. 17 4.1 優化實驗目的 .............................................................................................................. 17 4.2 優化實驗設計與數據 .................................................................................................. 17 4.2.1 向量化資料結構 .................................................................................................... 18 4.2.2 分流 ........................................................................................................................ 19 4.2.3 記憶體存取測試 .................................................................................................... 23 4.2.4 記憶體種類使用 .................................................................................................... 26 4.2.5 使用 host-memory .................................................................................................. 29 4.3 結果與分析 .................................................................................................................. 30. III.

(4) 第五章 Ray Tracer 優化設計與結果 ....................................................................... 31 5.1 演算法流程剖析 ........................................................................................................... 31 5.2 實做優化與實驗數據 ................................................................................................... 35 5.3 實驗結果與分析 .......................................................................................................... 40. 第六章結論與未來方向.......................................................................................... 42 6.1 結論 ............................................................................................................................... 42 6.2 未來研究方向 ............................................................................................................... 42. 參考文獻...................................................................................................................... 48. IV.

(5) 附表目錄表表表表表表. 1. OpenCL 各平台硬體限制資料舉例 ................................................. 12 2:向量化實驗數據 ................................................................................ 18 3.分流實驗數據結果............................................................................. 23 4.記憶體存取實驗數據結果 copy-host memory ................................ 25 5. 記憶體存取實驗數據結果 use-host memory ................................. 25 6.記憶體存取測試實驗數據................................................................. 28. 表表表表表. 7.host-device 傳輸時間表 ..................................................................... 29 8. CPU/GPU 運算能力理論值 GFLOPs 表 ......................................... 31 9. 2 bvh-nodes 10 triangles 原始版本 CPU 執行數據 ........................ 34 10. 106 bvh-nodes 399 triangles 原始版本 CPU 執行數據 ................. 34 11.優化變型設計執行結果 2nodes 10triangles.................................... 38. 表表表表. 12. 優化變型設計執行結果 106nodes 399triangles ........................... 38 13. Intersect 實驗數據 ........................................................................... 39 14. Occluded 實驗數據 ......................................................................... 39 15. Intersect 比較 CL 版本與 C 之速度 ............................................... 39. 表 16. Occluded 比較 CL 版本與 C 之速度 ............................................. 40. V.

(6) 附圖目錄圖 1.光追蹤演算法示意圖........................................................................... 4 圖 2. OpenCL Device Architecture Diagram [opencl-1-2-quick-reference-card] ......................................................... 9 圖 3.Intel 範例基礎程式 SimpleOptimizations.cl .................................... 18 圖 4.單分支-平均分流 ............................................................................... 20 圖 5.單分支-導向單一分流 ....................................................................... 21 圖圖圖圖圖. 6.雙分支-平均分流 ............................................................................... 22 7.大幅度 global 記憶體存取................................................................. 24 8.最低幅度 global 記憶體存取............................................................. 25 9.Global Memory ................................................................................... 27 10.Local Memory................................................................................... 27. 圖圖圖圖. 11.Private Memory................................................................................. 28 12. 2 bvh-nodes 10 triangles 場景圖 .................................................... 33 13. 106 bvh-nodes 399 triangles 場景圖 .............................................. 33 14. 三角形碰撞 Intersect 程式架構圖 ............................................... 35. 圖 15. 光影偵測 Occluded 程式架構圖 ................................................. 36. VI.

(7) 第一章緒論 1.1 研究背景：自從 NVidia 推出 CUDA 架構以來，計算機除了能使用傳統的 CPU 進行運算之外，還可以使用圖形處理器上的多核心進行運算，而 GPU 從原先僅止於圖形計算的功能轉型成為可以做為一般性運算的通用圖形處理器(GPGPU)。 OpenGL 很早以前非 GPGPU 的時候就已經在圖形處理器進行許多優化並能夠達到及時(real-time)的效果，因此大多數的強調及時的像是電玩遊戲都已經使用 OpenGL 技術為主，但是如果對於較為複雜的場景或是更為複雜的光學效果，光是只有 OpenGL 不一定是足夠的，也因此發展出像是光追蹤(Ray-Tracing)以及廣域照明(Global-Illumination)之技術，而這類的方法可以畫出更有質感的場景，但是該方法的最大缺點就是運算時間要很長，原因是這些方法需要打出很多個 samples 的光線去尋找物體的存在，而直接光源與間接光源可能也要取非常多的 samples 去模擬出該光源對於物體的效果，這些都是需要龐大的運算時間，所以傳統上 Ray-Tracing 光是使用 CPU 是較難達成 real-time 的速度。現今 GPGPU 的架構設計能夠讓 GPU 的平行運算能力充分應用，對於 Ray-Tracing 演算法的發射多條光線可以做為分散的單位，以此來做為平行處理計算的方針，目前已經有國內外學者研究在提到關於在 CUDA 上的 Ray-Tracing 的設計方針，都可以大幅加速其本身龐大的運算時間。而後 APPLE 推出 OpenCL 語言之後，各大圖形處理器製造商紛紛投入研發 GPGPU 的支援，也因此現今 OpenCL 幾乎可以在各個廠牌的 GPU 上運行，甚至. 1.

(8) 連 CPU 也可以支援 OpenCL 的平行運算處理，所以現今演算法也可以透過 OpenCL 程式語言可使用多種核心的功能，進而達到比較 CPU 與 GPU 的對於單一程式的運行效能以及分析各種演算法是否較於適合 CPU 或 GPU 去執行。. 1.2 研究目的：因為 OpenCL 跨平台的特性，我們可以較為方便的方法得 CPU 及 GPU 做平行運算的處理，也就是同一份程式碼可以只經過微調就可以讓多種平台使用，因此使用 OpenCL 來做為比較 Ray-Tracing 在 CPU/GPU 效能的基準並且能夠分析哪些適合 CPU 或 GPU 執行會比較有效的運用計算機的資源。透過各個平台的比較，能夠讓計算機的運算能力得到完整的發揮，並且探討程式編寫方式對於平行處理運算的影響以及各個廠牌的 CPU/GPU 在設計上有哪些指令架構的部份可以做為程式設計者應用的資訊，希望於計算機圖學以及平行程式設計領域能夠有些許貢獻。. 1.3 論文架構本論文總共六章，第一章為緒論，敘述本研究的背景與動機；第二章為文獻資料探討，介紹光追蹤的演算法與其適合平行化的分析；第三章對 OpenCL 做一個概略的介紹，並分析其最佳化的方式以及做法；第四章則實做關於 OpenCL 優化的實驗，藉以找出何種平台適合哪種寫法；第五章對我們實作出來的光追蹤演算法進行平行化撰寫實驗，並提供實驗結果與分析；第六章提出結論與未來方向。. 2.

(9) 第二章文獻探討 2.1 光跡追蹤 Ray Tracing 光線追蹤(Ray tracing)，是一個描述光線與物理表面所發生的互動行為而取得其路徑的方式。在早些年代光線追蹤並不是電腦圖學領域的專有名詞，像是一些光學鏡片產生的效果如相機，望遠鏡，眼鏡等都是運用其物理上的原理所構成。但是在電腦圖學上來說，這是一種成像構圖的演算法，藉由追蹤眼睛發射出來的光線，去掃描範圍內所放置的物體並顯示出來的數學模型，很特別的是光並不是從光源去追蹤，是因為中就成像的圖像就是眼睛的位置，為了確保每條光都是有意義的不會白費運算時間，使用反向的追蹤反而是比較有效率的做法，而光追蹤可以做出許多光學效果，許多高品質的效果都是 OpenGL 難以辦到的，因此要追求高畫質的三維影像常常使用這種方法來執行。現實生活中的光線，以物理學的觀點，是追蹤光在不同介質間傳播的情形，由於不同物體本身的屬性，使其可以產生各種的物理現象，例如說光被吸收，反射出去，或是改變行徑方向等。。。，當然在計算機圖學上不太可能用太複雜的物理公式去模擬光的行為，因此會簡化光線通過介質產生的反射折射陰影之類來解決許多較為複雜的狀況。因此在計算機圖學的光線追蹤運算方法，主要是把光線理想化成一個單一的向量線，並且透過光線追蹤期在各種不同物體上的反映情況做處理。首先打出一條光線，並且開始測試各個可能被光線射中的物體，取得最近距離的集中點位置以及其物體屬性，依據不同參數的設定決定是否會產生反射折射的現象而產生出. 3.

(10) 新的光線，取得其傳播路徑與其傳播圖中所經過的物體的顏色資訊，最後在計算出最終的顏色投射在視網膜上。. 圖 1.光追蹤演算法示意圖. 2.2 基礎演算法分析光線追蹤並不是一個非常新穎的技術，在此參考了被廣泛使用的光追蹤演算法的寫法如下,由此演算法為基礎來分析光追蹤演算法是否有機會使用平行化增強其效能.. 4.

(11) 經典光線追蹤演算法流程： For(對圖像中的每一個像素) { 產生一條的從眼睛穿越該像素的光線。初始化最近距離 T 為無限大，最近碰撞物體為空值。 For(對場景中的每一個物體) { 如果光線與物體相交 { 如果相交處的距離 t 小於最近距離 T { 把最近距離 T 為交點的 t 值把最近碰撞物體最近物體為新碰撞到的物體 } } } 如果最近物體為空值 { 用背景色設定顏色 } 否則 { 往每個光源檢查是否是在陰影中如果表面是反射面，生成反射光，遞迴光追蹤() 如果表面透明，生成折射光，遞迴光追蹤() 使用最近物體和最近距離 T 計算著色把最終著色顏色填回該像素 } } [1].經典光追蹤演算法. 光線追蹤在 OpenCL 實作的優點分析:. 在 OpenCL 架構中提供了不論是 CPU 或是 GPU 都提供了良好的平行運算環境，也就是說可以把光線追蹤的主要步驟對每一個圖像素產生的 Ray 做為平行化的 work-item，有多少條光線就分出多少個支流，每個光線由於獨立不受其他打出的光線影響，可以做為單獨運算，不會有相依性過高的問題，因此是光線追蹤很. 5.

(12) 適合做為平行處理。. 光線追蹤在 OpenCL 實作的缺點分析:. 然而 OpenCL 的 kernel 程式是無法執行遞迴處理，而且程式限制所有的分流都需要一開始決定好，不可以在 kernel 端中再 fork 出新的分流核心，因此在設計上不能使用以上的架構，如果要製作反射折射的設計就要另外做處理，而且雖然 OpenCL 提供了平行化的良好設計但是記憶體的流程反而是最重要的課題，由於記憶體與 host 端並不是完全可以互相使用，還是要經過額外傳輸或是宣告的處理才能使用，並且在大小限制上也比 C 程式還要多，但是還是有其使用上的限制會使其無法達到其最大的效能，是要在撰寫程式時需要有許多留意的地方。. 6.

(13) 第三章. OpenCL 優化架構. 3.1 開放式計算語言 OpenCL 介紹 NVidia 公司在於 GPGPU 運算很早就提出了 CUDA 架構，整合了 C 語言共用的程式撰寫架構，目的在於使用該廠牌的顯示卡提供平行運算之功能，並且這個架構可以轉換提供 OpenCL 使用其核心，並針對此架構下有許多更先進的設計，然而由於這算是完全屬於 NVidia 的架構，其他廠牌並沒辦法使用，為了能夠顧各個平台的通用性，我們選擇了 OpenCL 做為實驗的平行運算架構。 OpenCL(Open Computing Language)，是一個跨平台的撰寫程式的框架，其標準上允許多平台的支援並透過期內建的程式語言來撰寫裝置例如 GPU 與 CPU。OpenCL 本身擁有自己的程式語言，其主要是在選擇 CL 裝置下執行名為 kernel 端的程式碼，這個程式語言本身的是以 C99 為基底所構成，邏輯與設計規則幾乎跟 C 語言差沒有太大差別，但還是有一些規定與 C 不同，除此之外還有提供呼叫與控制的 API 給原本的程式語言使用，其最大的功能就是能夠做 SIMD(Single Instruction Multiple Data)的平行運算。並藉由一個低階(low-level)、高效能(high-performance)與可攜式抽象化(portable abstraction)提供嵌入式系統與消費性軟體的高效能解決方案。[2] 由於 OpenCL 是可以跨平台的緣故，所以目前已有許多廠牌的 CPU 及 GPU 支援 OpenCL 的架構設計，然而 CPU 與 GPU 在原先的設計上就已經 7.

(14) 有不小的差異，光是在計算機上的位置就會影響到資料傳輸的方式，CPU 可以直接運用主記憶體，在運算上直接以主記憶體為運算目標，經過優化可以不必浪費資料複製的時間，然而 GPU 可能是因為是外接的設計所以通常都有自己的獨立記憶體，並且在執行運算前後要把資料從主記憶體到顯示卡的記憶體之間做轉換，然而通常外接顯示卡可執行的運算單元是高於 CPU 的，也因此擁有計算量大的優勢，而又有另一種架構則是 CPU 與 GPU 做在同一塊晶片上，由於不需要獨立的記憶體，則是選擇直接使用主記憶體為運算的位置，不但有多運算單元的優點並且也有節省記憶體傳輸問題的功能，但由於坐在同一塊晶片的原故可能沒辦法把放太多核心，因此距離外接顯示卡的運算能力還是有一些差距的。有鑑於此，因為有各種不同的硬體架構，程式上如果想要發揮該裝置的最佳效能，必須有許多需要去注意的地方，同樣的一段演算法可能經過些調整可能比較適合某種架構的硬體設計這也是目前使用 OpenCL 一項要克服的難題。. 8.

(15) 圖 2. OpenCL Device Architecture Diagram [opencl-1-2-quick-reference-card]. OpenCL 的記憶體分配架構圖如上圖所示，OpenCL 在裝置執行核心 (Kernel)程式時，分成 Global， Constant， Local， Private 四種記憶體，這類的分類在於 CPU 上使用是沒有特別的意義的因為 CPU 使用的記憶體都是主記憶體，無法完全透過這些記憶體種類的差異達到優化的效果，而這些架構設計主要是針對外接顯示卡的設計，顯示卡的設計上記憶體的存取速度是 Private Memory 最快但是空間最少，並且只有自己單獨的執行緒續看的到，而 Local Memory 則是速度很快但是比 Private Memory 慢，空間稍大一點並且是同一個工作群看的到，並且處於分享的狀態，但是工作群外就看不到，而 Constant Memory 的是宣告為不會被修改類似 C 語言的 const 標籤，他會轉到 Constant Memory 中可以大幅提高記憶體的存取速度達到 Local Memory 速度的效果，而 Global Memory 是所有 memory 中最慢的， 9.

(16) 但是他的以讓所有的工作群都共用的資料。在 SIMD 的平行運算架構中，如果能夠將指令切割的非常乾淨，就會真的讓同一段程式碼同時運行完成，並且達到最大的運算效能，但是在 OpenCL 中，首先要考慮的是把資料複製到裝置上的記憶體，由其是外接顯示卡傳輸資料的時間可能會影響到整體程序的效能，因此記憶體傳輸的時間是首要考慮的課題，當資料量過大的時候也要注意是否能夠在裝置環境限制下能夠正確分配資料或是壓縮資料容量也是很重要的。. 3.2 OpenCL 使用條件與限制雖然 OpenCL 對於平行運算擁有很好的設計，使用上幾乎與 C 語言沒什麼不同，但是還是有很多條件與限制要注意的地方避免產生不良的設計導致整體效能低落，在此從幾個層面來探討，主要分成硬體限制與程式語言限制兩部分探討。. 3.2.1 硬體限制與需求 OpenCL 雖然是一個支援多核心的平行運算，但並不是說核心越多就可以達到使用單核心計算相同倍數的效果，其同時間可以處理的工作單元 (work-item)是該裝置本身提供的執行緒上限而定，例如說 Intel HD2500 核心 (core)數量與執行緒是相同的[3.]，但是在 NVidia 的機器中，雖然越高階的顯示卡會擁有更多的核心但是只有最多 32 個執行緒可以做分配[4.]。因為此原因不能直接單純的計算核心數量視為他們速度倍率的依據。. 10.

(17) 接下來是記憶體的限制，目前的 OpenCL 架構下的硬體設計都有限制各種記憶體的限制用量，在此的使用量相較於單純的 C 程式來說是少非常多的，現在關於運算裝置主要分兩種類，一種是獨立顯示卡通常會擁有自己的獨立記憶體，而另一種則是 CPU 與鑲嵌在 CPU 上的 GPU 的內嵌顯示卡，他們則是共享系統的主記憶體。在目前的 OpenCL 的架構中，要執行 Kernel 程式之前，需要先把要被計算的資料先傳輸到裝置上，經過 Kernel 計算後再將資料傳輸回來，因此這傳輸的過程當中會有額外的成本，而使用 CPU 及內嵌顯示卡的 GPU 時，可以允許使用資料原本在程式內的記憶體位置來免去傳輸的時間，但是這部分也是有限制大小的，如果遇到資料非常龐大的情形，無論是哪種平台還是會容易造成記憶體創建失敗並無法執行 kernel 程式。在 OpenCL 的記憶體架構下分成 global, local, private 與 constant 記憶體中使用特定的記憶體形式會有達到存取加速的效果，但是這些記憶體是有限制的，而且容量非常小，並不能放置太過龐大的資料，並且這些空間還要被每個工作群組(work-group)或是工作單元(work-item)分配，每個執行緒所能分配到的限制又更小了。. 11.

(18) 表 1. OpenCL 各平台硬體限制資料舉例另外就是 Command Queue 可以去分配執行不同的任務，但是如果一次分派太多任務的話會導致競爭的問題，避免同時讓忙碌中的裝置加入過多任務也是能夠讓其發揮最佳效能。由於其採用的是 SIMD 架構，能盡量做到的就是資料不要互相依賴為最好的方法，雖然有提供 global & local 等記憶體溝通方式，但是不同的執行緒不斷存取同一塊記憶體必然會發生存取競爭的問題，這部分在算是在平行運算架構下都會發生的問題。最後是流程方面，OpenCL 在 kernel 程式一次把所有事情做完，中途是不允許隨時跳回 C 程式端繼續執行，除非工作已經完全完成。. 3.2.2 程式語言限制需求. 12.

(19) OpenCL 的程式語言架構並沒有完全整合在 C 語言之中，而是需要把 CL kernel 端程式碼存在另外的檔案中在 runtime 時才會呼叫 OpenCL Compiler 編譯並執行，因此其資料結構是無法與 C 共用，需要另外在 kernel 端程式碼中也宣告出來，而當資料傳輸的時候他的作法是複製整塊記憶體的資料，因此只要在 kernel 端有宣告的容器即可裝載進去，但是 C 語言中的指標在傳輸後到 kernel 中是無效的，這部分也是需要注意的功能。目前 OpenCL 不支援動態記憶體配置，也就是說任何參數都需要在程式碼中宣告或是從外面傳進去後，內部不允許使用 memory allocate，因此在轉換程式碼實需要先把所有資料視為已知才能正確計算。不可以宣告非固定值的全域變數，只允許使用 constant 來宣告並且該值會屬於到 constant memory，存取速度幾乎接近於 private memory。其中增加執行緒等功能也不支援在 kernel 程式中，需要在 host 端(C 程式)透過 OpenCL 的 API 做控制才可以。在極度物件導向化的程式，其實並不太適合移植到 OpenCL 做平行化，因為其原本是以 C99 為藍本作為依據，並不支援複雜的資料結構與多形，若要改寫類似程式到 CL 端會導致非常多問題，如果不改寫原先的程式的話事實上會導致許多資料傳輸的問題發生。另外在 kernel 執行上有許多選項，可以讓 host 端不等待資料傳輸或是執行程式結束就可以繼續，但是要注意的是同步往往是一些重要的問題，特別是寫習慣單執行緒程式後需要注意的地方。. 3.3 OpenCL 優化方針. 13.

(20) 在 CUDA 中的優化方針主要因應記憶體的部分，OpenCL 在許多設計架構下都與 CUDA 類似，但近期的 CUDA 已經在硬體上的設計試圖去改善記憶體使用的問題，這部分 OpenCL 不見得能夠完全使用 CUDA 最新的架構，為了配合各種平台的限制通常設計上會採取有多種選擇的方法提供各廠商製造運算單元使用。以下這些都是主要介紹 OpenCL 以及如何優化的一些方法，當然其中最重要的還是演算法本身的設計，雖然不是每種方法都是完全適用，但是只要可以找出幾種正確的方法還是有許多改善程式效能的空間。. 3.3.1 通用性優化方針不論是 CPU 或是 CPU 都可以有一些通用的方式去加快其執行的效能，其中主要有幾個方法，首先就是要減少記憶體的複製，由於 OpenCL 核心程式需要把資訊傳到核心執行，並且之後還要把資料回傳回來，這部分常常佔有不少時間，甚至可能比 kernel 運算還要久，所以最小化記憶體的需求是在設計程式時需要仔細考慮的部分。其次避免使用太多的同步功能，在 OpenCL 中，可以針對所有的核心程式進行同步，同步在此指的是要等待其他的 kernel 執行到某一行之後才會做之後的程式，在資料有相依性的情況時，要使用這個指令避免程序先後的影響導致資料處理錯誤，然而等待等於是讓已經完成的核心空閒不處理資料，如果使用太頻繁會導致平行運算速度降低[5.]。另外一些技巧則是 OpenCL 內有許多已經進行優化過的 build-in Functions，可處理一些常見的運算。還有就是減少分支的數量(if 的數量)， 14.

(21) 因為在類似的指令常常會統整起來一併處理，如果是一個沒有分支的程式就會依序逐行執行會快速很多。 Use Host Pointer 在 CPU 或是 On-Chip 的 GPU 因為使用與主記憶體相同的 memory，可以使用該方式降低資料傳輸的成本，而外接顯示卡(NVidia) 則可以使用像是 pinned-memory 來達到相同的效果[8.]。. 3.3.2. CPU 優化架構. Intel 在 CPU 方面也是可以支援 OpenCL，然而因為使用主記憶體所以 GPU 方面的優化方法是沒有效果的，轉而要從其設計的指令集提供的功能，其中 Intel® Streaming SIMD Extensions (Intel® SSE) 和 Intel® Advanced Vector Extensions (Intel® AVX) instructions 的架構能夠將類似型態的資料自動處理群組排程並由單一指令去處理多筆資料[6.]，這樣儘管核心處理器數量沒有 GPU 多，但是單一 CPU 可以做的事情也就更多了。而在於程式碼撰寫方面其實該設計已經有 Implicit Vectorization，就是在於不加任何標籤下自動完成指令的群組處理，然而有時候因為演算法的原故該方法不見得會是最好的處理辦法，因此也可以在程式碼內調整標籤減少隱含的 Vectorization[5.]，這部分就是可能需要去研究調整的部分。而 OpenCL 有設計一種 Vector Data Types，像是 int2，int3，int4，float2， float4 等等。。。。。在於 GPU 的設計架構下只是為了方便一些多維元素的運算的設計並沒有特別針對其優化，然而在 CPU 中則會套用上述功能讓單一指令執行整個 Vector Data 的運算。因此如果要使用 CPU 版本的設計就要多使用這些架構來加速程式的運行。 15.

(22) 3.3.3. GPU 優化架構. 解決完 host 端到 kernel 的傳輸問題後，之後要考慮的就是每個執行緒的存取資料是不是獨立的，當所有執行緒都要存取同一筆資料時，就會有等待的效應，雖然現在 GPU 設計(NVidia)有針對這種等待處理做自動跳至另一筆資料避免存取相同的等待問題[7.]，但是這部分還是能夠避免少去存取相同的位置為最好的方式，因此可以使用的就是 Local Memory 與 Constant Memory 的方式，把一些資料複製至各個工作群的 Local Memory 就可以達到減少相同記憶體位置的存取次數，而把一些不會被改變的資料使用 Constant Memory 也可以達到讀取速度的有效提升，這些有效的優化方法主要適用於 GPU 為主，但是在以主記憶體執行的 CPU 則是效果沒有這麼顯著。. 16.

(23) 第四章相關技術實驗與結果. 4.1 優化實驗目的為了測是如何讓 OpenCL 程式最佳化，參考了各個廠牌的程式撰寫說明書與優化建議，撰寫一些程式來做為實驗，本章節盡量把問題單純化，為了避免其他額外的影響導致實驗失準，盡量採取最小的改動來估計其運算的誤差，得到的結果可作為程式撰寫參考，並評估哪些優化方法比較適合光追蹤演算法的撰寫。. 4.2 優化實驗設計與數據本實驗主要是針對 OpenCL Kernel 程式的撰寫使其產生不同的影響，採用的測試機器與環境設備為以下項目：. CPU： GPU：. Intel Nvidia Intel. I5-3550 Tesla C1060 HD 2500. (獨立顯示卡－使用獨立記憶體) (內嵌顯示卡－使用主記憶體). 作業系統：Microsoft Windows 7 編譯環境：Microsoft Visual Studio 2010 Professional OpenCL(v1.2)：Intel OpenCL SDK 2013 Nvidia CUDA SDK 4.2 以下實驗是根據 Intel OpenCL Example 的程式碼做為改寫藍本，並以此當作基礎範例做為實驗對照組，藉此觀察不同的程式撰寫方式對於執行效能的差異。. 實驗分成三個項目：1.向量化資料結構 2.分流 3.記憶體存取測試 4.記憶體種類使用 5.使用 host-memory. 17.

(24) 實驗測試資料都是的使用皆使用 64*1024*1024 大小的隨機生成浮點數陣列為輸入檔案數量，而 work-item size 則限制為 16。. 4.2.1 向量化資料結構在 OpenCL 中有設計向量型態可以一次存放一整組資料，這種結構是為了方便圖形處理以及多維運算會有很方便的功能，而根據 Intel CPU 的設計中，對於 OpenCL 的向量變數設計，可以使其一次使用一個指令處理四筆資料，在此的實驗則是使用正常的資料形態處理與使用 float4 資料型態的區別，程式碼如下所示： __kernel. void. SimpleKernel( const __global float *input, __global float *output) { size_t index = get_global_id(0); output[index] = rsqrt(fabs(input[index])); }. __kernel. void. SimpleKernel4( const __global float4 *input, __global float4 *output) { size_t index = get_global_id(0); output[index] = rsqrt(fabs(input[index])); }. 圖 3.Intel 範例基礎程式 SimpleOptimizations.cl. 這兩個程式的差異在於上方的程式會一個一個執行，而下方的是會存放在 float4 中群組執行。. 單位(ms). I5-3550. HD 2500. NV Tesla C1060. SimpleKernel. 48.46. 73.41. 24.69. SimpleKernel4. 44.58. 59.91. 6.7. 表 2:向量化實驗數據 18.

(25) 從表中可知道使用向量結構(int4, float4 , ...)對於整體的程序不論在任何平台都會有加速的效果，尤其是在 GPU 更有顯著的效果，這部分是因為 GPU 原先的功能是處理圖形繪圖資料，而其本身就是批次處理的設計，因此 GPU 處理向量化的能力是比 CPU 強很多的。而向量處理對於程式的加速都有很顯著的效果，但是其中的條件限制為最好是一次對整個向量型態做相同的計算，如果常常要提取向量型態的某個元素單獨計算，則使用這種方式效果有限，這是在設計資料結構時需要小心的事項。. 4.2.2 分流由於 SIMD 架構下會讓相似的指令一併執行的關係，使類似的指令更有效率，所以根據編程建議是越少分流越好，但寫程式不可能沒有分流的情況，因此設計此實驗來看分流對於各個裝置的影響，設計原則上是要讓所有的分流都執行一樣多的指令但是為了避免其 compiler 自動簡化指令產生優化效能，把許多指令執行的順序做些許修改，如同以下程式碼所示：. 19.

(26) __kernel void ExpKernel_0( const __global float *input, __global float *output) { size_t index = get_global_id(0); float in=input[index]; float out=0.0f; float data[30]={0}; if(index%2>0) { data[0]+=in+XXX;. //XXX為不規則整數數字. data[1]+=in+XXX;. //執行順序也會打亂不一定從[0]開始. data[2]+=in+XXX; ... data[29]+=in+XXX; out=data[0]+data[1]+data[2]+data[3]+data[4]+data[5]+... +data[29]; } else { data[0]+=in+XXX; //XXX為不規則整數數字 data[1]+=in+XXX; //執行順序也會打亂不一定從[0]開始 data[2]+=in+XXX; ... data[29]+=in+XXX; out=data[0]+data[1]+data[2]+data[3]+data[4]+data[5]+... +data[29]; } output[index]=out; }. 圖 4.單分支-平均分流. 這邊的設計主要讓程式執行會有一半走到上半部，一半走到下半部，平均處理的情況下執行 kernel 程式。. 20.

(27) __kernel void ExpKernel_1( const __global float *input, __global float *output) { size_t index = get_global_id(0); float in=input[index]; float out=0.0f; float data[30]={0}; if(index%2>=0) { data[0]+=in+XXX;. //XXX為不規則整數數字. data[1]+=in+XXX;. //執行順序也會打亂不一定從[0]開始. data[2]+=in+XXX; ... data[29]+=in+XXX; out=data[0]+data[1]+data[2]+data[3]+data[4]+data[5]+... +data[29]; } else { data[0]+=in+XXX; //XXX為不規則整數數字 data[1]+=in+XXX; //執行順序也會打亂不一定從[0]開始 data[2]+=in+XXX; ... data[29]+=in+XXX; out=data[0]+data[1]+data[2]+data[3]+data[4]+data[5]+... +data[29]; } output[index]=out; }. 圖 5.單分支-導向單一分流. 這部分雖然有分流，但是全部都會走向上方的分流，這樣一來就可以視為相同分流的情形。. 21.

(28) __kernel void ExpKernel_2( const __global float *input, __global float *output) {. size_t. index = get_global_id(0);. float in=input[index];float out=0.0f; float data[30]={0}; if(index%2) {. if(index%4) {. data[0]+=in+XXX;. //XXX為不規則整數數字. data[1]+=in+XXX;. //執行順序也會打亂不一定從[0]開始. ... data[29]+=in+XXX;} else {. data[0]+=in+XXX;. //XXX為不規則整數數字. data[1]+=in+XXX;. //執行順序也會打亂不一定從[0]開始. .. data[29]+=in+XXX;} out=data[0]+data[1]+data[2]+data[3]+data[4]+data[5]+... +data[29]; } else {. if((index+1)%4) {. data[0]+=in+XXX;. //XXX為不規則整數數字. data[1]+=in+XXX;. //執行順序也會打亂不一定從[0]開始. ... data[29]+=in+XXX;} else {. data[0]+=in+XXX;. //XXX為不規則整數數字. data[1]+=in+XXX;. //執行順序也會打亂不一定從[0]開始. ... data[29]+=in+XXX;} out=data[0]+data[1]+data[2]+data[3]+data[4]+data[5]+... +data[29]; } output[index]=out; }. 圖 6.雙分支-平均分流. 這個部分則是當分流呈現巢狀結構時會不會有更進一步的影響，裡面的執行會平均分配給四條巢狀的分支。. 22.

(29) 單位(ms) 分支-平均分流分支-單一分流雙重分支-平均. I5-3550 122.499 67.965 136.659. HD 2500 223.122 143.474 224.42. NV Tesla C1060 37.519 22.165 37.532. 表 3.分流實驗數據結果. 由這些實驗結果表示在單一導向分流時都會被各個平台視為沒有做分支處理，在程式本身就會各個執行單元做相同的指令，因此在執行時間上非常短，各個裝置在增加分支後皆會有產生降速的現象，而且影響效能非常嚴重，儘管他們只是在做類似的運算，還是會拖垮整體的效能由其從 Intel 裝置上(前兩項)更為嚴重，而雙重分支與一個分支來比較，則發現 GPU 影響是來的比 CPU 小，這部分其實有些不符合一般我們對 CPU 的概念，通常認為 CPU 在處理分支問題時應該會有較好的效能，但是我們經過多方測試包括避免 compiler 做了比預期還更多的優化，因此把可能的變因例如迴圈與暫存器處理的優化排除之後發現還是有不小的影響，這部分也許是因為要符合 SIMD 架構使其指令傾向相同會有較佳的設計有關，其實相關文獻也有提到 OpenCL 架構下 CPU 對於分支問題也是會有一定程度的影響，因此盡可能的在 OpenCL 程式的設計中，越少出現分流一定是比較好的選擇。. 4.2.3 記憶體存取測試在 OpenCL 架構中，是無法直接共用系統端的主記憶體，雖然有辦法提供一個使用 host-memory 的功能，但是其編排的虛擬地址是無法與 host-program 混唯一用的，因此每次的計算都要先把資料導向到 kernel 端才可執行，而這些傳進去的參數是否也跟一般的 C 程式一樣使用，這就是本實驗的重點，實驗的設計主要針對傳進去的 global-memory input 做運算最後. 23.

(30) 傳回 output 的記憶體中，兩種不同的撰寫方式的方式如下：. __kernel void ExpKernel_3( const __global float *input, __global float *output) { size_t index = get_global_id(0); output[index]=0.0f; output[index] += rsqrt(fabs(input[index])); output[index]+=input[index]+21; output[index]+=input[index]+22; output[index]+=input[index]+23; ...// (+=函式總共50次數字皆為不規則) output[index]+=input[index]+XXX; output[index]+=input[index]+XXX;. }. 圖 7.大幅度 global 記憶體存取. 這邊的設計在於所有的 input output 都採取直接存取的方式，任何計算以及複製都直接從傳進去的 global memory 來做。. 24.

(31) __kernel void ExpKernel_4( const __global float *input, __global float *output) { size_t index = get_global_id(0); float in=input[index]; float out=0.0f; out+= rsqrt(fabs(in)); out+=in+21; out+=in+22; out+=in+23; out+=in+34; ... (+=函式總共50次數字皆為不規則) out+=in+22; out+=in+23;. output[index]=out; }. 圖 8.最低幅度 global 記憶體存取. 這裡的設計則是在 kernel 程式內增加變數取代先複製一次原本要被計算的資料，途中的計算完全不會動到 input 與 output，直到計算完成後才把回傳值設回去。. 單位(ms) In out 高存取 In out 低存取. I5-3550 480.754 63.983. HD 2500 725.742 161.241. NV Tesla C1060 349.846 16.349. 表 4.記憶體存取實驗數據結果 copy-host memory. 單位(ms) In out 高存取 In out 低存取. I5-3550 223.355 63.315. HD 2500 723.626 156.424. NV Tesla C1060 349.806 16.346. 表 5. 記憶體存取實驗數據結果 use-host memory. 從實驗結果因為與記憶體有關特別使用了 OpenCL 的使用 host 端記憶體的功能，不過兩者的結果差異在於 CPU 在使用 host 端記憶體會比較快，但是在低存取 global memory 的情況卻沒有嚴重的差異，顯然記憶體擺放位. 25.

(32) 置盡量可以的話採用 host-memory 對於記憶體的存取會有比較快速的存取速度。在這邊有些比較特別的地方為什麼 CPU 都是用主記憶體還是有效能上的差異，經過幾個測試實驗後發現只要是從外部傳輸進來的 data 存取上都會比在 kernel 端宣告的記憶體慢上許多，原因是從外部傳輸進來的記憶體會放在動態區域存取時需要耗費更多時間，而在 kernel 端的參數宣告不允許動態因此所有變數基本上都是靜態且確認範圍的因此內部存取會快上許多。從中得知一項重要的訊息，越少存取傳進去的 global memory 會有更好的效能，由此可以接延伸到下一個實驗。. 4.2.4 記憶體種類使用從上一個實驗所知，記憶體並不像一般的 C 程式都是使用相同的主記憶體，在 OpenCL 中把記憶體分成幾種不同的型態，分別有各自的功能需求，越快的記憶體通常限制也就越大，這邊的實驗主要針對 global memory, local memory, private memory 的使用方式設計實驗，主要針對三種記憶體的使用方式改寫同樣的程式，並強迫程式進行記憶體存取競爭，也就是說每個 kernel 都要搶第 0~15 編號的 input 資料，其改寫發如下：. 26.

(33) __kernel void ExpKernel_5( const __global float *input, __global float *output) { size_t index = get_global_id(0); size_t index_y = get_local_id(0); float out=0.0f; for(int i=0;i<16;i++) { out += rsqrt(fabs(input[i]))+fabs(input[index]); } output[index]=out; }. 圖 9.Global Memory. 這個程式就是直接使用 global memory 來做為存取的標準。 __kernel void ExpKernel_6( const __global float *input, __global float *output) { size_t index = get_global_id(0); size_t index_y = get_local_id(0);. float out=0.0f; __local float inin[16]; inin[index_y]=input[index_y];. barrier(CLK_LOCAL_MEM_FENCE);. for(int i=0;i<16;i++) { out += rsqrt(fabs(inin[i]))+fabs(input[index]); } output[index]=out; }. 圖 10.Local Memory. 這邊的設計則是使用每個 work-group 個複製一份資料到 local memory ，最後在執行時會減少存取相同記憶體的次數。 27.

(34) __kernel void ExpKernel_7( const __global float *input, __global float *output) { size_t index = get_global_id(0); size_t index_y = get_local_id(0); float out=0.0f; float inin[16]; for(int i=0;i<16;i++) { inin[i]=input[i]; } for(int i=0;i<16;i++) { out += rsqrt(fabs(inin[i]))+fabs(input[index]); } output[index]=out; }. 圖 11.Private Memory. 這段程是採取的方法是直接每個 kernel 把該使用的記憶體複製並執行，但這不是很好的方法，因為在複製過程中就無形增加了存取競爭，若是將該資料直接設為一個固定參數及可解決，但是這就跟記憶體使用無關了，因此這段程式僅供參考。. 單位(ms) Global memory. I5-3550 187.42. HD 2500 962.33. NV Tesla C1060 243.82. Local memory. 192.37. 354.12. 37.99. Private memory but copy. 187.10. 962.33. 245.63. 表 6.記憶體存取測試實驗數據. 從實驗數據顯示這些針對記憶體種類做的調整對於 CPU 是沒有任何效果的，主要卻對於 GPU 就有顯著的幫助，在適當的分配下，可以讓每個 kernel. 28.

(35) 讀取一部分最後整個工作群組共用這些記憶體，結果使用 local 記憶體就有更快的效果。而使用 private-memory 則因為一開始的複製也會造成存取競爭，結果就沒有預期的有效果，不過從上一個實驗 4.2.3 中卻是有增快的，因此這算是撰寫的問題，最好的建議是讓比較多次被存取的小比資料複製到 private-memory，才會有加速的效果。. 4.2.5 使用 host-memory 這部分並不採取任何演算法而是單就測試資料的傳輸來做測試，在 OpenCL API 中有提供使用 host-memory 的功能在 CPU 與內嵌 GPU 上他們原先都是使用主記憶體，因此在硬體的設計上只要把一塊資料的記憶體框住並告知 kernel 程式即可直接達到，而在外接顯示卡則也使用類似的概念把記憶體釘住，也是直接使用主記憶體的方式，以下是針對各個裝置的資料傳輸測試，讀的資料為 64*1024*1024 的浮點數。. 單位(ms) Copy Host memory Use Host Memory. I5-3550 41.8 0.18. HD 2500 56.71 0.2. NV Tesla C1060 110.18 48.5. 表 7.host-device 傳輸時間表. 從表中的情況，只要是原先使用主記憶體的裝置設定為 use host memory 之後，其傳輸的時間原本是要複製一份現在則是幾乎沒有任何成本，而外接顯示卡，使用後雖然也變快了，但是其中就要傳輸還是要透過較慢的通道，因此無法跟內嵌裝置效果來的好。. 29.

(36) 4.3 結果與分析從以上這些實驗結果顯示，不見得所有方法都適合各種裝置的使用，但是也有許多共通的加速方針，例如說使用主記憶體，減少傳輸參數的存取次數，減少分流，使用向量化參數。而主要需要注意的是只要是 GPU 都可以對記憶體種類需要更有效的應用，三種記憶體架構雖然有其使用的限制，大小空間也是有限的，但是只要盡可能的減少競爭的問題，就可以發揮最好的效能。目前為止對於 OpenCL 提供的 profiling tool 功能相當有限，並不像是 CPU 已經有很完整的效能測試工具，因此實驗上主要也是由 OpenCL API 所提供的資訊作為測試依據，實驗數據中的時間都是從 OpenCL 中的提供的偵測功能，盡可能的達到較小誤差的情況下比較。由這些實驗結果我們可以做為設計 OpenCL 演算法以及資料結構的依據，但事實上不同的廠牌所做的硬體優化不盡相同，有時無法面面俱到，因此希望能在盡可能少的修正下能夠加速 OpenCL 的平行化程式。. 30.

(37) 第五章. Ray Tracer 優化設計與結果. 本實驗環境： CPU： GPU：. Intel Intel Nvidia Intel. Xeon E5506 I5-3550 Tesla C1060 HD 2500. (獨立顯示卡－使用獨立記憶體) (內嵌顯示卡－使用主記憶體). 作業系統：Microsoft Windows 7 編譯環境：Microsoft Visual Studio 2010 Professional OpenCL(v1.2)：Intel OpenCL SDK 2013 Nvidia CUDA SDK 4.2 這部分要稍微提到，Xeon E5506 與 Tesla C1060 是在同一台機器上，HD 2500 與 I5-3550 也是在同一台機器上，這些資料對於之後的實驗結果有些不小的影響，從表 1 的資料顯示中，官方提供各個運算單元的 single-precision GFLOPs (Giga FLoating-point Operations Per Second)是各有高低整理出來如下所示。. GFLOPs. Xeon E5506 34.128. I5-3550 105.6. Tesla C1060 933. HD 2500 55.2. 表 8. CPU/GPU 運算能力理論值 GFLOPs 表. 根據表中資訊，在運算能力上 Tesla C1060(GPU) > I5-3550(CPU) > HD 2500(GPU) > Xeon E5506(CPU)，而實驗結果比較是否與此相符，也是本次實驗所要觀察的項目之一。. 5.1 演算法流程剖析這邊採用的演算法是使用 Bounding volume hierarchy(BVH)-Tree 加速結構的. 31.

(38) 光追蹤演算法，其目的在於減少無意義的三角形碰撞的數量，加速整體的計算速度，這部分是與原本 C 版本的程式中就可以使用的光追蹤的加速方法，而實做到 OpenCL 後也保留了該加速結構，經由分析程式中的各個步驟占用時間，選擇適合平行化且佔用時間最多的部分。從中我們發現三角型相交與光影計算是程式中最消耗時間的部分，並且這些副程式都是以光線個體去做平行化的依據，因此在這次的實驗中我們就以兩個演算法"bvh-Intersect"(三角相交) 與" Occluded"(光影測試) 這兩個是主要我們嘗試以 OpenCL 撰寫的部分。我們把程式中主要的光追蹤演算法分成這幾個步驟 :. Part Initialization : 初始化必須使用到的大多數資料結構 Part A : bvh-Intersect 讓光線走訪 bvh-tree 取得碰撞點 Part B :整理碰撞後的資料，產生相應需求之作標位置與距離及其資料並整理備用，依據資料產生 sample 環境光源，帶入著色公式算出光線的顏色備用，並產生需要計算光影函式需要的光線。 Part C : Occluded 利用上步驟產生的光線計算是否被遮蔽。 Part D：依據結果判定最終顏色並預備回傳。 Part End : 將所有用到的動態記憶體釋放。. 使用該演算法計算以下場景: 32.

(39) 圖 12. 2 bvh-nodes 10 triangles 場景圖. 圖 13. 106 bvh-nodes 399 triangles 場景圖. 33.

(40) 以下是依照 C 程式分段的執行結果數據如下(未使用 OpenCL): 2 bvh-nodes 10 triangles. Xeon E5506. I5-3550. Part Initialization. 10.04. 7.145. 44.941. 22.559. 112.869. 60.831. 56.294. 30.206. Part D. 0.085. 0.044. Part End. 6.339. 4.153. 235.512. 127.871. Part A Intersect Part B Part C Occluded. Total. 表 9. 2 bvh-nodes 10 triangles 原始版本 CPU 執行數據. 2 bvh-nodes 10 triangles. Xeon E5506. I5-3550. Part Initialization. 10.101. 6.934. Part A Intersect. 59.113. 31.563. 119.115. 64.311. 84.883. 47.864. Part D. 0.082. 0.045. Part End. 6.195. 4.108. 284.237. 157.693. Part B Part C Occluded. Total. 表 10. 106 bvh-nodes 399 triangles 原始版本 CPU 執行數據. 從表中會發現 Part B 是最佔時間的項目，但是因為這個部分包含了許多步驟合在一起的關係事實上有許多細目沒有一一列舉，但其流程因為牽涉到許多比較不適合平行化以及嘗試平行化但導致效能低落的結果，因此把這些步驟都列入 Part B 而我們主要關注的部分是 Part A 與 Part C 的部分，這兩個步驟都是分散光線進行分散式的運算，本身處理起來相較於其他步驟是較為有意義的，且這兩個步驟的總和佔了總時間快要一半，因此我們決定採取針對 Part A 與 Part C 也就是 Intersect 與 Occluded 的功能為目標設計 OpenCL 平行化程式。. 34.

(41) 5.2 實做優化與實驗數據這部分主要為兩個演算法在 OpenCL 框架的實做與實驗，以下是兩個的程式 "bvh-Intersect"(三角相交) 與" Occluded"(光影測試)設計架構圖:. 分散每一個光線為一個 kernel 做計算。 Host Program. 走訪 BVH-Tree, 並把有可能碰撞的節點丟到 stack 反覆迴圈後整理出所有可能碰撞的子葉。 Kernel Program 對有符合條件的子葉中的三角片做三角形碰撞並記錄下來相關資訊。. Host Program 回傳到 host-program 進行後續處理。圖 14. 三角形碰撞 Intersect 程式架構圖. 35.

(42) 分散每一個從碰撞點產生往光源為路徑的光線傳入 kernel 做計算。 Host Program. 走訪 BVH-Tree, 並把有可能碰撞的節點丟到 stack 反覆迴圈後整理出所有可能碰撞的子葉。 Kernel Program 對有符合條件的子葉中的三角片做測試三角形碰撞是否碰撞，直到有碰撞或發現沒有碰撞為止。. Host Program 回傳到 host-program 進行最後顏色處理圖 15. 光影偵測 Occluded 程式架構圖. 從我們兩個架構圖來看，會發現這兩個演算法其實都是在做類似的事情，都是用光線去撞物體的流程，主要的差別在於圖 14.的算法需要把碰撞後相關資訊紀錄下來並且要跑完所有可能性確保碰撞到的三角型是最近的，而圖 15.則是在最後碰撞三角型時只要其中一次碰到即可回傳是否碰撞，兩者的設計架構原則上是類似的。依照之前第四章的經驗，我們盡可能的優化演算法，整體而言的設計方針如下：. 36.

(43) 1.In-Output 的記憶體在 kernel 中被存取最小化。 2.直接使用 host-memory 傳遞省去在 kernel 的創建記憶體與傳輸時間。 3.可能產生競爭存取之參數採取先複製到 private-memory 後再處理 4.盡可能減少分支與迴圈數量。 5.使用 BVH-tree 使走訪更有效率。 6.盡可能使用 OpenCL 內建的向量化功能一次處理多筆資料。以上就是這兩個修改的程式的優化方針，另外我們針對 Intersect CL program 做了另一點變形的設計，為了測試其中的 bvh-bodes 與 triangles 這兩個型態的資料陣列是否適合使用到 local 以及 private memory，我們做了如此的兩種設計變更：. local memory 版本: 將程式中的 bvh-bodes 與 triangles 在規劃下分散至整個 work-group 中的單位執行緒各自存取，並且分工複製到指定 local-memory 中再行計算。 private memory 版本: 將程式中每個 work-item 各自需要（有用到才複製並非全複製）的 bvh-bodes 與 triangles 到 private-memory 中再行計算。 37.

(44) 2nodes 10triangles 基本優化設計 Local Memory 複製 Private Memory 複製. Xeon E5506. I5-3550. Tesla C1060. HD 2500. 24.242. 9.144. 5.288. 44.639. 35.274. 10.554. 9.682. 47.757. 41.981. 9.356. 6.258. 58.134. I5-3550. Tesla C1060. HD 2500. 13.042. 9.992. 68.17. 14.108. 14.071. 95.336. 表 11.優化變型設計執行結果 2nodes 10triangles. Xeon E5506 106nodes 399triangles 基本優化設計 30.767 Private 50.737 Memory 複製. 表 12. 優化變型設計執行結果 106nodes 399triangles. 從結果中顯示，我們採取的程式變形並沒有達到有更進一步的加速，這部分經過程式碼的流程中我們發現，移動去其他記憶體的 bvh-nodes 與 triangles 的資料結構中，幾乎很少有被重複被存取超過兩次以上，這是因為在 BVH 樹型結構中幾乎會使三角形被存取次數降到最低，不同 kernel 在同一時間內會存取到同一筆資料的情況是很難發生的，且其之後的存取最多只有一次，因此並沒有賺到因為存取競爭所產生的時間節省。這是這些變形之所以沒有更進一步的加速的原因。另外由於 OpenCL 之使用記憶體容量都有一定限制，因此實驗中沒有再用更大的場景，這部分我們沒有特別把資料壓縮或是分批處理，但若是採取資料分批處理的程序，必定會導致有額外的資料傳輸成本，這部分會是將來在設計演算需要考慮的。而經過上述實驗結果，我們採取將兩個演算法的轉換成 OpenCL 的程式碼的策略方針都是使用基本優化方針，以下是兩種演算法的數據成果：. 38.

(45) Intersect Kernel Time Total CL API Time Kernel Time Total CL API Time. Xeon E5506 I5-3550 2 nodes 10 triangles 24.242 9.144. Tesla C1060. HD 2500. 5.288. 44.639. 41.207. 15.438. 49.31. 106 nodes 399 triangles 30.767 13.042. 9.992. 68.17. 48.127. 19.684. 73.745. Xeon E5506 I5-3550 2 nodes 10 triangles 29.697 8.381. Tesla C1060. HD 2500. 10.6. 55.563. 46.119. 21.133. 21.158. 59.881. 106 nodes 399 triangles 36.427 12.438. 50.307. 77.333. 55.432. 63.002. 82.13. 19.953. 24.592. 表 13. Intersect 實驗數據. Occluded Kernel Time Total CL API Time Kernel Time Total CL API Time. 24.642. 表 14. Occluded 實驗數據. Kernel Time 是透過 OpenCL API 中事件回傳的時間，而 Total CL API Time 則是包含執行以及傳輸資料，呼叫 OpenCL 的 API 與一些計算時間的參數。而我們可以把這些數據跟原始的 C 程式版本做比較，如下：. Intersect C OpenCL 加速比例 C OpenCL 加速比例. Xeon E5506 I5-3550 2 nodes 10 triangles 44.941 22.559 24.242 9.144 1.853 2.467 106 nodes 399 triangles 59.113 31.563. Tesla C1060. HD 2500. 44.941 5.288 8.498. 22.559 44.639 0.505. 59.113. 31.563. 30.767 1.921. 9.992 5.916. 68.17 0.463. 13.042 2.42. 表 15. Intersect 比較 CL 版本與 C 之速度. 39.

(46) Occluded C OpenCL 加速比例 C OpenCL 加速比例. Xeon E5506 I5-3550 2 nodes 10 triangles 56.294 30.206 29.697 8.381 1.895 3.604 106 nodes 399 triangles 84.883 47.864 36.427 12.438 2.33 3.848. Tesla C1060. HD 2500. 56.294 10.6 5.31. 30.206 55.563 0.539. 84.883 50.307 1.687. 47.864 77.333 0.618. 表 16. Occluded 比較 CL 版本與 C 之速度. 5.3 實驗結果與分析從以上圖表中得知，大部分的情況中使用 OpenCL 做平行化處理的結果通常是比較好的，以都是 CPU 的情況來說，使用 OpenCL 平行處理在適合平行化的條件下會使得在演算法速度相較使用單顆核心運算速度提升，但是 GPU 的情況就不一定如此，從之前的各項數據中發現 HD2500(GPU)在實驗中皆是 4 種平台中最慢的，而 C1060(GPU)卻是所有平台中最快的，這印證之前在一開始廠商提供的理論運算能力表 8. 中的資訊，而且該值居然比使用單核的 CPU 還要慢，我們的推論是因為 HD2500 的 GPU 並不是算是非常高階的，且並未做出許多硬體上的指令加速，在執行複雜的運算(多迴圈與分支)時，較難以達到理想的效能，且與其比較的 I5-3550 算是 Intel 中高端的 CPU，實驗數據皆顯示該平台的執行速度僅比 Nvidia Tesla C1060 慢一些。由於我們使用的 OpenCL API 設定為自動分配 work-item 數量，並且使用 host memory(主記憶體)減少記憶體傳輸的成本，因此已經使用建議上最佳的設定，雖然並不是所有的方法都適合本演算法中的優化，但若要有更高的效能則必須還是要從演算法設計著手，設計普通的 C 程式與設計平行計算的 OpenCL 程式在想法上就有顯著的差異，尤其是越是物件導向的資料結構架構下對於 OpenCL 下並. 40.

(47) 非很容易整合，許多資料都要先進行重新整理順序移除多型與重整指標動態宣告的部分才可以傳送給 OpenCL 使用，這無形中都會增加在程式中的執行成本，這部份由於我們只純粹考慮演算法兩種版本(C 與 OpenCL)因此暫時假設忽略這部分的整合成本，這算是要設計屬於 OpenCL 程式應該要在最初的資料架構下就要考慮周詳。而這邊採取的演算法都已經與原始程式碼多少有些更動，其撰寫法則大部分都跟 C 語言類似，但是主要是要以資料的存取同步做考量，並且 GPGPU 並不像 CPU 的設計已經有一段不小的歷史，許多針對指令的懮化是比 GPU 領先的，而 GPU 則是強項在於非常多的核心做較為單純的指令運算(尤其是圖形陣列處理)，各個平台所擅長的項目不盡相同，還是要依據演算法的性質決定適合的裝置。. 41.

(48) 第六章結論與未來方向. 6.1 結論現今使用 OpenCL 作平行化程式設計是一種泛用的程式架構，其目前已經支援各種不同平台，除了桌上型筆記型電腦 CPU、GPU 之外，現在也已經有嵌入式行動裝置的支援，也可以使用於各種作業系統，相信會是一個良好的協助平行計算的選擇之一。由前幾章的一些實驗配合 OpenCL 所的設計演算法優化方針，尤其是主要於記憶體的使用上要考慮的使用頻率與存取競爭問題，以及盡可能的資料相依性越低越好，越少的使用屏障同步對於程式的分工更有效率，這依然要看演算法的性質來決定，否則會無法達到加速的目的。. 6.2 未來研究方向由於本次的設計主要還是在於 OpenCL 記憶體使用限制下的場景，未來應可以設計針對當資料量極大時所採取的切割以及分流方法，但是依照目前為止 CPU 廠商則盡可能希望的 OpenCL 發展偏向直接使用主記憶體或是與 C 語言共用資料結構，也可能得以解決這些問題。而光追蹤演算法內也可透過變更演算法的設計以致於其他適合平行化的部分的方法，甚至可以加入次要光線追蹤的效果，這些是未來可能的研究方向。. 42.

(49) 附錄 1. Intersect OpenCL 程式碼 void triangle4_intersect(__global struct Triangle4 *triangle4, struct Ray *ray, __global struct Hit *hit) { float3 O = (float3) (ray->org[0], ray->org[1], ray->org[2]); float3 D = (float3) (ray->dir[0], ray->dir[1], ray->dir[2]); float4 C[3] = {triangle4->v0[0]-O.x, triangle4->v0[1]-O.y, triangle4->v0[2]-O.z}; float4 R[3] = SSE3F_FLOAT3_CROSS(C, -D); float4 det =. SSE3F_FLOAT3_DOT(triangle4->Ng, D);. float4 T = SSE3F_SSE3F_DOT(triangle4->Ng, C); float4 U = SSE3F_SSE3F_DOT(R, triangle4->e2); float4 V = SSE3F_SSE3F_DOT(R, triangle4->e1); float4 absDet = ABS(det); float4 signDet = select((float4)1.0, (float4)-1.0, signbit(det)); float4 _t = T * signDet;. float4 _u = U * signDet;. float4 _v = V * signDet;. float4 _w = absDet - _u - _v; int4 mask = (triangle4->id0 != -1) & (det != 0.f) & (_t >= absDet*ray->near) & (absDet*hit->t >= _t) & (min(min(_u,_v),_w) >= 0.f); if (! any(mask)) return; float4 rcpAbsDet = 1.0f / absDet; float4 t = _t * rcpAbsDet;. float4 u = _u * rcpAbsDet;. float4 v = _v * rcpAbsDet;. float4 __t = select(INFINITY, t, mask); float4 tmp = min(__t, __t.zwxy); float4 reducedMin = min(tmp, tmp.yxwz); mask = isequal(__t, reducedMin); int tri = round(dot(select((float4)0, (float4)(0,1,2,3), mask), (float4)1.0)); float tA[4] = {t.x, t.y, t.z, t.w}; float uA[4] = {u.x, u.y, u.z, u.w}; float vA[4] = {v.x, v.y, v.z, v.w}; int id0A[4] = {triangle4->id0.x, triangle4->id0.y, triangle4->id0.z, triangle4->id0.w}; int id1A[4] = {triangle4->id1.x, triangle4->id1.y, triangle4->id1.z, triangle4->id1.w}; hit->t = tA[tri];. hit->u = uA[tri];. hit->v = vA[tri];. hit->id0 = id0A[tri]; hit->id1 = id1A[tri]; return; }. 43.

(50) __kernel. void. BVH4Traverse(__global struct Ray *rays, __global struct Hit *hits,__global int. *bvh_root, __global struct BVH4_Node *bvh_nodes, __global struct Triangle4 *bvh_triangles) { __global struct BVH4_Node *p = bvh_nodes; int idx = get_global_id(0);. int nodeID, ofs, num;. struct Ray ray = rays[idx];. float4 tNearX, tNearY, tNearZ, tNear, tFarX, tFarY, tFarZ, tFar; int4 _hit;. int numHits=0;. int stackPtr = 1;. int hitList[4];. int stack[100]; stack[0] = *bvh_root;. const int nearX = ray.dir[0] >= 0 ? 0*sizeof(float4) : 1*sizeof(float4); const int nearY = ray.dir[1] >= 0 ? 2*sizeof(float4) : 3*sizeof(float4); const int nearZ = ray.dir[2] >= 0 ? 4*sizeof(float4) : 5*sizeof(float4); const int farX. = nearX ^ sizeof(float4);. const int farY. = nearY ^ sizeof(float4);. const int farZ. = nearZ ^ sizeof(float4);. hits[idx].t = ray.far;. while (stackPtr > 0) { nodeID = stack[--stackPtr]; if (nodeID >= 0) { p = node(bvh_nodes, nodeID); tNearX = (float4_ITEM(p,nearX) - ray.org[0]) * ray.rdir[0]; tNearY = (float4_ITEM(p,nearY) - ray.org[1]) * ray.rdir[1]; tNearZ = (float4_ITEM(p,nearZ) - ray.org[2]) * ray.rdir[2]; tNear = max( max( max(tNearX, tNearY), tNearZ ), ray.near ); tFarX = (float4_ITEM(p,farX) - ray.org[0]) * ray.rdir[0]; tFarY = (float4_ITEM(p,farY) - ray.org[1]) * ray.rdir[1]; tFarZ = (float4_ITEM(p,farZ) - ray.org[2]) * ray.rdir[2]; tFar = min( min( min(tFarX, tFarY), tFarZ ), ray.far ); _hit = tNear <= tFar; int _hitA[4] = {_hit.x, _hit.y, _hit.z, _hit.w}; float tNearA[4] = {tNear.x, tNear.y, tNear.z, tNear.w}; numHits = 0; for (int i=0; i<4; i++) { if (_hitA[i]) hitList[ numHits++ ] = i; } if (0 == numHits) continue; if (1 == numHits) { stack[stackPtr++] = p->child[ hitList[0] ]; continue; }. 44.

(51) if (2 == numHits) { if (tNearA[hitList[0]] < tNearA[hitList[1]]) { stack[stackPtr++] = p->child[ hitList[1] ]; stack[stackPtr++] = p->child[ hitList[0] ]; } else { stack[stackPtr++] = p->child[ hitList[0] ]; stack[stackPtr++] = p->child[ hitList[1] ]; } continue; } for (int i=0; i<numHits; i++) { int cMax = i; int dMax = tNearA[hitList[i]]; for (int j=i+1; j<numHits; j++) { int cur = hitList[j]; if (tNearA[cur] > dMax) { cMax = j;. dMax = tNearA[cur];. } } // for j if (cMax != i) { int tmp = hitList[i];. hitList[i] = hitList[cMax];. hitList[cMax] = tmp; } stack[stackPtr++] = p->child[ hitList[i] ]; } // for i } // if (nodeID >= 0) else { nodeID ^= 0x80000000;. ofs = nodeID >> 5;. num = nodeID & 0x1F;. if (!num) continue; for (int i=ofs; i<ofs+num; i++) { triangle4_intersect( &bvh_triangles[i], &ray, &hits[idx] ); } ray.far = hits[idx].t; } // else } // while (stackPtr > 0) }. 45.

(52) 2. Occluded OpenCL 程式碼 int triangle4_occluded(__global struct Triangle4 *triangle4, struct Ray *ray) { float3 O = (float3) (ray->org[0], ray->org[1], ray->org[2]); float3 D = (float3) (ray->dir[0], ray->dir[1], ray->dir[2]); float4 C[3] = {triangle4->v0[0]-O.x, triangle4->v0[1]-O.y, triangle4->v0[2]-O.z}; float4 R[3] = SSE3F_FLOAT3_CROSS(C, -D); float4 det =. SSE3F_FLOAT3_DOT(triangle4->Ng, D);. float4 T = SSE3F_SSE3F_DOT(triangle4->Ng, C); float4 U = SSE3F_SSE3F_DOT(R, triangle4->e2); float4 V = SSE3F_SSE3F_DOT(R, triangle4->e1); float4 absDet = ABS(det); float4 signDet = select((float4)1.0, (float4)-1.0, signbit(det)); float4 _t = T * signDet; float4 _u = U * signDet; float4 _v = V * signDet; float4 _w = absDet - _u - _v; int4 mask = (triangle4->id0 != -1) & (det != 0.f) & (_t >= absDet*ray->near) & (absDet*ray->far >= _t) & (min(min(_u,_v),_w) >= 0.f); return any(mask); } __kernel *bvh_root,. void BVH4Occluded(__global struct Ray *rays,. __global int *ishit,. __global int. __global struct BVH4_Node *bvh_nodes, __global struct Triangle4 *bvh_triangles). { __global struct BVH4_Node *p = bvh_nodes; int idx = get_global_id(0); ishit[idx]=0; int nodeID, ofs, num; struct Ray ray = rays[idx]; float4 tNearX, tNearY, tNearZ, tNear, tFarX, tFarY, tFarZ, tFar; int4 _hit; int numHits=0;. int hitList[4];. int stackPtr = 1;. int stack[100];. stack[0] = *bvh_root; const int nearX = ray.dir[0] >= 0 ? 0*sizeof(float4) : 1*sizeof(float4); const int nearY = ray.dir[1] >= 0 ? 2*sizeof(float4) : 3*sizeof(float4); const int nearZ = ray.dir[2] >= 0 ? 4*sizeof(float4) : 5*sizeof(float4);. 46.

(53) const int farX. = nearX ^ sizeof(float4);. const int farZ. = nearZ ^ sizeof(float4);. const int farY = nearY ^ sizeof(float4);. while (stackPtr > 0) { nodeID = stack[--stackPtr]; if (nodeID >= 0) { p = node(bvh_nodes, nodeID); tNearX = (float4_ITEM(p,nearX) - ray.org[0]) * ray.rdir[0]; tNearY = (float4_ITEM(p,nearY) - ray.org[1]) * ray.rdir[1]; tNearZ = (float4_ITEM(p,nearZ) - ray.org[2]) * ray.rdir[2]; tNear = max( max( max(tNearX, tNearY), tNearZ ), ray.near ); tFarX = (float4_ITEM(p,farX) - ray.org[0]) * ray.rdir[0]; tFarY = (float4_ITEM(p,farY) - ray.org[1]) * ray.rdir[1]; tFarZ = (float4_ITEM(p,farZ) - ray.org[2]) * ray.rdir[2]; tFar = min( min( min(tFarX, tFarY), tFarZ ), ray.far ); _hit = tNear <= tFar; int _hitA[4] = {_hit.x, _hit.y, _hit.z, _hit.w}; numHits = 0; for (int i=0; i<4; i++) { if (_hitA[i]) { stack[stackPtr++] = p->child[i]; } } } // if (nodeID >= 0) else { nodeID ^= 0x80000000; ofs = nodeID >> 5; num = nodeID & 0x1F; if (!num) continue; for (int i=ofs; i<ofs+num; i++) { if(triangle4_occluded(&bvh_triangles[i], &ray)){ ishit[idx]=1; return; } } } // else } // while (stackPtr > 0) return; }. 47.

(54) 參考文獻參考書目： [Du12]. P. Du, R. Weber, P. Luszczek, S. Tomov, G. Peterson, J. Dongarra: From CUDA to OpenCL: Towards a Performance-portable Solution for Multi-platform GPU Programming. Elsevier Science Publishers B. V. Amsterdam, 2012.. [GHKMS11] B. Gaster L. Howes, D. R. Kaeli, P. Mistry, D. Schaa: Heterogeneous Computing with OpenCL. Morgan Kaufmann, 2011. [KH12]. R. Karrenberg and S. Hack: Improving Performance of OpenCL on CPUs. Springer-Verlag Berlin, 2012.. [MGMF11]. A. Munshi, B. Gaster, T. G. Mattson, J. Fung, D. Ginsburg: OpenCL Programming Guide. Addison-Wesley, 2011.. [Rei08]. S. Reiter: Real–time Ray Tracing of Dynamic Scenes. Diploma Thesis, Institute for Graphics and Parallel Processing, Johannes Kepler University, 2008.. [SK11]. J. Sanders, E. Kandrot: Cuda by Example:An Introduction to General-purpose GPU Programming. Addison-Wesley, 2011.. [Sca11]. M. Scarpino: OpenCL in Action: How to Accelerate Graphics and Computations. Manning Publications, 2011.. [Shi09]. M. Shih, Y.-F. Chiu, Y.-C. Chen, C.-F. Chang: Real-Time Ray Tracing with CUDA. In ICA3PP ’09: Proceedings of the 9th International Conference on Algorithms and Architectures for Parallel Processing, pp. 327–337, 2009.. 48.

(55) [Suf07]. K. Suffern: Ray Tracing from the Ground Up. A K Peters , 2007.. [Wat99]. A. Watt: 3D Computer Graphics. Addison Wesley, 1999.. [Woo04]. S. Woop: A Ray Tracing Hardware Architecture for Dynamic Scenes. Computer Graphics Lab, Saarland University, 2004.. 參考網站： [1.] 經典光跡追蹤演算法 http://zh.wikipedia.org/wiki/%E5%85%89%E7%B7%9A%E8%BF%BD%E8%B9% A4 [2.]OpenCL Introduction http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/ [3.]Intel Ivy Bridge http://en.wikipedia.org/wiki/Ivy_Bridge_%28microarchitecture%29 [4.]NVidia CUDA http://en.wikipedia.org/wiki/CUDA [5.]Intel OpenCL Optimization-Guide http://software.intel.com/sites/landingpage/opencl/optimization-guide/ [6.] Intel Streaming_SIMD_Extensions http://en.wikipedia.org/wiki/Streaming_SIMD_Extensions [7.] OpenCL Programming for the CUDA Architecture http://www.nvidia.com/content/cudazone/download/OpenCL/NVIDIA_OpenCL_Prog rammingOverview.pdf [8.] CUDA Pinned memory http://www.orangeowlsolutions.com/archives/443. 49.

(56)