11

17 個のレジスタを使用するカーネルがあり、それを 16 個に減らすと占有率が 100% になります。私の質問は、別の方法でアルゴリズムを完全に書き直すことを除いて、使用される数またはレジスタを減らすために使用できる方法はありますか? 私は常に、コンパイラは私よりずっと賢いと思っていました。たとえば、わかりやすくするためだけに追加の変数を使用することがよくあります。この考え方は間違っていますか？

注: --max_registers (または構文が何であれ) フラグについては知っていますが、ローカル メモリの使用は、占有率が 25% 低下するよりも有害です (これをテストする必要があります)。

4

5 に答える 5

8

占有率は少し誤解を招く可能性があり、100% の占有率を主な目標にするべきではありません。グローバル メモリへの完全に結合されたアクセスを取得できる場合、ハイエンド GPU では 50% の占有率でグローバル メモリへのレイテンシを隠すのに十分です (float の場合、double の場合はさらに低くなります)。このトピックの詳細については、昨年の GTCのAdvanced CUDA Cプレゼンテーションをご覧ください。

あなたのケースでは、maxrregcount を 16 に設定した場合としない場合の両方でパフォーマンスを測定する必要があります。ローカル配列へのランダム アクセスを行わないと仮定すると、十分なスレッドがあるため、ローカル メモリへのレイテンシは隠されます (非合体アクセス）。

レジスターの削減に関する特定の質問に答えるには、より詳細な回答のコードを投稿してください! コンパイラが一般的にどのように機能するかを理解することは役立つかもしれませんが、nvcc は大きなパラメータ空間を持つ最適化コンパイラであるため、レジスタ数を最小限に抑えることは全体的なパフォーマンスとバランスを取る必要があることに注意してください。

于 2010-02-18T10:37:13.630 に答える
6

私の意見では、nvcc コンパイラはあまりスマートではありません。
たとえば、int の代わりに short を使用したり、参照によって変数を渡したり使用したり (eg&variable)、ループを展開したり、テンプレートを使用したり (C++ のように) など、明白なことを試すことができます。分割、超越関数が順番に適用されている場合は、それらをループとして作成してみてください。条件文を削除して、冗長な計算に置き換えるようにしてください。

コードを投稿すると、特定の回答が得られる場合があります。

于 2010-02-17T19:12:39.123 に答える
4

共有メモリをキャッシュとして利用すると、レジスタの使用量が減り、レジスタがローカル メモリに流出するのを防ぐことができます...

カーネルがいくつかの値を計算し、これらの計算された値がすべてのスレッドで使用されると考えてください。

__global__ void kernel(...) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int id0 = blockDim.x * blockIdx.x;

    int reg = id0 * ...;
    int reg0 = reg * a / x + y;


    ...

    int val =  reg + reg0 + 2 * idx;

    output[idx] = val > 10;
}

したがって、reg と reg0 をレジスタとして保持し、ローカル メモリ (グローバル メモリ) に流出させる可能性がある代わりに、共有メモリを使用することができます。

__global__ void kernel(...) {
    __shared__ int cache[10];

    int idx = threadIdx.x + blockDim.x * blockIdx.x;

    if (threadIdx.x == 0) {
      int id0 = blockDim.x * blockIdx.x;

      cache[0] = id0 * ...;
      cache[1] = cache[0] * a / x + y;
    }
    __syncthreads();


    ...

    int val =  cache[0] + cache[1] + 2 * idx;

    output[idx] = val > 10;
}

詳細については、このペーパーを参照してください。

于 2013-01-22T13:17:05.613 に答える
2

一般に、見当の圧力を最小限に抑えるのは適切な方法ではありません。コンパイラは、予測されるカーネル パフォーマンス全体を適切に最適化し、レジスタを含む多くの要因を考慮します。

レジスタを減らすと速度が遅くなった場合、どのように機能しますか

ほとんどの場合、コンパイラは不十分なレジスタ データを「ローカル」メモリにスピルする必要がありました。これは本質的にグローバル メモリと同じであるため、非常に遅くなります。

最適化の目的で、必要に応じて const や volatile などのキーワードを使用して、最適化フェーズでコンパイラを支援することをお勧めします。

いずれにせよ、CUDA カーネルの実行を遅くすることが多いのは、レジスタのような小さな問題ではありません。グローバル メモリ、アクセス パターン、可能であればテクスチャ メモリでのキャッシュ、PCIe 経由のトランザクションを使用して作業を最適化することをお勧めします。

于 2010-08-05T11:57:37.550 に答える
1

レジスタ使用量を下げると命令数が増える理由を簡単に説明します。コンパイラーは、レジスターを使用して、それらの値の再計算を避けるためにコードで複数回使用されるいくつかの操作の結果を格納している可能性があります。使用するレジスターを少なくする必要がある場合、コンパイラーは、レジスターに格納される値を再計算することを決定します。それ以外は。

于 2010-02-21T13:41:14.233 に答える