4

定数メモリ (グローバル変数) に配列があり、関数呼び出し cudaGetSymbolAddress によってその配列への参照を取得しました。グローバル変数を使用するのではなく、この参照を使用して定数データをフェッチすると、カーネルの実行が遅くなります。これの理由は何ですか？

__constant__ int g[2] = {1,2};
// __device__ int g[2] = {1,2};

// kernel: use by reference
__global__ void add_1( int *a, int *b, int *c, int *f )
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = f[0] * a[tid] + f[1] * b[tid];
}

// kernel: use global variable
__global__ void add_2( int *a, int *b, int *c, int *f )
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = g[0] * a[tid] + f[1] * b[tid];
}

int main()
{
    ......
    // a,b,c are large arrays in device memory of size 40960.

    int *f;
    cudaGetSymbolAddress( (void **)&f, (char *)&g);

    add_1 <<< 160, 256 >>> ( a, b, c, f );

    ......
}

これはサンプル コードで、ワープ内のすべてのスレッドが同じ場所を同時にロードします。コメント化されたコードは、定数メモリに直接アクセスすることによるものです

定数メモリ キャッシュを使用しない理由の説明（talonmiesによる）

その理由は、定数キャッシュの欠如です。キャッシュされたアクセスは、定数状態空間にあると明示的にマークされた変数に対してコンパイラが特定の PTX 命令 (ld.const) を発行した場合にのみ発生します。そして、コンパイラがこれを行うことを認識する方法は、変数が宣言されたときです。これは、__constant__コード生成に影響を与える静的なコンパイル時の属性です。実行時に同じプロセスが発生することはありません。

グローバル メモリにポインターを渡し、そのポインターが定数状態空間にあるとコンパイラが判断できない場合、定数キャッシュを介してそのメモリにアクセスするための正しい PTX が生成されません。その結果、アクセスが遅くなります。

未回答の問題

配列gが変数として宣言されている場合でも__device__、それを参照するとコードが遅くなるのはなぜですか。PTXグローバルメモリをレジスタにロードするためのコードを見ると、次のようになります。

  • ld.global.s32レジスタに 4 バイトをロードする2 つの命令が使用されます。(参照を使用するコード内)
  • ld.global.v2.s328 バイトを 2 つのレジスタにロードする1 つの命令が使用されます (グローバル変数を使用するコード内)。

違いは何ですか。ドキュメントの参照をいただければ幸いです。

4

1 に答える 1

2

グローバル メモリとは異なり、定数メモリへのアクセスは、それらが均一でない場合 ((コンピューティング機能 1.x の半分) ワープのすべてのスレッドが同じアドレスにアクセスする場合、シリアル化されます (複数のトランザクションに分割されます)。

そのため、アクセスが均一である可能性が高い場合にのみ、定数メモリを使用してください。

于 2012-10-17T12:10:11.773 に答える