memory - CUDA 常量内存分配是如何工作的？

Question

我想了解一下如何分配常量内存(使用 CUDA 4.2)。我知道总可用的常量内存是 64KB。但是这个内存是什么时候在设备上实际分配的呢？此限制适用于每个内核、cuda 上下文还是整个应用程序？

假设 .cu 文件中有多个内核，每个内核使用的常量内存都少于 64K。但总的恒定内存使用量超过 64K。是否可以按顺序调用这些内核？如果使用不同的流同时调用它们会发生什么？

如果有一个大型 CUDA 动态库，其中包含许多内核，每个内核都使用不同数量的常量内存？

如果有两个应用程序都需要一半以上的可用常量内存会怎样？第一个应用程序运行良好，但第二个应用程序何时会失败？在应用启动、cudaMemcpyToSymbol() 调用或内核执行时？

Answer 1

Parallel Thread Execution ISA Version 3.1第 5.1.3 节讨论常量库。

Constant memory is restricted in size, currently limited to 64KB whichcan be used to hold statically-sized constant variables. There is anadditional 640KB of constant memory, organized as ten independent 64KBregions. The driver may allocate and initialize constant buffers inthese regions and pass pointers to the buffers as kernel functionparameters. Since the ten regions are not contiguous, the drivermust ensure that constant buffers are allocated so that each bufferfits entirely within a 64KB region and does not span a regionboundary.

可以用一个简单的程序来说明常量内存的使用。

__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编译并执行cuobjdump -sass <executable or obj>拆卸你应该看到

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

我在 SASS 的右侧进行了注释。

在 sm30 上，您可以看到参数在从偏移量 0x140 开始的常量 bank 0 中传递。

用户定义__constant__变量在常量库 3 中定义。

如果你执行 cuobjdump --dump-elf <executable or obj>您可以找到其他有趣的常量信息。

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

每次启动都会对内核参数常量库进行版本控制，以便可以执行并发内核。编译器和用户常量是每个 CUmodule。开发人员有责任管理这些数据的一致性。例如，开发人员必须确保 cudaMemcpyToSymbol 以安全的方式更新。