这部分内容主要是介绍一些专门为CUDA应用程序设计的调试工具和方法。设计这些工具和方法的目的是让我们可以在代码运行的时候检查应用程序。接下来，代码检查将被分成两个独立但是相关的部分，分别是内核调试和内存调试。

内核调试是指在运行中检查内核执行的流和状态的能力。CUDA调试工具让我们能检查GPU上任何线程以及任何代码位置的任何变量的状态。在检查应用程序正确性的时候，这会变得非常有用。

内存调试专注于发现程序的怪异行为，如无效的内存访问、对同一内存地址的冲突访问以及其他具有未定义结果的行为。因为内存调试工具比内核调试工具更加自动化，所以在用内核调试工具进行更深入的探索前，它们为找出错误或判断应用程序的正确性提供了快捷的方法。

内 核 调 试

内核调试是通过检查一个或多个线程的执行和状态来确定内核的正确性。在CUDA里内核调试有3种主要的方法：cuda-gdb、printf和assert。

使用cuda-gdb

再用cuda-gdb调试CUDA应用程序之前，必须先用特定标志编译程序。这个过程与用gdb和类似的工具编译用于调试的主机程序是很接近的。只要添加两个标志nvcc中：-g和-G：$ nvcc -g -G foo.cu -o foo,这些标志嵌入到主机和设备代码的调试信息中，并且关闭了大多数优化以确保程序状态能被检查到。只要使用调试编译了应用程序，就能像用gdb那样用cuda-gdb启动一个CUDA应用程序。给定一个编译和链接应用程序foo，可以通过以下方法将可执行文件传给cuda-gdb:

$ cuda-gdb foo
...
(cuda-gdb)

提示符(cuda-gdb),意味着cuda-gdb加载了来自可执行文件的符号并准备执行。要想运行该程序，只要输入run命令即可。run命令或者set args命令设置后，它可以包含命令行参数。

通常来说，cuda-gdb完全支持gdb提供的很多功能，包括断点、观察点和检查程序状态的能力等。然而，cuda-gdb还提供CUDA特定的调试功能。接下来将简单总结这些扩展功能和它们使用的实例。

CUDA焦点
虽然CUDA程序能包含多个主机线程和许多CUDA线程，但是cuda-gdb调试会话一次只处理一个线程。为了能在同一应用程序中调试多个设备线程，cuda-gdb提供了一种功能，即可以指定被检查的上下文(即设备线程)。可以使用cuda-gdb报告当前的焦点信息，包括当前设备、当前块、当前线程等。

例如，如果cuda-gdb调试对话的当前焦点是设备上正在执行的CUDA线程，那么我们可以使用下面的语句检索该焦点的完整说明：

(cuda-gdb) cuda thread lane warp block sm grid device kernel

该命令示例输出如下：

kernel 1026, grid 1027, block(0,0,0), thread(64,0,0), device 0, sm 1, warp 2, lane 0

可以使用类似的语句将焦点改为不同的设备线程。除了线程属性，还能提供了一个特定的线程，例如，当前块中的第128个线程，用以下语句：

(cuda-gdb) cuda thread (128)

如果没有显式地设置焦点属性，那么cuda-gdb将重新使用当前焦点的属性。

使用gdb help命令可以获得CUDA焦点选项的更多信息：

(cuda-gdb) help cuda

检查CUDA内存
与gdb一样，cuda-gdb支持检查变量，在堆(即CUDA全局内存)和寄存器中使用print语句：

(cuda-gdb) print scalar
(cuda-gdb) print arr[0]
(cuda-gdb) print (*arr)[3]

cuda-gdb还能用于检查CUDA的共享内存。例如，可以用一下命令访问共享内存的第二个字：

(cuda-gdb) print *(@shared int*)0x4

请注意，因为对于每个SM来说共享内存都是本地的，所以这一语句可能不会堆每个焦点中相同内存单元进行评估。

因此，使用cuda-gdb能检查任何共享内存数据。

获取环境信息
我们能用gdb info命令检索当前CUDA环境和平台的相关信息。可以用以下语句查找完整的环境信息列表：

为了报告当前系统中所有设备的信息，可以使用info cuda devices子命令。两个Fermi M2090GPU的系统中其输出如下：

请注意，许多子命令复制了复制了前面“CUDA焦点”部分已描述过的功能，并且这些子命令相对于当前cuda-gdb焦点进行操作。然而，更多种类的元数据可通过info cuda子命令访问到。cuda命令和info cuda子命令都有各自的作用，这取决于要查找的信息的类型和数量。
CUDA调试可调参数
CUDA-gdb通过set子命令展示了许多可调参数，对cuda-gdb行为的调整是非常有用的：

下表总结了cuda-gdb提供的有用的可调参数。通过使用help set cuda 命令，可以得到每个参数附加的信息。这些参数能用set cuda< tunable-name >命令进行设置。对于许多命令，< value >的选择是on/off，但是对于其他一些命令，还可以选择非布尔值。

操作这些cuda-gdb参数能够获得更多的cuda-gdb调试会话信息，你可以自定义它的行为使之符合你的要求。

实践CUDA-GDB
要想得到使用cuda-gdb的实践经验，在Wrox.com上下载debug-segfault.cu文件，使用该文件在CUDA中调试一个无效的内存访问以进行试验。首先，用已给出的Makefile文件建立debug-segfault.cu，该过程会设置标志-g和-G。然后，将应用程序载入到cuda-gdb中：

$ cuda-gdb ./debug-segfault
...
(cuda-gdb)

用run命令启动debug-segfault,它不需要任何命令行参数：

(cuda-gdb) run

执行该命令后很可能会看到许多滚动的文本，它提供了CUDA上下文和内核事件的信息。最终，当内核中发生内存错误时，调试对话将恢复到提示符(cuda-gdb)。下面的输出显示了在文件debug-segfault.cu的第34行出现了错误：

可以使用list命令来检查上下文周围的代码：

非法地址访问行是第34行。这一行包含了对input arr的多次间接解引用。通过输出arr[tid]中的地址，来测试由偏移量tid引用的数组内容。

这看起来似乎不太正确，一个空的内存地址不能被解引用。可以尝试解除cuda-gdb调式会话内部地址的引用，来二次检查问题：

显然地址是无效的，这就意味着arr中的内容被写入了无效值，或者没有进行合适的初始化。再看最初的源代码，我们应该会发现没有cudaMemcpy真正填满了设备数组d_matrix。在内核启动前加入下面这一行代码能避免这个内存错误：

cudaMemcpy(d_matrix, d_ptrs, N * sizeof(int *), cudaMemcpyHostToDevice);

修正之后的版本可以从Wrox.com下载debug-segfault.fixed.cu文件中找到。如果我们仍然在那个cuda-gdb调试会话中，那么还可以检查GPU上其他线程的状态。使用cuda命令获取当前设备、块和线程，代码如下所示：

试着在一个设备里转换到其他的线程上，并检查不同线程的状态：

虽然有许多线程都存在内存错误。焦点很容易放在具有最低逻辑ID的线程中。输入quit和y可以退出cuda-gdb会话:

CUDA的printf

在主机调试时，可能会经常遇到printf输出主机应用程序的状态。如果能在GPU设备代码中使用printf来简单检查内部设备的状态那就太好了。但是，内核在有着上前线程的设备上运行着，要整理这些内核的输出是一个有趣的挑战。从CUDA4.0开始，NVIDIA在设备上支持printf功能。基于CUDA的printf接口，与我们在主机上C/C++研发中习惯使用的一样(甚至有着相同的头文件,stdion.h)，这使得我们能直接过渡到基于CUDA的printf中。

这里有一些基于CUDA的printf语句使用的说明。首先，它只能在计算能力是2.0或更高的版本中实现。第二，除非显式使用CUDA同步，否则在线程间没有输出顺序。第三，在内核上将执行的printf输出返回到主机显示前，需要使用一个固定大小的循环设备缓冲区临时存储该输出。因此，如果产生输出的速度比显式输出的速度快，那么缓冲区就会覆盖掉原有的输出。这个缓冲区的大小可以用cudaGetDeviceLimit检索，并用cudaSetDeviceLimit进行设置。

以下常见事件会导致固定大小的缓冲区转回到主机以用于显示：

1. 任何CUDA内核启动。
2. 用CUDA主机API的任何同步（例如，cudaDeviceSynchronize,cudaStreamSynchronize,cudaEventSynchronize等）。
3. 任何同步内存复制，如cudaMemcpy。

这为快速调试内核提供了一个可用且友好的方法。但是，谨防过多地使用printf。可以使用线程和块索引来限制输出调试信息的线程，比避免过多地输出线程，导致调试信息缓冲区超载。

CUDA的assert工具

另一个常见的主机错误工具是assert。assert能让我们声明某一的条件，在正确执行时该条件必须为真。如果assert失败，则应用程序执行有以下两种情况的一种：1)有assert失败的消息，立即中止；2）如果在cuda-gdb会话中运行，控制将会传到cuda-gdb，以便可以在assert失败的位置检查应用程序的状态。和printf一样，只有GPU计算能力为2.0及以上时才提供assert供。它依赖和主机相同的头文件，assert.h。

在GPU中使用assert在主机上使用assert有一点不同，一旦设备上有失败的assert(即任何包含表达式的计算结果为0的assert)，就会有一个CUDA线程将在存储失败的assert信息后立即退出。但是，这个信息只会显示到主机上下一个CUDA同步点的stderr中(例如，cudaDeviceSynchronize、cudaStreamSynchronize等)。这意味着在每一个同步点，信息将显示自上一个同步点开始有失败assert的线程。如果在检测到第一个assert失败后，使用任何CUDA主机API调用，那么应用程序都将返回CUDA错误代码cudaErrorAssert。

和printf一样，在内核里使用assert就和在主机里一样，如下所示：

与主机的assert一样，通过使用在包含assert.h头文件前定义的NDEBUG预处理宏编译，可以对代码发行版本禁用assert评估。

内 存 调 试

cuda-gdb对CUDA内核执行进行细粒度检查是很有效的，虽然printf/assert对于大量的错误检测来说是简单的机制，但是对于调试CUDA内存错误的主要工具是cuda-memcheck。cuda-memcheck的操作在用于交互方面更加自动化和粗粒度，但是对于CUDA内核中的内存错误，cuda-memcheck提供了更详细的数据。cuda-memcheck包含两个独立的工具：memcheck工具、racecheck工具。

memcheck工具用于检查CUDA内核中越界和未对齐的访问。racecheck工具用于检查共享内存的冲突访问，这些冲突访问会导致未定义的行为。这些工具用于调试不稳定内核行为是非常有用的，这些行为是因为线程读取或写入到意外的位置而引起的。

cuda-memcheck的编译

使用cuda-memcheck编译应用程序比使用cuda-gdb更为复杂。当使用-g -G建立应用程序后，这些选项会对性能起负面影响。当使用cuda-memcheck工具时，很重要的是，为保证错误可复写，应用程序的性能必须稳定。但是，一些编译标志对仔细分析cuda-memcheck信息和准确找到问题发生的位置是必需的。

有一些可用的编译选项对性能影响很小，却能彻底提升cuda-memcheck信息的可读性。首先，应该使用-lineinfo选项进行编译。这个标志把信息嵌入到可执行文件中，该可执行文件使用设备指令将文件名和行号联系起来。可执行文件应该总是使用符号信息进行编译。这可以使cuda-memcheck输出主机的堆栈踪迹，这些堆栈踪迹可以准确地找出内核的启动位置。包含符号信息的编译标志是平台特有的，它使用nvcc中的-Xcompiler选项将参数传递到主机编译器中。例如，在带有gcc的Linux系统中，会用到-Xcompiler-rdynamic;在Windows系统中，会用到-Xcompiler/Zi。

当使用这些编译标志时，会生成一个可执行文件，它包含了足够多的用于显示mem-check和racecheck帮助信息的元数据，这会使得性能特性与原始的应用程序非常接近。
memcheck工具

memcheck工具可以检查6种类型的错误：

• 内存访问错误：对全局内存、本地内存或共享内存的越界或未对齐访问。未对齐的原子操作会触发内存访问错误，但是这只有当引用全局内存时才会发生。
• 硬件异常：硬件报告错误。参考CUDAMEMCHECK指南的附录B，其中包括每一个可能的硬件错误的详细信息。
• malloc/free错误：使用CUDA内核里的CUDA动态内存分配时，memcheck能找到malloc和free API调用的非正常使用。
• CUDA API错误：任何由CUDA API调用返回的错误代码。
• cudaMalloc内存泄漏：任何被应用程序使用cudaMalloc的内存分配，再执行完成前没有被释放。
• 设备堆内存泄露：使用CUDA内核中的CUDA动态内存分配时，memcheck会找到未释放的分配。

因为用cuda-gdb调试的debug-segfault程序显示了内存访问错误，所以可以对来自于memcheck工具的诊断信息和来自cuda-gdb的诊断信息进行比较。

假设想检查一个名为app的应用程序的内存错误。app能正确编译以维持性能，但仍会报告堆栈和行信息，memcheck能用以下语句调用：

$ cuda-memcheck [memcheck_options] app [app_options]

在debug_segfault上使用默认的选项运行memcheck会产生下面的输出：

memcheck工具不仅在debug-segfault.cu的第25行指出了一个无效的内存访问，还提供了无效访问的方向(写)、被写入( __ global __ )的内存空间、写入的大小(4字节)、执行写操作的线程以及造成无效引用的具体地址。相比于cuda-gdb,memcheck工具需要的手动工作比较少，而且为debug-segfault提供了更详细和精确的内存错误信息。

我们也能注意到memcheck也报告了第二次错误，调用cudaMemcpy返回的是CUDA错误4.回想可知，memcheck处理的错误类型之一是由CUDA API调用返回的错误代码。参考cuda.h。CUDA错误4是CUDA_ERROR_DEINITIALIZED,这表明CUDA驱动器处于关闭过程中。这个错误可能是由之前的内存访问错误引起的：该驱动器正在从以外的设备行为种恢复过来。
racecheck工具

racecheck用于识别共享内存中存储数据的冲突访问(一般称为冲突)。另一方面，racecheck在同一线程块种寻找多个线程，这些线程块引用共享内存中的同一位置，这些共享内存是不同步的，这些引用中至少有一个引用对这个位置进行写操作。调试共享内存的正确性是非常重要的，理由如下：

• 首先，因为共享内存在片上的且被一个线程块共享，所以它常被用作多线程间的低延迟通信通道。如果不合理地同步那些多线程访问，那么就可能发生冲突。因此，需要一个工具来处理这种常见情况，因为共享内存更容易被误用，导致冲突访问。
• 第二，共享内存的正确性不能直接通过主机的应用程序来检查。全局内存的调试被简化了，因为主机有立即检查全局状态的能力，但共享内存不存在这样的直接通道。支持这种性能首先需要将这种状态传输到全局内存，然后再返回主机。racecheck工具帮我们做了这些事。

考虑一个简单的且使用共享内存和本地同步的单一线程块的并行归约问题。为了研究racecheck的效果，下面的例子去掉了本地同步，因此再冲突访问存在时可以观察由race-check产生的诊断。从Wrox.com上可以下载到源代码debug-hazards.cu中。

用前面已讨论过的编译选项编译debug-hazards.cu：

$ nvcc -arch=sm_20 -lineinfo -Xcompiler -rdynamic -o debug-hazards debug-hazards.cu

在运行debug-hazards前，要知道racecheck在应用程序执行中会生成一个大的被后处理的转储文件。racecheck也会在命令行的终端生成一个详细的报告，所以为了以后分析可以将终端输出保存成文件。钙离子使用–save CLI参数将转储文件的位置设置在有几百MV可用磁盘空间的地方。对于较大的应用程序，转储文件将占用更多的磁盘空间。这个例子将终端输出转向日志文件，用于以后的检查。

现在，可以使用以下命令运行racecheck，以分析debug-hazards:

$ cuda-memecheck --tool racecheck --save racecheck.dump ./debug-hazards > log

检查日志文件，将看到很多重叠的部分，类似于以下代码：

从第一行开始看：

这一行表明了3个重要的事情。首先，检查到一个潜在的冲突！这是否是一个好的(或可怕的)开始，取决于你的看法。

第二，这一行报告了Read-After-Write(RAW)冲突。这意味着两个线程没有按照任何顺序访问了相同地址，一个执行读操作，一个执行写操作。因为没有顺序，所以读线程应该在写线程之前还是之后加载数值是未定义的。这个未定义行为是不可取的，因此造成了冲突。

第三，这一行指出了哪个线程块存在冲突(共享内存上的风险只能发生在单一线程块上)。这个信息是否有用取决于应用程序。因为这个应用程序的每个块都在做相同的工作，它可能对调试没有帮助。

现在，看下一行：

这一行提供了线程上的信息，这个线程在RAW风险中执行写操作。它指出了线程ID(31,0,0)，正在执行的指令地址(oxc8)和正在执行的源代码行。

下一行在读线程上提供了相同的信息：

下一行：

指明存储在冲突位置的当前值。

剩下的行显示了来自于主机位置的堆栈跟踪，这里启动的内核引起了该冲突。

现在，为读线程和写线程提供的信息可以用于分析冲突。回想一下可知，写线程在文件debug-hazards.cu的第50行执行了simple-reduction函数。读线程在同一个文件和同一个函数的第66行。这里为了方便突出一下相关行：

冲突发生在写入local_mem[local_tid]和读取local_mem[i]之间。读线程正在扫描共享内存的每个条目，同时写线程正在填充这些单元中的一个。在这个应用程序中，期望的行为是确保所有写线程都在读线程开始扫描共享内存前完成，并且在扫描完成前没有其他的写操作被执行。因为错误报告中显示为先写后读冲突，这是指，在作为扫描的一部分被执行的读操作和对本地内存的写操作之间，同步丢失了。因此有一个条件是在内存地址上读操作先于写操作。第一次尝试消除这种竞争条件，尝试不对54行的__synthreads进行注释，然后用相同的命令重新建立并重新运行。__synthreads将确保在线程0开始扫描前所有的写操作都已完成，避免了写后读的冲突。

看一下新的输出日志，出现了一个新的警告。注意这个冲突发生在代码中两个相同的位置之前，但是这次是读后写冲突。这说明现在写操作发生在扫描线程完成之前。如果一些线程在扫描线程完成前进入外层循环的下一个迭代，那么会发生这种情况。内核必须保证，对内存位置进行任何额外的写操作完成前，扫描线程必须读完它的当前值。为了避免这种情况，必须在73行插入另一个同步点，使所有线程等待扫描完成。重新建立并运行，看看是否阻止了读后写冲突。

现在应该有一个日志文件报告以下内容：

现在racecheck在这个程序里找不到冲突了。虽然这不能保证程序里没有冲突了，但它是指示共享内存中没有冲突的强大指示器。