Come funziona l'allocazione di memoria costante CUDA?

Question

Mi piacerebbe avere un'idea di come viene allocata la memoria costante (utilizzando CUDA 4.2). So che la memoria costante disponibile totale è 64 KB. Ma quando questa memoria è effettivamente allocata sul dispositivo? Questo limite si applica a ogni kernel, contesto di cuda o per l'intera applicazione?

Diciamo che ci sono diversi kernel in un file .cu, ognuno con meno di 64 K di memoria costante. Ma l'utilizzo di memoria costante totale è più di 64 KB. È possibile chiamare questi kernel in modo sequenziale? Cosa succede se vengono chiamati contemporaneamente utilizzando flussi diversi?

Cosa succede se c'è una grande libreria dinamica CUDA con un sacco di kernel che utilizzano diverse quantità di memoria costante?

Cosa succede se ci sono due applicazioni ognuna delle quali richiede più della metà della memoria costante disponibile? La prima applicazione funziona bene, ma quando la seconda app fallirà? All'inizio dell'app, alle chiamate cudaMemcpyToSymbol() o all'esecuzione del kernel?

Answer 1

Parallel Thread Execution ISA Version 3.1 la sezione 5.1.3 tratta le banche costanti.

La memoria costante è di dimensioni limitate, attualmente limitata a 64 KB che è possibile utilizzare per contenere variabili costanti di dimensioni statiche. Vi è un 640KB aggiuntivi di memoria costante, organizzati come dieci regioni indipendenti da 64 KB . Il driver può allocare e inizializzare i buffer costanti in queste regioni e passare i puntatori ai buffer come parametri della funzione kernel . Poiché le dieci regioni non sono contigue, il driver deve garantire che i buffer costanti siano allocati in modo che ciascun buffer si adatti interamente a una regione di 64 KB e non contenga un limite di regione .

Un semplice programma può essere utilizzato per illustrare l'uso della memoria costante.

__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; 
} 

compilare questo per compute_30, sm_30 ed eseguire cuobjdump -sass <executable or obj> di smontare si dovrebbe vedere

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; 
      ..................................... 

ho annotato alla destra del Sass.

Su sm30 è possibile vedere che i parametri vengono passati nel banco costante 0 a partire dall'offset 0x140.

definiti dall'utente costanti variabili sono definite in costante banca 3.

Se si esegue cuobjdump --dump-elf <executable or obj> si possono trovare altre informazioni costante interessante.

32bit elf: abi=6, sm=30, flags = 0x1e011e 
Sections: 
Index Offset Size ES Align Type Flags Link  Info Name 
    1  34 142 0 1 STRTAB  0 0  0 .shstrtab 
    2 176 19b 0 1 STRTAB  0 0  0 .strtab 
    3 314  d0 10 4 SYMTAB  0 2  a .symtab 
    4 3e4  50 0 4 CUDA_INFO  0 3  b .nv.info._Z9constantsPiPsPcPd 
    5 434  30 0 4 CUDA_INFO  0 3  0 .nv.info 
    6 464  90 0 4 CUDA_INFO  0 3  a .nv.info._Z10parametersiscdPiPsPcPd 
    7 4f4 160 0 4 PROGBITS  2 0  a .nv.constant0._Z10parametersiscdPiPsPcPd 
    8 654 150 0 4 PROGBITS  2 0  b .nv.constant0._Z9constantsPiPsPcPd 
    9 7a8  30 0 8 PROGBITS  2 0  0 .nv.constant3 
    a 7d8  c0 0 4 PROGBITS  6 3 a00000b .text._Z10parametersiscdPiPsPcPd 
    b 898  c0 0 4 PROGBITS  6 3 a00000c .text._Z9constantsPiPsPcPd 

.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  a  .text._Z10parametersiscdPiPsPcPd 
    2   0  0  3  0  7  .nv.constant0._Z10parametersiscdPiPsPcPd 
    3   0  0  3  0  b  .text._Z9constantsPiPsPcPd 
    4   0  0  3  0  8  .nv.constant0._Z9constantsPiPsPcPd 
    5   0  0  3  0  9  .nv.constant3 
    6   0  4  1  0  9  kd_p1 
    7   4  2  1  0  9  kd_p2 
    8   6  1  1  0  9  kd_p3 
    9   8  8  1  0  9  kd_p4 
    10   16  32  1  0  9  kd_floats 
    11   0  192  12  10  a  _Z10parametersiscdPiPsPcPd 
    12   0  192  12  10  b  _Z9constantsPiPsPcPd 

La banca di parametri del kernel è versionata per avvio in modo da poter eseguire i kernel contemporanei. Il costitutore e le costanti dell'utente sono per modulo CUmodule. È responsabilità dello sviluppatore gestire la coerenza di questi dati. Ad esempio, lo sviluppatore deve assicurarsi che un cudaMemcpyToSymbol sia aggiornato in modo sicuro.