CUDA定数のメモリ割り当てはどのように機能しますか？

Question

私は一定のメモリがどのように割り当てられるかについていくつかの洞察を得たいと思います（CUDA 4.2を使用）。私は利用可能な合計定数メモリが64KBであることを知っています。しかし、このメモリが実際にデバイスに割り当てられるのはいつですか？この制限は各カーネル、Cudaコンテキスト、またはアプリケーション全体に適用されますか？

.cuファイルに複数のカーネルがあり、それぞれ64K未満の定数メモリを使用しているとします。しかし、全体の定数メモリ使用量は64K以上です。これらのカーネルを順番に呼び出すことは可能ですか？異なるストリームを使用して同時に呼び出されるとどうなりますか？

さまざまな量の定数メモリを使用する多数のカーネルを持つ大きなCUDA動的ライブラリがあるとどうなりますか？

利用可能な定数メモリの半分以上を必要とする2つのアプリケーションがあるとどうなりますか？最初のアプリケーションは正常に動作しますが、2番目のアプリケーションはいつ失敗しますか？アプリ起動時に、cudaMemcpyToSymbol()の呼び出し時またはカーネル実行時？

Answer 1

Parallel Thread Execution ISA Version 3.1セクション5.1.3では、定数バンクについて説明します。

固定メモリのサイズは現在制限されていますが、現時点では64KBに制限されており、固定サイズの定数を保持するために使用できます。 追加の640KBの定数メモリがあり、10個の独立した64KB 領域として編成されています。ドライバは、これらの領域の に定数バッファを割り当てて初期化し、カーネル関数 のパラメータとしてバッファへのポインタを渡すことができます。 10個の領域は連続していないため、ドライバ は、各バッファ が64KBの領域内に完全に収まり、領域 の境界にまたがらないように、定数バッファが割り当てられるようにする必要があります。

単純なプログラムを使用して定数メモリの使用を説明することができます。

__constant__ int kd_p1; 
__constant__ short kd_p2; 
__constant__ char kd_p3; 
__constant__ double kd_p4; 

__constant__ float kd_floats[8]; 

__global__ void parameters(int p1, short p2, char p3, double p4, int* pp1, short* pp2, char* pp3,  double* pp4) 
{ 
    *pp1 = p1; 
    *pp2 = p2; 
    *pp3 = p3; 
    *pp4 = p4; 
    return; 
} 

__global__ void constants(int* pp1, short* pp2, char* pp3, double* pp4) 
{ 
    *pp1 = kd_p1; 
    *pp2 = kd_p2; 
    *pp3 = kd_p3; 
    *pp4 = kd_p4; 
    return; 
} 

、compute_30のためにこれをコンパイルしsm_30、あなたは私がSASSの右側に注釈を付け

Fatbin elf code: 
================ 
arch = sm_30 
code version = [1,6] 
producer = cuda 
host = windows 
compile_size = 32bit 
identifier = c:/dev/constant_banks/kernel.cu 

    code for sm_30 
      Function : _Z10parametersiscdPiPsPcPd 
    /*0008*/  /*0x10005de428004001*/  MOV R1, c [0x0] [0x44];  // stack pointer 
    /*0010*/  /*0x40001de428004005*/  MOV R0, c [0x0] [0x150];  // pp1 
    /*0018*/  /*0x50009de428004005*/  MOV R2, c [0x0] [0x154];  // pp2 
    /*0020*/  /*0x0001dde428004005*/  MOV R7, c [0x0] [0x140];  // p1 
    /*0028*/  /*0x13f0dc4614000005*/  LDC.U16 R3, c [0x0] [0x144]; // p2 
    /*0030*/  /*0x60011de428004005*/  MOV R4, c [0x0] [0x158];  // pp3 
    /*0038*/  /*0x70019de428004005*/  MOV R6, c [0x0] [0x15c];  // pp4 
    /*0048*/  /*0x20021de428004005*/  MOV R8, c [0x0] [0x148];  // p4 
    /*0050*/  /*0x30025de428004005*/  MOV R9, c [0x0] [0x14c];  // p4 
    /*0058*/  /*0x1bf15c0614000005*/  LDC.U8 R5, c [0x0] [0x146]; // p3 
    /*0060*/  /*0x0001dc8590000000*/  ST [R0], R7;     // *pp1 = p1 
    /*0068*/  /*0x0020dc4590000000*/  ST.U16 [R2], R3;    // *pp2 = p2 
    /*0070*/  /*0x00415c0590000000*/  ST.U8 [R4], R5;    // *pp3 = p3 
    /*0078*/  /*0x00621ca590000000*/  ST.64 [R6], R8;    // *pp4 = p4 
    /*0088*/  /*0x00001de780000000*/  EXIT; 
    /*0090*/  /*0xe0001de74003ffff*/  BRA 0x90; 
    /*0098*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a8*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b8*/  /*0x00001de440000000*/  NOP CC.T; 
      ........................................... 


      Function : _Z9constantsPiPsPcPd 
    /*0008*/  /*0x10005de428004001*/  MOV R1, c [0x0] [0x44];  // stack pointer 
    /*0010*/  /*0x00001de428004005*/  MOV R0, c [0x0] [0x140];  // p1 
    /*0018*/  /*0x10009de428004005*/  MOV R2, c [0x0] [0x144];  // p2 
    /*0020*/  /*0x0001dde428004c00*/  MOV R7, c [0x3] [0x0];  // kd_p1 
    /*0028*/  /*0x13f0dc4614000c00*/  LDC.U16 R3, c [0x3] [0x4]; // kd_p2 
    /*0030*/  /*0x20011de428004005*/  MOV R4, c [0x0] [0x148];  // p3 
    /*0038*/  /*0x30019de428004005*/  MOV R6, c [0x0] [0x14c];  // p4 
    /*0048*/  /*0x20021de428004c00*/  MOV R8, c [0x3] [0x8];  // kd_p4 
    /*0050*/  /*0x30025de428004c00*/  MOV R9, c [0x3] [0xc];  // kd_p4 
    /*0058*/  /*0x1bf15c0614000c00*/  LDC.U8 R5, c [0x3] [0x6];  // kd_p3 
    /*0060*/  /*0x0001dc8590000000*/  ST [R0], R7; 
    /*0068*/  /*0x0020dc4590000000*/  ST.U16 [R2], R3; 
    /*0070*/  /*0x00415c0590000000*/  ST.U8 [R4], R5; 
    /*0078*/  /*0x00621ca590000000*/  ST.64 [R6], R8; 
    /*0088*/  /*0x00001de780000000*/  EXIT; 
    /*0090*/  /*0xe0001de74003ffff*/  BRA 0x90; 
    /*0098*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00a8*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b0*/  /*0x00001de440000000*/  NOP CC.T; 
    /*00b8*/  /*0x00001de440000000*/  NOP CC.T; 
      ..................................... 

が表示されるはず分解しcuobjdump -sass <executable or obj>を実行します。

sm30では、オフセット0x140から始まる定数バンク0でパラメータが渡されることがわかります。あなたが他の興味深い一定の情報を見つけることができcuobjdump --dump-elf <executable or obj>を実行した場合

ユーザー定義定数変数が一定のバンク3

で定義されています。

カーネルパラメータ定数バンクは、並列カーネルを実行できるように起動ごとにバージョン管理されます。コンパイラとユーザ定数はCUmodule単位です。このデータの一貫性を管理するのは開発者の責任です。たとえば、開発者は、cudaMemcpyToSymbolが安全に更新されていることを確認する必要があります。