TERCOM算法 - 在CUDA中從單線程切換到多線程

Question

我目前正致力於使用僅使用1個線程的端口TERCOM algorithm來使用多線程。簡而言之，TERCOM算法接收5個測量值和航向，並將這些測量結果與預先存儲的地圖進行比較。該算法將選擇最佳匹配，即最低平均絕對差（MAD），並返回該位置。

該代碼完美工作與一個線程和for循環，但是當我嘗試使用多個線程和塊它返回錯誤的答案。看起來多線程版本不像「單線程」版本那樣「運行」計算。有誰知道我做錯了什麼？

下面是使用for循環的代碼

__global__ void kernel (int m, int n, int h, int N, float *f, float heading, float *measurements) 
{ 
    //Without threads 
    float pos[2]={0}; 
    float theta=heading*(PI/180); 
    float MAD=0; 

    // Calculate how much to move in x and y direction 
    float offset_x = h*cos(theta); 
    float offset_y = -h*sin(theta); 

    float min=100000; //Some High value 

    //Calculate Mean Absolute Difference 
    for(float row=0;row<m;row++) 
    { 
     for(float col=0;col<n;col++) 
     { 
      for(float g=0; g<N; g++) 
      { 
       f[(int)g] = tex2D (tex, col+(g-2)*offset_x+0.5f, row+(g-2)*offset_y+0.5f); 
       MAD += abs(measurements[(int)g]-f[(int)g]); 
      } 
      if(MAD<min) 
      { 
       min=MAD; 
       pos[0]=col; 
       pos[1]=row; 
      } 
      MAD=0;     //Reset MAD 
     } 
    } 

    f[0]=min; 
    f[1]=pos[0]; 
    f[2]=pos[1]; 
} 

這是我嘗試使用多線程

__global__ void kernel (int m, int n, int h, int N, float *f, float heading, float *measurements) 
{ 
    // With threads 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    int idy = blockIdx.y * blockDim.y + threadIdx.y; 
    float pos[2]={0}; 
    float theta=heading*(PI/180); 
    float MAD=0; 

    // Calculate how much to move in x and y direction 
    float offset_x = h*cos(theta); 
    float offset_y = -h*sin(theta); 

    float min=100000; //Some High value 

    if(idx < n && idy < m) 
    { 
     for(float g=0; g<N; g++) 
     { 
      f[(int)g] = tex2D (tex, idx+(g-2)*offset_x+0.5f, idy+(g-2)*offset_y+0.5f); 
      MAD += abs(measurements[(int)g]-f[(int)g]); 
     } 

     if(MAD<min) 
     { 
      min=MAD; 
      pos[0]=idx; 
      pos[1]=idy; 
     } 
     MAD=0;     //Reset MAD 
    } 
    f[0]=min; 
    f[1]=pos[0]; 
    f[2]=pos[1]; 
} 

到這裏進入內核

dim3 dimBlock(16,16); 
dim3 dimGrid; 
dimGrid.x = (n + dimBlock.x - 1)/dimBlock.x; 
dimGrid.y = (m + dimBlock.y - 1)/dimBlock.y; 

kernel <<< dimGrid,dimBlock >>> (m, n, h, N, dev_results, heading, dev_measurements); 

Answer 1

的基本問題是，你在代碼中有一個內存競賽，圍繞使用f作爲某種線索d本地暫存空間和輸出變量。每個併發線程將嘗試同時將值寫入f中的相同位置，這會產生未定義的行爲。

是最好的，我可以告訴大家，使用f作爲暫存空間甚至沒有必要在所有與內核的主要計算部分可以寫成這樣：

if(idx < n && idy < m) 
{ 
    for(float g=0; g<N; g++) 
    { 
     float fval = tex2D (tex, idx+(g-2)*offset_x+0.5f, idy+(g-2)*offset_y+0.5f); 
     MAD += abs(measurements[(int)g]-fval); 
    } 
    min=MAD; 
    pos[0]=idx; 
    pos[1]=idy; 
} 

[免責聲明：寫在瀏覽器中，使用風險自擔]

在該計算的最後，每個線程有自身min和pos值。至少這些必須存儲在唯一的全局內存中（即輸出必須有足夠的空間用於每個線程結果）。然後，您需要執行某種簡化操作，以從線程局部值集合中獲取全局最小值。這可能在主機或設備代碼中，或者兩者的組合中。有許多代碼可用於CUDA並行縮減，您應該可以通過搜索和/或查看隨CUDA工具包提供的示例來查找這些代碼。將它們調整到您需要保留位置以及最小值的指定情況應該是微不足道的。