Нахождение минимума между потоками

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); 

Но эту программу распечатку: Min = -207521258711807190000000000000000000000,00 позы = 0

I Мы рассмотрели множество примеров сокращения, но, похоже, у каждого они имеют значения, хранящиеся в массиве, а не в каждом отдельном потоке.

Так на вопросы:

1. Можно ли сделать функцию atomicMin вернуть расположение?
2. Может ли кто-нибудь дать мне подсказку о том, как решить проблему с помощью библиотеки тяги?

Answer 1

Для вашего кода тяги вы пишете dev_MAD [0], но вычислите результаты, как если бы вы написали весь массив.

IIUC, вы пытаетесь найти минимальное значение и соответствующее местоположение, у вас есть значения как переменные в каждом потоке, но не сохранены в памяти.

Есть несколько простых способов, которые я могу сделать для этого, но оба они связаны с хранением значений в памяти и вычислением минимума/местоположения во втором проходе.

Во-первых, вы можете просто использовать Thrustmin_element, как вы уже пробовали, но сохраните значения в device_vector в своем ядре и затем вызовите thrust :: min_element независимо.

Во-вторых, вы могли бы сэкономить некоторое пространство памяти и пропускную способность, вычислив минимум/местоположение в потоковом блоке сначала (а затем используйте thrust :: min_element впоследствии). Для этого вы можете использовать сокращение CUB с помощью пользовательского оператора сокращения (сравнение по значению, datum - {value, index}).