0

現在、 1 つのスレッドのみを使用して複数のスレッドを使用するTERCOM アルゴリズムの移植に取り組んでいます。簡単に説明すると、TERCOM アルゴリズムは 5 つの測定値と方位を受け取り、この測定値を事前に保存されたマップと比較します。アルゴリズムは、最適な一致、つまり最小の平均絶対差 (MAD) を選択し、位置を返します。

コードは 1 つのスレッドと 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);
4

1 に答える 1

1

fここでの基本的な問題は、ある種のスレッド ローカル スクラッチ スペースと出力変数の両方としての使用を中心に、コード内でメモリ競合が発生していることです。すべての並行スレッドは、同時に同じ場所に値を書き込もうとする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ます。少なくとも、これらは一意のグローバル メモリに格納する必要があります (つまり、出力には、各スレッドの結果に対して十分なスペースが必要です)。次に、一連のスレッド ローカル値からグローバル最小値を取得するために、何らかのリダクション操作を実行する必要があります。これは、ホスト、デバイス コード、またはその 2 つの組み合わせにある可能性があります。CUDA ツールキットで提供されている例を検索および/または参照することで、CUDA 並列リダクションに使用できる多くのコードが既に利用可能です。最小値とともに位置を保持する必要がある特定のケースにそれらを適応させるのは簡単です。

于 2013-07-30T12:31:34.033 に答える