2

Nvidia ドライバーを使用して小さな OpenCL ベンチマークを実行しています。私のカーネルは 1024 個のヒューズ乗算加算を実行し、結果を配列に格納します。

#define FLOPS_MACRO_1(x)    { (x) = (x) * 0.99f + 10.f; } // Multiply-add
#define FLOPS_MACRO_2(x)    { FLOPS_MACRO_1(x) FLOPS_MACRO_1(x) }
#define FLOPS_MACRO_4(x)    { FLOPS_MACRO_2(x) FLOPS_MACRO_2(x) }
#define FLOPS_MACRO_8(x)    { FLOPS_MACRO_4(x) FLOPS_MACRO_4(x) }
// more recursive macros ...
#define FLOPS_MACRO_1024(x) { FLOPS_MACRO_512(x) FLOPS_MACRO_512(x) }

__kernel void ocl_Kernel_FLOPS(int iNbElts, __global float *pf)
{
   for (unsigned i = get_global_id(0); i < iNbElts; i += get_global_size(0))  
   {
      float f = (float) i;
      FLOPS_MACRO_1024(f)
      pf[i] = f;
    }   
}

しかし、生成された PTX を見ると、次のように表示されます。

    .entry ocl_Kernel_FLOPS(
    .param .u32 ocl_Kernel_FLOPS_param_0,
    .param .u32 .ptr .global .align 4 ocl_Kernel_FLOPS_param_1
)
{
    .reg .f32   %f<1026>; // 1026 float registers !
    .reg .pred  %p<3>;
    .reg .s32   %r<19>;    

    ld.param.u32    %r1, [ocl_Kernel_FLOPS_param_0];
    // some more code unrelated to the problem
    // ...

BB1_1:
    and.b32     %r13, %r18, 65535;
    cvt.rn.f32.u32  %f1, %r13;
    fma.rn.f32  %f2, %f1, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f3, %f2, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f4, %f3, 0f3F7D70A4, 0f41200000;
    fma.rn.f32  %f5, %f4, 0f3F7D70A4, 0f41200000;
    // etc
    // ...

私が正しければ、PTX は1026個の float レジスタを使用して 1024 個の演算を実行し、2 つのレジスタのみを使用してすべての乗加算演算を実行できたとしても、レジスタを 2 回再利用することはありません。1026 は、(仕様によると) スレッドが持つことができるレジスタの最大数をはるかに超えているため、これはメモリの流出につながると思います。

それはコンパイラのバグですか、それとも完全に何かが欠けていますか?

Quadro K2000 GPU で nvcc バージョン 6.5 を使用しています。

編集

実際、私は仕様で何かを見逃していました:

「PTX は仮想レジスタをサポートしているため、コンパイラ フロントエンドが多数のレジスタ名を生成することは非常に一般的です。すべての名前の明示的な宣言を必要とするのではなく、PTX は、共通のプレフィックス文字列が追加された一連の変数を作成するための構文をサポートしています。整数サフィックス。たとえば、プログラムが %r0、%r1、...、%r99 という名前の .b32 変数を多数、たとえば 100 個使用しているとします。

4

1 に答える 1

5

PTX ファイル形式は、仮想マシンと命令セットのアーキテクチャを記述することを目的としています。

PTX は、汎用の並列スレッド実行用に仮想マシンと ISA を定義します。PTX プログラムは、インストール時にターゲット ハードウェアの命令セットに変換されます。PTX-to-GPU トランスレータおよびドライバにより、NVIDIA GPU をプログラム可能な並列コンピュータとして使用できます。

そのため、取得している PTX 出力は「GPU アセンブラー」の形式ではありません。これは、事実上あらゆる形式の並列計算を記述できるようにすることを目的とした中間表現にすぎません。

次に、PTX 表現は、それぞれのターゲット GPU の実際のバイナリにコンパイルされます。これは、実際のアーキテクチャから抽象化できるようにするために重要です-具体的には、例に関して:特定のターゲットマシンで使用可能なレジスタの数に関係なく、プログラムの同じPTX 表現を使用できるはずです。表示される 1026 個の「レジスタ」は「仮想」レジスタであり、最終的には、実際に使用可能な (少数の) 実際のハードウェア レジスタにマップされる可能性があります。--ptxas-options=-vコンパイル中に引数を NVCC に追加して、レジスタの使用に関する追加情報を取得できます。

(これは、 LLVMの背後にあるものとほぼ同じ考え方です。つまり、元のソース コードと実際のターゲット アーキテクチャの両方から抽象化して、最適化および議論できる表現を持つことです)。

于 2015-02-27T14:47:47.253 に答える