0

多くの操作と少数の分岐がある CUDA カーネルがあります。のように見えます

__global__
void kernel(Real *randomValues, Real mu, Real sigma)
{
    int row = blockDim.y * blockIdx.y + threadIdx.y;
    int col = blockDim.x * blockIdx.x + threadIdx.x;

    if ( row >= cnTimeSteps || col >= cnPaths ) return;

    Real alphaLevel = randomValues[row*cnPaths+col];
    Real q = 0.0;
    Real x = 0.0;

    if ( alphaLevel < p_low)
    {
        q = sqrt( -2*log( alphaLevel ) );
        x = (((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6) / ((((d1*q+d2)*q+d3)*q+d4)*q+1);
    }
    else if ( alphaLevel < p_high )
    {
        q = alphaLevel-0.5;
        Real r = q*q;
        x= (((((a1*r+a2)*r+a3)*r+a4)*r+a5)*r+a6)*q / (((((b1*r+b2)*r+b3)*r+b4)*r+b5)*r+1);
    }
    else
    {
        q = sqrt( -2*log( 1.0-alphaLevel ) );
        x = -(((((c1*q+c2)*q+c3)*q+c4)*q+c5)*q+c6) / ((((d1*q+d2)*q+d3)*q+d4)*q+1);
    }

    randomValues[row*cnPaths+col] = sigma * x + mu;
}

ここで、a's、b's、c's、d's はすべて定数値です (デバイス定数メモリ内)。

static __device__ __constant__ Real a1 = 1.73687;
static __device__ __constant__ Real a2 = 1.12321100;

等々。

カーネルのプロファイリングを行った後、理論上の占有率は 100% であることがわかりましたが、60% を超えていません。

カーネルを最適化するために、thisとthis GTC talks を実行しました。

一方では、IPC が平均 1.32 の発行された命令と 0.62 の実行された命令を報告していることを確認しました。命令のシリアル化は約 50% ですが、SM アクティビティはほぼ 100% です。一方、約 38 のアクティブなワープがありますが、8 つは次の命令を実行する資格がありますが、ワープ発行効率では、サイクルの約 70% に適格なワープがないことがわかります。logストールの理由は「その他」として報告されますが、これは と の計算に関係していると思いますsqrt。

  1. ほとんどのサイクルに適格なワープがない場合、SM アクティビティが 99.82% になるにはどうすればよいでしょうか?
  2. ストールを減らすにはどうすればよいですか？
  3. ワープ内のスレッドは同じブランチに入らない可能性があるため、定数メモリへの要求はおそらくシリアライズされますが、これは本当ですか? これらの定数をグローバル メモリに配置する必要がありますか (共有メモリも使用する可能性があります)。

Nsight Visual Studio を使用するのは初めてなので、すべてのパフォーマンス分析の意味を理解しようとしています。ところで、私のカードは Quadro K4000 です。

4

3 に答える 3

3

1) ほとんどのサイクルに適格なワープがない場合、SM アクティビティが 99.82% になるにはどうすればよいですか?

レジスタとワープ スロットがワープに割り当てられている場合、ワープはアクティブです。SM で少なくとも 1 つのワープがアクティブな場合、SM はアクティブです。

SM アクティビティを効率と混同しないでください。

2) どうすれば失速を減らすことができますか?

上記のコードの場合、倍精度実行ユニットが使用可能になるのを待ってワープが停止します。Quadro K4000 は、倍精度操作で 8 スレッド/サイクルのスループットを実現します。

この問題の解決策は次のとおりです。倍精度演算の数を減らします。たとえば、単精度浮動小数点のスループットは倍精度スループットの 24 倍であるため、連続演算を float に移動すると、パフォーマンスが大幅に向上する可能性があります。b. GK10x の 8 倍の倍精度スループットを持つ GK110 でカーネルを実行します。

達成占有率を増やしても、K4000 でのこのカーネルのパフォーマンスは向上しない場合があります。提供された情報が不十分で、達成された占有率が理論上の占有率よりも​​大幅に少ない理由を判断できません。

達成された FLOPs の実験を使用して、カーネルのパフォーマンスが倍精度のスループットによって制限されているかどうかを確認できます。

3) ワープ内のスレッドは同じブランチに入らない可能性があるため、定数メモリへの要求はおそらくシリアライズされますが、これは本当ですか? これらの定数をグローバル メモリに配置する必要がありますか (共有メモリも使用する可能性があります)。

コードには、一定のメモリ ロードでメモリ アドレスの相違はありません。ワープ制御フローの分岐は、各リクエストでスレッドの一部がアクティブになることを意味します。

初期グローバル ロードは結合されない場合があります。誰かがレビューできるように、cnPaths の値を提供する必要があります。また、メモリの実験やソース相関の実験も参照できます。

if ステートメントと else ステートメントをより効率的な方法でコーディングして、コンパイラーが分岐分岐の代わりに述語を使用できるようにすることができます。

于 2013-08-28T19:15:44.763 に答える
1

あなたの Real データ型は float の typedef だと思います。使用される定数値に f サフィックスを追加して、コンパイラが不要なキャストを追加するのを防ぐことができます。

例えば

q = alphaLevel-0.5;

定数 0.5 は double 値で、alphaLevel は real=float 値です。alphaLevel は double にキャストされます。q は float 型です。減算の結果は、再びフロートにダウンキャストする必要があります。

Real が dobule の typedef である場合、すべての計算で double と float が混在し、同じアップ キャストとダウン キャストになります。

于 2013-08-28T10:30:34.483 に答える