OpenACC 中parallel 和kernels的區別

Kernels構件

Kernels構件源于PGI Accelerator模型的region構件。嵌套kernels構件里的循環可能會被編譯器轉換成能在GPU上高效并行的部分。在這個過程中有三步。?

1：判斷并行中遇到的循環。

2：把抽象的并行轉換成硬件上的并行。對于NVIDIA CUDA GPU， 它會把并行的循環映射到grid層次(blockIdx) 或 thread層次(threadIdx)。OpenACC申明， gang 對應grid， vector 對應thread。編譯器可能會通過strip-mining(一種拆分循環利用緩存的技術)把一層的循環映射到多層。

3：編譯器生成并優化代碼。

?

在kernels構件中，編譯器用自動并行技術識別并行的循環。這個識別能被指令(directives)和條款(clauses)增強，比如說loop independent。使用-Minfo標志可以讓PGI 編譯器顯示編譯信息。你能看到類似Loop is parallelizable 如果編譯器認為這個循環可以被并行化。。

?

在第二步中，PGI編譯器使用目標硬件的模型去選擇循環被映射成vector并行還是gang。比如說，循環中有多個步進為1的數組時，更多的會被映射成vector(thread)并行。在NVIDIA GPU上，這種映射更傾向于生成能同時運行的代碼。同樣能用 loop gang 明確生成grid上的并行，或者用loop vector(64)生成64個thread的block。但這會造成移植上的困難。

第三步底層代碼生成和優化。

Parallel 構件

Parallel構件源于OpenMp的parallel構件。OpenMP會立即產生很多多余的線程，當運行到一個循環是，一個線程運行一部分。

同樣的，OpenACC的parallel構件也會產生多余的gangs。不同的是OpenACC的parallel在結束時沒有同步。就像帶了nowait參數的OpenMP一樣。

Kernels和parallel主要的不同是， 整個parallel構件會被變成一個kernel。比如說在CUDA中，會變成一個kernel<<<grid, block>>>();

?

舉例分析：

單個循環

???????? 在單個循環中kernels和parallel幾乎一樣。

#pragma acc kernels loopfor( i = 0; i < n; ++i )a[i] = b[i] + c[i];

等價于

 #pragma acc kernels{for( i = 0; i < n; ++i )a[i] = b[i] + c[i];}

如果a,b,c是指針的話，編譯器沒法消除指針的歧義，可能不會生成并行代碼。可以用restrict屬性/加-Msafeptr(不推薦)/jia independent子語

#pragma acc kernels loop independentfor( i = 0; i < n; ++i )a[i] = b[i] + c[i];

如果用parallel構件，就相當于告訴編譯器兩件事

1. 把接下來的循環映射成kernel
2. 把每一個步長分到gangs(grid)上

#pragma acc parallel loopfor( i = 0; i < n; ++i )a[i] = b[i] + c[i];

這會造成每一個block里都只有一個thread的情況。

下面一種和上面完全一致

 #pragma acc parallel{#pragma acc loopfor( i = 0; i < n; ++i )a[i] = b[i] + c[i];}

但是如果把里面的loop子語去掉

#pragma acc parallel{for( i = 0; i < n; ++i )a[i] = b[i] + c[i];}

循環就會在所有gang中運行。這是另一種浪費。

???????? 嵌套循環：

  !$acc kernels loopdo j = 1, mdo i = 1, na(j,i) = b(j,i) * alpha + c(i,j) * betaenddo
enddo

編譯器發現j被更多的數組用作stride-1索引，他會選擇把j循環作為vector并行，把i循環作為gangs并行。

但是，編譯器在這里有點自由。他可能會生成二維的grid，把i，j作為在gangs中多個block的索引。或者生成二維的block，把i，j作為vector中多個threads的索引。但是他總是尋找好性能的實現。

 !$acc parallel loopdo j = 1, mdo i = 1, na(j,i) = b(j,i) * alpha + c(i,j) * betaenddoenddo

或者

!$acc paralleldo j = 1, m!$acc loopdo i = 1, na(j,i) = b(j,i) * alpha + c(i,j) * betaenddoenddo

這樣編譯器能把i或j循環轉成vector并行。

 !$acc parallel loopdo i = 1, ndo j = 1, ma(j,i) = b(j,i) * alpha + c(i,j) * betaenddoenddo

這樣編譯器更可能會把i作為gangs并行，j作為vector并行。

不緊湊的嵌套循環(Non-tight Nested Loop)

#pragma acc kernels loopfor( i = 0; i < nrows; ++i ){double val = 0.0;int nstart = rowindex[i];int nend = rowindex[i+1];#pragma acc loop vector reduction(+:val)for( n = nstart; n < nend; ++n )val += m[n] * v[colndx[n]];r[i] = val;}

編譯器會為外循環生成gang并行，內循環vector并行

#pragma acc kernels loopfor( i = 0; i < nrows; ++i ){double val = 0.0;int nstart = rowindex[i];int nend = rowindex[i+1];for( n = nstart; n < nend; ++n )val += m[n] * v[colndx[n]];r[i] = val;}

如果沒有loop導語， 編譯器可能會把外循環在gangs和vector上并行， 或者自動把內循環vector化。

 #pragma acc parallel loopfor( i = 0; i < nrows; ++i ){double val = 0.0;int nstart = rowindex[i];int nend = rowindex[i+1];#pragma acc loop vector reduction(+:val)for( n = nstart; n < nend; ++n )val += m[n] * v[colndx[n]];r[i] = val;}

這樣，編譯器就不需要分析了。或者可以去掉loop vector reduction而讓編譯器自動尋找并行的機會。

相鄰的循環

 !$acc kernelsdo i = 1, na(i) = mob(i)*charge(i)enddodo i = 2, n-1c(i) = 0.5*(a(i-1) + a(i+1))enddo
!$acc end kernels

這兩個循環會被獨自的分析，調度和編譯。A會在第二個循環開始前結束。這和下面一樣。

 !$acc kernels loopdo i = 1, na(i) = mob(i)*charge(i)enddo!$acc kernels loopdo i = 2, n-1c(i) = 0.5*(a(i-1) + a(i+1))enddo

但是，寫成parallel構件就會很不同。

!$acc parallel!$acc loopdo i = 1, na(i) = mob(i)*charge(i)enddo!$acc loopdo i = 2, n-1c(i) = 0.5*(a(i-1) + a(i+1))enddo!$acc end parallel

兩個循環會同時被并行化，而且沒有同步操作。A可能還沒完成，第二個循環就開始了。

同樣的問題可能會在reduction中出現

 sum = 0.0!$acc kernels!$acc loop reduction(+:sum)do i = 1, nsum = sum + v(i)enddodo i = 1,nv(i) = v(i) / sumenddo!$acc end kernels

因為用的是kernel，第一個規約會在第二個循環開始前結束。

如果寫成parallel，這個依賴關系就又破了

   sum = 0.0!$acc parallel!$acc loop reduction(+:sum)do i = 1, nsum = sum + v(i)enddo!$acc loopdo i = 1,nv(i) = v(i) / sumenddo!$acc end parallel

事實上，多數教程都建議不要把reduction 和acc loop parallel 一起用。

總結：

???????? Kernels 和parallel構件都用于解決一樣的問題。區別是 kernels構件更隱含一些，給編譯器更多的自由性選擇。Parallel更加嚴格，需要程序員更多的分析。