Source code should be compiled with nvcc using the -G option in order to provide non-optimized code with debug symbols and linetable information. It is also possible to compile with -linetable in order to debug with symbolics; however, this will optimize SASS and will result in erratic source stepping and source breakpoint behavior.

When stepping, all warps except the debugger focus warp are frozen. However, when stepping over __syncthreads(), all warps in the current block are allowed to make progress. This allows the focus warp to get past the __syncthreads() function.

Stepping out of a global function will resume the application (same result as pressing F5, or selecting the menu command Debug > Continue).

Stepping over the last line of a kernel function will also resume application execution.

Stepping over the last line of an inline function has the same result as stepping out of the inline function.

To experience the best stepping behavior when debugging, we recommend using curly braces when possible in your code. Place the curly braces on their own line of source code. For example, because of the way that the compiler produces scope information, the following coding styles result in different stepping behavior.

These code snippets results in less than optimal stepping behavior:

for () statement; 
for ()        statement; 
for () {        statement;
} 

This code snippet results in optimal stepping behavior:

for ()     
{
     statement;     
} 

In general, breaking up computation by declaring temporary variables results in code that is easier to debug. The chances of a variable being "live" over a certain set of lines increases. This means that as you step through code, the values of variables will be easier to identify.

1. Start a debugging session.

2. Select a thread you would like to single step using the Nsight CUDA Debug Location toolbar.

   See Specify Debugger Context for more information.

3. From the Debug menu, choose Step Into.

   As an alternative, press F11 on the keyboard.

4. The program counter for the current thread will advance one source line.

To run the target application up to the cursor location:

1. In the Document window containing the source, right-click on the desired line of source code.

2. From the Context menu, choose Run To Cursor...

3. The program counter will advance to the specified line of source code.

From the Debug menu, choose Stop Debugging.

Alternatively, you can also press the Stop icon on the debugger toolbar, or type SHIFT+F5.

The target application will now stop.

Source code should be compiled with nvcc using the -G option in order to provide non-optimized code with debug symbols and linetable information. It is also possible to compile with -linetable in order to debug with symbolics; however, this will optimize SASS and will result in erratic source stepping and source breakpoint behavior.

You can use NVIDIA Nsight™ VSCE to set source breakpoints in CUDA code. Source breakpoints allow you to halt your application at the specified line of CUDA source code. You can also constrain your breakpoints to break only when a particular block and thread are executing. Constraining breakpoints to specific blocks and threads is a critical capability in a massively parallel environment where you can have hundreds or thousands of threads running simultaneously.

You set and configure CUDA breakpoints with the same familiar dialog boxes, keyboard shortcuts, and tool windows used for other languages supported by Visual Studio Code.

In a massively parallel environment with hundreds or thousands of threads, it is critical to be able to narrow your breakpoints to just the areas of interest. The CUDA Debugger supports setting conditional breakpoints for GPU threads with arbitrary expressions. Expressions may use program variables, the intrinsics blockIdx and threadIdx, and a few short-hand macros to make it easier to set conditions based on blockIdx and threadIdx.

To set a block or thread condition on a CUDA C breakpoint:

1. Set a breakpoint on a line of source code.

2. Right-click on the breakpoint.

3. From the drop-down menu, select Edit Condition...

4. Type an expression you want evaluated by the debugger at this breakpoint.

NVIDIA Nsight™ VSCE supports function breakpoints. To set a function breakpoint, use one of the following methods:

• From the Run and Debug tab Breakpoints section, select Add Function Breakpoint.

• From the Run menu, select New Breakpoint > Function Breakpoint...

The user can enter the function of the CUDA kernel.

The CUDA debugger has options to halt execution when certain runtime activities occur, such as the launch of a kernel or when API errors occur.

The CUDA debugger can automatically insert an internal breakpoint on the launch function. Enable the Break on Launch feature by setting the breakOnLaunch property in the launch configuration to true. For more information on launch configurations see Create a Launch Configuration.

When enabled, execution will be halted when any __global__ kernel is launched. This is useful to when setting function breakpoints is tedious. As shown below, there is no user breakpoint shown, but execution is halted on the kernel entry point.

The settings to control behavior when CUDA API errors are configured via the onAPIError property in the launch configuration. The valid values are:

• hide will not report any error of any kind.

• ignore will emit a warning but continue the execution of the application.

• stop will emit an error and stop the application.