有没有办法在CUDA中访问常量内存库的值

Question

我一直在尝试调试使用内联 PTX 程序集的 cuda 程序。具体来说，我在指令级别进行调试，并试图确定指令的参数值。有时，反汇编包括对常量内存的引用。我试图让 gdb 打印这个常量内存的值，但没有找到任何说明如何执行此操作的文档。例如，反汇编包括 IADD R0, R0, c[0x0] [0x148]

我想确定如何让 gdb 打印 c[0x0] [0x148] 的值。我试过使用 print * (@constant) ... 但这似乎不起作用（我在这里传递 0x148 并且它没有打印出任何内容）。这可以在 cuda-gdb 中实现吗？

我试图通过在编译期间传递编译器选项 --disable-optimizer-constants 来避免这种情况，但这不起作用。

Answer 1

这样做的方法是 print *(void * @parameter *) 地址

其中 addr 是应该打印的常量组 0 中的地址。

例子

假设我们在名为 foo.cu:

的文件中有一个简单的内核

#include <cuda.h>
#include <stdio.h>
#include <cuda_runtime.h>

__global__ void myKernel(int a, int b, int *d)
{
    *d = a + b;
}

int main(int argc, char *argv[]) {
   if (argc < 3) {
       printf("Requires inputs a and b to be specified\n");
       return 0;
   }

   int * dev_d;
   int d;
   cudaMalloc(&dev_d, sizeof(*dev_d));
   myKernel<<<1, 1>>>(atoi(argv[1]), atoi(argv[2]), dev_d);
   cudaMemcpy(&d, dev_d, sizeof(d), cudaMemcpyDeviceToHost);
   cudaFree(dev_d);
   printf("D is: %d\n", d);
   return 0;
}

通过

编译

$ nvcc foo.cu -o foo.out

接下来，假设我们有兴趣反汇编这个程序，所以我们用命令行为我们的程序执行cuda-gdb：

$ cuda-gdb --args ./foo.out 10 15

在 cuda-gdb 中，我们通过键入

进入内核

(cuda-gdb) set cuda break_on_launch application
(cuda-gdb) start
Temporary breakpoint 1, 0x000055555555b12a in main ()

(cuda-gdb) cont

内核内部，查看我们感兴趣调试的反汇编：

(cuda-gdb) x/15i $pc
=> 0x555555b790a8 <_Z8myKerneliiPi+8>:  MOV R1, c[0x0][0x20]
   0x555555b790b0 <_Z8myKerneliiPi+16>: MOV R0, c[0x0][0x144]
   0x555555b790b8 <_Z8myKerneliiPi+24>: MOV R2, c[0x0][0x148]
   0x555555b790c0 <_Z8myKerneliiPi+32>:
   0x555555b790c8 <_Z8myKerneliiPi+40>: MOV R3, c[0x0][0x14c]
   0x555555b790d0 <_Z8myKerneliiPi+48>: IADD R0, R0, c[0x0][0x140]
   0x555555b790d8 <_Z8myKerneliiPi+56>: STG.E [R2], R0
   0x555555b790e0 <_Z8myKerneliiPi+64>:
   0x555555b790e8 <_Z8myKerneliiPi+72>: NOP
   0x555555b790f0 <_Z8myKerneliiPi+80>: NOP
   0x555555b790f8 <_Z8myKerneliiPi+88>: NOP
   0x555555b79100 <_Z8myKerneliiPi+96>:
   0x555555b79108 <_Z8myKerneliiPi+104>:        EXIT
   0x555555b79110 <_Z8myKerneliiPi+112>:        BRA 0x70
   0x555555b79118 <_Z8myKerneliiPi+120>:        NOP

传递给 IADD 指令的第二个参数在常量内存库之一中。让我们看看它的实际价值是多少。我们前进到 IADD 指令：

(cuda-gdb) stepi 4
0x0000555555b790d0 in myKernel(int, int, int*)<<<(1,1,1),(1,1,1)>>> ()
(cuda-gdb) x/i $pc
=> 0x555555b790d0 <_Z8myKerneliiPi+48>: IADD R0, R0, c[0x0][0x140]

我们现在可以得到c[0x0][0x140]的内容如下：

(cuda-gdb) print (int) *(void * @parameter *) 0x140
 = 10

在这里，我们知道参数应该有 32 位，所以我们将其转换为（32 位）int。如果我们不这样做，我们会得到太多的位，例如：

(cuda-gdb) print *(void * @parameter *) 0x140
 = 0xf0000000a

注意在print命令后加/x可以保留十六进制格式：

(cuda-gdb) print/x (int) *(void * @parameter *)0x140
 = 0xa

有没有办法在CUDA中访问常量内存库的值

Is there a way to access value of constant memory bank in CUDA

cuda

cuda-gdb

ptx

例子