The behavior of nvbit_get_related_function seems to have changed

[keram88]

Hello,

I noticed a possible issue with the nvbit_get_related_functions function in the 1.7.6 release. It seems to return other kernels as “related functions,” which I didn’t expect.

Here’s a minimal example:

...
__global__ void fp64_nan_to_mem(double a, double b, double* res) {
    res[threadIdx.x] = a+b;
}

__global__ void fp32_nan_to_mem(float a, float b, float* res) {
    res[threadIdx.x] = a+b;
}
...
int main() {
    ...
    fp64_nan_to_mem<<<1, 2*warpSize>>>(INFINITY, -INFINITY, d_res);
    ...
    fp32_nan_to_mem<<<1, 2*warpSize>>>(INFINITY, -INFINITY, d_res32);
    ...
}

I wrote a small NVBit tool (test_related_functions.tgz) to print the related functions of each kernel launch. With version 1.7.6, I see:

$ LD_PRELOAD=./test_related_functions.so test/test
...
Instrumenting kernel fp64_nan_to_mem(double, double, double*)
Instrumenting related function fp64_nan_to_mem(double, double, double*)
Instrumenting related function fp32_nan_to_mem(float, float, float*)
Instrumenting kernel fp32_nan_to_mem(float, float, float*)
Instrumenting related function fp32_nan_to_mem(float, float, float*)
Instrumenting related function fp64_nan_to_mem(double, double, double*)

With version 1.7.5, the output is:

$ LD_PRELOAD=./test_related_functions.so test/test
...
Instrumenting kernel fp64_nan_to_mem(double, double, double*)
Instrumenting kernel fp32_nan_to_mem(float, float, float*)

The 1.7.5 behavior makes sense to me since the kernels don’t call any other functions. Could you clarify what the expected behavior should be in 1.7.6?

Thanks,
Mark Baranowski

[x-y-z]

Which CUDA version are you using?

I got

$ LD_PRELOAD=./test_related_functions.so test/test
...
Instrumenting kernel fp64_nan_to_mem(double, double, double*)
Instrumenting kernel fp32_nan_to_mem(float, float, float*)

on CUDA 13 using your example.

There is no behavior change from 1.7.5 to 1.7.6 in terms of related function.

[keram88]

I am currently working with CUDA 12.8, driver version 570.172.08 and NVCC version 12.3. My GPUs are an RTX 3090 and an RTX 3060. Thanks for looking into this.

[x-y-z]

Can you upload your test.cu binary? Thanks.

[keram88]

Here are my compiled binaries, including my test program:

trf_binaries.tgz

[x-y-z]

It seems that .nv.callgraph section is in a different cubin than fp*_nan_to_mem's cubin. NVBit, in this case, assumes every function calls each other to avoid missing related functions. I will need to figure out how to read .nv.callgraph in this case.

The behavior change you see is due to a bugfix. 1.7.6 fixed an issue in related function callgraph.

I wonder if you can compile your binary with a single GPU arch to fix the issue, since that is what I did in #158 (comment).

[keram88]

I get the same behavior when I compile for sm_86, which matches my hardware. When I compile for sm_80, I also get the following message:
cuobjdump info : No ELF file found to extract from '/tmp/nvbit_tool_tmpdir.9MWWSQ/test_related_functions.so'. You may try with -all option

[x-y-z]

What I meant was to recompile test.cu with just sm_86 not your tool. Sorry for the confusion. If that still fails, I will need more time to work on NVBit side to fix this.

[x-y-z]

Or maybe you can try to use a newer nvcc version.

[keram88]

For clarification, I compiled my test program for sm_86. I will try a different nvcc in the interim.