0

私はCUDAで（確かにローカルメモリを集中的に使用する）コードを書いており、開発段階を過ぎて加速しています。コマンド ライン プロファイラーは、私の占有率が (私が思うに) 非常に低い (主要なカーネルでは 0.083 - 0.417) ことを示しており、これを改善したいと考えています。残念ながら、計算作業には大量の__shared__メモリ (128 スレッド ブロックあたり 16 ～ 20 kB) といくつかのレジスタ (主要なカーネルでは 63 と報告されていますが、実際にこれだけ多く使用しているかどうかはわかりません...) が必要です。

しかし、私の本当の質問は、cmem[2]. 次に例を示します。

8 つのレジスタを使用、40 バイト cmem[0]、51584 バイト cmem[2]、368 バイト cmem[14]、4 バイト cmem[16]

すべてのカーネルがこの大量の を使用しcmem[2]ているようで、それが何であるかさえわかりません。通常の呼び出しでデバイスに大量のメモリを保存し、呼び出しを介しcudaMallocていくつかの double とポインタをメモリに保存します (ただし、50 kB にはほど遠いですが)、それだけです。したがって、私の質問は次のとおりです。これをすべて正確に使用するにはどうすればよいですか?カーネルの占有率を制限していますか?__constant__cudaMemcpyToSymbolcmem[2]

また、CUDA 4.2 および Ubuntu 10.04.4 64 ビットを搭載した GTX 550 Ti で実行しています。コードも MPI 並列化されているため、nvcc で作成された実行可能ファイルは mpirun でラップされています。

4

1 に答える 1

3

Fermiアーキテクチャでは、CUDAドライバーは定数変数の格納にcmem[2]を使用します。同じモジュール内のすべての関数は同じ定数を共有します。この一定のバンクのサイズは、理論上のSM占有率に影響を与えたり、起動オーバーヘッドを増加させたりすることはありません。最大サイズ（64KB）を超えると、コンパイラエラーが発生します。

CUDAバイナリユーティリティcuobjdumpを使用して、割り当てをデバッグできます。

次の定数を持つファイルsm20.cuがある場合

__constant__ float k_float_array[] = { 0.f, 1.f, 2.f, 3.f };
__constant__ double k_double_array[] = { 0.0, 1.0, 2.0, 3.0 };
__constant__ int k_int_array[] = { 0, 1, 2, 3 };

__global__ void empty_kernel(float* a)
{
    return;
}

あなたが実行することができます

cuobjdump.exe -elf sm20.cu.obj

実行可能ファイルを引数として使用することもできます。cuobjdump --helpを実行して、オプションのリストを取得します。

このコマンドは次の出力を生成します

Fatbin elf code:
================
arch = sm_20
code version = [1,5]
producer = cuda
host = windows
compile_size = 32bit
identifier = c:/dev/constant/sm20.cu

32bit elf: abi=5, sm=20, flags = 0x140114
Sections:
Index Offset   Size ES Align   Type   Flags Link     Info Name
    1     34     a6  0  1    STRTAB       0    0        0 .shstrtab
    2     da     e9  0  1    STRTAB       0    0        0 .strtab
    3    1c4     80 10  4    SYMTAB       0    2        6 .symtab
    4    244     18  0  4 CUDA_INFO       0    3        0 .nv.info
    5    25c     20  0  4 CUDA_INFO       0    3        8 .nv.info._Z12empty_kernelPf
    6    27c     24  0  4  PROGBITS       2    0        8 .nv.constant0._Z12empty_kernelPf
    7    2a0     40  0  8  PROGBITS       2    0        0 .nv.constant2
    8    2e0     10  0  4  PROGBITS       6    3  2000007 .text._Z12empty_kernelPf

elfセクション.nv.constant2には、cmem[2]の内容が含まれています。このサイズは0x40==64バイトで、私の期待に一致します。

.nv.constant2はセクションインデックス7です。

.section .strtab

.section .shstrtab

.section .symtab
 index     value     size      info    other  shndx    name
   0          0        0        0        0      0     (null)
   1          0        0        3        0      8     .text._Z12empty_kernelPf
   2          0        0        3        0      6     .nv.constant0._Z12empty_kernelPf
   3          0        0        3        0      7     .nv.constant2
   4          0       16        1        0      7     k_float_array
   5         16       32        1        0      7     k_double_array
   6         48       16        1        0      7     k_int_array
   7          0       16       12       10      8     _Z12empty_kernelPf

.symtabには、定義された3つの定数配列が含まれています。shndx==7のすべてのシンボルは.nv.constant2にあります

.nv.constant0._Z12empty_kernelPf
0x00000000  0x00000000  0x00000000  0x00000000  0x00000000
0x00000000  0x00000000  0x00000000  0x00000000

.nv.constant2
0x00000000  0x3f800000  0x40000000  0x40400000  0x00000000
0x00000000  0x00000000  0x3ff00000  0x00000000
0x40000000  0x00000000  0x40080000  0x00000000
0x00000001  0x00000002  0x00000003

.nv.constant2は、セクションのバイナリデータを定義します。これは、宣言された定数変数と一致します。定数が多い場合、.symtabセクションは各シンボルのオフセットとサイズを識別します。

// skipping .nv.info and .text

cuobjdumpは、PTXおよびSASSコードをダンプするために使用できます。PTXおよびSASSアセンブリを使用して、定数を使用しているカーネルを判別できます。

于 2012-10-17T22:52:17.063 に答える