Address value changing

[AKKamath]

For the below code snippet, when instrumented using the mem_printf example, I notice some odd behaviour.

__device__ int dummy;
__global__ void test(int *tests)
{
    dummy = tests[0];
    tests[0] = 5;
}

int main()
{
    int *tests;
    cudaMalloc((void **)&tests, sizeof(int));
    test<<<1,1>>>(tests);
    cudaDeviceSynchronize();
}

The output I get is: OPCODE MOV MAPS TO ID 0 OPCODE LDG.E.SYS MAPS TO ID 1 OPCODE STG.E.SYS MAPS TO ID 2 0x00007f3d02e00000 - opcode_id 1 0x00007f3d06600300 - opcode_id 2 0x00007f3d22000200 - opcode_id 2

The order of operations is as expected, first, we load tests[0], then we store into dummy, then we store into tests[0].

However, notice the addresses these operations are mapped to. The first and third operation should have the same address, the address of tests[0], but they differ (0x00007f3d02e00000 and 0x00007f3d22000200).

I tested this for other programs and observed the same behavior; the first operation's address differs from subsequent operations for the same memory location.

[x-y-z]

Have you tried to use nvdisasm to dump the actual SASS code of your program and see if the SASS code matches your expectation?

Your dummy variable is written but not used, so it could be the case that compiler optimized the dummy = tests[0] away.

[ovilla]

Thanks for reporting this. Which architecture specifically?

We have recently discovered a problem in nvbit when printf is used in the injected device code on Turing and Ampere and we are evaluating possible solutions.

[AKKamath]

Yes, I'm using Turing as well, GTX 1660Ti specifically. I dumped the SASS and it matches expectations.

I'll try noting down the addresses in an array and dumping it at the end of the kernel, see if its specifically because of printf.

[x-y-z]

You can use mem_trace to do that. Also, could you upload the binary? Just to make sure we can debug the same binary as you have. Thanks.

[AKKamath]

address.zip I double-checked against mem_trace, it was indeed because of printf; mem_trace gives expected output. I've attached the binary and source file anyway, in case it helps.

[AKKamath]

This seems to be fixed in the latest version, NVBit 1.5.2. Closing.