2012-10-17 54 views
3

我試圖使用共享內存來緩存OpenACC的東西。與OpenACC共享內存的使用

基本上就是我工作的是一個矩陣乘法,和我有什麼是這樣的:

typedef float ff; 

// Multiplies two square row-major matrices a and b, puts the result in c. 
void mmul(const restrict ff* a, 
      const restrict ff* b, 
      restrict ff* c, 
      const int n) { 
#pragma acc data copyin(a[0:n*n], b[0:n*n]) copy(c[0:n*n]) 
{ 

#pragma acc region 
{ 

#pragma acc loop independent vector(16) 
    for (int i = 0; i < n; ++i) { 
#pragma acc loop independent vector(16) 
    for (int j = 0; j < n; ++j) { 
     ff sum = 0; 
     for (int k = 0; k < n; ++k) { 
     sum += a[i + n * k] * b[k + n * j]; 
     } 
     c[i + n * j] = sum; 
    } 
    } 

} 
} 
} 

我想這樣做的是使用共享內存對矩陣的高速緩存片「一個'和'b'用於計算'c',這與CUDA mmul算法的作用類似。

基本上在CUDA我也知道我的塊的確切大小,將能夠:

  • 聲明一個共享內存塊
  • 複製的的「相關」部分的大小數據塊
  • 使用這個數據

我明白我可以使用

#pragma acc cached 

指令,我可以指定塊大小矢量幫派選項,但我有一些麻煩理解如何將映射到CUDA體系結構。

有沒有辦法實現與OpenACC類似的東西?有沒有關於使用緩存指令的好教程/資源,或者有關如何將CUDA的一些共享內存的功能映射到OpenACC?

+1

PGI加速器編譯器可能已經在使用共享內存。您是否使用-Minfo開關檢查了輸出?這[教程](http://developer.nvidia.com/cuda/openacc-example-part-1)可能是有趣的。 –

+1

是的,但Minfo開關只告訴我我的實現正在使用多少共享內存。儘管這很有用,但我更感興趣的是知道是否有明確的方式**操縱這些內存。能夠看到高水平的cuda生成是非常有用的,但。 – leo

+0

@leo你找到了你的問題的答案?你能否在OpenACC中明確定義共享內存? – Millad

回答

4

如果您正在使用PGI加速器編譯器,你可以轉儲在所產生的PTX文件,看到什麼是在執行的下屬怎麼回事:

pgcc -acc -fast -Minfo -ta=nvidia,cc13,keepptx matrixMult.c -o matrixMult 

生成的PTX將被保存在當前目錄。

編輯:你可能更喜歡看到高級代碼(C或Fortran的CUDA)。所以請使用以下-ta=nvidia,cc13,keepptx,keepgpu