0

数日前、私は同じアルゴリズムの非常に単純な置換と推力の実装を実行する私のコードのパフォーマンスを比較していました。Thrustを支持して1桁（！）の不一致を発見したので、デバッガーをコードに「サーフ」させて、魔法がどこで発生するかを発見し始めました。

驚いたことに、ファンクターのすべてを取り除き、核心に迫った後、私の非常に単純な実装が実際には彼らの実装と非常に似ていることを発見しました。Thrustには、ブロックサイズとグリッドサイズの両方を決定する賢い方法があることがわかりました（ところで、正確には、どのように機能するのですか？！）。そのため、それらの設定を取得して、コードを再度実行しました。私は数マイクロ秒を得ましたが、ほとんど同じ状況です。そして、結局、理由はわかりませんが、「試してみる」ために、カーネルとBINGOの後にcudaThreadSynchronize（）を削除しました。私はギャップをゼロにし（そしてそれ以上）、実行時間の桁違いを獲得しました。配列の値にアクセスすると、期待どおりの配列であることがわかりました。正しく実行されました。

ここでの質問は次のとおりです。cudaThreadSynchronize（et similia）をいつ削除できますか？なぜそれがそのような巨大なオーバーヘッドを引き起こすのですか？スラスト自体が最後に同期しないことがわかります（マクロ__THRUST_SYNCHRONOUSが定義されておらず、定義されていない場合、NOPであるsynchronize_if_enabled（const char *メッセージ））。詳細とコードは次のとおりです。

// my replace code
template <typename T>
__global__ void replaceSimple(T* dev, const int n, const T oldval, const T newval)
{
    const int gridSize = blockDim.x * gridDim.x;
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    while(index < n)
    {
        if(dev[index] == oldval)
            dev[index] = newval;
        index += gridSize;
    }
}

// replace invocation - not in main because of cpp - cu separation
template <typename T>
void callReplaceSimple(T* dev, const int n, const T oldval, const T newval)
{       
    replaceSimple<<<30,768,0>>>(dev,n,oldval,newval);
    cudaThreadSynchronize();
}

// thrust replace invocation
template <typename T>
void callReplace(thrust::device_vector<T>& dev, const T oldval, const T newval)
{
    thrust::replace(dev.begin(), dev.end(), oldval, newval);
}

パラメータの詳細：配列：n = 10,000,000要素を2に設定、oldval = 2、newval = 3

  • 推力呼び出しを実行する時間置換（推力）：0.057ミリ秒
  • 同期を使用してcallReplaceSimpleを実行する時間：0.662ミリ秒
  • 同期なしでcallReplaceSimpleを実行する時間：0.011ミリ秒

スラストを含むCUDA5.0を使用しました。カードは、GeForce GTX 570で、クアッドコアQ9550 2.83 GHz、2GBRAMを搭載しています。

4

1 に答える 1

6

カーネルの起動は非同期です。呼び出しを削除するcudaThreadSynchronize()と、カーネルの起動時間のみが測定され、カーネルが完了するまでの時間は測定されません。

于 2013-03-26T01:38:43.747 に答える