NVIDIA电脑消毒剂是一个强大的工具，可以节省您的时间和精力，同时提高CUDA应用程序的可靠性和性能。
在我们之前的帖子中，高效的CUDA调试：如何使用NVIDIA Compute Sanizer搜索Bug，我们探索了并行编程领域中的高效调试。我们讨论了如何在CUDA环境中调试代码既富有挑战性又耗时，尤其是在处理数千个线程时，以及Compute Sanitizer如何帮助完成这一过程。
这篇文章继续我们对高效CUDA调试的探索。它重点介绍了其他一些Compute Sanitizer工具，并介绍了几个示例。
NVIDIA电脑消毒剂
Compute Sanitizer是一套工具，可以对代码的功能正确性执行不同类型的检查。有四种主要工具：
内存检查：内存访问错误和泄漏检测赛马检查：共享内存数据访问危险检测工具初始检查：未初始化的设备全局内存访问检测工具同步检查：线程同步危险检测
除这些工具外，Compute Sanitizer功能还包括：
本文的重点是调试代码并使用初始化检查、和同步使用同步检查。有关使用的详细信息内存检查发现内存泄漏和赛马检查有关查找竞争条件的信息，请参见高效的CUDA调试：如何使用NVIDIA计算机消毒剂搜索错误.
初始化检查
初始化检查帮助您识别和解决CUDA代码中未初始化的内存访问错误。未初始化的内存访问可能会导致CUDA应用程序中出现不可预测的行为和错误的结果。
初始化检查可以检测设备代码中对全局内存的未初始化内存访问。它提供了有关访问位置和时间的详细信息，以及访问线程的堆栈跟踪。这有助于揭示问题的根本原因并解决问题。
为了提供一个示例，以下代码示例受益于初始化检查：
#包括<stdio.h>#定义螺纹32#定义块2__全局__void addToVector（float*v）{int tx=线程Idx.x+块Dim.x*块Idx.x；v[tx]+=tx；}int main（int argc，char**argv）{float*d_vec=空；float*h_vec=空；h_vec=（float*）malloc（BLOCKS*THREADS*sizeof（float））；cudaMalloc（（void**）&d_vec，sizeof（float）*BLOCKS*THREADS）（浮动）*块*螺纹）；cudaMemset（d_vec，0，块*螺纹）；//将阵列归零添加到矢量<<<块，线程>>>（d_vec）；cudaMemcpy（h_vec，d_vec，BLOCKS*THREADS*sizeof（float），cudaMemcpyDeviceToHost）；cudaDeviceSynchronize（）；printf（“之后：向量0，1..N-1:%f%f..%f\N”，h_vec[0]，h_vec[1]，h_vesc[BLOCKS*THREADS-1]）；cudaFree（d_vec）；游离（hvec）；退出（0）；}
此代码包含一个名为添加到矢量它对向量中的每个元素执行简单的值相加，结果写回相同的元素。乍一看，它看起来很好：在设备上用库达·马洛克，然后用归零cudaMemset卡，然后在内核中执行计算。它甚至打印出正确的答案：
$nvcc-lineinfo initcheck_example.cu-o initchecke_example$ ./初始化检查示例之后：向量0，1。。N-1:0.000000 1.000000。。63
但代码中包含一个小错误。（如果你能发现，就得20分。）
使用初始化检查检查对设备全局内存中向量的任何访问是否试图读取未初始化的值。
$compute-sanitizer--工具initcheck/初始化检查示例=====计算机消毒器=====大小为4字节的未初始化__global__内存读取=====在/home/pgraham/Code/BlogExamples/initcheck_example.cu:8:addToVector（float*）中的0x70处====~====通过块（0,0,0）中的线程（16,0.0）. . .=====大小为4字节的未初始化__global__内存读取=====在/home/pgraham/Code/BlogExamples/initcheck_example.cu:8:addToVector（float*）中的0x70处====~====通过块（0,0,0）中的线程（17,0.0）。=========之后：矢量0，1。。N-1:0.000000 1.000000。。63=====错误摘要：48个错误
这应该会打印出大量信息（为了简洁起见，对显示的输出进行了编辑），但有些地方不正确。大量输出是回溯信息，可以使用--show-backtrace否选项：
$compute消毒程序——工具initcheck——显示回溯no/初始化检查示例
查看输出，您可以看到总共48个错误。报告显示他们都是大小为4字节的未初始化__global__内存读取.
每条消息都表示试图从全局设备内存中读取某些内容，并且该内容的大小为4个字节。一个合理的猜测是，错误是指试图访问向量的元素，这些元素由每个4个字节的浮点组成。
查看第一个错误，消息的下一部分指示导致错误的线程和线程块。在本例中，它是块0中的线程16。在设置内核时，每个线程都会访问向量的不同元素，即向量的元素17，d_vec[16]，未初始化。
在输出中，您可能会看到第一个导致错误的线程是另一个线程。GPU可以按照它认为合适的顺序安排扭曲（32个线程组成的组）。但是检查输出的其余部分，并确信导致错误的向量中最低的元素是元素17（块0中的线程16）。
接下来，查看初始化（或应该初始化）数组的代码行：
cudaMemset（d_vec，0，块*螺纹）；//将阵列归零
检查的定义cudaMemset卡，它需要三个参数：
- 指向要设置的设备内存的指针(
d_vec在这种情况下）
- 该内存区域中每个字节应设置的值（在本例中为0）
- 要设置的字节数(
垫块*螺纹)
现在问题开始变得更加清楚。矢量包含64个元素，由垫块*螺纹，但每个元素都是一个浮点值，因此整个向量的长度为256字节。cuda内存集只初始化了前64个字节 （前16个元素），这意味着剩余192个字节（相当于48个元素）未初始化。这48个元素对应48个错误。
这与元素17（线程16，块0）是第一个导致错误的观察结果有关。答对了，发现问题了。
要解决问题，请更改cuda内存集呼叫：
cudaMemset（d_vec，0，sizeof（float）*块*螺纹）；
并检查以确保消毒剂满意。
检查未使用的内存
的另一个功能初始化检查该工具识别到应用程序结束时尚未访问的已分配设备内存。在某些程序中，这可能需要慎重考虑——例如，使用大型静态缓冲区来处理一系列潜在的问题大小。但当这更可能是一个错误导致错误时，请使用初始化检查，如以下代码示例所示：
#包括<stdio.h>#定义N 10__global__void initArray（float*数组，float值）{int threadGlobalID=threadIdx.x+blockIdx.x*blockDim.x；if（threadGlobalID<N）数组[threadGlobalID]=值；回报；}整型main（）{float*数组；const int numThreadsPerBlock=4；const int numBlocks=2；cudaMalloc（（void**）&数组，sizeof（float）*N）；initArray（数组，3.0）；cudaDeviceSynchronize（）；cudaFree（阵列）；退出（0）；}
这个基本的代码示例将揭示潜在的错误。它正在初始化数组，但线程数和块数是硬编码的。执行配置<<< … >>>启动一个由八个线程组成的网格，而数据集有10个元素。最后两个元素未使用。
使用轨用存储器选项。所需的语法取决于使用的CUDA版本。对于12.3之前的版本，请提供参数对使用以下代码示例：
从12.3版本开始，您不必提供参数：
$nvcc-o unused-lineinfo unused.cu$compute-sanitizer--tool initcheck--track-unused-memory/未使用的=====计算机消毒器=====分配0x7fe0a7200000中的未使用内存，大小为40字节======偏移量0x20（0x7fe0a7200020）处未写入8个字节=====20%的分配未使用。==============错误摘要：1个错误
显然，轨用存储器指示40字节数组（10 x 4字节浮点）包括8个未写入的字节。使用数组地址（第一个长0x…数字）和偏移量（0 x 20，即32位十进制，因此32个字节或8个浮点）查看哪些字节未使用。正如预期的那样，未使用数组中的浮点9和浮点10。
要解决此问题，请使用N个定义num块以下为：
const int numBlocks=（N+numThreadsPerBlock-1）/numThreadsPerBlock；
这个--轨用存储器选项用于分配给的设备内存库达·马洛克。该功能不适用于统一内存(cudaMalloc管理分配的内存）。
同步检查
除了块和扭曲之外，同步不同级别线程的功能是一个强大的CUDA功能，由协作组编程模型支持。协作组具有用于定义、分区和同步线程组的设备代码API，与标准组相比具有更大的灵活性和控制能力同步线程函数，它同步块中的所有线程。有关更多信息，请参阅合作小组：灵活的CUDA线程编程.
然而，这种能力带来了更多引入错误的机会。这里就是同步检查可以帮助识别和解决CUDA代码中的同步错误。同步检查可以确定CUDA应用程序是否正确使用同步原语及其协作组API对应项。
同步的一个有趣用途是将掩码应用于线程扭曲。设置扭曲，使某些线程为true，其他线程为false，从而使每个线程能够根据该属性单独执行不同的操作。有关更多信息，请参阅使用CUDA扭曲级别基本体.
帮助实现这一点的一个有用函数是__气球同步定义如下：
unsigned int __ballot_sync（无符号int掩码，int谓词）；
面具是初始掩码，通常创建时所有位都设置为1，表示扭曲中的所有线程最初都处于活动状态。谓词是每个线程评估的条件，其中谓词计算结果为真的（非零）或假（零）。
这个投票函数计算warp中每个线程的谓词，并返回表示该线程结果的掩码。它还提供了一个同步点。经纱上的所有线都必须达到这个高度__选票同步在他们中的任何人能够继续前进之前。
例如，设置一个掩码，使扭曲中的偶数线程为真，奇数线程为假：
__ballot_sync（0xffffff，线程ID%2==0）；
初始遮罩0xffffff（0xffff）是十六进制表示形式，计算结果为11111111111111111111111111111111二进制格式。这样可以确保所有32个线程都参与投票。
投票结果是一个面具，0xaaaaaaaa，以二进制表示为10101010101010101010101010101010。偶数线程（线程ID 0、2、4…）设置为true，奇数线程设置为false。
选票通常与__同步扭曲，它可以根据提供的掩码同步扭曲中的线程。
以下示例同时使用_气球同步和_同步扭曲以下为：
静态constexpr int NumThreads=32；__shared__ int smem[NumThreads]；__全局__void sumValues（int*sum_out）{int threadID=线程Idx.x；无符号整数掩码=__ballot_sync（0xffffffff，threadID<（NumThreads/2））；if（threadId<=（线程数/2））{smem[threadId]=线程ID；__syncwarp（掩码）；if（线程ID==0）{*sum_out=0；for（int i=0；i<（NumThreads/2）++i）*sum_out+=最小值[i]；}}__syncThreads（）；}整型main（）{int*sum_out=nullptr；cudaMallocManaged（（void**）&sum_out，sizeof（int））；sumVaules<<<1，NumThreads>>>（sum_out）；cudaDeviceSynchronize（）；printf（“总和=%d\n”，*Sum_out）；cudaFree（sum_out）；返回0；}
在进一步阅读之前，请看一看代码，并根据您对投票和同步扭曲功能。看看你能不能找出问题所在。（这一次得50分——这更具挑战性。）
此代码的目的是让每个线程为共享内存分配一个值，然后将所有值相加以得到一个答案。然而，这只适用于一半可用线程。通过执行配置设置32个线程的单个扭曲<<<1，numThreads>>>执行内核sumValues（总和值）.
在该内核中，使用__选票同步具有threadID<NumThreads/2作为谓词，在曲速的前半部分求值为true，其中螺纹ID<16（线程0、1、..15）。
对于这16个线程，指定一个值(线程ID)共享内存，并执行__syncwarp（掩码）同步这些线程，以确保它们都已写入共享内存。然后更新汇总（_O）基于这些值的全局和。
接下来，尝试编译并运行以下代码：
$nvcc-o ballot_example-lineinfo ballot_example.cu$ ./选票_示例总和=0
答案是零，不正确。它应该是120（15+14+13+…+2+1+0）。
你发现错误了吗？代码的条件部分是使用执行的if（threadId<=（NumThreads/2））。此代码使用<=而不是<作为比较器，意味着前17个线程执行。
当线程17尝试调用时会发生什么同步扭曲当它不包含在掩码中时？它会导致整个内核停止运行，因此永远无法进行总和计算。因此，输出为零。
所有这些都会以静默方式失败，只有不正确的输出表示出现了问题。在更复杂的代码中，这可能是一场噩梦。
使用同步检查提供了以下内容：
$compute-sanitizer--工具同步检查--show-backtrace否/选票_示例=====计算机消毒器=====检测到屏障错误。无效参数======在/home/pgraham/Code/devblog/NCS_Part2/ballot_example.cu:32:sumValues（int*）中的0x220处======通过块（0,0,0）中的线程（0,0，0）=========. . .============检测到屏障错误。无效的参数======在/home/pgraham/Code/devblog/NCS_Part2/ballot_example.cu:32:sumValues（int*）中的0x220处====~====通过块（0,0,0）中的线程（16,0.0）=========总和=0=====错误摘要：17个错误
关于这17个“无效参数”错误同步检查文档声明如果不是所有线程都达到__同步扭曲在面具参数。
在这种情况下，螺纹17或螺纹（16,0,0）不是 在掩码中处于活动状态，因此不应调用同步扭曲。这会导致所有其他线程调用同步扭曲也注册错误。他们正在单独呼叫同步扭曲，但因为其中一个导致它失败，所有其他同步扭曲呼叫也必须失败。这是一个集体操作，共导致17个错误。
结论
本文向您介绍了几个示例，说明如何使用初始化检查和同步检查中的功能NVIDIA电脑消毒剂。要开始使用Compute Sanitizer，请下载CUDA工具包。
要了解更多信息，请访问/NVIDIA/电脑消毒剂样本在GitHub上，并阅读计算消毒剂文档。在中加入对话NVIDIA开发者论坛专用于消毒工具。
祝你找虫子好运！
