GPGPU + SIMD、大改進的 OpenMP 4.0

| | 0 Comments| 08:40
Categories:

很久以前,Heresy 曾經介紹過 OpenMP 這個超簡單的程式平行化的標準(參考);當時 Heresy 並沒有特別去管他的版本,不過實際上以 VisualStudio 來說,就算到了最新的 2013 版,也都還是只支援到 OpenMP 2.0(MSDN)。而實際上,OpenMP 在 2008 年就已經把標準更新到 3.0 了,而最新的版本,則是 2011 年的 3.1 版(維基百科)。

而在今年,OpenMP 則是發表了最新版的 OpenMP 4.0,做了相當大的改進!雖然現在似乎還沒有可以完整支援的編譯器(官方編譯器支援列表,不過似乎沒有更新的很快),不過這邊還是稍微來整理一下好了。

OpenMP 的官方網站是:http://www.openmp.org/,目前最新版的 OpenMP C API 文件也已經可以下載了(PDF),另外也有快速參考卡可以下載(PDF)。如果以 Intel 的《What will be new in OpenMP 4.0?》這份投影片(PDF)來看,OpenMP 4.0 最主要的新功能,基本上有五個:

  • 使用者定義的 Reduction(User-defined Reductions)
  • 支援 SIMD 的平行化(Support for SIMD parallelism)
  • 強化 Task 的功能(Task Extensions)
  • 支援加速裝置以及協同處理器(Support for Accelerators and Coprocessors)
  • Thread-affinity Support

其中,Heresy 個人覺得最重要的,應該還是針對 SIMD 以及加速裝置的支援了!也就是說,以後如果開發環境有支援的話,就可以直接用 OpenMP 來做 GPGPU、用顯示卡、或是 Intel Xeon Phi 這類的計算卡來做大量平行化的加速了!


加速裝置

在以往的 OpenMP 來說,都是採用「shared memory」、所有的執行序都共用同一塊記憶體空間的配置,而從 OpenMP 4.0 開始,他已經可以在單台電腦上,去操作多張加速計算卡(顯卡、或是協同處理器)、讓他們使用各自獨立的記憶體空間,來做計算了~

透過新的 target 這個新的 directive,可以指定要使用哪個裝置來做計算、並配置相關的記憶體對應,之後就可以用指定的裝置來計算了~而再搭配 teams,也可以再去指定執行序的群組、來做更進一步的最佳化。在《OpenMP for Accelerators》這份投影片(PDF)裡面,有提供了一些簡單的範例:

double B[N] = ...;
#pragma omp target device(0) map(tofrom:B)
#pragma omp parallel for
for( i = 0; i < N; i )
B[i] = sin( B[i] );

在上面的程式碼裡面,除了本來用來把迴圈平行化的 omp parallel for 外,這邊還加入了一行新的、以 target 為首的 directive,這邊就是用來指定要使用第一個 device、並且把 B 這個變數對應到這個 device 上。而根據投影片的說法,這樣理論上是可以在 Intel Xeon Phi 這樣的協同處理器上運作的~

而如果是在顯示卡這種有做核心群組(例如 CUDA 的 thread block、nVIDIA GPU 架構內的 SMX)的話,則還可以透過 teams 來針對硬體的配置來做最佳化;修改後的程式可能會像下面這樣:

double B[N] = ...;
#pragma omp target device(0) map(tofrom:B)
#pragma omp teams num_teams(num_blocks) num_threads(bsize)
#pragma omp distribute
for( i = 0; i < N; i = num_blocks )
#pragma omp parallel for
for( b = i; b < i num_blocks; b )
B[b] = sin( B[b] );

理論上,這樣在 nVIDIA 的 GPU 上,應該會有更好的效能。

這樣的程式雖然看來比本來的 OpenMP 複雜了些,但是相較於直接去寫 nVIDIA CUDA、或是 OpenCL 程式比起來,算是相當簡單的了!基本上,這種 GPGPU 的程式開發方法,應該也是當初 OpenACC 想做的;如果之後主流開發環境都能支援 OpenMP 4.0 的話,應該就不需要特別去找支援 OpenACC 的編譯器了。


SIMD

除了可以使用顯示卡等計算卡來做計算外,Heresy 認為另一個最重要的新功能,則是他可以透過 SIMD(Single instruction, multiple data)來做平行化(或者應該說是向量化?);使用 SIMD 的形式來做計算,主要是在同一個時間處理迴圈裡的多組資料,基本上是可以透過處理器的特殊指令,來做加速的~理論上這樣在符合條件的情況下,也會有相當程度的效能增進的~

要使用的話,最基本的使用方法也相當簡單,只要在 for 迴圈前加上

#pragma omp simd

就可以了。這樣的好處,是不必去特別研究 SIMD 的指令怎麼用,就可以簡單地把程式裡的迴圈向量化了。

而 SIMD 的向量化處理也可以和本來的平行化並用,只要寫成:

#pragma omp parallel for simd

這樣就好了~如此一來,程式除了會把迴圈做向量化處理外,也會建立多個執行序、來做平行化的計算。

另外,透過新的 directive,也可以指定自定義的函式可以在 SIMD 迴圈內呼叫,不過這邊就先不細講了。


至於其他像是可以自定義 reduction 的計算方法,或是在 task 方面的諸多強化,其實在需要時也都是相當有用的~不過由於現在 Heresy 手邊沒有可以拿來玩的編譯器,所以就先不介紹了。

而實際上,就算上面的 device 和 SIMD,也都僅只是看著文件寫的粗淺介紹,等到有可以玩的編譯器後,Heresy 應該會再認真玩看看、然後再來做分享的。


參考:

附註:

  • gcc 4.9 開始應該是有支援 OpenMP 4.0 了,不過基本上 target 和 simd 似乎都不能實際運作(參考
  • Intel® Composer XE 2013 SP1 應該是有支援 OpenMP 4.0 了,不過不確定支援到什麼程度(官網

Leave a Reply

發佈留言必須填寫的電子郵件地址不會公開。 必填欄位標示為 *