nvbit_add_call_arg_mref_addr64 for LDGSTS

[mahmoodn]

Hi Regarding the following function

/* The 64-bit memory reference address accessed by this instruction
  Typically memory instructions have only 1 MREF so in general id = 0 */
void nvbit_add_call_arg_mref_addr64(const Instr* instr, int id = 0,
                                    bool is_variadic_arg = false);

I assume that it is saying for LDG and STS, it sould be nvbit_add_call_arg_mref_addr64(instr, 0);. Can you please explain what really happens when the id is 0? Because when I change it to 1, the traces LDG instruction has always 0x0 value in its argument list.

Also, what about LDGSTS? Should we call it two times?

nvbit_add_call_arg_mref_addr64(instr, 0);
nvbit_add_call_arg_mref_addr64(instr, 0);

Does that matter to push multiple zeros?

[ovilla]

In general void nvbit_add_call_arg_mref_addr64(const Instr* instr, int id = 0, bool is_variadic_arg = false) adds the 64 bit address used by the instruction as argument to the instrumentation call, "int id" is used to select which address to add (if the instruction uses more than 1 address).

if the instruction has no memory references, for instance an ADD, calling nvbit_add_call_arg_mref_addr64(instr, 0); will add some unspecified value, so it should not be done.

if the instruction has 1 memory reference, for instance LDG, STS, etc, calling nvbit_add_call_arg_mref_addr64(instr, 0); will add the one and only address used by the instruction. You can add it as many times as you want, but the "0" index address is always the same, so it will show up as the same value for the corresponding different instrumentation function parameters.

if the instruction has 2 memory references, like LDGSTS, calling nvbit_add_call_arg_mref_addr64(instr, 0); will pass the first and calling nvbit_add_call_arg_mref_addr64(instr, 1); will add the second.

lets' assume my instrumentation function is this:

__device__ void foo(uint64_t addr1, uint64_t addr0) {

}

if instr = ADD R3, R1, R2

nvbit_insert_call(instr, "foo", IPONT_BEFORE);
nvbit_add_call_arg_mref_addr64(instr, 0);  // addr1 of foo will have unspecified value
nvbit_add_call_arg_mref_addr64(instr, 1); // addr2 of foo will have unspecified value

if instr = LDG R2, [R3]

nvbit_insert_call(instr, "foo", IPONT_BEFORE);
nvbit_add_call_arg_mref_addr64(instr, 0);  // addr1 of foo will have the address in memory generated using R3
nvbit_add_call_arg_mref_addr64(instr, 1); // addr2 of foo will have unspecified value

if instr = LDGSTS [R2], [R3]

nvbit_insert_call(instr, "foo", IPONT_BEFORE);
nvbit_add_call_arg_mref_addr64(instr, 0);  // addr1 of foo will have the address in memory generated using R3
nvbit_add_call_arg_mref_addr64(instr, 1); // addr2 of foo will have the shared memory address in memory generated using R2

[mahmoodn]

Thanks a lot for the explanation. It is clear now.