1

Mark Harris のリダクションの例を見て、スレッドにリダクション操作なしで中間値を格納できるかどうかを確認しようとしています。

たとえば、CPU コード:

for(int i = 0; i < ntr; i++)
{
    for(int j = 0; j < pos* posdir; j++)
    {
        val = x[i] * arr[j];
        if(val > 0.0)
        {
            out[xcount] = val*x[i];
            xcount += 1;
        }
    }
}

同等の GPU コード:

const int threads = 64; 
num_blocks = ntr/threads;

__global__ void test_g(float *in1, float *in2, float *out1, int *ct, int posdir, int pos)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    __shared__ float t1[threads];
    __shared__ float t2[threads];

    int gcount  = 0;

    for(int i = 0; i < posdir*pos; i += 32) {
        if (threadIdx.x < 32) {
            t1[threadIdx.x] = in2[i%posdir];
        }
       __syncthreads();

        for(int i = 0; i < 32; i++)
        {
            t2[i] = t1[i] * in1[tid];
                if(t2[i] > 0){
                    out1[gcount] = t2[i] * in1[tid];
                    gcount = gcount + 1;
                }
        }
    }        
    ct[0] = gcount;
}

ここでやろうとしていることは、次の手順です。

(1)共有メモリ変数t1にin2の値を32個格納し、

(2) i と in1[tid] の各値について、t2[i] を計算し、

(3) if t2[i] > 0i の特定の組み合わせについてt2[i]*in1[tid]、out1[gcount]

しかし、私の出力はすべて間違っています。t2[i] が 0 より大きい回数をすべて取得することさえできません。

各 i および tid の gcount の値を保存する方法に関する提案はありますか?? デバッグしていると、ブロック (0,0,0) とスレッド (0,0,0) の t2 の値が更新されていることがわかります。CUDA カーネルがフォーカスをブロック (0,0,0) とスレッド (32,0,0) に切り替えた後、out1[0] の値が再び書き直されます。各スレッドの out1 の値を取得/保存し、それを出力に書き込むにはどうすればよいですか?

これまでに 2 つのアプローチを試しました: (NVIDIA フォーラムで @paseolatis が提案)

(1) 定義済みoffset=tid*32; and replace out1[gcount] with out1[offset+gcount]、

(2)定義

__device__ int totgcount=0; // this line before main()
atomicAdd(&totgcount,1);
out1[totgcount]=t2[i] * in1[tid];

int *h_xc = (int*) malloc(sizeof(int) * 1);
cudaMemcpyFromSymbol(h_xc, totgcount, sizeof(int)*1, cudaMemcpyDeviceToHost);
printf("GPU: xcount = %d\n", h_xc[0]); // Output looks like this: GPU: xcount = 1928669800

助言がありますか？前もって感謝します ！

4

2 に答える 2

2

OK、コードが何をすべきかについてのあなたの説明をあなたが投稿したものと比較しましょう（これはラバーダックデバッグと呼ばれることもあります）。

  1. in2の32個の値を共有メモリ変数に格納しますt1

    カーネルには次のものが含まれています。

    if (threadIdx.x < 32) {
    t1[threadIdx.x] = in2[i%posdir];
}

    これは、からすべての値に同じ値を効果的にロードしています。私はあなたがこのようなものがもっと欲しいと思う：in2t1

    if (threadIdx.x < 32) {
    t1[threadIdx.x] = in2[i+threadIdx.x];
}

  2. iとの値ごとにin1[tid]、を計算t2[i]します。

    この部分は問題ありませんが、なぜt2共有メモリに必要なのですか？これは、内部反復が完了した後に破棄できる中間結果にすぎません。あなたは簡単に次のようなものを持つことができます：

    float inval = in1[tid];
.......
for(int i = 0; i < 32; i++)
{
     float result = t1[i] * inval;
     ......

  3. iのそのt2[i] > 0特定の組み合わせについては t2[i]*in1[tid]、out1[gcount]

    ここから問題が実際に始まります。ここでこれを行います：

            if(t2[i] > 0){
            out1[gcount] = t2[i] * in1[tid];
            gcount = gcount + 1;
        }

    これは記憶の競争です。gcountはスレッドローカル変数であるため、各スレッドは、異なる時間に、指定out1[gcount]されたものを独自の値で上書きします。このコードが記述どおりに正しく機能するために必要なのはgcount、グローバルメモリ変数として持つ必要があり、アトミックメモリの更新を使用して、各スレッドが値をgcount出力するたびに一意の値を使用するようにすることです。ただし、アトミックメモリアクセスを頻繁に使用すると非常にコストがかかることに注意してください（これが、コメントでカーネルの起動ごとに出力ポイントがいくつあるかを尋ねた理由です）。

結果のカーネルは次のようになります。

__device__ int gcount; // must be set to zero before the kernel launch

__global__ void test_g(float *in1, float *in2, float *out1, int posdir, int pos)
{
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    __shared__ float t1[32];

    float ival = in1[tid];

    for(int i = 0; i < posdir*pos; i += 32) {
        if (threadIdx.x < 32) {
            t1[threadIdx.x] = in2[i+threadIdx.x];
        }
        __syncthreads();

        for(int j = 0; j < 32; j++)
        {
            float tval = t1[j] * ival;
            if(tval > 0){
                int idx = atomicAdd(&gcount, 1);
                out1[idx] = tval * ival
            }
        }
    }        
}

免責事項：ブラウザで記述されており、コンパイルまたはテストされていないため、自己責任で使用してください。

書き込みctもメモリ競合でしたが、gcountがグローバル値になったことで、カーネルの後で、を必要とせずに値を読み取ることができることに注意してくださいct。

gcount編集：カーネルを実行する前にゼロ化に問題があるようです。これを行うには、cudaMemcpyToSymbolまたはのようなものを使用する必要がcudaGetSymbolAddressありますcudaMemset。次のようになります。

const int zero = 0;
cudaMemcpyToSymbol("gcount", &zero, sizeof(int), 0, cudaMemcpyHostToDevice);

繰り返しになりますが、通常の免責事項：ブラウザで記述され、コンパイルまたはテストされたことがなく、自己責任で使用してください。

于 2012-04-23T20:40:07.390 に答える
1

あなたがしていることを行うためのより良い方法は、各スレッドに独自の出力を与え、それを独自にインクリメントしてcount値を入力させることです-このように、double-forループは任意の順序で並行して発生する可能性があります。これはGPUが行うことです良い。スレッドがout1配列を共有しているため、出力が間違っているため、すべてが上書きされます。

また、共有メモリにコピーするコードを、__syncthreads()後を付けて別のループに移動する必要があります。ループから__syncthreads()外れると、パフォーマンスが向上するはずです。つまり、共有配列はin2のサイズである必要があります。これが問題になる場合は、この回答の最後にこれに対処するためのより良い方法があります。

threadIdx.x < 32また、チェックを外側に移動する必要があります。したがって、コードは次のようになります。

if (threadIdx.x < 32) {
    for(int i = threadIdx.x; i < posdir*pos; i+=32) {
        t1[i] = in2[i];
    }
}
__syncthreads();

for(int i = threadIdx.x; i < posdir*pos; i += 32) {
    for(int j = 0; j < 32; j++)
    {
         ...
    }
}

次に、、__syncthreads()アトミック加算gcount += count、およびローカル出力配列からグローバル出力配列へのコピーを配置します。この部分はシーケンシャルであり、パフォーマンスを低下させます。可能であれば、各ローカル配列の配列へのポインターのグローバルリストを作成し、それらをCPUにまとめます。

もう1つの変更点は、t2に共有メモリが必要ないことです。これは役に立ちません。そして、これを行う方法では、単一のブロックを使用している場合にのみ機能するようです。ほとんどのNVIDIAGPUから優れたパフォーマンスを引き出すには、これを複数のブロックに分割する必要があります。これは、共有メモリの制約に合わせて調整できます。もちろん、__syncthreads()ブロック間にはないので、各ブロックのスレッドは、内側のループと外側のループのパーティションの全範囲を通過する必要があります。

于 2012-04-23T20:14:01.007 に答える