Control GPU Execution
In this section, learn more about how to control GPU execution, set GPU breakpoints, change debugger focus, and configure how to break on runtime activities.
1. Control GPU Execution
In this section, learn more about how to control GPU execution, set GPU breakpoints, change debugger focus, and configure how to break on runtime activities.
2. Control Execution
The CUDA Debugger allows you to control the execution of your CUDA program during the debugging session. The CUDA Debugger supports stepping into a line of code, and running the program to the cursor.
In order to use any of the below features, the debugged CUDA program must be halted as a result of stopping at a breakpoint or from a previous execution control statement.
2.1. Prerequisites
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.
2.2. Stepping 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.
2.3. Best Practices
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.
2.4. Single Stepping Your Program
-
Start a debugging session.
-
Select a thread you would like to single step using the Nsight CUDA Debug Location toolbar.
See Specify Debugger Context for more information.
-
From the Debug menu, choose Step Into.
As an alternative, press F11 on the keyboard.
-
The program counter for the current thread will advance one source line.
Note: |
Using Step Into causes the debugger to execute the next instruction of the currently selected thread. Stepping occurs on the current focus thread. |
2.5. Running to the Cursor
To run the target application up to the cursor location:
-
In the Document window containing the source, right-click on the desired line of source code.
-
From the Context menu, choose Run To Cursor...
-
The program counter will advance to the specified line of source code.
2.6. Stopping the CUDA Debugger
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.
3. Set GPU Breakpoints
NVIDIA Nsight Visual Studio Code Edition's debugger can control execution, based on GPU breakpoints.
3.1. Prerequisites
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.
3.2. GPU Breakpoints
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.
3.2.1. Creating a Source Breakpoint
You can set breakpoints on any line of executable CUDA source code.
To set a breakpoint on a line of source code:
-
In a source code window, move the insertion point to the line where you want the application to break.
To set a breakpoint in a statement that spans two or more lines, move the insertion point to the last line of the statement.
-
From the Debug menu, choose Toggle Breakpoint.
Alternatively, press the F9 key, or left-click in the left margin of the source code document.
A red breakpoint glyph appears in the left margin of the source code document.
The color of the breakpoint glyph indicates the status of the breakpoint. A red filled glyph, for example, indicates a normal, enabled breakpoint. A white filled glyph indicates that the breakpoint is disabled.
For more information on breakpoints, see the Breakpoints section in the Visual Studio Code documentation.
3.2.2. Managing Breakpoints
Use the standard Visual Studio Code Breakpoints tool window to enable, disable, add and delete breakpoints.
For assistance in locating sample applications, see Working with Samples.
-
From the Run And Debug tab, select the Breakpoints section.
-
Enable or disable a breakpoint by marking the checkbox next to a breakpoint.
-
Delete a breakpoint by selecting a breakpoint and then clicking the Delete button in the toolbar.
-
Delete all breakpoints by clicking the Remove All Breakpoints button.
-
Navigate to the source code location of the breakpoint double-clicking on the breakpoint.
3.3. Conditional Breakpoints
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:
-
Set a breakpoint on a line of source code.
-
Right-click on the breakpoint.
-
From the drop-down menu, select Edit Condition...
-
Type an expression you want evaluated by the debugger at this breakpoint.
3.4. Function Breakpoints
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.
4. Break on Activity
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.
4.1. Break on Launch
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.
4.2. Break on API Error
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.
Notices
Notice
NVIDIA® Nsight™ Application Development Environment for Heterogeneous Platforms, Visual Studio Code Edition 2021.1.0 User GuideSend Feedback
THE INFORMATION IN THIS GUIDE AND ALL OTHER INFORMATION CONTAINED IN NVIDIA DOCUMENTATION REFERENCED IN THIS GUIDE IS PROVIDED “AS IS.” NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO THE INFORMATION FOR THE PRODUCT, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE. Notwithstanding any damages that customer might incur for any reason whatsoever, NVIDIA’s aggregate and cumulative liability towards customer for the product described in this guide shall be limited in accordance with the NVIDIA terms and conditions of sale for the product.
THE NVIDIA PRODUCT DESCRIBED IN THIS GUIDE IS NOT FAULT TOLERANT AND IS NOT DESIGNED, MANUFACTURED OR INTENDED FOR USE IN CONNECTION WITH THE DESIGN, CONSTRUCTION, MAINTENANCE, AND/OR OPERATION OF ANY SYSTEM WHERE THE USE OR A FAILURE OF SUCH SYSTEM COULD RESULT IN A SITUATION THAT THREATENS THE SAFETY OF HUMAN LIFE OR SEVERE PHYSICAL HARM OR PROPERTY DAMAGE (INCLUDING, FOR EXAMPLE, USE IN CONNECTION WITH ANY NUCLEAR, AVIONICS, LIFE SUPPORT OR OTHER LIFE CRITICAL APPLICATION). NVIDIA EXPRESSLY DISCLAIMS ANY EXPRESS OR IMPLIED WARRANTY OF FITNESS FOR SUCH HIGH RISK USES. NVIDIA SHALL NOT BE LIABLE TO CUSTOMER OR ANY THIRD PARTY, IN WHOLE OR IN PART, FOR ANY CLAIMS OR DAMAGES ARISING FROM SUCH HIGH RISK USES.
NVIDIA makes no representation or warranty that the product described in this guide will be suitable for any specified use without further testing or modification. Testing of all parameters of each product is not necessarily performed by NVIDIA. It is customer’s sole responsibility to ensure the product is suitable and fit for the application planned by customer and to do the necessary testing for the application in order to avoid a default of the application or the product. Weaknesses in customer’s product designs may affect the quality and reliability of the NVIDIA product and may result in additional or different conditions and/or requirements beyond those contained in this guide. NVIDIA does not accept any liability related to any default, damage, costs or problem which may be based on or attributable to: (i) the use of the NVIDIA product in any manner that is contrary to this guide, or (ii) customer product designs.
Other than the right for customer to use the information in this guide with the product, no other license, either expressed or implied, is hereby granted by NVIDIA under this guide. Reproduction of information in this guide is permissible only if reproduction is approved by NVIDIA in writing, is reproduced without alteration, and is accompanied by all associated conditions, limitations, and notices.
Trademarks
NVIDIA, the NVIDIA logo, and cuBLAS, CUDA, CUDA-GDB, CUDA-MEMCHECK, cuDNN, cuFFT, cuSPARSE, DIGITS, DGX, DGX-1, DGX Station, NVIDIA DRIVE, NVIDIA DRIVE AGX, NVIDIA DRIVE Software, NVIDIA DRIVE OS, NVIDIA Developer Zone (aka "DevZone"), GRID, Jetson, NVIDIA Jetson Nano, NVIDIA Jetson AGX Xavier, NVIDIA Jetson TX2, NVIDIA Jetson TX2i, NVIDIA Jetson TX1, NVIDIA Jetson TK1, Kepler, NGX, NVIDIA GPU Cloud, Maxwell, Multimedia API, NCCL, NVIDIA Nsight Compute, NVIDIA Nsight Eclipse Edition, NVIDIA Nsight Graphics, NVIDIA Nsight Integration, NVIDIA Nsight Systems, NVIDIA Nsight Visual Studio Edition, NVIDIA Nsight Visual Studio Code Edition, NVLink, nvprof, Pascal, NVIDIA SDK Manager, Tegra, TensorRT, Tesla, Visual Profiler, VisionWorks and Volta are trademarks and/or registered trademarks of NVIDIA Corporation in the United States and other countries. Other company and product names may be trademarks of the respective companies with which they are associated.