Need help debugging out of bounds errors using -lineinfo

[russellmatt66]

Hi, I’ve run into some out of bounds errors on an application that I’ve written for an RTX2060 6GB. The output from compute-sanitizer --tool memcheck ./imhd-debug > sanitizer.log looks like the below, repeating some 100 odd times,

Starting iteration 0
Evolving fluid interior and boundary
========= Invalid __global__ read of size 4 bytes
=========     at FluidAdvance(float *, float *, float *, float *, float *, float *, float *, float *, const float *, const float *, const float *, const float *, const float *, const float *, const float *, const float *, float *, float *, float *, float *, float *, float *, float *, float *, float, float, float, float, float, int, int, int)+0xfc50
=========     by thread (0,2,0) in block (0,0,0)
=========     Address 0x74fafcffc110 is out of bounds
=========     and is 16,112 bytes before the nearest allocation at 0x74fafd000000 of size 1,048,576 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time

I read online that you can add information to the output of compute-sanitizer by including the -lineinfo flag in the compilation process.

Using my novice understanding of CMake, I cobbled together a build file for the project, and added the following to the CMakeLists.txt file:

target_compile_options(imhd-debug PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:
                        -v
                        -g
                        -lineinfo
                        >)

However, when I build the project with

cmake ../src
make clean
make
compute-sanitizer --tool memcheck ./imhd-debug > sanitizer.log

there is no difference in the output of compute-sanitizer, so I need some human assistance to resolve this, thank you.

[Robert_Crovella]

Why is compute-sanitizer not reporting lineinfo like I've asked it too?

[Robert_Crovella]

Here is the text at the linked article:

I have a CUDA C/C++ application that I wrote for an RTX2060, which I am now debugging. The first step that I took was to run the executable with compute-sanitizer, and discovered there were some out-of-bounds accesses occurring. Essentially, a very similar issue to the one happening in this thread: Unspecified launch failure on Memcpy. Here is what one looks like:

Starting iteration 0
Evolving fluid interior and boundary
========= Invalid __global__ read of size 4 bytes
=========     at FluidAdvance(float *, float *, float *, float *, float *, float *, float *, float *, const float *, const float *, const float *, const float *, const float *, const float *, const float *, const float *, float *, float *, float *, float *, float *, float *, float *, float *, float, float, float, float, float, int, int, int)+0xfc50
=========     by thread (0,2,0) in block (0,0,0)
=========     Address 0x74fafcffc110 is out of bounds
=========     and is 16,112 bytes before the nearest allocation at 0x74fafd000000 of size 1,048,576 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time

Using my novice understanding of CMake I had a working build system functioning, and in order to determine where the out of bounds accesses are occurring, I added the following line to CMakeLists.txt:

target_compile_options(imhd-debug PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:
                        -v
                        -g
                        -lineinfo
                        >)

However, when I rebuild the project, and run compute-sanitizer, I do not find any information added to the output that indicates the lines in the source where the errors are occurring.

Why is this, and what do I need to do to fix it? Thank you for any help.

EDIT: Question is solved! imhd-debug represented the main executable, but the function where the out of bounds access was occurring is defined in a different library that I needed to add a target_compile_options( ... -lineinfo) for. After doing this, I found the line information that I needed! :)

did you verify, via CMake verbose output, that the -lineinfo switch is actually present on the nvcc compile command line for the file/module in question?

– Robert Crovella

21 hours ago

• As far as I can tell, the build is not actually calling nvcc, here is the output from the build: pastebin.com/uVxN1Yu5. There is $ fatbinary ... that includes the flag --generate-line-info

– Ashamandarei

20 hours ago

• 3

You should edit it into the question, not via external link. This is essentially a CMake question , even though the title suggests that compute-sanitizer is somehow not behaving. And you may have misinterpreted my suggestion about “CMake verbose”. I am not suggesting a CUDA verbose setting (which seems to be how you got to looking at fatbinary ...) but instead a CMake verbose output, so it will show you the command lines it is issuing rather than just a progress statement.

– Robert Crovella

19 hours ago

• Thank you for the corrections. Yes, my lack of proper understanding about CMake is essentially the crux of the matter here, and I apologize for any misinterpretation on my end. I followed your advice, and added set( CMAKE_VERBOSE_MAKEFILE on ) at the beginning of my CMakeLists.txt. Consequentially, I realized that in CMakeLists.txt, I was failing to provide the -lineinfo option to the library where the function with the out-of-bounds access was occurring. After rectifying this, I am now seeing the line information that I need. Thank you very much for your help!