Error after instrumenting CALL.REL.NOINC

[Lucindaaaa]

Issue

When we instrument any instruction with CALL.REL.NOINC (CALL.ABS.NOINC is fine), there will be an error msg like

Cuda error in function 'cudaDeviceSynchronize()' file 'instr_count.cu' in line 240 : an illegal memory access was encountered.

Inside instr_count, if I added the if(!strcmp(i->getOpcode(), "CALL.REL.NOINC")) continue; inside instrument_function_if_needed() to just skip CALL.REL.NOINC, it runs to completion with no obvious problem.

Setup

GPU: A100
Driver Version: 555.42.02
CUDA Version: 12.5
Compute Capability: 8.0
NVBit version: 1.7.4

I just want to report this in case it's not already known, since I accidentally encountered it.

[x-y-z]

Do you have a program we can test it? Thanks.

[cesar-avalos3]

Maybe they have a simpler workload, we saw it happen running BERT (using the accel-sim mlperf_inference / reference gpu-app-collection version). More specifically the attention kernels of pytorch (fmha_cutlassF_[...]_sm80)

[Lucindaaaa]

Sorry for the delay. I used quicksilver application, this link (https://gitlab.com/yumengliu/quicksilver_build.git) has a build script and Makefile to build the application, and I put some info in README to help run it (very brief). Please let me know if you need more information. Thank you!

[rodhuega]

Hello,

I am having the same issue with multiple apps in SM_86 architecture. All of them can be observed using the tool instr_count.

• CNS_all from proxy-apps: Kernel D_energy_4q. cns_all_proxy_apps.tar.gz
• ISPASS-LIB from ISPASS: Kernel Pathcalc_Portfolio_KernelGPU2. ISPASS_LIB.tar.gz
• Parboil-sad from parboil: Kernel mb_sad_calc. (The compilation of this application is not as easy/clean as the other two. If needed, ask, and I will upload it.)

cns_all from proxy-apps-doe and parboil-sad work fine with Turing and Blackwell, but they do not have any CALL instructions with those architectures. However, ISPASS-LIB works in both architectures (Blackwell has this other problem #148) despite having CALL.REL.NOINC instructions, while it does not in Ampere and crashes.

I am using NVBIT 1.7.4 and CUDA 12.8

[rodhuega]

Tested in the new 1.7.5 release, and it works at least for my workloads.