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=-20752125871180719000000000000000000000.00 pos=0

私は多くのリダクションの例を見てきましたが、個々のスレッドではなく、配列に値が格納されているようです。

質問に：

  1. atomicMin 関数に場所を返すようにすることは可能ですか?
  2. スラストライブラリの問題を解決する方法についてヒントを教えてもらえますか?
4

1 に答える 1

0

Thrust コードでは、dev_MAD[0] に書き込みますが、配列全体に書き込んだかのように結果を計算します。

IIUC、最小値と対応する場所を見つけようとしています。値は各スレッドに変数としてありますが、メモリには保存されていません。

これを行うために考えられる簡単な方法がいくつかありますが、どちらも値をメモリに保存し、2 番目のパスで最小値/位置を計算する必要があります。

まず、すでに試したようにThrustのmin_elementを使用することもできますが、値をカーネルの device_vector に保存してから、 throw::min_element を個別に呼び出します。

第二に、最初にスレッドブロック内の最小/場所を計算することで、メモリスペースと帯域幅を節約できます (その後、thrust::min_element を使用します)。このために、カスタムの reduce 演算子を使用してCUBの削減を使用できます (値で比較、データムは {value,index} です)。

于 2013-08-12T10:44:23.230 に答える