CUDA 代码中意外的大 cmem[2] 使用

Question

我正在 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 调用在设备上存储了大量内存，并通过cudaMemcpyToSymbol 调用在__constant__ 内存中存储了一些双精度值和指针（虽然远不及50 kB），但仅此而已。所以我的问题是：我究竟如何才能使用所有这些cmem[2]，它是否限制了我的内核占用？

另外，我正在使用 CUDA 4.2 和 Ubuntu 10.04.4 64 位的 GTX 550 Ti 上运行。 nvcc 生成的可执行文件包含在 mpirun 中，因为代码也是 MPI 并行化的。

Answer 1

在 Fermi 架构上，CUDA 驱动程序使用 cmem[2] 来存储 constant 变量。同一模块中的所有函数共享相同的常量。这个常量库的大小不会影响您的理论 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 程序集可用于确定哪些内核正在使用常量。