0

这个问题与我几周前发布的一个现有问题有关：TERCOM 算法 - 从单线程更改为 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 pos=0

我看过很多减少示例，但似乎每个人都将值存储在数组中，而不是在每个单独的线程中。

所以对于问题：

  1. 是否可以让 atomicMin 函数返回位置？
  2. 谁能给我一个关于如何解决推力库问题的提示？
4

1 回答 1

0

对于您的 Thrust 代码，您正在写入 dev_MAD[0]，但计算结果就像您已写入整个数组一样。

IIUC，您试图找到最小值和相应的位置，您将这些值作为每个线程中的变量但不存储在内存中。

我可以想到几种简单的方法来做到这一点，但都涉及将值存储到内存并在第二遍中计算最小值/位置。

首先，您可以像已经尝试过的那样使用Thrust的min_element ，但是您可以将这些值存储到内核中的 device_vector 中，然后独立调用thrust::min_element 。

其次，您可以通过首先计算线程块中的最小值/位置来节省一些内存空间和带宽（然后使用推力::min_element）。为此，您可以将CUB的归约与自定义归约运算符一起使用（比较值，基准为 {value,index}）。

于 2013-08-12T10:44:23.230 回答