How to set breakpoints Nsight 5.3/CUDA 9.0/Visual Studio 2017?

I have a very basic question on debugging CUDA code.

Visual Studio, NSight and the device (GeForce GT 720) are all on the same Windows 7 machine.

I have created a .cu file containing a kernel

__global__ void Generate_OnCUDA(void)

, and done a debug build of my project, which contains many C++ source files as well.

I am unable to set a breakpoint in this kernel. Visual Studio disallows setting a breakpoint, with the message “The breakpoint will not currently be hit. No executable code of the debugger’s target code is associated with this line”. Same for every line of the kernel.

I can set breakpoints in my C++ code outside the kernel, including in my C++ code in the .cu file.

How do I set breakpoints, examine variables, etc in the kernel, when the executable is being run under Visual Studio?

The nvcc command generated by Visual Studio:

1>C:\MyProject>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\bin\nvcc.exe" -gencode=arch=compute_30,code=\"sm_30,compute_30\" --use-local-env --cl-version 2017 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio017\Community\VC\Tools\MSVC4.11.25503\bin\HostX86\x64" -x cu  -IC:\MyProject\ ... -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\include"  -G  --keep --keep-dir x64\Debug -maxrregcount=0  --machine 64 --compile -cudart static  -g   -DWIN64 -D_WINDOWS -D_DEBUG  -Xcompiler "/EHsc /W3 /nologo /Od /FS /Zi /RTC1 /MTd " -o x64\Debug\MyKernel.cu.obj "C:\MyProject\MyKernel.cu"

CUDA 9.0 runtime, NSight 5.3, Visual Studio Community 2017 v15.3.2, 64-bit Windows 7

Hi,

GT 720 is not supported by nsight yet, please check the supported gpu list.

@harryz: Tnx. I have ordered a card on the supported list. Never imagined such an old product would not be supported in a current NSight release.

Actually, I suggest that you should buy the maxwell (GTX 9 series) or pascal card (GTX 10 series), kepler cards will be dropped soon, we only run the tests for GK104 and GK107 now.

@harryz: I have now installed a GTX 1050 Ti and updated to drivers v385.41. But I still cannot set breakpoints in the kernel w/Visual Studio. The best I can do is set a breakpoint on the opening brace of the global kernel function. VS stops here, but stepping (F10) continues without stopping.

So, I have to repeat my original question: how do you use Visual Studio 2017/NSight 5.3/CUDA runtime 9.0 to do the basic operations of debugging the kernel (breakpoints, stepping and examining variables).

Hi,

Do you use the original c++ debugger or the “start CUDA debugging” to debug the app? Looks like you were trying to debug the cuda kernel coda via visual c++ debugger. Also could you try the latest nsight 5.4?

Best Regards
Harry

@harryz: I am using the Visual Studio 2017 debugger, necessary for C++. My project is a large C++ pgm; the CUDA part is just one subroutine. In order to start it, and to check the results, I need to use the VS debugger.

I haven’t tried NSight 5.4, but I will. Strange it didn’t notify me of the availability of a new version–all the “Notify” options are on.

I read some postings that said you had to have another instance of Visual Studio running to debug CUDA. That would be really awkward, if it’s even possible.

FYI, the GPU and VS are on the same machine, and O/S is 64-bit.

You cannot debug the cuda code via visual c++ debugger, nsight also cannot debug the code on cpu, sorry for the inconvenience.

There is a workaround solution but it’s really not easy for you to use, you can start your app with “NSIGHT_CUDA_DEBUGGER=1”, then use your visual c++ debugger to attach to it, then you can debug the code on cpu. Detach the c++ debugger and use nsight to attach to you app, then you can debug your gpu code.

@harryz: Bummer. I suppose, then, I can debug both codes simultaneously with two instances of VS2017?

I hope NVidia will consider integrating the NSight capabilities into VS. MS has done it with their C++ Accelerated Massive Parallelism (C++ AMP); Intel has done it with their Fortran, at least for the code that runs on the CPU.

Two instances of visual studio also should work.

I remember there is a internal bug, but dev thinks it’s not critical, sorry for the inconvenience.

I have successfully set a breakpoint in my kernel code. For other beginners, here is what I did (Visual Studio 2017):

  1. Start my app using the “Start CUDA debugging” button on the toolbar (also in the nSight menu).
  2. Set a breakpoint by clicking a line in the kernel code.
  3. The CUDA debugger ignores all breakpoints set in the C++ code, but stops at the breakpoint in the CUDA code.

Both the GPU and the debugger are on the same system.

I haven’t tries harryz’s recommendation of two instances of Visual Studio, because I read that they will interfere with each other.

Hi Barry,
I am a beginner. I want to debug the cuda code in visual studio. I tried a lot of things but unable to debug the program. I am getting an error for both debuggers,

  1. CUDA Nsight Next-Gen debugger: could not initialize driver for debugging.
  2. CUDA Nsight Legacy debugger: cuda grid launch failed.

can you help me here?

Thanks in advance.

UPDATE TO Harry’s comment from 2017:
Note that as of NsightVSE 5.5, you can debug native c++ code (as well as CUDA kernels), using the ‘Next-Gen CUDA Debugger’ (on Pascal and later GPUs). The ‘Legacy Debugger’ still does not support debugging native c++ code (only CUDA kernels) and only works with Pascal and earlier GPUs.

For Next-Gen/Legacy Debugger supported GPU configurations (GPU, OS, DriverMode, DriverVersion), see
https://developer.nvidia.com/nsight-visual-studio-edition-supported-gpus-full-list#SupportedComputeConfigs

Vijay (or to comment on Vijay’s question),
Please see https://devtalk.nvidia.com/default/topic/1044822/nsight-visual-studio-edition/cuda-c-code-debugging-is-failing-in-visual-studio-/?offset=2#5301166