0

さまざまなシナリオでスループットをテストして、いくつかのGPUをテストする必要があります。

これには、単純な64b乗算が含まれます。

__device__ void add(unsigned int *data, bool flag){
unsigned int index = threadIdx.x;
unsigned int result;

asm ("{\n\t"
     "add.cc.u32    %1, %1, %1;\n\t"
     "addc.u32  %0, 0, 0;\n\t"
     "}"
     : "=r"(result), "+r"(index): );

if(flag)
    data[threadIdx.x] = result;
}

64bモジュロ：

__device__ void mod(){
    asm ("{\n\t"
         ".reg .u64 t1;\n\t"
         "cvt.u64.u32   t1, %0;\n\t"
         "rem.u64       t1, t1, t1;\n\t" 
         "}"
         : : "r"(index));
}

および64bmul+ mod：

__device__ void mulmod
    asm ("{\n\t"
         ".reg .u64 t1;\n\t"
         ".reg .u64 t2;\n\t"
         "mul.wide.u32  t1, %0, %0;\n\t"
         "cvt.u64.u32   t2, %0;\n\t"
         "rem.u64       t1, t1, t2;\n\t"
         "}"
         : : "r"(index));
}

私の意図では、メモリアクセスはまったく役に立たないと思います。次に、スレッドのインデックス変数を入力として使用したいと思います。

そして、レジスターなしで書き込むので、レジスターの使用法を気にする必要はなく、可能な限り多くのスレッドを起動できます（各GPUで許可されます）

私が知りたいのですが：

  • これが適切な方法である場合

  • 最高のスループットを取得できるスレッド数を最大化する以外に、特定のスレッド構成はありますか？

4

1 に答える 1

2

最初の「サブ質問」への答えはノーです。これは適切な方法ではありません。作成した関数はコンパイラによって発行されないためです。

上記にリンクした質問の私の回答で詳細を確認できますが、短いバージョンでは、C コンパイラ レベルのデッド コードの最適化により、メモリに書き込まれる値に関与しないコードがすべて削除されます。したがって、これらの関数に値を返させる必要があり、コンパイラがデバイス関数の呼び出しが冗長であると推測してそれを排除できないような方法で戻り値を使用する必要があります。

さらに、SM ごとに十分な数のアクティブなワープを用意して、アーキテクチャ内のすべての命令スケジューリング レイテンシを償却し、命令スケジューラとパイプラインのレイテンシではなく、デバイス機能の命令スループットを確実に測定する必要があります。

于 2013-02-03T19:12:26.063 に答える