3

多くのアルゴリズムは、特定の収束基準に達するまで繰り返されます（たとえば、特定のマトリックスの安定性）。多くの場合、反復ごとに1つのCUDAカーネルを起動する必要があります。私の質問は、最後のカーネル呼び出しの過程で行列が変更されたかどうかを効率的かつ正確に判断するにはどうすればよいでしょうか。同様に満足できないと思われる3つの可能性があります。

  • カーネル内でマトリックスが変更されるたびにグローバルフラグを書き込みます。これは機能しますが、非常に非効率的であり、技術的にスレッドセーフではありません。
  • アトミック操作を使用して上記と同じことを行います。繰り返しますが、最悪のシナリオではスレッドごとに1つのグローバル書き込みが発生するため、これは非効率的です。
  • 削減カーネルを使用して、行列のパラメーター（合計、平均、分散など）を計算します。これは場合によっては高速かもしれませんが、それでもやり過ぎのようです。また、行列が変更されたが、合計/平均/分散が変更されていない（たとえば、2つの要素が交換された）場合を想像することもできます。

上記の3つのオプションのいずれか、またはベストプラクティスと見なされる、および/または一般的により効率的な代替手段はありますか？

4

2 に答える 2

4

また、2012年に投稿したであろう回答に戻りますが、ブラウザーがクラッシュします。

基本的な考え方は、ワープ投票命令を使用して単純で安価な削減を実行し、ブロックごとにゼロまたは1つのアトミック操作を使用して、カーネルの起動ごとにホストが読み取ることができる固定されたマップ済みフラグを更新できることです。マップされたフラグを使用すると、カーネルの起動ごとに転送をホストする明示的なデバイスが不要になります。

これには、カーネルのワープごとに1ワードの共有メモリが必要です。これは小さなオーバーヘッドであり、テンプレートパラメータとしてブロックごとのワープの数を指定すると、いくつかのテンプレートトリックでループ展開が可能になります。

完全に機能する試験版（C ++ホストコードを使用しているため、現時点では機能するPyCUDAインストールにアクセスできません）は次のようになります。

#include <cstdlib>
#include <vector>
#include <algorithm>
#include <assert.h>

__device__ unsigned int process(int & val)
{
    return (++val < 10);
}

template<int nwarps>
__global__ void kernel(int *inout, unsigned int *kchanged)
{
    __shared__ int wchanged[nwarps];
    unsigned int laneid = threadIdx.x % warpSize;
    unsigned int warpid = threadIdx.x / warpSize;

    // Do calculations then check for change/convergence 
    // and set tchanged to be !=0 if required
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int tchanged = process(inout[idx]);

    // Simple blockwise reduction using voting primitives
    // increments kchanged is any thread in the block 
    // returned tchanged != 0
    tchanged = __any(tchanged != 0);
    if (laneid == 0) {
        wchanged[warpid] = tchanged;
    }
    __syncthreads();

    if (threadIdx.x == 0) {
        int bchanged = 0;
#pragma unroll
        for(int i=0; i<nwarps; i++) {
            bchanged |= wchanged[i];
        }
        if (bchanged) {
            atomicAdd(kchanged, 1);
        }
    }
}

int main(void)
{
    const int N = 2048;
    const int min = 5, max = 15;
    std::vector<int> data(N);
    for(int i=0; i<N; i++) {
        data[i] = min + (std::rand() % (int)(max - min + 1));
    }

    int* _data;
    size_t datasz = sizeof(int) * (size_t)N;
    cudaMalloc<int>(&_data, datasz);
    cudaMemcpy(_data, &data[0], datasz, cudaMemcpyHostToDevice);

    unsigned int *kchanged, *_kchanged;
    cudaHostAlloc((void **)&kchanged, sizeof(unsigned int), cudaHostAllocMapped);
    cudaHostGetDevicePointer((void **)&_kchanged, kchanged, 0);

    const int nwarps = 4;
    dim3 blcksz(32*nwarps), grdsz(16);

    // Loop while the kernel signals it needs to run again
    do {
        *kchanged = 0;
        kernel<nwarps><<<grdsz, blcksz>>>(_data, _kchanged);
        cudaDeviceSynchronize(); 
    } while (*kchanged != 0); 

    cudaMemcpy(&data[0], _data, datasz, cudaMemcpyDeviceToHost);
    cudaDeviceReset();

    int minval = *std::min_element(data.begin(), data.end());
    assert(minval == 10);

    return 0;
}

これkchangedは、カーネルがホストに対して再度実行する必要があることを通知するために使用するフラグです。カーネルは、入力の各エントリがしきい値を超えるまでインクリメントされるまで実行されます。各スレッドの処理の最後に、ワープ投票に参加します。その後、各ワープの1つのスレッドが、投票結果を共有メモリにロードします。1つのスレッドがワープの結果を減らしてから、kchanged値をアトミックに更新します。ホストスレッドは、デバイスが終了するまで待機し、マップされたホスト変数から結果を直接読み取ることができます。

アプリケーションが必要とするものにこれを適応させることができるはずです

于 2014-05-29T19:22:11.030 に答える
3

元の提案に戻ります。関連する質問を自分の答えで更新しましたが、これは正しいと思います。

グローバルメモリにフラグを作成します。

__device__ int flag;

各反復で、

  1. フラグをゼロに初期化します（ホストコード内）：

    int init_val = 0;
cudaMemcpyToSymbol(flag, &init_val, sizeof(int));

  2. カーネルデバイスコードで、マトリックスに変更が加えられた場合は、フラグを1に変更します。

    __global void iter_kernel(float *matrix){

...
  if (new_val[i] != matrix[i]){
    matrix[i] = new_val[i];
    flag = 1;}
...
}

  3. カーネルを呼び出した後、反復の最後に（ホストコードで）、変更をテストします。

    int modified = 0;
cudaMemcpyFromSymbol(&modified, flag, sizeof(int));
if (modified){
  ...
  }

別々のブロックまたは別々のグリッドにある複数のスレッドがflag値を書き込んでいる場合でも、同じ値（この場合は1）を書き込むことだけが行われている限り、危険はありません。書き込みは「失われる」ことはなく、スプリアス値はflag変数に表示されません。

この方法での同等性のテストfloatまたはdouble量には疑問がありますが、それはあなたの質問のポイントではないようです。「変更」を宣言するための好ましい方法がある場合は、代わりにそれを使用してください（おそらく、許容範囲内の同等性のテストなど）。

このメソッドの明らかな機能強化は、スレッドごとに1つの（ローカル）フラグ変数を作成し、各スレッドに、変更ごとではなく、カーネルごとに1回グローバルフラグ変数を更新させることです。これにより、カーネルごとのスレッドごとに最大で1つのグローバル書き込みが発生します。別のアプローチは、共有メモリ内のブロックごとに1つのフラグ変数を保持し、すべてのスレッドにその変数を更新させることです。ブロックの完了時に、グローバルフラグを更新するために（必要に応じて）グローバルメモリに1回の書き込みが行われます。この場合、カーネル全体でブール結果が1つしかないため、複雑な削減に頼る必要はありません。また、すべてのスレッドが同じ書き込みを行っている限り、共有変数またはグローバル変数への複数のスレッドの書き込みを許容できます。価値。

アトミックを使用する理由や、それがどのように役立つかはわかりません。

少なくとも最適化されたアプローチの1つ（たとえば、ブロックごとの共有フラグ）と比較すると、削減カーネルはやり過ぎのように見えます。また、CRCに満たないものや、同様に複雑な計算によって2つの異なるマトリックス結果が「同じ」とエイリアスされる可能性があるなど、前述の欠点があります。

于 2014-05-28T16:51:59.357 に答える