CUDA「卷積」與OpenMP版本一樣緩慢

Question

我試圖用另一個modelWidth * modelHeight * 31多維數據集「卷積」一個featWidth * featHeight * 31立方體。問題在於這個內核速度很慢（當然，我的速度比順序的CPU代碼快，但和OpenMP版本一樣慢）。我使用的是Quadro FX 1800（是的，64個CUDA內核......）。

__constant__ float d_model[31*22*22]; 
#define IMUL(a,b) (__mul24((a), (b))) 
#define IMAD(a,b,c) (__mul24((a), (b)) + (c)) 
__global__ void dMatch(float *score, const int featWidth, const int featHeight, const int modelWidth, const int modelHeight, const int scoreWidth, const int scoreHeight) 
{ 
    const int x = IMAD(blockIdx.x, blockDim.x, threadIdx.x); 
    const int y = IMAD(blockIdx.y, blockDim.y, threadIdx.y); 
    if(x < scoreWidth && y < scoreHeight) 
    { 
    const int scoreIdx = IMAD(x, scoreHeight, y); 
    score[scoreIdx] = 0.f; 
    const int baseFeatIdx = IMUL(x,scoreHeight) + IMAD(modelHeight-1, x, y); 
    for(int z = 0; z < 31; ++z) 
    { 
    // Index positionning 
    int featIdx = IMAD(z, IMUL(featWidth,featHeight), baseFeatIdx); 
    int modelIdx = IMUL(z, IMUL(modelWidth,modelHeight)); 

    float value = 0.f; 

    // filter 
    for(int xx=0; xx<modelWidth; xx++) 
    { 
     const int xxmodelIdx = IMAD(xx, modelHeight, modelIdx); 
     const int xxfeatIdx = IMAD(xx, featHeight, featIdx); 
     for(int yy=0; yy<modelHeight; yy++) 
     { 
     value += d_model[xxmodelIdx+yy] * tex1Dfetch(texFeatures,xxfeatIdx+yy); 
     } 
    } 
    score[scoreIdx] += value; 
    } 
} 
} 

無論如何，我推出這個內核8*8線程塊，並用的(scoreWidth/8)*(scoreHeight/8)（scoreWidth和scoreHeight是結果矩陣大小）的網格尺寸。 我想知道你是否有任何線索有什麼問題，或者我的代碼中速度很慢。

編輯：

更快的版本（150毫秒下降了480毫秒的過程！）由於TERA：

__global__ void dMatch(float *score, const int featWidth, const int featHeight, const int modelWidth, const int modelHeight, const int scoreWidth, const int scoreHeight) 
{ 
    const int y = IMUL(4,IMAD(blockIdx.x, blockDim.x, threadIdx.x)); 
    const int x = IMAD(blockIdx.y, blockDim.y, threadIdx.y); 
    if(x < scoreWidth && y < scoreHeight) 
    { 
    const int scoreIdx = IMAD(x, scoreHeight, y); 
    const int baseFeatIdx = IMUL(x,scoreHeight) + IMAD(modelHeight-1, x, y); 
    float value=0.f, value1 = 0.f, value2 = 0.f, value3 = 0.f; 
    float feat,feat1,feat2,feat3; 

    // Index positionning 
    int featIdx = 0; 
    int modelIdx = 0; 
    int xxmodelIdx; 
    int xxfeatIdx; 
    float val; 
    for(int z = 0; z < 31; ++z) 
    { 
     featIdx = IMAD(z,IMUL(featWidth,featHeight),baseFeatIdx); 
     modelIdx = IMUL(z,IMUL(modelWidth,modelHeight)); 

     // filter 
     for(int xx=0; xx<modelWidth; xx++) 
     { 
      xxmodelIdx = IMAD(xx, modelHeight, modelIdx); 
      xxfeatIdx = IMAD(xx, featHeight, featIdx); 
      feat=tex1Dfetch(texFeatures,xxfeatIdx+0); 
      feat1=tex1Dfetch(texFeatures,xxfeatIdx+1); 
      feat2=tex1Dfetch(texFeatures,xxfeatIdx+2); 
      feat3=tex1Dfetch(texFeatures,xxfeatIdx+3); 
      for(int yy=0; yy<modelHeight; yy++) 
      { 
       val = d_model[xxmodelIdx+yy]; 
       value += val * feat; 
       value1 += val * feat1; 
       value2 += val * feat2; 
       value3 += val * feat3; 
       feat = feat1; 
       feat1 = feat2; 
       feat2 = feat3; 
       feat3 = tex1Dfetch(texFeatures,xxfeatIdx+yy+4); 
      } 
     } 
    } 
    score[scoreIdx] = value; 
    if(y+1 < scoreHeight) 
     score[scoreIdx+1] = value1; 
    if(y+2 < scoreHeight) 
     score[scoreIdx+2] = value2; 
    if(y+3 < scoreHeight) 
     score[scoreIdx+3] = value3; 
} 

與此dim3 threads(16,16); dim3 grid(divup(scoreHeight,64), divup(scoreWidth,16));推出。

Answer 1

分析器說什麼？ NVidia NSight（Windows上的Visual Studio和Linux上的Eclipse的插件）允許您查看兩個檔位的位置，並提供各種提示以優化性能。

我的猜測（沒有在profiler上查看）是你的塊太小。在warp內部有32個線程，這是基本的調度元素。 NVIDIA GPU能夠快速運行，因爲它可以通過在其他線程上運行而隱藏延遲，而目前的線程正在執行前面的指令。雖然每個SM（特斯拉和費米）或16（開普勒）上可能有8個塊，但峯值可能會很小（我可能是錯的，但啓動塊有一定的延遲）。我會考慮使用更大的塊。

的紋理獲取是次優的，如果我理解正確的代碼 - 在經線程在baseFeatId，因此在featIdx和xxfeatIdx相差modelHeight - 1。因此，紋理提取完全是隨機的，並且不會利用數據局部性。反轉x和y會使它更有效率。

但是，好的規則是與分析器一起檢查 - 如果您的問題在GPU上受到計算限制，那麼您應該專注於計算方面。如果你的問題是內存綁定，你應該看看內存訪問模式。可能還有其他幾個部分看起來像優化的地方，但直到看到瓶頸時纔會知道。一旦你知道它，你可能想閱讀best practices guide的特定章節。