<!DOCTYPE html
PUBLIC "-//W3C//DTD XHTML 1.0 Transitional//EN" "http://www.w3.org/TR/xhtml1/DTD/xhtml1-transitional.dtd">
<html xmlns="http://www.w3.org/1999/xhtml" lang="en-us" xml:lang="en-us">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8"></meta>
<meta http-equiv="X-UA-Compatible" content="IE=edge"></meta>
<meta name="copyright" content="(C) Copyright 2005"></meta>
<meta name="DC.rights.owner" content="(C) Copyright 2005"></meta>
<meta name="DC.Type" content="concept"></meta>
<meta name="DC.Title" content="CUDA-MEMCHECK"></meta>
<meta name="abstract" content="The user manual for CUDA-MEMCHECK."></meta>
<meta name="description" content="The user manual for CUDA-MEMCHECK."></meta>
<meta name="DC.Coverage" content="Tools"></meta>
<meta name="DC.subject" content="CUDA MEMCHECK, CUDA MEMCHECK features, CUDA MEMCHECK tools, CUDA MEMCHECK supported OS, CUDA MEMCHECK supported devices, CUDA MEMCHECK error, CUDA MEMCHECK racecheck, CUDA MEMCHECK backtrace, CUDA MEMCHECK hardware exception, CUDA MEMCHECK memory access"></meta>
<meta name="keywords" content="CUDA MEMCHECK, CUDA MEMCHECK features, CUDA MEMCHECK tools, CUDA MEMCHECK supported OS, CUDA MEMCHECK supported devices, CUDA MEMCHECK error, CUDA MEMCHECK racecheck, CUDA MEMCHECK backtrace, CUDA MEMCHECK hardware exception, CUDA MEMCHECK memory access"></meta>
<meta name="DC.Format" content="XHTML"></meta>
<meta name="DC.Identifier" content="abstract"></meta>
<link rel="stylesheet" type="text/css" href="../common/formatting/commonltr.css"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/site.css"></link>
<title>CUDA-MEMCHECK :: CUDA Toolkit Documentation</title>
<!--[if lt IE 9]>
<script src="../common/formatting/html5shiv-printshiv.min.js"></script>
<![endif]-->
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.min.js"></script>
<script type="text/javascript" charset="utf-8" src="../common/formatting/jquery.ba-hashchange.min.js"></script>
<link rel="canonical" href="http://docs.nvidia.com/cuda/cuda-memcheck/index.html"></link>
<link rel="stylesheet" type="text/css" href="../common/formatting/qwcode.highlight.css"></link>
</head>
<body>
<article id="contents">
<div id="release-info">CUDA-MEMCHECK
(<a href="../../pdf/CUDA_Memcheck.pdf">PDF</a>)
-
CUDA Toolkit v5.5
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated
July 19, 2013
-
<a href="mailto:cudatools@nvidia.com?subject=CUDA Tools Documentation Feedback: cuda-memcheck">Send Feedback</a></div>
<div class="topic nested0" id="abstract"><a name="abstract" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#abstract" name="abstract" shape="rect">CUDA-MEMCHECK</a></h2>
<div class="body conbody"></div>
</div>
<div class="topic concept nested0" id="introduction"><a name="introduction" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#introduction" name="introduction" shape="rect">1. Introduction</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="about-cuda-memcheck"><a name="about-cuda-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#about-cuda-memcheck" name="about-cuda-memcheck" shape="rect">1.1. About CUDA-MEMCHECK</a></h3>
<div class="body conbody">
<p class="p">CUDA-MEMCHECK is a functional correctness checking suite included in the
CUDA toolkit. This suite contains multiple tools that can perform different
types of checks. The <dfn class="term">memcheck</dfn> tool is capable of precisely detecting and
attributing out of bounds and misaligned memory access errors in CUDA applications.
The tool also reports hardware exceptions encountered by the GPU.
The <dfn class="term">racecheck</dfn> tool can report shared memory data access hazards
that can cause data races.This document describes the usage of these tools.
</p>
<p class="p">CUDA-MEMCHECK can be run in <dfn class="term">standalone mode</dfn> where the user's
application is started under CUDA-MEMCHECK. The <dfn class="term">memcheck</dfn> tool can also
be enabled in <dfn class="term">integrated mode</dfn> inside CUDA-GDB.
</p>
</div>
</div>
<div class="topic concept nested1" id="why-cuda-memcheck"><a name="why-cuda-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#why-cuda-memcheck" name="why-cuda-memcheck" shape="rect">1.2. Why CUDA-MEMCHECK?</a></h3>
<div class="body conbody">
<p class="p">NVIDIA allows developers to easily harness the power of GPUs to solve problems in
parallel using CUDA. CUDA applications often run thousands of threads in parallel.
Every programmer invariably encounters memory access errors and thread ordering
errors that are hard to detect and time consuming to debug. The number of
such errors increases substantially when dealing with thousands of threads.
The CUDA-MEMCHECK suite is designed to detect such errors in your CUDA application.
Using the memcheck tool, CUDA-MEMCHECK can identify memory access errors as well
as hardware reported program errors. The racecheck tool in CUDA-MEMCHECK
can identify hazards caused by race conditions in the CUDA program.
</p>
</div>
</div>
<div class="topic concept nested1" id="how-to-get-cuda-memcheck"><a name="how-to-get-cuda-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#how-to-get-cuda-memcheck" name="how-to-get-cuda-memcheck" shape="rect">1.3. How to Get CUDA-MEMCHECK</a></h3>
<div class="body conbody">
<p class="p">CUDA-MEMCHECK is installed as part of the CUDA toolkit.</p>
</div>
</div>
<div class="topic concept nested1" id="cuda-memcheck-tools"><a name="cuda-memcheck-tools" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#cuda-memcheck-tools" name="cuda-memcheck-tools" shape="rect">1.4. CUDA-MEMCHECK tools</a></h3>
<div class="body conbody">
<div class="p">
Tools allow use the basic CUDA-MEMCHECK infrastructure to provide different
checking mechanisms. Currently, the supported tools are :
<ul class="ul">
<li class="li"><dfn class="term">Memcheck</dfn> - The memory access error and leak detection tool.
See <a class="xref" href="index.html#memcheck-tool" shape="rect">Memcheck Tool</a></li>
<li class="li"><dfn class="term">Racecheck</dfn> - The shared memory data access hazard detection tool.
See <a class="xref" href="index.html#racecheck-tool" shape="rect">Racecheck Tool</a></li>
</ul>
</div>
<div class="tablenoborder"><a name="cuda-memcheck-tools__supported-modes-by-tool" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="cuda-memcheck-tools__supported-modes-by-tool" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 1. Supported Modes by CUDA-MEMCHECK tool</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e177" rowspan="1" colspan="1">Tool Name</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e180" rowspan="1" colspan="1">Standalone Mode</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e183" rowspan="1" colspan="1">Integrated Mode</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e177" rowspan="1" colspan="1">Memcheck</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e180" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e183" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e177" rowspan="1" colspan="1">Racecheck</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e180" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e183" rowspan="1" colspan="1">No</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="using-cuda-memcheck"><a name="using-cuda-memcheck" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#using-cuda-memcheck" name="using-cuda-memcheck" shape="rect">2. Using CUDA-MEMCHECK</a></h2>
<div class="body conbody">
<div class="p">CUDA-MEMCHECK tools can be invoked by running the <samp class="ph codeph">cuda-memcheck</samp>
executable as follows:
<pre class="pre screen" xml:space="preserve">
<strong class="ph b">cuda-memcheck [options] app_name [app_options] </strong>
</pre></div>
<p class="p">
For a full list of options that can be specified to memcheck and their default values, see <a class="xref" href="index.html#command-line-options" shape="rect">Command Line Options</a>.
</p>
</div>
<div class="topic concept nested1" id="command-line-options"><a name="command-line-options" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#command-line-options" name="command-line-options" shape="rect">2.1. Command Line Options</a></h3>
<div class="body conbody">
<p class="p"> Command line options can be specified to <samp class="ph codeph">cuda-memcheck</samp>.
With some exceptions, the options to memcheck are usually of
the form <samp class="ph codeph">--option value</samp>. The option list can be terminated
by specifying <samp class="ph codeph">--</samp>. All subsequent words on the command line are
treated as the application being run and its arguments.
</p>
<p class="p"> The table below describes the supported
options in detail. The first column is the option name as passed to CUDA-MEMCHECK.
Some options have a one character short form, which is given in parentheses.
These options can be invoked using a single hypen. For example, the help option
can be invoked as <samp class="ph codeph">-h</samp>. The options that have a short form do not
take a value.
</p>
<p class="p"> The second column contains the permissible values for the option. In case the
value is user defined, this is shown below in braces {}. An option
that can accept any numerical value is represented as <dfn class="term"> {number} </dfn>.
Blank entries indicate that the value is not present.
</p>
<p class="p"> The third column contains the default value of the option. Some options have
different default values depending on the architecture they are being run on.
</p>
<div class="tablenoborder"><a name="command-line-options__cuda-memcheck-command-line-options" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="command-line-options__cuda-memcheck-command-line-options" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 2. CUDA-MEMCHECK Command line options</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="20%" id="d54e299" rowspan="1" colspan="1">Option</th>
<th class="entry" valign="top" width="20%" id="d54e302" rowspan="1" colspan="1">Values</th>
<th class="entry" valign="top" width="20%" id="d54e305" rowspan="1" colspan="1">Default</th>
<th class="entry" valign="top" width="40%" id="d54e308" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">demangle</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">full, simple, no</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">full</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Enables demangling of device function names. For more information, see <a class="xref" href="index.html#name-demangling" shape="rect">Name Demangling</a>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">destroy-on-device-error</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">context,kernel</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">context</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">This controls how the application proceeds on hitting a memory access error. For more information, see <a class="xref" href="index.html#error-actions" shape="rect">Error Actions</a>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">error-exitcode</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">{number}</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">0</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">The exit code memcheck will return if the original application succeeded but there were any memcheck detected errors. This
is meant to allow CUDA-MEMCHECK to be integrated into automated test suites
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">flush-to-disk</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">yes,no</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Forces every disk write to be flushed to disk. When enabled, this will make CUDA-MEMCHECK tools much slower. </td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">force-blocking-launches</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">yes,no</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">no on SM 2.0+, yes on SM 1.x</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">This forces all host kernel launches to be sequential. When enabled, the number and precision of memcheck reported errors
will decrease. This option only has effect on SM 2.0 and higher GPUs
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">help (h)</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Displays the help message</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">prefix</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">{string}</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">========</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">The string prepended to CUDA-MEMCHECK output lines</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">print-level</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">info,warn,error,fatal</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">info</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">The minimum level print level of messages from CUDA-MEMCHECK.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">racecheck-report</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">hazard,analysis,all</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">hazard</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Controls how racecheck reports information. For more information, see <a class="xref" href="index.html#racecheck-report-modes" shape="rect">Racecheck report modes</a>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">read</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">{filename}</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">The input memcheck file to read data from. This can be used in conjunction with the save option to allow processing records
after a run.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">save</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">{filename}</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Filename where CUDA-MEMCHECK will save the output from the current run.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">show-backtrace</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">yes,host,device,no</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1"> Displays a backtrace for most types of errors. No disables all backtraces, Yes enables all backtraces. Host enables only
host side backtraces. Device enables only device side backtraces. For more information, see <a class="xref" href="index.html#stack-backtraces" shape="rect">Stack Backtraces</a>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">tool</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">memcheck, racecheck</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">memcheck</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Controls which CUDA-MEMCHECK tool is actively running</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e299" rowspan="1" colspan="1">version (V)</td>
<td class="entry" valign="top" width="20%" headers="d54e302" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="20%" headers="d54e305" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="40%" headers="d54e308" rowspan="1" colspan="1">Prints the version of cuda-memcheck</td>
</tr>
</tbody>
</table>
</div>
<div class="tablenoborder"><a name="command-line-options__memcheck-tool-command-line-options" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="command-line-options__memcheck-tool-command-line-options" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 3. <dfn class="term">Memcheck</dfn> Tool Command line options</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="20%" id="d54e565" rowspan="1" colspan="1">Option</th>
<th class="entry" valign="top" width="20%" id="d54e568" rowspan="1" colspan="1">Values</th>
<th class="entry" valign="top" width="20%" id="d54e571" rowspan="1" colspan="1">Default</th>
<th class="entry" valign="top" width="40%" id="d54e574" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e565" rowspan="1" colspan="1">check-device-heap</td>
<td class="entry" valign="top" width="20%" headers="d54e568" rowspan="1" colspan="1">yes,no</td>
<td class="entry" valign="top" width="20%" headers="d54e571" rowspan="1" colspan="1">yes on SM 2.0+, no on SM 1.x</td>
<td class="entry" valign="top" width="40%" headers="d54e574" rowspan="1" colspan="1">Enable checking of device heap allocations. This applies to both error checking and leak checking. This option only has effect
on SM 2.0 and higher GPUs
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e565" rowspan="1" colspan="1">leak-check</td>
<td class="entry" valign="top" width="20%" headers="d54e568" rowspan="1" colspan="1">full,no</td>
<td class="entry" valign="top" width="20%" headers="d54e571" rowspan="1" colspan="1">no</td>
<td class="entry" valign="top" width="40%" headers="d54e574" rowspan="1" colspan="1">Prints information about all allocations that have not been freed via cudaFree at the point when the context was destroyed.
For more information, see <a class="xref" href="index.html#leak-checking" shape="rect">Leak Checking</a>.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="20%" headers="d54e565" rowspan="1" colspan="1">report-api-errors</td>
<td class="entry" valign="top" width="20%" headers="d54e568" rowspan="1" colspan="1">yes,no</td>
<td class="entry" valign="top" width="20%" headers="d54e571" rowspan="1" colspan="1">yes</td>
<td class="entry" valign="top" width="40%" headers="d54e574" rowspan="1" colspan="1">Report errors if any CUDA API call fails. For more information, see <a class="xref" href="index.html#api-error-checking" shape="rect">CUDA API Error Checking</a>.
</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="supported-operating-systems"><a name="supported-operating-systems" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#supported-operating-systems" name="supported-operating-systems" shape="rect">2.2. Supported Operating Systems</a></h3>
<div class="body conbody">
<p class="p">The standalone CUDA-MEMCHECK binary is supported on all CUDA supported platforms
i.e. Windows, Mac OS X, and supported Linux distributions.
CUDA-MEMCHECK can interoperate with CUDA-GDB on Linux and Mac OS X.
</p>
</div>
</div>
<div class="topic concept nested1" id="supported-devices"><a name="supported-devices" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#supported-devices" name="supported-devices" shape="rect">2.3. Supported Devices</a></h3>
<div class="body conbody">
<p class="p">The CUDA-MEMCHECK tool suite is supported on all CUDA capable GPUs with SM versions
1.1 and above. Individual tools may support a different subset of GPUs.
The table below contains the list of current tools and their supported GPUs
</p>
<div class="tablenoborder"><a name="supported-devices__supported-devices-by-tool" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="supported-devices__supported-devices-by-tool" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 4. Supported Devices by CUDA-MEMCHECK tool</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e682" rowspan="1" colspan="1">Tool Name</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e685" rowspan="1" colspan="1">SM 1.x</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e688" rowspan="1" colspan="1">SM 2.x</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e691" rowspan="1" colspan="1">SM 3.0</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e694" rowspan="1" colspan="1">SM 3.5</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e682" rowspan="1" colspan="1">Memcheck</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e685" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e688" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e691" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e694" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e682" rowspan="1" colspan="1">Racecheck</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e685" rowspan="1" colspan="1">No</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e688" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e691" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e694" rowspan="1" colspan="1">Yes</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="compilation-options"><a name="compilation-options" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#compilation-options" name="compilation-options" shape="rect">2.4. Compilation Options</a></h3>
<div class="body conbody">
<p class="p">
The CUDA-MEMCHECK tools <dfn class="term">memcheck</dfn> and <dfn class="term">racecheck</dfn> do not
need any special compilation flags to function.
</p>
<p class="p">
The output displayed by the CUDA-MEMCHECK tools is more useful with some extra
compiler flags. The <samp class="ph codeph">-G</samp> option to nvcc forces the compiler to
generate debug information for the CUDA application. To generate line number
information for applications without affecting the optimization level of
the output, the <samp class="ph codeph">-lineinfo</samp> option to nvcc can be used.
The CUDA-MEMCHECK tools fully support both of these options
and can display source attribution of errors for applications
compiled with line information.
</p>
<p class="p"> The stack backtrace feature of the CUDA-MEMCHECK tools is more useful
when the application contains function symbol names. For the host backtrace,
this varies based on the host OS. On Linux, the host compiler
must be given the <samp class="ph codeph">-rdynamic</samp> option to retain function
symbols. On Windows, the application must be compiled for debugging,
i.e. the <samp class="ph codeph">/Zi</samp> option. When using nvcc, flags to the host
compiler can be specified using the <samp class="ph codeph">-Xcompiler</samp> option.
For the device backtrace, the full
frame information is only available when the application is compiled
with device debug information. The compiler can skip generation of
frame information when building with optimizations.
</p>
<div class="p">
Sample command line to build with function symbols and device side line
information on linux:
<pre class="pre screen" xml:space="preserve">
nvcc -Xcompiler -rdynamic -lineinfo -o out in.cu
</pre></div>
</div>
</div>
</div>
<div class="topic concept nested0" id="memcheck-tool"><a name="memcheck-tool" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#memcheck-tool" name="memcheck-tool" shape="rect">3. Memcheck Tool</a></h2>
<div class="body conbody">
<p class="p"></p>
</div>
<div class="topic concept nested1" id="what-is-memcheck"><a name="what-is-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#what-is-memcheck" name="what-is-memcheck" shape="rect">3.1. What is Memcheck ?</a></h3>
<div class="body conbody">
<p class="p">
The <dfn class="term">memcheck</dfn> tool is a run time error detection tool for
CUDA applications. The tool can
precisely detect and report out of bounds and misaligned memory accesses to
global, local, shared and global atomic instructions in CUDA
applications. It can also detect and report hardware reported error
information. In addition, the memcheck tool can detect and report memory
leaks in the user application.
</p>
<p class="p"></p>
</div>
</div>
<div class="topic concept nested1" id="supported-error-detection"><a name="supported-error-detection" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#supported-error-detection" name="supported-error-detection" shape="rect">3.2. Supported Error Detection</a></h3>
<div class="body conbody">
<p class="p">
The errors that can be reported by the memcheck tool are summarized in the table
below. The location column indicates whether the report originates from the
host or from the device. The precision of an error is explained in the paragraph
below.
</p>
<div class="tablenoborder"><a name="supported-error-detection__memcheck-error-types" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="supported-error-detection__memcheck-error-types" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 5. Memcheck reported error types</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="16.666666666666664%" id="d54e850" rowspan="1" colspan="1">Name</th>
<th class="entry" valign="top" width="33.33333333333333%" id="d54e853" rowspan="1" colspan="1">Description</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e856" rowspan="1" colspan="1">Location</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e859" rowspan="1" colspan="1">Precision</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e862" rowspan="1" colspan="1">See also</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e850" rowspan="1" colspan="1"><dfn class="term">Memory access error</dfn></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e853" rowspan="1" colspan="1">
Errors due to
out of bounds or misaligned accesses to memory by a global,
local, shared or global atomic access.
</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e856" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e859" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e862" rowspan="1" colspan="1"><a class="xref" href="index.html#memory-access-error-reporting" shape="rect">Memory Access Error Reporting</a></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e850" rowspan="1" colspan="1"><dfn class="term">Hardware exception</dfn></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e853" rowspan="1" colspan="1">
Errors that are reported
by the hardware error reporting mechanism.
</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e856" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e859" rowspan="1" colspan="1">Imprecise</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e862" rowspan="1" colspan="1"><a class="xref" href="index.html#hardware-exception-reporting" shape="rect">Hardware Exception Reporting</a></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e850" rowspan="1" colspan="1"><dfn class="term">Malloc/Free errors</dfn></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e853" rowspan="1" colspan="1">
Errors that occur due to incorrect
use of <samp class="ph codeph">malloc()/free()</samp>
in CUDA kernels.
</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e856" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e859" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e862" rowspan="1" colspan="1"><a class="xref" href="index.html#device-side-allocation-checking" shape="rect">Device Side Allocation Checking</a></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e850" rowspan="1" colspan="1"><dfn class="term">CUDA API errors</dfn></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e853" rowspan="1" colspan="1">
Reported when a CUDA API call in the application
returns a failure.
</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e856" rowspan="1" colspan="1">Host</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e859" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e862" rowspan="1" colspan="1"><a class="xref" href="index.html#api-error-checking" shape="rect">CUDA API Error Checking</a></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e850" rowspan="1" colspan="1"><dfn class="term">cudaMalloc memory leaks</dfn></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e853" rowspan="1" colspan="1">
Allocations of device memory using <samp class="ph codeph">cudaMalloc()</samp>
that have not been freed by the application.
</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e856" rowspan="1" colspan="1">Host</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e859" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e862" rowspan="1" colspan="1"><a class="xref" href="index.html#leak-checking" shape="rect">Leak Checking</a></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e850" rowspan="1" colspan="1"><dfn class="term">Device Heap Memory Leaks</dfn></td>
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e853" rowspan="1" colspan="1">
Allocations of device memory using <samp class="ph codeph">malloc()</samp>
in device code that have not been freed by the application.
</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e856" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e859" rowspan="1" colspan="1">Imprecise</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e862" rowspan="1" colspan="1"><a class="xref" href="index.html#device-side-allocation-checking" shape="rect">Device Side Allocation Checking</a></td>
</tr>
</tbody>
</table>
</div>
<p class="p">
The memcheck tool reports two classes of errors
<dfn class="term">precise</dfn> and <dfn class="term">imprecise</dfn>.
</p>
<p class="p"><dfn class="term">Precise</dfn> errors in memcheck are those that the tool can uniquely
identify and gather all information for.
For these errors, memcheck can report the block and thread coordinates
of the thread causing the failure, the PC of the instruction performing the
access, as well as the address being accessed and its size and type. If the CUDA
application contains line number information (by either being compiled with device
side debugging information, or with line information), then the tool will also
print the source file and line number of the erroneous access.
</p>
<p class="p"><dfn class="term">Imprecise</dfn> errors are errors reported by the hardware
error reporting mechanism that could not be precisely attributed to a particular
thread. The precision of the error varies based on the type of the error
and in many cases, memcheck may not be able to attribute the cause
of the error back to the source file and line. Imprecise error reporting
is only supported on SM 2.0 and higher GPUs.
</p>
</div>
</div>
<div class="topic concept nested1" id="using-memcheck"><a name="using-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#using-memcheck" name="using-memcheck" shape="rect">3.3. Using Memcheck</a></h3>
<div class="body conbody">
<div class="p">
The memcheck tool is enabled by default when running the
CUDA-MEMCHECK application. It can also be explicitly enabled by using
the <samp class="ph codeph">--tool memcheck</samp> option.
<pre class="pre screen" xml:space="preserve">
<strong class="ph b">cuda-memcheck [memcheck_options] app_name [app_options]</strong>
</pre></div>
<p class="p">
When run in this way, the memcheck tool will look for precise, imprecise,
malloc/free and CUDA API errors. The reporting of device leaks must be explictly
enabled.
Errors identified by the memcheck tool are displayed on the screen after
the application has completed execution. See
<a class="xref" href="index.html#understanding-memcheck-errors" shape="rect">Understanding Memcheck Errors</a> for more information about
how to interpret the messages printed by the tool.
</p>
</div>
</div>
<div class="topic concept nested1" id="understanding-memcheck-errors"><a name="understanding-memcheck-errors" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#understanding-memcheck-errors" name="understanding-memcheck-errors" shape="rect">3.4. Understanding Memcheck Errors</a></h3>
<div class="body conbody">
<p class="p">
The memcheck tool can produce a variety of different errors. This is a short
guide showing some samples of errors and explaining how the information
in each error report can be interpreted.
</p>
<ol class="ol">
<li class="li">
<div class="p"><dfn class="term">Memory access error</dfn>: Memory access errors are generated for errors
that the memcheck tool can correctly attribute and identify the erroneous
instruction. Below is an example of a precise memory access error
<pre class="pre screen" xml:space="preserve">
========= Invalid __global__ write of size 4
========= at 0x00000060 in memcheck_demo.cu:6:unaligned_kernel(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x400100001 is misaligned
</pre></div>
<div class="p">
Let us examine this error line by line :
<pre class="pre screen" xml:space="preserve">Invalid __global__ write of size 4</pre>
The first line shows the memory segment, type and size being accessed.
The memory segment is one of :
<ul class="ul">
<li class="li">__global__ : for device global memory</li>
<li class="li">__shared__ : for per block shared memory</li>
<li class="li">__local__ : for per thread local memory</li>
</ul>
In this case, the access was to device global memory.
The next field contains information about the type of access,
whether it was a read or a write. In this case, the access is a write.
Finally, the last item is the size of the access in bytes. In this
example, the access was 4 bytes in size.
</div>
<div class="p"><pre class="pre screen" xml:space="preserve">at 0x00000060 in memcheck_demo.cu:6:unaligned_kernel(void)</pre>
The second line contains the PC of the instruction, the source file
and line number (if available) and the CUDA kernel name.
In this example, the instruction causing the access was at
PC 0x60 inside the <samp class="ph codeph">unaligned_kernel</samp> CUDA kernel.
Additionally, since the application was compiled with line number
information, this instruction corresponds to line 6 in the memcheck_demo.cu
source file.
</div>
<div class="p"><pre class="pre screen" xml:space="preserve">by thread (0,0,0) in block (0,0,0)</pre>
The third line contains the thread indices and block indices of the
thread on which the error was hit.
In this example, the thread doing the erroneous access belonged to
the first thread in the first block.
</div>
<div class="p"><pre class="pre screen" xml:space="preserve">Address 0x400100001 is misaligned</pre>
The fourth line contains the memory address being accessed and the type
of of access error. The type of access error can either be out of bounds
access or misaligned access.
In this example, the access was to address 0x400100001 and the access
error was because this address was not aligned correctly.
</div>
<p class="p"></p>
</li>
<li class="li">
<p class="p"><dfn class="term">Hardware exception</dfn>: Imprecise errors are generated for errors that
the hardware reports to the memcheck tool. Hardware exceptions have a variety
of formats and messages. Typically, the first line will provide some information
about the type of error encountered.
</p>
<p class="p"></p>
</li>
<li class="li">
<div class="p"><dfn class="term">Malloc/free error</dfn>: Malloc/free errors refer to the errors in the
invocation of device side <samp class="ph codeph">malloc()/free()</samp> in CUDA kernels.
An example of a malloc/free error :
<pre class="pre screen" xml:space="preserve">
========= Malloc/Free error encountered : Double free
========= at 0x000079d8
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x400aff920
</pre></div>
<div class="p">
We can examine this line by line.
<pre class="pre screen" xml:space="preserve">Malloc/Free error encountered : Double free</pre>
The first line indicates that this is a malloc/free error, and contains
the type of error. This type can be :
<ul class="ul">
<li class="li">Double free : This indicates that the thread called
<samp class="ph codeph">free()</samp> on an allocation that
has already been freed.
</li>
<li class="li">Invalid pointer to free : This indicates that <samp class="ph codeph">free</samp>
was called
on a pointer that was not returned
by <samp class="ph codeph">malloc()</samp></li>
<li class="li">Heap corruption : This indicates generalized heap corruption,
or cases where the state of the heap
was modified in a way that memcheck
did not expect
</li>
</ul>
In this example, the error is due to calling <samp class="ph codeph">free()</samp>
on a pointer which had already been freed.
</div>
<div class="p"><pre class="pre screen" xml:space="preserve">at 0x000079d8</pre>
The second line gives the PC on GPU where the error was reported.
This PC is usually inside of system code, and is not interesting
to the user. The device frame backtrace will contain the location
in user code where the <samp class="ph codeph">malloc()/free()</samp> call was
made.
</div>
<div class="p"><pre class="pre screen" xml:space="preserve">by thread (0,0,0) in block (0,0,0)</pre>
The third line contains the thread and block indices of the thread
that caused this error. In this example, the thread has threadIdx = (0,0,0)
and blockIdx = (0,0,0)
</div>
<div class="p"><pre class="pre screen" xml:space="preserve">Address 0x400aff920</pre>
This line contains the value of the pointer passed to <samp class="ph codeph">free()</samp>
or returned by <samp class="ph codeph">malloc()</samp></div>
<p class="p"></p>
</li>
<li class="li">
<div class="p"><dfn class="term">Leak errors</dfn>: Errors are reported for allocations
created using cudaMalloc and for allocations on the device heap that were
not freed when the CUDA context was destroyed.
An example of a cudaMalloc allocation leak report follows :
<pre class="pre screen" xml:space="preserve">
========= Leaked 64 bytes at 0x400200200
</pre>
The error message reports information about the size of the allocation
that was leaked as well as the address of the allocation on the device.
</div>
<div class="p">
A device heap leak message will be explicitly identified as such:
<pre class="pre screen" xml:space="preserve">
========= Leaked 16 bytes at 0x4012ffff6 on the device heap
</pre></div>
<p class="p"></p>
</li>
<li class="li">
<div class="p"><dfn class="term">CUDA API error</dfn>: CUDA API errors are reported for CUDA
API calls that return an error value. An example of a CUDA API error:
<pre class="pre screen" xml:space="preserve">
========= Program hit error 11 on CUDA API call to cudaMemset
</pre>
The message contains the returned value of the CUDA API call, as well as
the name of the API function that was called.
</div>
</li>
</ol>
</div>
</div>
<div class="topic concept nested1" id="integrated-mode"><a name="integrated-mode" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#integrated-mode" name="integrated-mode" shape="rect">3.5. Integrated Mode</a></h3>
<div class="body conbody">
<div class="p">You can execute the memcheck tool from within CUDA-GDB by using the
following option before running the application:
<pre class="pre screen" xml:space="preserve">(cuda-gdb) <strong class="ph b">set cuda memcheck on</strong></pre></div>
<p class="p">
In integrated mode, the memcheck tool improves the precision of error
reporting by CUDA-GDB. The memory access checks are enabled, allowing
identification of the thread that may be causing a warp or
device level exception.
</p>
</div>
</div>
<div class="topic concept nested1" id="api-error-checking"><a name="api-error-checking" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#api-error-checking" name="api-error-checking" shape="rect">3.6. CUDA API Error Checking</a></h3>
<div class="body conbody">
<p class="p"> The memcheck tool supports reporting an error if a CUDA API call made by the user
program returned an error. The tool supports this detection for both
CUDA run time and CUDA driver API calls. In all cases, if the API function
call has a nonzero return value, CUDA-MEMCHECK will print an error message
containing the name of the API call that failed and the return value of the
API call.
</p>
<p class="p">
CUDA API error reports do not terminate the application, they merely provide
extra information. It is up to the application to check the
return status of CUDA API calls and handle error conditions appropriately.
</p>
</div>
</div>
<div class="topic concept nested1" id="device-side-allocation-checking"><a name="device-side-allocation-checking" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#device-side-allocation-checking" name="device-side-allocation-checking" shape="rect">3.7. Device Side Allocation Checking</a></h3>
<div class="body conbody">
<p class="p">
On SM 2.0 and higher GPUs, the <dfn class="term">memcheck</dfn> tool checks accesses to
allocations in the device heap.
</p>
<p class="p">
These allocations are created by calling <samp class="ph codeph">malloc()</samp> inside a kernel.
This feature is implicitly enabled and can be disabled by specifying the
<samp class="ph codeph">--check-device-heap no</samp> option. This
feature is only activated for kernels in the application that call
<samp class="ph codeph">malloc()</samp>.
</p>
<p class="p">
The current implementation does not require space on the device heap, and so
the heap allocation behavior of the program with and without memcheck
should remain similar. The <dfn class="term">memcheck</dfn> tool does require space in
device global memory to track these heap allocations and will print an internal
error message if it is not able to allocate this space in device global memory.
</p>
<div class="p">
In addition to access checks, the <dfn class="term">memcheck</dfn> tool can now perform
libc style checks on the <samp class="ph codeph">malloc()/free()</samp> calls. The tool
will report an error if the application calls a <samp class="ph codeph">free()</samp> twice
on a kernel, or if it calls <samp class="ph codeph">free()</samp> on an invalid pointer.
<div class="note note"><span class="notetitle">Note:</span> Make sure to look at the device side backtrace to find the location
in the application where the <samp class="ph codeph">malloc()/free()</samp> call was
made
</div>
</div>
</div>
</div>
<div class="topic concept nested1" id="leak-checking"><a name="leak-checking" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#leak-checking" name="leak-checking" shape="rect">3.8. Leak Checking</a></h3>
<div class="body conbody">
<p class="p">The <dfn class="term">memcheck</dfn> tool can detect leaks of allocated memory.
</p>
<p class="p">Memory leaks are device side allocations that have not been freed by the time
the context is destroyed. The <dfn class="term">memcheck</dfn> tool tracks device memory
allocations created using the CUDA driver or runtime APIs. Starting in CUDA 5,
allocations that are created dynamically on the device heap by calling
<samp class="ph codeph">malloc()</samp> inside a kernel are also tracked.
</p>
<p class="p">For an accurate leak checking summary to be generated, the application's
CUDA context must be destroyed at the end. This can be done explicitly by
calling <samp class="ph codeph">cuCtxDestroy()</samp> in applications using the CUDA driver
API, or by calling
<samp class="ph codeph">cudaDeviceReset()</samp> in applications programmed against the CUDA
run time API.
</p>
<p class="p">The <samp class="ph codeph">--leak-check full</samp> option must be specified to enable
leak checking.
</p>
</div>
</div>
</div>
<div class="topic concept nested0" id="racecheck-tool"><a name="racecheck-tool" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#racecheck-tool" name="racecheck-tool" shape="rect">4. Racecheck Tool</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="what-is-racecheck"><a name="what-is-racecheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#what-is-racecheck" name="what-is-racecheck" shape="rect">4.1. What is Racecheck ?</a></h3>
<div class="body conbody">
<p class="p">
The <dfn class="term">racecheck</dfn> tool is a run time shared memory data access hazard
detector. The primary use of this tool is to help identify memory access
race conditions in CUDA applications that use shared memory.
</p>
<p class="p">
In CUDA applications, storage declared with the <samp class="ph codeph">__shared__</samp>
qualifier is placed in on chip <dfn class="term">shared memory</dfn>. All threads in a
thread block can access this per block shared memory. Shared memory goes out of
scope when the thread block completes execution. As shared memory is on chip,
it is frequently used for inter thread communication and as a temporary
buffer to hold data being processed. As this data is being accessed by
multiple threads in parallel, incorrect program assumptions may result
in data races.
Racecheck is a tool built to identify these hazards and help users write
programs free of shared memory races.
</p>
<p class="p">
Currently, this tool only supports detecting accesses to on-chip shared memory.
For supported architectures, see <a class="xref" href="index.html#supported-devices" shape="rect">Supported Devices</a>.
</p>
</div>
</div>
<div class="topic concept nested1" id="what-are-hazards"><a name="what-are-hazards" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#what-are-hazards" name="what-are-hazards" shape="rect">4.2. What are Hazards?</a></h3>
<div class="body conbody">
<p class="p">
A <dfn class="term">data access hazard</dfn> is a case where two threads
attempt to access the same location in memory resulting in nondeterministic
behavior, based on the relative order of the two accesses.
These hazards cause <dfn class="term">data races</dfn> where the behavior or the output
of the application depends on the order in which all parallel threads are
executed by the hardware. Race conditions manifest as intermittent application
failures or as failures when attempting to run a working application on a
different GPU.
</p>
<div class="p">
The racecheck tool identifies three types of canonical hazards in a program.
These are :
<ul class="ul">
<li class="li"> Write-After-Write (<dfn class="term">WAW</dfn>) hazards
<p class="p"> This hazard occurs when two threads attempt to
write data to the same memory location. The resulting value
in that location depends on the relative order of the two
accesses.
</p>
</li>
<li class="li"> Read-After-Write (<dfn class="term">RAW</dfn>) hazards
<p class="p"> This hazard occurs when two threads access the same memory location,
with one thread performing a read and another a write. In
this case, the writing thread is ordered before the reading
thread and the value returned to the reading thread is
not the original value at the memory location.
</p>
</li>
<li class="li"> Write-After-Read (<dfn class="term">WAR</dfn>) hazards
<p class="p"> This hazard occurs when two threads access the same memory
location, with one thread performing a read and the other a write.
In this case, the reading thread reads the value before the
writing thread commits it.
</p>
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested1" id="using-racecheck"><a name="using-racecheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#using-racecheck" name="using-racecheck" shape="rect">4.3. Using Racecheck</a></h3>
<div class="body conbody">
<div class="p">
The racecheck tool is enabled by running the CUDA-MEMCHECK application
with the <samp class="ph codeph">--tool racecheck</samp> option.
<pre class="pre screen" xml:space="preserve">
<strong class="ph b">cuda-memcheck --tool racecheck [memcheck_options] app_name [app_options]</strong>
</pre></div>
<p class="p">
Once racecheck has identified a hazard, the user can make program modifications
to ensure this hazard is no longer present.
In the case of Write-After-Write hazards, the program should be modified
so that multiple writes are not happening to the same location.
In the case of Read-After-Write and Write-After-Read hazards, the reading
and writing locations should be deterministically ordered. In CUDA kernels,
this can be achieved by inserting a <samp class="ph codeph">__syncthreads()</samp> call
between the two accesses.
</p>
<div class="note note"><span class="notetitle">Note:</span> The racecheck tool does not perform any memory access error checking.
It is recommended that users first run the memcheck tool to ensure the
application is free of errors
</div>
</div>
</div>
<div class="topic concept nested1" id="racecheck-report-modes"><a name="racecheck-report-modes" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#racecheck-report-modes" name="racecheck-report-modes" shape="rect">4.4. Racecheck report modes</a></h3>
<div class="body conbody">
<div class="p">
The racecheck tool can produce two types of output :
<ul class="ul">
<li class="li"><dfn class="term">Hazard</dfn> reports
<p class="p"> These reports contain detailed information about one particular hazard.
Each hazard report is byte accurate and represents information about
conflicting accesses between two threads that affect this byte of
shared memory.
</p>
</li>
<li class="li"><dfn class="term">Analysis</dfn> reports
<p class="p"> These reports contain a post analysis set of reports. These reports
are produced by the racecheck tool by analysing multiple hazard
reports and examining active device state.
For more about analysis reports, see
</p>
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested1" id="understanding-racecheck-analysis-reports"><a name="understanding-racecheck-analysis-reports" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#understanding-racecheck-analysis-reports" name="understanding-racecheck-analysis-reports" shape="rect">4.5. Understanding Racecheck Analysis Reports</a></h3>
<div class="body conbody">
<p class="p">
In <dfn class="term">analysis</dfn> reports, the racecheck tool produces a series of high level
messages that identify the source locations of a particular race, based on observed
hazards and other machine state
</p>
<div class="p">
A sample racecheck analysis report is below:
<pre class="pre screen" xml:space="preserve">
========= Race reported between Write access at 0x00000050 in raceGroupBasic.cu:53:WAW(void)
========= and Write access at 0x00000050 in raceGroupBasic.cu:53:WAW(void)
</pre></div>
<p class="p">
The analysis record contains high level information about the hazard that is conveyed to
the end user. Each line contains information about a unique location in the application
which is participating in the race.
</p>
<div class="p">
The first line contains the type of access. The access can be either a :
<ul class="ul">
<li class="li">Read</li>
<li class="li">Write</li>
</ul>
The next item on the line is the PC of the location where the access happened from. In
this case, the PC is 0x50. If the application was compiled was compiled with line number
information, this line will also contain the file name and line number of the access.
Finally, the line contains the kernel name of the kernel containing the access.
</div>
<p class="p">
A given analysis report will always contain at least one line which is performing a write
access. A common strategy to eliminate races which contain only write accesses is to
ensure that the write access is performed by only one thread. In the case of races
with multiple readers and one writer, introducing explicit program ordering
via a <samp class="ph codeph">__syncthreads()</samp> call can avoid the race condition.
</p>
</div>
</div>
<div class="topic concept nested1" id="understanding-racecheck-hazard-reports"><a name="understanding-racecheck-hazard-reports" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#understanding-racecheck-hazard-reports" name="understanding-racecheck-hazard-reports" shape="rect">4.6. Understanding Racecheck Hazard Reports</a></h3>
<div class="body conbody">
<p class="p">
In <dfn class="term">hazard</dfn> reporting mode, the racecheck tool produces a series of messages
detailing information about hazards in the application. The tool is byte accurate and
produces a message for each byte on which a hazard was detected. Additionally, when enabled,
the host backtrace for the launch of the kernel will also be displayed.
</p>
<div class="p">
A sample racecheck hazard is below:
<pre class="pre screen" xml:space="preserve">
========= ERROR: Potential WAW hazard detected at __shared__ 0x0 in block (0, 0, 0) :
========= Write Thread (0, 0, 0) at 0x00000088 in raceWAW.cu:18:WAW(void)
========= Write Thread (1, 0, 0) at 0x00000088 in raceWAW.cu:18:WAW(void)
========= Current Value : 0, Incoming Value : 2
</pre></div>
<p class="p">
The hazard records are dense and capture a lot of interesting information.
In general terms, the first line contains information about the hazard
its severity, type and address, as well as information about the thread
block where it occurred.
The next 2 lines contain detailed information about the two threads that were
in contention. These two lines are ordered chronologically, so the first entry
is for the access that occurred earlier and the second for the access that
occurred later. The final line is printed for some hazard types and captures
the actual data that was being written.
</p>
<div class="p">
Examining this line by line, we have :
<pre class="pre screen" xml:space="preserve">ERROR: Potential WAW hazard detected at __shared__ 0x0 in block (0, 0, 0)</pre></div>
<div class="p">
The first word on this line indicates the severity of this hazard. This can be
<ul class="ul">
<li class="li"><dfn class="term">INFO</dfn> : The lowest level of severity. This is for hazards
that have no impact on program execution and hence are not contributing to
data access hazards. It is still a good idea to find and eliminate
such hazards
</li>
<li class="li"><dfn class="term">WARNING</dfn> : Hazards at this level of severity are determined
to be programming model hazards, however may be intentionally created
by the programmer. An example of this are hazards due to warp level
programming that make the assumption that threads are proceeding
in groups. Such hazards are typically only encountered by advanced
programmers. In cases where a beginner programmer encounters such
errors, he should treat them as sources of hazards.
</li>
<li class="li">ERROR : The highest level of severity. Correspond to hazards that are
very likely candidates for causing data access races. Programmers
would be well advised to examine errors at this level of severity.
</li>
</ul>
In this case, the message is at the ERROR level of severity.
</div>
<div class="p">
The next piece of information here is the type of hazard. The racecheck tool
detects three types of hazards:
<ul class="ul">
<li class="li">WAW or Write-After-Write hazards</li>
<li class="li">WAR or Write-After-Read hazards</li>
<li class="li">RAW or Read-After-Write hazards</li>
</ul>
The type of hazard indicates the accesses types of the two threads that were in
contention. In this example, the hazard is of Write-After-Write type.
</div>
<p class="p">
The next piece of information is the address in shared memory that was being
accessed. This is the offset in per block shared memory that was being accessed
by both threads. Since the racecheck tool is byte accurate, the message is only
for the byte of memory at given address. In this example, the byte being accessed
is byte 0x0 in shared memory.
</p>
<p class="p">
Finally, the first line contains the block index of the thread block to which
the two racing threads belong.
</p>
<div class="p">
The second line contains information about the first thread to write to this
location.
<pre class="pre screen" xml:space="preserve">Write Thread (0, 0, 0) at 0x00000088 in raceWAW.cu:18:WAW(void)</pre>
The first item on this line indicates the type of access being performed by
this thread to the shared memory address. In this example, the thread was
writing to the location.
The next component is the index of the thread the thread block. In this case,
the thread is at index (0,0,0). Following this, we have the byte offset of the
instruction which did the access in the kernel. In this example, the offset
is 0x88. This is followed by the source file and line number
(if line number information is available). The
final item on this line is the name of the kernel that was being executed.
</div>
<p class="p">
The third line contains similar information about the second thread
which was causing this hazard. This line has an identical format to
the previous line.
</p>
<div class="p">
The fourth line contains information about the data in the two accesses.
<pre class="pre screen" xml:space="preserve">Current Value : 0, Incoming Value : 2</pre>
If the second thread in the hazard was performing a write access, i.e.
the hazard is a Write-After-Write (WAW) or a Write-After-Read (WAR)
this line contains the value after the access by the first thread
as the <dfn class="term">Current Value</dfn> and the value that will be written
by the second access as the <dfn class="term">Incoming Value</dfn>.
In this case, the first thread wrote the value 0 to the shared memory location.
The second thread is attempting to write the value 2.
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="cuda-memcheck-features"><a name="cuda-memcheck-features" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#cuda-memcheck-features" name="cuda-memcheck-features" shape="rect">5. CUDA-MEMCHECK Features</a></h2>
<div class="topic concept nested1" id="nonblocking-mode"><a name="nonblocking-mode" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#nonblocking-mode" name="nonblocking-mode" shape="rect">5.1. Nonblocking Mode</a></h3>
<div class="body conbody">
<p class="p"> By default, on SM 2.0 and higher GPUs the standalone CUDA-MEMCHECK tool will launch
kernels in nonblocking mode. This allows the tool to support error reporting in
applications running concurrent kernels
</p>
<p class="p">
To force kernels to execute serially, a user can use the
<samp class="ph codeph">--force-blocking-launches yes</samp> option. Blocking launch mode is
always enabled on Mac OS X 10.6 and on Windows XP. This flag has no effect on
GPUs less than SM 2.0. One side effect is that when in blocking mode, only the
first thread to hit an error in a kernel will be reported.
</p>
</div>
</div>
<div class="topic concept nested1" id="stack-backtraces"><a name="stack-backtraces" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#stack-backtraces" name="stack-backtraces" shape="rect">5.2. Stack Backtraces</a></h3>
<div class="body conbody">
<p class="p">In standalone mode, CUDA-MEMCHECK can generate backtraces when given
<samp class="ph codeph">--show-backtrace</samp> option. Backtraces usually consist of
two sections - a saved host backtrace that leads upto the CUDA driver
call site, and a device backtrace at the time of the error. Each backtrace
contains a list of function calls showing the state of the stack at the
time the backtrace was created.
</p>
<p class="p">To get function names in the host backtraces, the user application must be
built with support for symbol information in the host application. For more
information, see <a class="xref" href="index.html#compilation-options" shape="rect">Compilation Options</a></p>
<p class="p">
In CUDA 5, the host stack backtrace will show a maximum of 61 frames.
</p>
<p class="p">
Backtraces are printed for most CUDA-MEMCHECK tool outputs, and the information
generated varies depending on the type of output. The table below explains the
kind of host and device backtrace seen under different conditions.
</p>
<div class="tablenoborder"><a name="stack-backtraces__cuda-memcheck-stack-backtrace-information" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="stack-backtraces__cuda-memcheck-stack-backtrace-information" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 6. CUDA-MEMCHECK Stack Backtrace Information</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="25%" id="d54e1753" rowspan="1" colspan="1">Output Type</th>
<th class="entry" valign="top" width="37.5%" id="d54e1756" rowspan="1" colspan="1">Host Backtrace</th>
<th class="entry" valign="top" width="37.5%" id="d54e1759" rowspan="1" colspan="1">Device Backtrace</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">Memory access error</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Kernel launch on host</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">Precise backtrace on device</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">Hardware exception</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Kernel launch on host</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">Imprecise backtrace on device
<a name="fnsrc_1" href="#fntarg_1" shape="rect"><sup>1</sup></a></td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">Malloc/Free error</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Kernel launch on host</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">Precise backtrace on device</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">cudaMalloc allocation leak</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Callsite of cudaMalloc</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">N/A</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">CUDA API error</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Callsite of CUDA API call</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">N/A</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">CUDA-MEMCHECK internal error</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Callsite leading to internal error</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">N/A</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">Device heap allocation leak</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">N/A</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="25%" headers="d54e1753" rowspan="1" colspan="1">Shared memory hazard</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1756" rowspan="1" colspan="1">Kernel launch on host</td>
<td class="entry" valign="top" width="37.5%" headers="d54e1759" rowspan="1" colspan="1">N/A</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested1" id="name-demangling"><a name="name-demangling" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#name-demangling" name="name-demangling" shape="rect">5.3. Name Demangling</a></h3>
<div class="body conbody">
<p class="p">
The CUDA-MEMCHECK suite now supports displaying mangled and demangled names for
CUDA kernels and CUDA device functions.
By default, tools display the fully demangled name, which contains the name
of the kernel as well as its prototype information. In the simple demangle
mode, the tools will only display the first part of the name. If demangling
is disabled, tools will display the complete mangled name of the kernel.
</p>
<p class="p"></p>
</div>
</div>
<div class="topic concept nested1" id="dynamic-parallelism"><a name="dynamic-parallelism" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#dynamic-parallelism" name="dynamic-parallelism" shape="rect">5.4. Dynamic Parallelism</a></h3>
<div class="body conbody">
<p class="p">The CUDA-MEMCHECK tool suite supports dynamic parallelism. The
<dfn class="term">memcheck</dfn> tool supports precise error reporting of out of bounds
and misaligned accesses on global, local and shared memory accesses as well as
on global atomic instructions for applications using dynamic parallelism.
In addition, the imprecise hardware exception reporting mechanism is also fully
supported. Error detection on applications using dynamic parallelism requires
significantly more memory on the device and as a result, in memory constrained
environments, <dfn class="term">memcheck</dfn> may fail to initialize with an internal out
of memory error.
</p>
<p class="p">
For limitations, see <a class="xref" href="index.html#known-issues" shape="rect">Known Issues</a>.
</p>
</div>
</div>
<div class="topic concept nested1" id="error-actions"><a name="error-actions" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#error-actions" name="error-actions" shape="rect">5.5. Error Actions</a></h3>
<div class="body conbody">
<p class="p">
On encountering an error, CUDA-MEMCHECK behavior depends on the type of error.
The default behavior of CUDA-MEMCHECK is to continue execution on purely host
side errors. Hardware exceptions detected by the memcheck tool cause the
CUDA context to be destroyed. Precise errors (such as memory access and
malloc/free errors) detected by the memcheck tool cause the kernel to be terminated.
This terminates the kernel without running any subsequent instructions and the
application continues launching other kernels in the CUDA context.
The handling of memory access and malloc/free errors detected by the memcheck tool
can be changed using the <samp class="ph codeph">--destroy-on-device-error</samp> option.
</p>
<p class="p">
For racecheck detected hazards, the hazard is reported, but execution is
not affected.
</p>
<div class="p">
For a full summary of error action, based on the type of the error see the
table below. The error action <dfn class="term">terminate kernel</dfn> refers to the
cases where the kernel is terminated early, and no subsequent instructions
are run. In such cases, the CUDA context is not destroyed and other kernels
continue execution and CUDA API calls can still be made.
<div class="note note"><span class="notetitle">Note:</span>
When kernel execution is terminated early, the application may not have
completed its computations on data. Any subsequent kernels that depend
on this data will have undefined behavior.
</div>
The action <dfn class="term">terminate CUDA context</dfn> refers to the cases where the
CUDA context is forcibly terminated. In such cases, all outstanding work for
the context is terminated and subsequent CUDA API calls will fail.
The action <dfn class="term">continue application</dfn> refers to cases where the
application execution is not impacted, and the kernel continues executing
instructions.
</div>
<p class="p"></p>
<div class="tablenoborder"><a name="error-actions__cuda-memcheck-error-action" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="error-actions__cuda-memcheck-error-action" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 7. CUDA-MEMCHECK Error Actions</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="30.76923076923077%" id="d54e1963" rowspan="1" colspan="1">Error Type</th>
<th class="entry" valign="top" width="15.384615384615385%" id="d54e1966" rowspan="1" colspan="1">Location</th>
<th class="entry" valign="top" width="23.076923076923077%" id="d54e1969" rowspan="1" colspan="1">Action</th>
<th class="entry" valign="top" width="30.76923076923077%" id="d54e1972" rowspan="1" colspan="1">Comments</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">Memory access error</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Terminate kernel</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">User can choose to instead terminate the CUDA context</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">Hardware exception</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Terminate CUDA context</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">Subsequent calls on the CUDA context will fail</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">Malloc/Free error</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Terminate kernel</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">User can choose to instead terminate the CUDA context</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">cudaMalloc allocation leak</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Host</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Continue application</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">Error reported. No other action taken.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">CUDA API error</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Host</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Continue application</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">Error reported. No other action taken.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">CUDA-MEMCHECK internal error</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Host</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Undefined</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">The application may behave in an undefined fashion</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">Device heap allocation leak</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Continue application</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">Error reported. No other action taken.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1963" rowspan="1" colspan="1">Shared memory hazard</td>
<td class="entry" valign="top" width="15.384615384615385%" headers="d54e1966" rowspan="1" colspan="1">Device</td>
<td class="entry" valign="top" width="23.076923076923077%" headers="d54e1969" rowspan="1" colspan="1">Continue application</td>
<td class="entry" valign="top" width="30.76923076923077%" headers="d54e1972" rowspan="1" colspan="1">Error reported. No other action taken.</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="cuda-memcheck-tool-examples"><a name="cuda-memcheck-tool-examples" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#cuda-memcheck-tool-examples" name="cuda-memcheck-tool-examples" shape="rect">6. CUDA-MEMCHECK Tool Examples</a></h2>
<div class="topic concept nested1" id="example-use-of-memcheck"><a name="example-use-of-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#example-use-of-memcheck" name="example-use-of-memcheck" shape="rect">6.1. Example Use of Memcheck</a></h3>
<div class="body conbody">
<div class="section">
<div class="p">
This section presents a walk-through of running the memcheck tool from
CUDA-MEMCHECK on a simple application called <samp class="ph codeph">memcheck_demo</samp>.
<div class="note note"><span class="notetitle">Note:</span> Depending on the SM type of your GPU, your system output may vary.
</div>
</div>
</div>
<div class="example">
<h3 class="title sectiontitle">memcheck_demo.cu source code</h3><pre xml:space="preserve">#include <stdio.h>
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__device__</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> x;
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__global__</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> unaligned_kernel(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>) {
*(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>*) ((<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">char</span>*)&x + 1) = 42;
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__device__</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> out_of_bounds_function(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>) {
*(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span>*) 0x87654320 = 42;
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">__global__</span> <span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> out_of_bounds_kernel(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>) {
out_of_bounds_function();
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> run_unaligned(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>) {
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Running unaligned_kernel\n"</span>);
unaligned_kernel<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute"><<<</span>1,1<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">>>></span>();
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Ran unaligned_kernel: %s\n"</span>,
cudaGetErrorString(cudaGetLastError()));
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Sync: %s\n"</span>, cudaGetErrorString(cudaThreadSynchronize()));
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span> run_out_of_bounds(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>) {
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Running out_of_bounds_kernel\n"</span>);
out_of_bounds_kernel<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute"><<<</span>1,1<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-attribute">>>></span>();
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Ran out_of_bounds_kernel: %s\n"</span>,
cudaGetErrorString(cudaGetLastError()));
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Sync: %s\n"</span>, cudaGetErrorString(cudaThreadSynchronize()));
}
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> main() {
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">int</span> *devMem;
printf(<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-string">"Mallocing memory\n"</span>);
cudaMalloc((<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">void</span>**)&devMem, 1024);
run_unaligned();
run_out_of_bounds();
cudaDeviceReset();
cudaFree(devMem);
<span xmlns:xslthl="http://xslthl.sf.net" class="xslthl-keyword">return</span> 0;
}</pre></div>
<div class="section">
<div class="p">
This application is compiled for release builds as :
<pre class="pre screen" xml:space="preserve">
nvcc -gencode arch=compute_20,code=sm_20 -o memcheck_demo memcheck_demo.cu
</pre></div>
</div>
</div>
<div class="topic concept nested2" id="memcheck-demo-output"><a name="memcheck-demo-output" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memcheck-demo-output" name="memcheck-demo-output" shape="rect">6.1.1. <samp class="ph codeph">memcheck_demo</samp> Output
</a></h3>
<div class="body conbody">
<p class="p">When a CUDA application causes access violations, the kernel launch may
terminate with an error code of unspecified launch failure or a subsequent
cudaThreadSynchronize call which will fail with an error code of unspecified
launch failure.
</p>
<p class="p">This sample application is causing two failures but there is no way to detect
where these kernels are causing the access violations, as illustrated in the
following output:
</p><pre class="pre screen" xml:space="preserve"><strong class="ph b">$ ./memcheck_demo</strong>
Mallocing memory
Running unaligned_kernel
Ran unaligned_kernel: no error
Sync: unspecified launch failure
Running out_of_bounds_kernel
Ran out_of_bounds_kernel: unspecified launch failure
Sync: unspecified launch failure
</pre></div>
</div>
<div class="topic concept nested2" id="memcheck-demo-output-with-memcheck-release-build"><a name="memcheck-demo-output-with-memcheck-release-build" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memcheck-demo-output-with-memcheck-release-build" name="memcheck-demo-output-with-memcheck-release-build" shape="rect">6.1.2. <samp class="ph codeph">memcheck_demo</samp> Output with Memcheck (Release Build)</a></h3>
<div class="body conbody">
<p class="p">In this case, since the application is built in release mode, the
CUDA-MEMCHECK output contains only the kernel names from the application causing
the access violation. Though the kernel name and error type are detected, there
is no line number information on the failing kernel. Also included in the output
are the host and device backtraces for the call sites where the functions were
launched. In addition, CUDA API errors are reported, such as the invalid
<samp class="ph codeph">cudaFree()</samp> call in the application.
</p><pre class="pre screen" xml:space="preserve">
<strong class="ph b">$ cuda-memcheck ./memcheck_demo</strong>
========= CUDA-MEMCHECK
Mallocing memory
Running unaligned_kernel
Ran unaligned_kernel: no error
Sync: no error
Running out_of_bounds_kernel
Ran out_of_bounds_kernel: no error
Sync: no error
========= Invalid __global__ write of size 4
========= at 0x00000028 in unaligned_kernel(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x400100001 is misaligned
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/local/lib/libcuda.so (cuLaunchKernel + 0x3ae) [0xddbee]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xcd27]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaLaunch + 0x1bb) [0x3778b]
========= Host Frame:memcheck_demo [0xdfc]
========= Host Frame:memcheck_demo [0xc76]
========= Host Frame:memcheck_demo [0xc81]
========= Host Frame:memcheck_demo [0xb03]
========= Host Frame:memcheck_demo [0xc27]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0x9b9]
=========
========= Invalid __global__ write of size 4
========= at 0x00000010 in out_of_bounds_kernel(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0xffffffff87654320 is out of bounds
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/local/lib/libcuda.so (cuLaunchKernel + 0x3ae) [0xddbee]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xcd27]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaLaunch + 0x1bb) [0x3778b]
========= Host Frame:memcheck_demo [0xdfc]
========= Host Frame:memcheck_demo [0xca0]
========= Host Frame:memcheck_demo [0xcab]
========= Host Frame:memcheck_demo [0xbbc]
========= Host Frame:memcheck_demo [0xc2c]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0x9b9]
=========
========= Program hit error 17 on CUDA API call to cudaFree
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/usr/local/lib/libcuda.so [0x28f850]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaFree + 0x20d) [0x364ed]
========= Host Frame:memcheck_demo [0xc3d]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0x9b9]
=========
========= ERROR SUMMARY: 3 errors
</pre></div>
</div>
<div class="topic concept nested2" id="memcheck-demo-output-with-memcheck-debug-build"><a name="memcheck-demo-output-with-memcheck-debug-build" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#memcheck-demo-output-with-memcheck-debug-build" name="memcheck-demo-output-with-memcheck-debug-build" shape="rect">6.1.3. <samp class="ph codeph">memcheck_demo</samp> Output with Memcheck (Debug Build)</a></h3>
<div class="body conbody">
<div class="p"> The application is now built with device side debug information and function
symbols as :
<pre class="pre screen" xml:space="preserve">
nvcc -G -Xcompiler -rdynamic -gencode arch=compute_20,code=sm_20 -o memcheck_demo memcheck_demo.cu
</pre></div>
<p class="p">Now run this application with CUDA-MEMCHECK and check the output. By default, the
application will run so that the kernel is terminated on memory access errors
but other work in the CUDA context can still proceed.
</p>
<p class="p"> In the output below the first kernel no longer reports an unspecified launch failure
as its execution has been terminated early after CUDA-MEMCHECK detected the error.
The application continued to run the second kernel. The error detected in the
second kernel causes it to terminate early. Finally, the application calls
<samp class="ph codeph">cudaDeviceReset()</samp>, which destroys the CUDA context and then
attempts to call <samp class="ph codeph">cudaFree()</samp>. This call returns an API error
that is caught and displayed by memcheck.
</p><pre class="pre screen" xml:space="preserve">
<strong class="ph b">$ cuda-memcheck ./memcheck_demo</strong>
========= CUDA-MEMCHECK
Mallocing memory
Running unaligned_kernel
Ran unaligned_kernel: no error
Sync: no error
Running out_of_bounds_kernel
Ran out_of_bounds_kernel: no error
Sync: no error
========= Invalid __global__ write of size 4
========= at 0x00000028 in memcheck_demo.cu:6:unaligned_kernel(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x400100001 is misaligned
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/local/lib/libcuda.so (cuLaunchKernel + 0x3ae) [0xddbee]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xcd27]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaLaunch + 0x1bb) [0x3778b]
========= Host Frame:memcheck_demo (_Z10cudaLaunchIcE9cudaErrorPT_ + 0x18) [0x11a4]
========= Host Frame:memcheck_demo (_Z35__device_stub__Z16unaligned_kernelvv + 0x1d) [0x101d]
========= Host Frame:memcheck_demo (_Z16unaligned_kernelv + 0x9) [0x1028]
========= Host Frame:memcheck_demo (_Z13run_unalignedv + 0x76) [0xeaa]
========= Host Frame:memcheck_demo (main + 0x28) [0xfce]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xd79]
=========
========= Invalid __global__ write of size 4
========= at 0x00000028 in memcheck_demo.cu:10:out_of_bounds_function(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x87654320 is out of bounds
========= Device Frame:memcheck_demo.cu:15:out_of_bounds_kernel(void) (out_of_bounds_kernel(void) : 0x10)
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/local/lib/libcuda.so (cuLaunchKernel + 0x3ae) [0xddbee]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xcd27]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaLaunch + 0x1bb) [0x3778b]
========= Host Frame:memcheck_demo (_Z10cudaLaunchIcE9cudaErrorPT_ + 0x18) [0x11a4]
========= Host Frame:memcheck_demo (_Z39__device_stub__Z20out_of_bounds_kernelvv + 0x1d) [0x1047]
========= Host Frame:memcheck_demo (_Z20out_of_bounds_kernelv + 0x9) [0x1052]
========= Host Frame:memcheck_demo (_Z17run_out_of_boundsv + 0x76) [0xf63]
========= Host Frame:memcheck_demo (main + 0x2d) [0xfd3]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xd79]
=========
========= Program hit error 17 on CUDA API call to cudaFree
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/usr/local/lib/libcuda.so [0x28f850]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaFree + 0x20d) [0x364ed]
========= Host Frame:memcheck_demo (main + 0x3e) [0xfe4]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xd79]
=========
========= ERROR SUMMARY: 3 errors
</pre></div>
</div>
<div class="topic concept nested2" id="leak-checking-in-cuda-memcheck"><a name="leak-checking-in-cuda-memcheck" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#leak-checking-in-cuda-memcheck" name="leak-checking-in-cuda-memcheck" shape="rect">6.1.4. Leak Checking in CUDA-MEMCHECK</a></h3>
<div class="body conbody">
<p class="p">To print information about the allocations that have not been freed at the time
the CUDA context is destroyed, we can specify the <samp class="ph codeph">--leak-check full</samp>
option to CUDA-MEMCHECK.
</p>
<p class="p">When running the program with the leak check option, the user is presented with
a list of allocations that were not destroyed, along with the size of the allocation
and the address on the device of the allocation. For allocations made on the host,
each leak report will also print a backtrace corresponding to the saved host stack
at the time the allocation was first made. Also presented is a summary of the total
number of bytes leaked and the corresponding number of allocations.
</p>
<p class="p">In this example, the program created an allocation using
<samp class="ph codeph">cudaMalloc()</samp> and has not called <samp class="ph codeph">cudaFree()</samp>
to release it, leaking memory. Notice that CUDA-MEMCHECK still prints errors
it encountered while running the application.
</p><pre class="pre screen" xml:space="preserve">
<strong class="ph b">$ cuda-memcheck --leak-check full memcheck_demo</strong>
========= CUDA-MEMCHECK
Mallocing memory
Running unaligned_kernel
Ran unaligned_kernel: no error
Sync: no error
Running out_of_bounds_kernel
Ran out_of_bounds_kernel: no error
Sync: no error
========= Invalid __global__ write of size 4
========= at 0x00000060 in memcheck_demo.cu:6:unaligned_kernel(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x400100001 is misaligned
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/local/lib/libcuda.so (cuLaunchKernel + 0x3ae) [0xddbee]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xcd27]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaLaunch + 0x1bb) [0x3778b]
========= Host Frame:memcheck_demo (_Z10cudaLaunchIcE9cudaErrorPT_ + 0x18) [0x122c]
========= Host Frame:memcheck_demo (_Z35__device_stub__Z16unaligned_kernelvv + 0x1d) [0x10a6]
========= Host Frame:memcheck_demo (_Z16unaligned_kernelv + 0x9) [0x10b1]
========= Host Frame:memcheck_demo (_Z13run_unalignedv + 0x76) [0xf33]
========= Host Frame:memcheck_demo (main + 0x28) [0x1057]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xde9]
=========
========= Invalid __global__ write of size 4
========= at 0x00000028 in memcheck_demo.cu:10:out_of_bounds_function(void)
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x87654320 is out of bounds
========= Device Frame:memcheck_demo.cu:15:out_of_bounds_kernel(void) (out_of_bounds_kernel(void) : 0x10)
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame:/usr/local/lib/libcuda.so (cuLaunchKernel + 0x3ae) [0xddbee]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xcd27]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaLaunch + 0x1bb) [0x3778b]
========= Host Frame:memcheck_demo (_Z10cudaLaunchIcE9cudaErrorPT_ + 0x18) [0x122c]
========= Host Frame:memcheck_demo (_Z39__device_stub__Z20out_of_bounds_kernelvv + 0x1d) [0x10d0]
========= Host Frame:memcheck_demo (_Z20out_of_bounds_kernelv + 0x9) [0x10db]
========= Host Frame:memcheck_demo (_Z17run_out_of_boundsv + 0x76) [0xfec]
========= Host Frame:memcheck_demo (main + 0x2d) [0x105c]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xde9]
=========
========= Leaked 1024 bytes at 0x400200000
========= Saved host backtrace up to driver entry point at cudaMalloc time
========= Host Frame:/usr/local/lib/libcuda.so (cuMemAlloc_v2 + 0x236) [0xe9746]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0x26dd7]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 [0xb37b]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaMalloc + 0x17a) [0x36e6a]
========= Host Frame:memcheck_demo (main + 0x23) [0x1052]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xde9]
=========
========= Program hit error 17 on CUDA API call to cudaFree
========= Saved host backtrace up to driver entry point at error
========= Host Frame:/usr/local/lib/libcuda.so [0x28f850]
========= Host Frame:/usr/local/lib/libcudart.so.5.0 (cudaFree + 0x20d) [0x364ed]
========= Host Frame:memcheck_demo (main + 0x3e) [0x106d]
========= Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xfd) [0x1eb1d]
========= Host Frame:memcheck_demo [0xde9]
=========
========= LEAK SUMMARY: 1024 bytes leaked in 1 allocations
========= ERROR SUMMARY: 3 errors
</pre></div>
</div>
</div>
<div class="topic concept nested1" id="integrated-cuda-memcheck-example"><a name="integrated-cuda-memcheck-example" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#integrated-cuda-memcheck-example" name="integrated-cuda-memcheck-example" shape="rect">6.2. Integrated CUDA-MEMCHECK Example</a></h3>
<div class="body conbody">
<p class="p">This example shows how to enable CUDA-MEMCHECK from within CUDA-GDB and how
to detect errors within the debugger so you can access the line number
information and check the state of the variables
</p>
<p class="p">In this example the unaligned kernel has a misaligned memory access in
block 1 lane 1, which gets trapped as an illegal lane address at line 6
from within CUDA-GDB. Note that CUDA-GDB displays the address and
that caused the bad access.
</p><pre class="pre screen" xml:space="preserve">
(cuda-gdb) <strong class="ph b">set cuda memcheck on</strong>
(cuda-gdb) <strong class="ph b">run</strong>
Starting program: memcheck_demo
[Thread debugging using libthread_db enabled]
Mallocing memory
[New Thread 0x7ffff6fe1710 (LWP 7783)]
[Context Create of context 0x6218a0 on Device 0]
[Launch of CUDA Kernel 0 (memset32_post<<<(1,1,1),(64,1,1)>>>) on Device 0]
Running unaligned_kernel
[Launch of CUDA Kernel 1 (unaligned_kernel<<<(1,1,1),(1,1,1)>>>) on Device 0]
Memcheck detected an illegal access to address (@global)0x400100001
Program received signal CUDA_EXCEPTION_1, Lane Illegal Address.
[Switching focus to CUDA kernel 1, grid 2, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
0x000000000078b8b0 in unaligned_kernel<<<(1,1,1),(1,1,1)>>> () at memcheck_demo.cu:6
6 *(int*) ((char*)&x + 1) = 42;
(cuda-gdb) <strong class="ph b">print &x</strong>
$1 = (@global int *) 0x400100000
(cuda-gdb) <strong class="ph b">continue</strong>
Continuing.
[Termination of CUDA Kernel 1 (unaligned_kernel<<<(1,1,1),(1,1,1)>>>) on Device 0]
[Termination of CUDA Kernel 0 (memset32_post<<<(1,1,1),(64,1,1)>>>) on Device 0]
Program terminated with signal CUDA_EXCEPTION_1, Lane Illegal Address.
The program no longer exists.
(cuda-gdb)
</pre></div>
</div>
</div>
<div class="topic concept nested0" id="memory-access-error-reporting"><a name="memory-access-error-reporting" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#memory-access-error-reporting" name="memory-access-error-reporting" shape="rect">A. Memory Access Error Reporting</a></h2>
<div class="body conbody">
<div class="p">The memcheck tool will report memory access errors when run standalone
or in integrated mode with CUDA-GDB. The table below describes the types of
accesses that are checked and the SM version where such checks happen
<div class="tablenoborder"><a name="memory-access-error-reporting__memcheck-memory-access-error-detection-support" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="memory-access-error-reporting__memcheck-memory-access-error-detection-support" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 8. Memcheck memory access error detection support</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="33.33333333333333%" id="d54e2326" rowspan="1" colspan="1">Error Type</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e2329" rowspan="1" colspan="1">SM 1.x</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e2332" rowspan="1" colspan="1">SM 2.x</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e2335" rowspan="1" colspan="1">SM 3.0</th>
<th class="entry" valign="top" width="16.666666666666664%" id="d54e2338" rowspan="1" colspan="1">SM 3.5</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2326" rowspan="1" colspan="1">Global</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2329" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2332" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2335" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2338" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2326" rowspan="1" colspan="1">Shared</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2329" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2332" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2335" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2338" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2326" rowspan="1" colspan="1">Local</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2329" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2332" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2335" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2338" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2326" rowspan="1" colspan="1">Global Atomic</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2329" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2332" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2335" rowspan="1" colspan="1">Yes</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2338" rowspan="1" colspan="1">Yes</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="33.33333333333333%" headers="d54e2326" rowspan="1" colspan="1">Load through texture</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2329" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2332" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2335" rowspan="1" colspan="1">N/A</td>
<td class="entry" valign="top" width="16.666666666666664%" headers="d54e2338" rowspan="1" colspan="1">Yes</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
</div>
<div class="topic concept nested0" id="hardware-exception-reporting"><a name="hardware-exception-reporting" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#hardware-exception-reporting" name="hardware-exception-reporting" shape="rect">B. Hardware Exception Reporting</a></h2>
<div class="body conbody">
<p class="p">The CUDA‐MEMCHECK tool will report hardware exceptions when run as a
standalone or as part of CUDA‐GDB. The table below enumerates the
supported exceptions, their precision and scope, as well as a brief
description of their cause. For more detailed information, see the
documentation for CUDA‐GDB.
</p>
<div class="tablenoborder"><a name="hardware-exception-reporting__cuda-exception-codes" shape="rect">
<!-- --></a><table cellpadding="4" cellspacing="0" summary="" id="hardware-exception-reporting__cuda-exception-codes" class="table" frame="border" border="1" rules="all">
<caption><span class="tablecap">Table 9. CUDA Exception Codes</span></caption>
<thead class="thead" align="left">
<tr class="row">
<th class="entry" valign="top" width="28.57142857142857%" id="d54e2470" rowspan="1" colspan="1">Exception code</th>
<th class="entry" valign="top" width="14.285714285714285%" id="d54e2473" rowspan="1" colspan="1">Precision of the Error</th>
<th class="entry" valign="top" width="21.428571428571427%" id="d54e2476" rowspan="1" colspan="1">Scope of the Error</th>
<th class="entry" valign="top" width="35.714285714285715%" id="d54e2479" rowspan="1" colspan="1">Description</th>
</tr>
</thead>
<tbody class="tbody">
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_1 : "Lane Illegal Address"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Per lane/thread error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when a thread accesses an illegal (out of bounds) global address.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_2 : "Lane User StackOverflow"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Per lane/thread error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when a thread exceeds its stack memory limit.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_3: "Device Hardware Stack Overflow" </td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Global error on the GPU</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when the application triggers a global hardware stack overflow. The main cause of this error is large amounts
of divergence in the presence of function calls.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_4: "Warp Illegal Instruction"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Warp error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread within a warp has executed an illegal instruction.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_5: "Warp Out-of-range Address"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Warp error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread within a warp accesses an address that is outside the valid range of local or shared memory regions.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_6: "Warp Misaligned Address"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Warp error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread within a warp accesses an address in the local or shared memory segments that is not correctly
aligned.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_7: "Warp Invalid Address Space"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Warp error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread within a warp executes an instruction that accesses a memory space not permitted for that instruction.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_8: "Warp Invalid PC"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Warp error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread within a warp advances its PC beyond the 40-bit address space.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_9: "Warp Hardware Stack Overflow"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Warp error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread in a warp triggers a hardware stack overflow. This should be a rare occurrence.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_10: "Device Illegal Address"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Global error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when a thread accesses an illegal (out of bounds) global address.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_11: "Lane Misaligned Address"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Per lane/thread error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when a thread accesses a global address that is not correctly aligned.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_12: "Warp Assert"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Per warp</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when any thread in the warp hits a device side assertion.</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">CUDA_EXCEPTION_13: "Lane Syscall Error"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Per lane</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">This occurs when a particular thread causes an syscall error, such as calling <samp class="ph codeph">free()</samp> in a kernel on an already free'd pointer.
</td>
</tr>
<tr class="row">
<td class="entry" valign="top" width="28.57142857142857%" headers="d54e2470" rowspan="1" colspan="1">"Unknown Exception"</td>
<td class="entry" valign="top" width="14.285714285714285%" headers="d54e2473" rowspan="1" colspan="1">Not precise</td>
<td class="entry" valign="top" width="21.428571428571427%" headers="d54e2476" rowspan="1" colspan="1">Global error</td>
<td class="entry" valign="top" width="35.714285714285715%" headers="d54e2479" rowspan="1" colspan="1">The precise cause of the exception is unknown. Potentially, this may be due to Device Hardware Stack overflows or a kernel
generating an exception very close to its termination.
</td>
</tr>
</tbody>
</table>
</div>
</div>
</div>
<div class="topic concept nested0" id="release-notes"><a name="release-notes" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#release-notes" name="release-notes" shape="rect">C. Release Notes</a></h2>
<div class="body conbody"></div>
<div class="topic concept nested1" id="new-features-5.5"><a name="new-features-5.5" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#new-features-5.5" name="new-features-5.5" shape="rect">C.1. New Features in 5.5</a></h3>
<div class="body conbody">
<ul class="ul">
<li class="li"> Analysis mode in racecheck tool. For more information,
see <a class="xref" href="index.html#racecheck-tool" shape="rect">Racecheck Tool</a></li>
<li class="li"> Support for racecheck on SM 3.5 GPUs.
</li>
</ul>
</div>
</div>
<div class="topic concept nested1" id="new-features-5.0"><a name="new-features-5.0" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#new-features-5.0" name="new-features-5.0" shape="rect">C.2. New Features in 5.0</a></h3>
<div class="body conbody">
<ul class="ul">
<li class="li">Reporting of data access hazards in shared memory accesses.
This is supported on Fermi SM 2.x and Kepler SM 3.0 GPUs. This functionality
is not supported on Windows XP.
For more information, see <a class="xref" href="index.html#racecheck-tool" shape="rect">Racecheck Tool</a>.
</li>
<li class="li">Support for SM 3.0 and SM 3.5 GPUs. For more information, see
<a class="xref" href="index.html#supported-devices" shape="rect">Supported Devices</a>.
</li>
<li class="li">Support for dynamic parallelism. All memory access error detection
is supported for applications using dynamic parallelism.
For more information, see <a class="xref" href="index.html#dynamic-parallelism" shape="rect">Dynamic Parallelism</a>.
</li>
<li class="li">Precise error detection for local loads/stores, shared loads/stores,
global atomics/reductions. On SM 3.5, added precise memory access error
detection for noncoherent global loads through the texture unit.
For more information, see <a class="xref" href="index.html#memory-access-error-reporting" shape="rect">Memory Access Error Reporting</a>.
</li>
<li class="li">Error detection in device side malloc()/free(), such
as double free() or invalid free() on the GPU.
For more information, see <a class="xref" href="index.html#device-side-allocation-checking" shape="rect">Device Side Allocation Checking</a>.
</li>
<li class="li">Leak checking for allocations on the device heap.
For more information, see <a class="xref" href="index.html#leak-checking" shape="rect">Leak Checking</a>.
</li>
<li class="li">Display of a saved stack backtrace on the host
and captured backtrace on the device for different errors.
For more information, see <a class="xref" href="index.html#stack-backtraces" shape="rect">Stack Backtraces</a>.
</li>
<li class="li">Reporting of CUDA API errors in the user's application.
For more information, see <a class="xref" href="index.html#api-error-checking" shape="rect">CUDA API Error Checking</a>.
</li>
<li class="li">Added display of mangled, demangled, and full prototype of the kernel.
For more information, see <a class="xref" href="index.html#name-demangling" shape="rect">Name Demangling</a>.
</li>
<li class="li">Increased functionality in integrated mode with CUDA-GDB. Added reporting of
the address and address space being accessed that caused a precise exception.
Added checking of device side malloc() and free() when in integrated mode.
For more information, see <a class="xref" href="index.html#integrated-mode" shape="rect">Integrated Mode</a>.
</li>
<li class="li">Support for applications compiled separately that use the device
side linker.
</li>
<li class="li">Support for applications compiled with the -lineinfo flag.
</li>
<li class="li">New style of command line options.
For more information, see <a class="xref" href="index.html#command-line-options" shape="rect">Command Line Options</a>.
</li>
<li class="li">Changed default behavior. CUDA-MEMCHECK will
display backtraces by default and will report API errors by default.
For more information, see <a class="xref" href="index.html#command-line-options" shape="rect">Command Line Options</a>.
</li>
</ul>
</div>
</div>
</div>
<div class="topic concept nested0" id="known-issues"><a name="known-issues" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#known-issues" name="known-issues" shape="rect">D. Known Issues</a></h2>
<div class="body conbody">
<p class="p">The following are known issues with the current release.</p>
<ul class="ul">
<li class="li">Applications run much slower under CUDA-MEMCHECK tools. This may cause
some kernel launches to fail with a launch timeout error when running with
CUDA-MEMCHECK enabled.
</li>
<li class="li">On Windows XP, the standalone CUDA-MEMCHECK tools will always run in blocking
launch mode.
</li>
<li class="li">When running CUDA-MEMCHECK tools in integrated mode with CUDA-GDB,
only the <dfn class="term">memcheck</dfn> tool is enabled. Also, the following
features are disabled:
<ul class="ul">
<li class="li">Nonblocking launches</li>
<li class="li">Leak checking</li>
<li class="li">API error checking</li>
</ul>
</li>
<li class="li">
CUDA-MEMCHECK tools do not support CUDA/Direct3D interop.
</li>
<li class="li">
The memcheck tool does not support CUDA API error checking for
API calls made on the GPU using dynamic parallelism.
</li>
<li class="li">
The racecheck tool does not support dynamic parallelism.
</li>
</ul>
</div>
</div>
<div class="topic concept nested0" id="notices-header"><a name="notices-header" shape="rect">
<!-- --></a><h2 class="title topictitle1"><a href="#notices-header" name="notices-header" shape="rect">Notices</a></h2>
<div class="topic reference nested1" id="notice"><a name="notice" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#notice" name="notice" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Notice</h3>
<p class="p">ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND
SEPARATELY, "MATERIALS") ARE BEING PROVIDED "AS IS." NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE
WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS
FOR A PARTICULAR PURPOSE.
</p>
<p class="p">Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the
consequences of use of such information or for any infringement of patents or other rights of third parties that may result
from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications
mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information
previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems
without express written approval of NVIDIA Corporation.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="trademarks"><a name="trademarks" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#trademarks" name="trademarks" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Trademarks</h3>
<p class="p">NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation
in the U.S. and other countries. Other company and product names may be trademarks of
the respective companies with which they are associated.
</p>
</div>
</div>
</div>
<div class="topic reference nested1" id="copyright-past-to-present"><a name="copyright-past-to-present" shape="rect">
<!-- --></a><h3 class="title topictitle2"><a href="#copyright-past-to-present" name="copyright-past-to-present" shape="rect"></a></h3>
<div class="body refbody">
<div class="section">
<h3 class="title sectiontitle">Copyright</h3>
<p class="p">© <span class="ph">2007</span>-<span class="ph">2013</span> NVIDIA
Corporation. All rights reserved.
</p>
</div>
</div>
</div>
</div>
<div class="fn"><a name="fntarg_1" href="#fnsrc_1" shape="rect"><sup>1</sup></a> In some cases, there may be no device backtrace
</div>
<hr id="contents-end"></hr>
<div id="release-info">CUDA-MEMCHECK
(<a href="../../pdf/CUDA_Memcheck.pdf">PDF</a>)
-
CUDA Toolkit v5.5
(<a href="https://developer.nvidia.com/cuda-toolkit-archive">older</a>)
-
Last updated
July 19, 2013
-
<a href="mailto:cudatools@nvidia.com?subject=CUDA Tools Documentation Feedback: cuda-memcheck">Send Feedback</a></div>
</article>
<header id="header"><span id="company">NVIDIA</span><span id="site-title">CUDA Toolkit Documentation</span><form id="search" method="get" action="search">
<input type="text" name="search-text"></input><fieldset id="search-location">
<legend>Search In:</legend>
<label><input type="radio" name="search-type" value="site"></input>Entire Site</label>
<label><input type="radio" name="search-type" value="document"></input>Just This Document</label></fieldset>
<button type="reset">clear search</button>
<button id="submit" type="submit">search</button></form>
</header>
<nav id="site-nav">
<div class="category closed"><span class="twiddle">▷</span><a href="../index.html" title="The root of the site.">CUDA Toolkit</a></div>
<ul class="closed">
<li><a href="../cuda-toolkit-release-notes/index.html" title="The Release Notes for the CUDA Toolkit from v4.0 to today.">Release Notes</a></li>
<li><a href="../eula/index.html" title="The End User License Agreements for the NVIDIA CUDA Toolkit, the NVIDIA CUDA Samples, the NVIDIA Display Driver, and NVIDIA NSight (Visual Studio Edition).">EULA</a></li>
<li><a href="../cuda-getting-started-guide-for-linux/index.html" title="This guide discusses how to install and check for correct operation of the CUDA Development Tools on GNU/Linux systems.">Getting Started Linux</a></li>
<li><a href="../cuda-getting-started-guide-for-mac-os-x/index.html" title="This guide discusses how to install and check for correct operation of the CUDA Development Tools on Mac OS X systems.">Getting Started Mac OS X</a></li>
<li><a href="../cuda-getting-started-guide-for-microsoft-windows/index.html" title="This guide discusses how to install and check for correct operation of the CUDA Development Tools on Microsoft Windows systems.">Getting Started Windows</a></li>
<li><a href="../cuda-c-programming-guide/index.html" title="This guide provides a detailed discussion of the CUDA programming model and programming interface. It then describes the hardware implementation, and provides guidance on how to achieve maximum performance. The Appendixes include a list of all CUDA-enabled devices, detailed description of all extensions to the C language, listings of supported mathematical functions, C++ features supported in host and device code, details on texture fetching, technical specifications of various devices, and concludes by introducing the low-level driver API.">Programming Guide</a></li>
<li><a href="../cuda-c-best-practices-guide/index.html" title="This guide presents established parallelization and optimization techniques and explains coding metaphors and idioms that can greatly simplify programming for CUDA-capable GPU architectures. The intent is to provide guidelines for obtaining the best performance from NVIDIA GPUs using the CUDA Toolkit.">Best Practices Guide</a></li>
<li><a href="../kepler-compatibility-guide/index.html" title="This application note is intended to help developers ensure that their NVIDIA CUDA applications will run effectively on GPUs based on the NVIDIA Kepler Architecture. This document provides guidance to ensure that your software applications are compatible with Kepler.">Kepler Compatibility Guide</a></li>
<li><a href="../kepler-tuning-guide/index.html" title="Kepler is NVIDIA's next-generation architecture for CUDA compute applications. Applications that follow the best practices for the Fermi architecture should typically see speedups on the Kepler architecture without any code changes. This guide summarizes the ways that an application can be fine-tuned to gain additional speedups by leveraging Kepler architectural features.">Kepler Tuning Guide</a></li>
<li><a href="../parallel-thread-execution/index.html" title="This guide provides detailed instructions on the use of PTX, a low-level parallel thread execution virtual machine and instruction set architecture (ISA). PTX exposes the GPU as a data-parallel computing device.">PTX ISA</a></li>
<li><a href="../optimus-developer-guide/index.html" title="This document explains how CUDA APIs can be used to query for GPU capabilities in NVIDIA Optimus systems.">Developer Guide for Optimus</a></li>
<li><a href="../video-decoder/index.html" title="This document provides the video decoder API specification and the format conversion and display using DirectX or OpenGL following decode.">Video Decoder</a></li>
<li><a href="../video-encoder/index.html" title="This document provides the CUDA video encoder specifications, including the C-library API functions and encoder query parameters.">Video Encoder</a></li>
<li><a href="../inline-ptx-assembly/index.html" title="This document shows how to inline PTX (parallel thread execution) assembly language statements into CUDA code. It describes available assembler statement parameters and constraints, and the document also provides a list of some pitfalls that you may encounter.">Inline PTX Assembly</a></li>
<li><a href="../cuda-runtime-api/index.html" title="The CUDA runtime API.">CUDA Runtime API</a></li>
<li><a href="../cuda-driver-api/index.html" title="The CUDA driver API.">CUDA Driver API</a></li>
<li><a href="../cuda-math-api/index.html" title="The CUDA math API.">CUDA Math API</a></li>
<li><a href="../cublas/index.html" title="The CUBLAS library is an implementation of BLAS (Basic Linear Algebra Subprograms) on top of the NVIDIA CUDA runtime. It allows the user to access the computational resources of NVIDIA Graphical Processing Unit (GPU), but does not auto-parallelize across multiple GPUs.">CUBLAS</a></li>
<li><a href="../cufft/index.html" title="The CUFFT library user guide.">CUFFT</a></li>
<li><a href="../curand/index.html" title="The CURAND library user guide.">CURAND</a></li>
<li><a href="../cusparse/index.html" title="The CUSPARSE library user guide.">CUSPARSE</a></li>
<li><a href="../npp/index.html" title="NVIDIA NPP is a library of functions for performing CUDA accelerated processing. The initial set of functionality in the library focuses on imaging and video processing and is widely applicable for developers in these areas. NPP will evolve over time to encompass more of the compute heavy tasks in a variety of problem domains. The NPP library is written to maximize flexibility, while maintaining high performance.">NPP</a></li>
<li><a href="../thrust/index.html" title="The Thrust getting started guide.">Thrust</a></li>
<li><a href="../cuda-samples/index.html" title="This document contains a complete listing of the code samples that are included with the NVIDIA CUDA Toolkit. It describes each code sample, lists the minimum GPU specification, and provides links to the source code and white papers if available.">CUDA Samples</a></li>
<li><a href="../cuda-compiler-driver-nvcc/index.html" title="This document is a reference guide on the use of the CUDA compiler driver nvcc. Instead of being a specific CUDA compilation driver, nvcc mimics the behavior of the GNU compiler gcc, accepting a range of conventional compiler options, such as for defining macros and include/library paths, and for steering the compilation process.">NVCC</a></li>
<li><a href="../cuda-gdb/index.html" title="The NVIDIA tool for debugging CUDA applications running on Linux and Mac, providing developers with a mechanism for debugging CUDA applications running on actual hardware. CUDA-GDB is an extension to the x86-64 port of GDB, the GNU Project debugger.">CUDA-GDB</a></li>
<li><a href="../cuda-memcheck/index.html" title="CUDA-MEMCHECK is a suite of run time tools capable of precisely detecting out of bounds and misaligned memory access errors, checking device allocation leaks, reporting hardware errors and identifying shared memory data access hazards.">CUDA-MEMCHECK</a></li>
<li><a href="../nsight-eclipse-edition-getting-started-guide/index.html" title="Nsight Eclipse Edition getting started guide">Nsight Eclipse Edition</a></li>
<li><a href="../profiler-users-guide/index.html" title="This is the guide to the Profiler.">Profiler</a></li>
<li><a href="../cuda-binary-utilities/index.html" title="The application notes for cuobjdump and nvdisasm.">CUDA Binary Utilities</a></li>
<li><a href="../floating-point/index.html" title="A number of issues related to floating point accuracy and compliance are a frequent source of confusion on both CPUs and GPUs. The purpose of this white paper is to discuss the most common issues related to NVIDIA GPUs and to supplement the documentation in the CUDA C Programming Guide.">Floating Point and IEEE 754</a></li>
<li><a href="../incomplete-lu-cholesky/index.html" title="In this white paper we show how to use the CUSPARSE and CUBLAS libraries to achieve a 2x speedup over CPU in the incomplete-LU and Cholesky preconditioned iterative methods. We focus on the Bi-Conjugate Gradient Stabilized and Conjugate Gradient iterative methods, that can be used to solve large sparse nonsymmetric and symmetric positive definite linear systems, respectively. Also, we comment on the parallel sparse triangular solve, which is an essential building block in these algorithms.">Incomplete-LU and Cholesky Preconditioned Iterative Methods</a></li>
<li><a href="../libnvvm-api/index.html" title="The libNVVM API.">libNVVM API</a></li>
<li><a href="../libdevice-users-guide/index.html" title="The libdevice library is an LLVM bitcode library that implements common functions for GPU kernels.">libdevice User's Guide</a></li>
<li><a href="../nvvm-ir-spec/index.html" title="NVVM IR is a compiler IR (internal representation) based on the LLVM IR. The NVVM IR is designed to represent GPU compute kernels (for example, CUDA kernels). High-level language front-ends, like the CUDA C compiler front-end, can generate NVVM IR.">NVVM IR</a></li>
<li><a href="../cupti/index.html" title="The CUPTI API.">CUPTI</a></li>
<li><a href="../debugger-api/index.html" title="The CUDA debugger API.">Debugger API</a></li>
<li><a href="../gpudirect-rdma/index.html" title="A tool for Kepler-class GPUs and CUDA 5.0 enabling a direct path for communication between the GPU and a peer device on the PCI Express bus when the devices share the same upstream root complex using standard features of PCI Express. This document introduces the technology and describes the steps necessary to enable a RDMA for GPUDirect connection to NVIDIA GPUs within the Linux device driver model.">RDMA for GPUDirect</a></li>
</ul>
<div class="category"><span class="twiddle">▼</span><a href="index.html" title="CUDA-MEMCHECK">CUDA-MEMCHECK</a></div>
<ul>
<li><a href="#introduction">1. Introduction</a><ul>
<li><a href="#about-cuda-memcheck">1.1. About CUDA-MEMCHECK</a></li>
<li><a href="#why-cuda-memcheck">1.2. Why CUDA-MEMCHECK?</a></li>
<li><a href="#how-to-get-cuda-memcheck">1.3. How to Get CUDA-MEMCHECK</a></li>
<li><a href="#cuda-memcheck-tools">1.4. CUDA-MEMCHECK tools</a></li>
</ul>
</li>
<li><a href="#using-cuda-memcheck">2. Using CUDA-MEMCHECK</a><ul>
<li><a href="#command-line-options">2.1. Command Line Options</a></li>
<li><a href="#supported-operating-systems">2.2. Supported Operating Systems</a></li>
<li><a href="#supported-devices">2.3. Supported Devices</a></li>
<li><a href="#compilation-options">2.4. Compilation Options</a></li>
</ul>
</li>
<li><a href="#memcheck-tool">3. Memcheck Tool</a><ul>
<li><a href="#what-is-memcheck">3.1. What is Memcheck ?</a></li>
<li><a href="#supported-error-detection">3.2. Supported Error Detection</a></li>
<li><a href="#using-memcheck">3.3. Using Memcheck</a></li>
<li><a href="#understanding-memcheck-errors">3.4. Understanding Memcheck Errors</a></li>
<li><a href="#integrated-mode">3.5. Integrated Mode</a></li>
<li><a href="#api-error-checking">3.6. CUDA API Error Checking</a></li>
<li><a href="#device-side-allocation-checking">3.7. Device Side Allocation Checking</a></li>
<li><a href="#leak-checking">3.8. Leak Checking</a></li>
</ul>
</li>
<li><a href="#racecheck-tool">4. Racecheck Tool</a><ul>
<li><a href="#what-is-racecheck">4.1. What is Racecheck ?</a></li>
<li><a href="#what-are-hazards">4.2. What are Hazards?</a></li>
<li><a href="#using-racecheck">4.3. Using Racecheck</a></li>
<li><a href="#racecheck-report-modes">4.4. Racecheck report modes</a></li>
<li><a href="#understanding-racecheck-analysis-reports">4.5. Understanding Racecheck Analysis Reports</a></li>
<li><a href="#understanding-racecheck-hazard-reports">4.6. Understanding Racecheck Hazard Reports</a></li>
</ul>
</li>
<li><a href="#cuda-memcheck-features">5. CUDA-MEMCHECK Features</a><ul>
<li><a href="#nonblocking-mode">5.1. Nonblocking Mode</a></li>
<li><a href="#stack-backtraces">5.2. Stack Backtraces</a></li>
<li><a href="#name-demangling">5.3. Name Demangling</a></li>
<li><a href="#dynamic-parallelism">5.4. Dynamic Parallelism</a></li>
<li><a href="#error-actions">5.5. Error Actions</a></li>
</ul>
</li>
<li><a href="#cuda-memcheck-tool-examples">6. CUDA-MEMCHECK Tool Examples</a><ul>
<li><a href="#example-use-of-memcheck">6.1. Example Use of Memcheck</a><ul>
<li><a href="#memcheck-demo-output">6.1.1. memcheck_demo Output </a></li>
<li><a href="#memcheck-demo-output-with-memcheck-release-build">6.1.2. memcheck_demo Output with Memcheck (Release Build)</a></li>
<li><a href="#memcheck-demo-output-with-memcheck-debug-build">6.1.3. memcheck_demo Output with Memcheck (Debug Build)</a></li>
<li><a href="#leak-checking-in-cuda-memcheck">6.1.4. Leak Checking in CUDA-MEMCHECK</a></li>
</ul>
</li>
<li><a href="#integrated-cuda-memcheck-example">6.2. Integrated CUDA-MEMCHECK Example</a></li>
</ul>
</li>
<li><a href="#memory-access-error-reporting">A. Memory Access Error Reporting</a></li>
<li><a href="#hardware-exception-reporting">B. Hardware Exception Reporting</a></li>
<li><a href="#release-notes">C. Release Notes</a><ul>
<li><a href="#new-features-5.5">C.1. New Features in 5.5</a></li>
<li><a href="#new-features-5.0">C.2. New Features in 5.0</a></li>
</ul>
</li>
<li><a href="#known-issues">D. Known Issues</a></li>
</ul>
</nav>
<nav id="search-results">
<h2>Search Results</h2>
<ol></ol>
</nav>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/formatting/common.min.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/omniture/s_code_us_dev_aut1-nolinktrackin.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/omniture/omniture.js"></script>
<noscript><a href="http://www.omniture.com" title="Web Analytics"><img src="http://omniture.nvidia.com/b/ss/nvidiacudadocs/1/H.17--NS/0" height="1" width="1" border="0" alt=""></img></a></noscript>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-write.js"></script>
<script language="JavaScript" type="text/javascript" charset="utf-8" src="../common/scripts/google-analytics/google-analytics-tracker.js"></script>
</body>
</html>
