How to use shared memory in the inject device function? about nvbit HOT 5 CLOSED

Ohh, I just found this limitation in the paper.

Injected functions may not use shared and constant memory because that memory can be used by the application itself.

But in my inject_kernel, I just want to read the shared memory of the origin kernel instead of modifying it. That is also not possible?

from nvbit.

We currently don't have an interface for reading shared memory within an injected function.
We will keep it in mind for next version, but unfortunately we will not be able to explore this in short time as it requires non trivial amount of plumbing.

from nvbit.

Can I pass the c[0x0][0x18] and c[0x0][0x1c] as the base address of the shared memory? (BTW, I'm using GV100).

It works for me in this simple test:

extern "C" __device__ __noinline__ void inject_kernel(unsigned int r2, unsigned int r3) {
    unsigned long x = ((unsigned long)r3 << 32) | r2;
    long *shared_mem_ptr = (long *)x;
    printf("%ld\n", *shared_mem_ptr);
}

from nvbit.

We never tried that explicitly, but it seems promising.

From the earlier post I thought you were looking for an API to dump the entire content of shmem (without knowing a particular address), but if you know the exact address I guess you can do the way you are doing above... it should be safe.

from nvbit.

Thanks for your explanation.

from nvbit.