CUDA memcheck 地址 - 如何确定代码中的位置？答案

【问题标题】：CUDA memcheck address - how to determine location in code?CUDA memcheck 地址 - 如何确定代码中的位置？
【发布时间】：2011-06-25 12:03:32
【问题描述】：

cuda-memcheck 正在报告发布模式 CUDA 内核的此信息：

========= Error: process didn't terminate successfully
========= Invalid __global__ read of size 4
=========     at 0x000002c8 in xx_kernel
=========     by thread (0,0,0) in block (0,0)
=========     Address 0x10101600014 is out of bounds
=========
========= ERROR SUMMARY: 1 error

此故障仅在释放模式下发生。在 cuda-gdb 下运行时也不会发生这种情况。

如何获取 0x000002c8 地址并确定导致故障的代码？我查看了缓存的中间文件（.ptx、.cubin 等），没有发现明显的方法可以确定错误的源代码。

这是在带有 CUDA 3.2 的 x86_64 Linux 上。

更新：原来这是 3.2 中的编译器错误。升级到 4.0 会使 memcheck 错误消失。另外，我可以从 4.0 开始使用 cuobjdump 对 CUBIN 进行反汇编，但由于它是发布模式并经过优化，因此很难将反汇编与源代码匹配。

【问题讨论】：

你能把你的内核代码贴出来，让我们看看为什么这个线程会访问一个越界区域吗？
不幸的是它的专有源代码，所以我不能发布实际代码。谢谢。

标签： cuda gpgpu gpu

【解决方案1】：

下载CUDA Toolkit 4.0 from the NVIDIA Developer Zone。使用支持 2.x cubin 的新 cuobjdump。

cuobjdump -sass /path/to/your/cubin > /path/to/dump.txt.

示例输出（在 sm_20 cubin 上测试，代码版本 2.3）

    ...
/*6018*/     /*0xe00100075003ff9a*/     CAL 0x46d8;
/*6020*/     /*0x10001de428000000*/     MOV R0, R4;
/*6028*/     /*0x00001de428000000*/     MOV R0, R0;
/*6030*/     /*0x40011de428000000*/     MOV R4, R16;
    ...

【讨论】：

我将下载 4.0 并尝试一下。谢谢。
移动到 4.0 修复了这个错误。谢谢。

【解决方案2】：

内核中的这种错误与内存访问有关，而内存访问不仅基于线程标识符。

考虑到您使用的每个内存区域都已为 GPU 正确分配，因此应该仅基于 threadIdx.x 之类的内容进行访问不会引起任何问题。因此：

要么你有一个错误的索引计算（例如data[blockDim.y * blockDim.x * threadIdx.z + blockDim.x * threadIdx.y + threadIdx.x] 这样的表达式很常见）
或您在索引计算中使用了另一个变量，这使其超出了您的数组范围（例如data[threadIdx.x + offset]）

---- 编辑（跟随 cmets） ----
有关设备> 2.x 的 cuobjdump 的补充信息，请参阅@Cicada 的答案

【讨论】：

你错了@Cicada！拥有thread (0,0,0) 的block (0,0) 并不意味着只能基于这个threadIdx 进行访问。无论thread (0,0,0) 是否尝试基于其唯一标识符访问内存位置，CUDA memcheck 输出都是相同的。
这可能与索引计算有关：再看看我的例子，它在线程标识符和 blockDim 上。我考虑两种情况：与内核配置相关的情况（threadIdx、blockDim、gridDim、blockdIdx...）和与其他来源相关的情况（变量、常量、参数...）
我需要的是一个 cubin 反汇编程序，但 cuobjdump 不适用于 2.x cubin。有其他选择吗？