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

Question

I have been trying to debug cuda programs that use inline PTX assembly. Specifically, I am debugging at the instruction level, and am trying to determine the values of arguments to the instructions. Occasionally, the disassembly includes a reference to constant memory. I am trying to have gdb print the value of this constant memory, but have not found any documentation that shows how to do this. For instance, a disassembly includes IADD R0, R0, c[0x0] [0x148]

I want to determine how to have gdb print the value of c[0x0] [0x148]. I have tried using print * (@constant) ... but this does not seem to work (I pass 0x148 here and it prints out nothing). Is this possible to do in cuda-gdb?

I have tried to avoid this by passing the compiler option --disable-optimizer-constants during compilation, but this does not work.

You can use `cuobjdump --dumpelf` and look a the constant sections (`.nv.constant0.*` for `c[0x0]` and `.nv.constant2.*` for `c[0x2]`. — njuffa, Apr 01 '21 at 21:31
@njuffa The constant bank can change on kernel launch, for instance, to pass parameters in. What I am trying to do is create a script that logs the values of operands to instructions before they are executed. The elf file won't help me with this. — Benjie, Apr 01 '21 at 22:47
If I knew how to do what you want to in `gdb`, I would have written answer. *One* of the constant banks is used to pass kernel arguments, correct. And you can copy data to `__constant__` data, which uses a different constant bank. Most apps do not update `__constant__` data. The compiler puts literal constant into yet another constant bank. Looking at the `elf` sections lets you at least look at all constant bank data that does not change dynamically, which should be better than no visibility at all. — njuffa, Apr 01 '21 at 23:09
Since those places are constant, can't you just print the values before the point where you're stepping through individual instructions? — einpoklum, Apr 02 '21 at 22:13
@einpoklum I need to be able to get an on-the-fly mapping between the constant bank references and the system variables before hand. It's not clear how to access these mappings from within gdb. — Benjie, Apr 05 '21 at 22:26
Why do you need any on-the-fly mappings? If you're not sure which constant will get used, print a bunch of them (well, unless the printing affects what gets used later on). — einpoklum, Apr 06 '21 at 19:31

Benjie · Accepted Answer · 2021-08-25T19:56:42.933

The way to do this is to print *(void * @parameter *) addr

where addr is the address inside the constant bank 0 that should be printed.

Example

Suppose we have a simple kernel in a file called 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;
}

which is compiled via

$ nvcc foo.cu -o foo.out

Next, suppose we are interested in disassembling this program, so we execute cuda-gdb with a command-line for our program:

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

Inside cuda-gdb, we get to the kernel by typing

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

(cuda-gdb) cont

Inside the kernel, we view the disassembly we are interested in debugging:

(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

The second argument being passed to the IADD instruction is in one of the constant memory banks. Let's find out what its value actually is. We advance go to the IADD instruction:

(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]

We can now obtain the contents of c[0x0][0x140] as follows:

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

Here, we knew the argument should have 32 bits, so we cast it as an (32-bit) int. If we hadn't done this, we would get too many bits, e.g.:

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

Note the hexadecimal format can be retained by adding /x after the print command:

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

I don't understand what you're suggesting here. Remember, even though you're answering yourself - you have to target your answer at other people who don't have the same context as you. — einpoklum, Aug 25 '21 at 15:08
What do you mean? If an instruction references c[0x0][0x140] then to print the value being passed to the instruction, you use print *(void * @parameter *) 0x140 — Benjie, Aug 25 '21 at 15:13
Please provide a concrete example with PTX instruction and the sequence of commands and outputs with cuda-gdb. — einpoklum, Aug 25 '21 at 15:20
@einpoklum I've edited the answer with a concrete example and the sequence of commands and outputs in cuda-gdb. Does it look better as an answer? — Benjie, Aug 25 '21 at 15:58
Yes. I've tweaked it a little. However, the example is a much better answer than the first few paragraphs. I would retain perhaps just the second line of the answer and say the example explains what I mean. If you like, I could make that edit. Anyway, now you have two upvotes from me. — einpoklum, Aug 25 '21 at 16:25

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

1 Answers1

Example

Linked