CUDA, 用于大量数据的超级运算：第14节

http://www.ddj.com/development-tools/220601124
Rob Farber
Debugging CUDA and using CUDA-GDB
Rob Farber 是西北太平洋国家实验室（Pacific Northwest National Laboratory）的高级科研人员。他在多个国家级的实验室进行大型并行运算的研究，并且是几个新创企业的合伙人。大家可以发邮件到rmfarber@gmail.com与他沟通和交流。

在关于CUDA的系列文章第13节CUDA,用于大量数据的超级运算: 里，我讨论了本系列文章第11节就探讨的“纹理内存”，并且介绍了诸如最近新推出的CUDA Toolkit 2.2（增加了向GPU上与2D纹理绑定的全局内存写入的能力）相关信息。本小节关注的是调试技巧和CUDA-GDB是怎样被利用以有效诊断和调试CUDA代码的－重点在怎样在查阅大量的数据，线程语法和语意区别的时候提高这一过程的速度，以调试在GPU上运行的内核。
首先，你需要了解有关CUDA-GDB现在情况的两个重要观点：
它仅仅在基于UNIX的系统上运行。微软用户特别要注意参照最近发布的Visual Studio debugger，我将在下一篇专栏文章里讨论。
X11不能在用于调试的GPU上运行。使用多GPU系统或者删除X11，通过ssh, VNC或其它方式远程读取单-GPU系统
本文章也提供了一个精心设计的样例。我推荐每个人都通读一遍－不管他们之前的GDP体验－原因如下：
每一个尚不了解CUDA-GDB的人：样例指令显示如何使用新的CUDA线程语法，强调允许调试和通过在CUDA启动的图形处理器上的内核的语义变化；
初学者：该样例显示了如何设定断点（象征性地和以行数为单位），以及执行其它基本的调试操作如运行程序，逐行通过源代码，持续的程序执行，退出调试器。
专家：GPU的快捷意味着调试CUDA应用程序需要挖掘大量数据以发现问题。调试策略，包括应用调试器助手功能和GDP模拟数组 可以有效节省调试时间（和个人的精力）。
实际上，CUDA-GDB是GNU GDB调试版本6.6的一个端口。使用过GDB的程序员们会发现他们已经熟悉了CUDA-GDB，并且应该通读这篇文章以了解与CUDA相关的一些建议。GDB新手可使用本文以开始马上调试他们的软件，但是应该参阅网上的指导教程和参考文件以了解更多有关这个强大的调试工具特点的信息。可以从GNU documentation for GDB（GDB GNU文件）。不论技术等级如何，所有的CUDA开发人员都应该至少阅读下CUDA-GDB: The NVIDIA CUDA Debugger（CUDA-GDB: NVIDIA CUDA调试器）的最新版本。
CUDA-GDB创建之前的调试方法
在创建CUDA-GDB之前，调试CUDA程序最简单和至少是不错的一个方法就是，向源代码添加打印语句，编译以在模拟程序中运行（通过将-deviceemu或-device-emulation 标识传递给nvcc编译器，进行初始化）。因为模拟程序在主处理器上运行（而非GPU上），打印语句可被编译和连接，这样程序员可以检查什么样的程序值可能会有重要意义。这个谨慎的一个方法就是看，在CUDA的早期，CUDA程序里正在进行什么。（这个调试方法在2008年4月，本系列文章的第一节里有所讨论）。我再次提及这个方法是因为它，作为最后的一个补救方法－可以帮助在代码里找到程序错误bug。基本上，如果你不相信GPU正在进行的内容－就尝试着在模拟程序上运行吧。如果代码还是不成功，你就知道它不是GPU。记住，模拟程序不会精确地再现GPU上所进行的操作，这就意味着在GPU上的程序错误bug和行为（包括竞态条件）可能不会在模拟环境里发生。
还有一个方法就是，利用GPU，并允许检查GPU计算结果，利用cudaMemcpy() 以从GPU转移变量到主机上从头开始的位置。基于主机的方法（包括GDB和/或打印语句）然后可以被用来检查开头位置的信息，以诊断问题。在下面的文章里，我将使用这个技巧以演示怎样在向大CUDA向量写入的时候，使用CUDA-GDB的简单助手调试功能发现错误。
映像内存和回归测试的重要性
此刻，值得提及的是，在CUDA2.2版本里提到的新的映像内存能力提供了一个重要，便捷和全新的能力，以在向CUDA移植原有集成程序代码时，促进回归测试。毫无疑问，回归测试是一个重要的软件行为。这个技巧在创建和验证正确工作的软件方面有非常重要的作用，怎么强调都不为过！
在使用原有软件时，程序员就已经有了一个工作代码基础，可用来与GPU生成结果对比，以帮助确认错误。映像内存（第12节”CUDA2.2 改变数据移动范式里有所讨论）通过透明地维持主机和设备内存空间之间的数据同步版本，极大地促进了这一过程。带着小心，程序员可以利用这一透明的同步性，在整个移植项目中，保持原有软件的功能性。其结果就是，会有一个已知的工作版本可以被用来对比所有的GPU结果和中间结果。
基本上，新的CUDA内核会被合并入遗产代码，你不需仔细考虑将数据从主机移动到GPU上，这样就可以在新的CUDA内核和相应的原有主代码之间轻松转换了。新的GPU版本然后在一个或多个案例中被评估，以观察它是否生成了正确结果。如果错误被确定，那么该计算阶段的原有主代码－加上中间结果－可被用来快速确认在GPU代码里出现的第一个错误。最后，足够的计算会在GPU上进行，因为不再需要保持与主机的同步性，映像被禁用或移除－因此创建遗产代码的GPU版本（仅适用于GPU），全速运行，无PCI限制。
使用大量数据运用CUDA-GDB的小贴士

因为多核构架，大量线程模式，和高性能，图形处理器是处理大量数据的出色平台。反之，手动方式在这如海的数据里发现错误会非常耗费时间和精力。
为了让调试变得轻松些，我想在代码里包含一些简单的助手调试程序，可从GDP或CUDA-GDB调用，或被用来作为测试套件的一部分，通过使用断言进行安全检查。当运行时，这些调试程序不会被调用（如必要，可用#ifdef语句消除），这样它们不会引发额外的任何内存或处理器运行时间。在运行GDB（或CUDA-GDB）时，互动调用任一数量的调试程序的能力就提供了一个便捷的方法，在程序执行时，寻找错误－而无需修改源代码或重新编译！
下面AssignScaleVectorWithError.cu样例程序里的函数sumOnHost()硕民给了一个调试程序。在这个例子中，在将向量从设备移动到主机上后，sumOnHost() 计算大向量的浮点数。很容易想象怎样将这个理念拓展，提供关于大型数据结构的信息。计算浮点数非常有用，因为它会产生一个单个数，以了解数据的意义，确认NaN (不是数字) 问题，执行其它安全检查。许多网络和碟子系统使用类似技巧，计算检查码（CHECKSUM）（或其它综合量），以确认数据错误。
使用一个数做为对比的手段，以发现数据差别，在这种情况下会非常有用：当已知的，基于主机的软件被用来对比中间结果和基于CUDA的内核时。不用盲目地浏览大量数字，以发现差别，程序员们可利用CUDA-GDB的力量第一时间隔离遗产主机和GPU内核结果和/或中间值之间的任何差别。
根据经验，这是个非常出色的节省时间的方法。注意，当对比浮点结果时，可能会有些变化，因为浮点仅仅是个近似表示。主机或GPU按照算法运算的次序，执行算术甚至合理的变化时，会出现微小差别（简单的改变，如使用不同的编译器转换或改变优化等级，就可导致微小差别），因而在甚至正确的工作代码结果里都会有细微的变动。
CUDA Toolkit的大部分基于UNIX的操作系统都包括CUDA-GDB，因为你只需打出：
cuda-gdb

就可以启动调试器了。（如果是另外一种情况，则请参考CUDA-GDB版本给出的注意事项及NVIDIA CUDA论坛，看看其他人是如何让调试器在OS操作系统上运行的）。
CUDA-GDB接受的自变量和选项和GDP完全一致。通常，CUDA-GDB从一个明确指定可执行程序以进行调试（如cuda-gdb a.out）的自变量开始。将程序的流程ID（PID）添加到指令行（如cuda-gdb a.out pid），CUDA-GDB也可以被用来调试已经运行的程序。
为了以用户友好形式调试程序，编译器需要为CUDA-GDB生成额外的调试信息，说明每个变量的数据类型或函数及在源行数和地址之间的一致性（以可执行模式）。为了让编译器生成该信息，当你运行nvcc编译器时，-g 和 -G 选项必须被指明。
下面的指令行是编译程序AssignScaleVectorWithError.cu以进行调试的：

nvcc -G -g AssignScaleVectorWithError.cu -o AssignScaleVectorWithError

那么，这些指令行选项都是做什么呢？
-G 选项明确说明用于CUDA内核的生成调试信息，并
强制编译-O0 (大部分是未优化的)
将所有变量都存入本地内存（并且可能会使得程序的执行变缓）

-g 选项通知nvcc为主机代码生成调试信息，并以可执行模式包含象征性调试信息；
最后，-o 选项通知编译器将可执行程序写入AssignScaleVectorWithError。
注:现在，当使用-cubin选项编译时，不可能生成调试信息。
下面是AssignScaleVectorWithError.cu的源代码：

#include <stdio.h>
#include <assert.h>
// A simple example program to illustrate
// debugging with cuda-gdb
// Vector on the device
float *a_d;
// Print a message if a CUDA error occurred
void checkCUDAError(const char *msg) {
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err) {
fprintf(stderr, "Cuda error: %s: %s./n", msg, cudaGetErrorString( err) );
exit(EXIT_FAILURE);
}
}
// Zero the vector
__global__ void zero(float *v_d, int n)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < n)
v_d[tid] = 0.f;
}
// Assign the vector with consecutive values starting with zero
__global__ void assign(float *v_d, int n)
{
int tid = threadIdx.x;
if(tid < n)
v_d[tid] = ((float) tid);
}
// Scale the vector
__global__ void scale(float *v_d, int n, float scaleFactor)
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if(tid < n)
v_d[tid] *= scaleFactor;
}
// Move the vector to the host and sum
float sumOnHost(const float *v_d, int n)
{
float sum=0.f;
int i;

// create space on the host for the device data
float *v_h = (float*)malloc(n*sizeof(float));
// check if the malloc succeeded
assert(v_h != NULL);

// copy the vector from the device to the host
cudaMemcpy(v_h,v_d, n*sizeof(float), cudaMemcpyDeviceToHost);

for(i=0; i<n; i++) sum += v_h[i];

// free the vector on host
free(v_h);
return(sum);
}
int main()
{
int nBlocks = 32;
int blockSize = 256;
int n = nBlocks*blockSize;
float scaleFactor = 10.f;

// create the vector a_d on the device and zero it
cudaMalloc((void**)&a_d, n*sizeof(float));
checkCUDAError("Create and zero vector");
// fill the vector with zeros
zero<<<nBlocks, blockSize>>>(a_d, n);
// assign the vector
assign<<<nBlocks, blockSize>>>(a_d, n);
// scale the vector by scaleFactor
scale<<<nBlocks, blockSize>>>(a_d, n, scaleFactor);
// calculate the sum of the vector on the host
float dSum = sumOnHost(a_d, n);
checkCUDAError("calculating dSum");
// Check if both host and GPU agree on the result
float hSum=0.f;
for(int i=0; i < n; i++) hSum += ((float)i)*scaleFactor;

if(hSum != dSum) {
printf("TEST FAILED!/n");
} else {
printf("test succeeded!/n");
}

// free the vector on the device
cudaFree(a_d);
}

总而言之，该程序创建了一个设备上的向量，a_d, 该向量被内核用零填满，zero()。向量a_d 然后被连续赋予不断上升的值，从零开始，assign()。最后，向量a_d 通过内核乘以一个标度，scale()。主程序sumOnHost() 被调用计算a_d内的值，然后放置在dSum中，并与变量 hSum内包含的主机生成数对比。如果值相同，我们会得到一条声明测试成功的消息。否则，程序显示测试失败。
正如我们在下面看到的，运行未修改程序生成一个失败信息，该信息暗示存在代码错误：

$ ./AssignScaleVectorWithError
TEST FAILED!

下面的指令行启动CUDA-GDB 以用来调试程序：

$cuda-gdb AssignScaleVectorWithError

你可以看到输出类似于：

NVIDIA (R) CUDA Debugger
BETA release
Portions Copyright (C) 2008,2009 NVIDIA Corporation
GNU gdb 6.6
Copyright (C) 2006 Free Software Foundation, Inc.
GDB is free software, covered by the GNU General Public License, and you are
welcome to change it and/or distribute copies of it under certain conditions.
Type "show copying" to see the conditions.
There is absolutely no warranty for GDB.  Type "show warranty" for details.
This GDB was configured as "x86_64-unknown-linux-gnu"...
Using host libthread_db library "/lib/libthread_db.so.1".

我们使用缩略的指令l (列表) 以查看源代码第81行附近的指令行：

(cuda-gdb) l 81
76        checkCUDAError("Create and zero vector");
77
78        // fill the vector with zeros
79        zero<<<nBlocks, blockSize>>>(a_d, n);
80        // assign the vector
81        assign<<<nBlocks, blockSize>>>(a_d, n);
82        // scale the vector by scaleFactor
83        scale<<<nBlocks, blockSize>>>(a_d, n, scaleFactor);
84
85        // calculate the sum of the vector on the host

现在，我们在开始执行assign() 内核前，使用指令在第81行设定断点（使用单个字母b 表示断点）：

(cuda-gdb) b 81
Breakpoint 1 at 0x40f216: file AssignScaleVectorWithError.cu, line 81.

如下指令所示，断点也可被象征性地设定，该指令每当内核scale() 被调用时，就会设置断点：

(cuda-gdb) b scale
Breakpoint 2 at 0x40f4e3: file AssignScaleVectorWithError.cu, line 38.

我们现在在调试器里运行该程序，使用字母r ，而无需打出整个指令run。 (注：有些输出可能会有所不同，如流程ID)。

(cuda-gdb) r
Starting program: /home/XXX/DDJ/Part14/AssignScaleVectorWithError
[Thread debugging using libthread_db enabled]
[New process 16805]
[New Thread 140439601190656 (LWP 16805)]
[Switching to Thread 140439601190656 (LWP 16805)]
Breakpoint 1, main () at AssignScaleVectorWithError.cu:81
81        assign<<<nBlocks, blockSize>>>(a_d, n);
Current language:  auto; currently c++

使用 p指令 （"print"缩写) ，我们使用自变量调用主机函数sumOnHost(), 以恰当地将GPU数组a_d 里的所有数据移到主机上，并计算值的总和。可以看到，对内核zero()的调用似乎运行正确，因为向量被零浮点值填满：

(cuda-gdb) p sumOnHost(a_d, n)
$1 = 0

我们使用next 指令(缩写为n) ，运行该程序的下一行指令。在本例中，程序在GPU上运行assign() 内核。
请注意，与一般的执行不同，在CUDA-GDB内对内核的调用同步发生。（通常情况下，内核异步启动。
因此，在输入next指令后，控制仅在assign()内核在GPU上运行到完成后，才返回。
正如CUDA-GDB手册第4.4章节所指出的，调试器支持以warp粒度步进GPU代码。这就意味着单个的线程不会被推进，但是在warp里的所有线程会前进。特别情况是，跳过线程阻碍调用，__syncThreads(), 这样所有的线程都会通过阻碍前进。此外，不可能跳过子程序，因为编译器现在内联代码。因此，仅可能步进入一个子程序。
我们再来看看使用print指令从sumOnHost() 返回的和值：

(cuda-gdb) n
83        scale<<<nBlocks, blockSize>>>(a_d, n, scaleFactor);
(cuda-gdb) p sumOnHost(a_d, n)
$2 = 32640

在本例中，返回值32640看上去很小，为所有整数（从[0 到 nBlocks*BlockSize)的和。因此我们选择“continue”（缩写为c），直到我们切入下一个断点，这个断点正好是设定在CUDA内核scale()里。（注：现在，我们忽略说明“Current CUDA Thread”指令行的含义）。

(cuda-gdb) c
Continuing.
[Current CUDA Thread <<<(0,0),(0,0,0)>>>]
Breakpoint 2, scale () at AssignScaleVectorWithError.cu:38
38      int tid = blockIdx.x * blockDim.x + threadIdx.x;

调试器允许我们检查GPU上的值。当在scale() 内，断点被设定以停止程序运行时，情况就是如此。
a_d 的地址，通过v_d函数自变量，被传递到这个内核。使用打印指令（缩写为p), 我们可以通过使用模拟数组的GNU概念，可以检查GPU内存里的向量值的逐次值。在下面指令的输出中可以看到，向量的前10个值（在指令中显示为syntax @10） 由assign() 内核正确设置：

(cuda-gdb) p *v_d@10
$3 = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}

然而，我们看到大于255的向量元素仍然设置为零，这就意味着在assign() 内核里有个问题。于是，使用指令语法(v_d+250)告知调试器从v_dpointer的开始就抵消250个元素：

(cuda-gdb) p *(v_d+250)@10
$4 = {250, 251, 252, 253, 254, 255, 0, 0, 0, 0}

我们打出quit退出CUDA-GDB，然后进一步查看代码：

(cuda-gdb) quit
The program is running.  Exit anyway? (y or n) y
rmfarber@k2:~/DDJ/Part14$

结果是，assign()的问题就是变量线程号仅被设置到threadIdx.x。当多个块被使用时，这是不正确的。
让我们通过使用CUDA-GDB拓展（可让我们查看warp内的单个线程）来进行验证。
再次启动CUDA-GDB：

$cuda-gdb AssignScaleVectorWithError

现在，我们在第31行设定一个断点，它在tid之后在assign() 内核里被初始化，运行调试器知道断点被击中：

(cuda-gdb) b 31
Breakpoint 1 at 0x40f4cf: file AssignScaleVectorWithError.cu, line 31.
(cuda-gdb) r
Starting program: /home/XXXX/DDJ/Part14/AssignScaleVectorWithError
[Thread debugging using libthread_db enabled]
[New process 22405]
[New Thread 139839913080576 (LWP 22405)]
[Switching to Thread 139839913080576 (LWP 22405)]
[Current CUDA Thread <<<(0,0),(0,0,0)>>>]
Breakpoint 1, assign () at AssignScaleVectorWithError.cu:31
31    if(tid < n)
Current language:  auto; currently c++

使用"info cuda threads" 指令, 我们看到如下输出：

(cuda-gdb) info cuda threads
<<<(0,0),(0,0,0)>>> ... <<<(31,0),(255,0,0)>>> assign ()
at AssignScaleVectorWithError.cu:31

CUDA 线程信息以以下格式显示：

<<<(BX,BY),(TX,TY,TZ)>>>

BX 和BY 是X 和Y 块索引，TX, TY, 和TZ 分别是线程X, Y, 和 Z索引。因此，我们可以看到assign()内核的块的索引范围从(0,0,0)到(255,0,0)。这就正确地表示了使用32块配置来在GPU上运行的内核，每个块包括256个线程。
下面的指令行显示调试器现在被设置，以检查第一个块的第一个线程：

[Current CUDA Thread <<<(0,0),(0,0,0)>>>]

输入tid ，显示就该线程"，它被正确地设置为零：

(cuda-gdb) p tid
$1 = 0

使用CUDA线程语法，我们使用缩写的语法转换到块31和线程255，以节省输入：

(cuda-gdb) thread <<<(31),(255)>>>
Switching to <<<(31,0),(255,0,0)>>> assign ()
at AssignScaleVectorWithError.cu:31
31    if(tid < n)

tid 变量的值显示，它被错误地设定为255.

(cuda-gdb) p tid
$2 = 255

我们现在知道assign() 内核错误地使用以下语句将threadIdx.x 赋值到tid:

int tid = threadIdx.x;

使用编辑器，将tid索引的赋值更改如下：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

在保存，重新编译和运行修订的程序之后，我们看到程序现在报告测试成功：

test succeeded!

启动CUDA-GDB，重复之前的调试步骤，我们现在看到线程<<<(31,0),(255,0,0)>>>里的tid正确地包含了值8191。

(cuda-gdb) thread <<<(31),(255)>>>
Switching to <<<(31,0),(255,0,0)>>> assign () at AssignVector.cu:31
31    if(tid < n)
(cuda-gdb) p tid
$1 = 8191

额外的CUDA-GDB调试拓展和语义
CUDA-GDB提供了诸多CUDA专用的指令：
thread – 显示现有重点主机和CUDA线程；
thread <<<(TX,TY,TZ)>>> - 转换到在指定座标的CUDA线程
thread <<<(BX,BY),(TX,TY,TZ)>>> -转换到在指定座标的CUDA块和线程
info cuda threads – 显示GPU上现有所有CUDA线程的全面总结
info cuda threads all – 显示GPU上现有线程的列表。这个列表可能非常之长。
info cuda state – 显示有关现有CUDA状态的信息。
Next and step 指令的特殊语义：
在warp 等级，执行被推进，在同一个warp里的所有线程，被当作现有CUDA线程，将进行运算。
一个特殊情况就是：步入线程阻碍调用，__syncthreads(), 这就导致在阻碍之后立即会设定一个隐式断点。所有的线程都被延续到这个__syncthreads()之后的断点。
关注事项和已知事项
如下所示，CUDA2.3版本需要关注的事项和已知事项：
主机系统的文字大小不再是我们的关注问题，因为CUDA-GDB（如CUDA2.2测试版本）支持32和64字节系统；
在论坛上有报道称，CUDA-GDB有时会挂住。注意，这是GDB的一个复杂端口，而NVIDIA似乎在处理问题方面做得很出色。
当使用调试器时，下列任何情况都可能影响程序行为或性能：
X11不能在用来调试的GPU上运行。建议的变通方案包括：
对单个GPU（VNC, ssh, 等）的远程读取；
使用两个GPU，X11仅在一个图形处理器上工作。
如CUDA 2.2, CUDA驱动器会自动将运行X11的设备排除，以避免被正在调试的应用程序选定；
使用-G 选项进行编译，这样变量流入到本地内存，极大地降低程序性能。（如第五节所示，本地内存可以比寄存器或共享内存慢150倍）。
内核启动不再异步，因为调试器拦阻内核启动。
不支持范围映射。这意味着如果变量在一个和外部范围内的变量同名的内部范围里被引入，仅可看到外部范围的值。AssignArray.cu样例显示了这个局限性。
调试器必须在内核内被停止，以检查设备内存（通过cudaMalloc()被分配），因为设备内存在内核函数之外是不可见的。
被分配了cudaMallocHost()的主机内存在CUDA-GDB里不可见。
不支持多GPU应用程序。
不是所有的非法程序行为都会在调试器里被“抓到”，如界外内存访问或被零除情况。
现在，不可能跳过设备代码里的子程序。
这些专栏文章注重的是利用运行时间界面。任何使用设备驱动器API的程序都不能使用CUDA-GDB进行调试，因为设备驱动器API不受到支持。
总结
已被验证的软件开发策略如断言和回归测试对开发正确和无bug的软件来说，非常重要。当移植遗产软件时，查阅CUDA里有用的映像内存特征。当bug显现时，CUDA-GDB可被用来追踪这些bug。因为GPU问题通常操控大量数据，因此在你的代码里有几个简单的，可以互动地从CUDA-GDB调用的程序以及仿真数组可以有效提高调试过程的速度。