查找線程

Question

這個問題是涉及到由我發佈了幾個星期前的現有問題之間的最小：TERCOM algorithm - Changing from single thread to multiple threads in CUDA

簡要說明，每個內核線程的計算MAD價值，我想知道最小的和它的位置。

我試着使用atomicMin這樣

__global__ void kernel (int m, int n, int h, int N, int *f, float heading, float *measurements, int *global_min) 
{ 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    int idy = blockIdx.y * blockDim.y + threadIdx.y; 

    float MAD=0; 
    float pos[2]; 
    float theta=heading*(PI/180); 
    float fval = 0; 

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

    //Calculate Mean Absolute Difference 
    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); 
     } 
    } 
    cuPrintf("%.2f \n",MAD); 

    atomicMin(global_min, MAD); 
    pos[0]=idx; 
    pos[1]=idy; 

    f[0]=*global_min; 
    f[1]=pos[0]; 
    f[2]=pos[1]; 
} 

而且它產生正確的結果，但atomicMin是無法找到最低的位置。

我還試圖用推力庫

__global__ void kernel (int m, int n, int h, int N, int *f, float heading, float *measurements, int *global_min, float *dev_MAD) 
{ 
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    int idy = blockIdx.y * blockDim.y + threadIdx.y; 

    float theta=heading*(PI/180); 
    float fval = 0; 

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

    //Calculate Mean Absolute Difference 
    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); 
      *dev_MAD += abs(measurements[(int)g]-fval); 
     } 
    } 
    cuPrintf("%.2f \n",MAD); 
} 

並調用內核這樣

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

thrust::device_ptr<float> dev_ptr(dev_MAD); 
thrust::device_ptr<float> min_pos = thrust::min_element(dev_ptr, dev_ptr + n*m); 
int abs_pos = min_pos - dev_ptr; 
float min_val=min_pos[0]; 

cudaMemcpy(&min_val, dev_MAD+abs_pos, sizeof(float), cudaMemcpyDeviceToHost); 

// Print out the result 
printf("Min=%.2f pos=%d\n",min_val,abs_pos); 

但這個方案打印出來：最小= -207521258711807190000000000000000000000.00 POS = 0

我我們看過很多縮減的例子，但是似乎在每個人中他們都將值存儲在一個數組中，而不是在每個單獨的線程中。

所以對這些問題：

1. 是否有可能使atomicMin函數返回的位置？
2. 任何人都可以給我一個關於如何解決推力庫問題的提示嗎？

Answer 1

對於您的Thrust代碼，您正在寫入dev_MAD [0]，但計算結果就像您已寫入整個陣列一樣。

IIUC，你試圖找到最小值和相應的位置，你有值作爲每個線程中的變量，但沒有存儲在內存中。

我可以考慮的一些簡單的方法是這樣做的，但都涉及將值存儲到內存並在第二遍中計算最小值/位置。首先，您可以使用Thrust的min_element，因爲您已經嘗試過了，但是您應該將值存儲到內核中的device_vector，然後獨立調用thrust :: min_element。其次，你可以通過首先計算線程塊內的最小值/位置（然後使用thrust :: min_element）來節省一些內存空間和帶寬。爲此，您可以使用CUB的自定義縮減運算符（比較值，基準值爲{value，index}）。