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

Question

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

memorycudaconstantsnvidiagpu-constant-memory

14

我希望了解一下在CUDA 4.2中如何分配常量内存。我知道总可用的常量内存为64KB，但是这个内存实际上是在设备上分配的吗？这个限制是针对每个核函数、CUDA上下文还是整个应用程序的？

假设一个.cu文件中有几个核函数，每个核函数使用的常量内存少于64K。但是总的常量内存使用量超过了64K。能够按顺序调用这些核函数吗？如果它们使用不同的流同时调用会发生什么情况？

如果有一个大型的CUDA动态库，其中包含许多核函数，每个核函数使用不同数量的常量内存，会发生什么情况？

如果有两个应用程序，每个应用程序需要超过可用常量内存的一半，第一个应用程序正常运行，但是第二个应用程序会在什么时候失败？在应用程序启动时，在cudaMemcpyToSymbol()调用时还是在核函数执行时？

- hthms

1

常量内存是上下文的属性，不是特定内核的属性。在较新的硬件上，内核在其参数列表之外不会“使用”常量内存，而该内存始终限制为最大4kb。 - talonmies

@talonmies... 常量内存不是64 KB吗？ - sgarizvi

1

@sgar91：是的，没错。但我并没有说反对。我所说的是，在Fermi/Kepler上，内核参数驻留在常量内存中，并且它们每个内核的最大限制为4kb。 - talonmies

这个问题和答案提供了关于常量内存的有趣信息。然而，它并没有说明当你尝试使用超过64KB时会发生什么。 - Roger Dahl

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Greg Smith · Accepted Answer

Parallel Thread Execution ISA Version 3.1 5.1.3讨论常量存储器。

目前，常量存储器的大小受限于64KB，可用于保存静态大小的常量变量。还有640KB的常量存储器，组织成十个独立的64KB区域。驱动程序可以在这些区域中分配和初始化常量缓冲区，并将指向缓冲区的指针作为内核函数参数传递。由于这十个区域不是连续的，因此驱动程序必须确保常量缓冲区分配到每个缓冲区都完全位于一个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进行编译，并执行cuobjdump -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;
            .....................................

我在SASS的右侧做了注释。

在sm30上，可以看到参数从0x140偏移开始传递给常量银行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以安全的方式更新。