cuda-memcheck，如何从地址得到源代码？

Question

我成功地使用cuda-memcheck来检测错误的内存访问。在编译cuda代码时，使用-g -G参数可以得到类似以下的良好源代码位置：

========= Error: process didn't terminate successfully
========= Invalid __global__ read of size 1
=========     at 0x00000710 in /some/path/somefile.cu:117:some_function
=========     by thread (0,14,0) in block (1,16,0)
=========     Address 0x00abac20 is out of bounds

现在我尝试使用-l开关来获取内存泄漏信息。然而，在这里，我只能获得地址：

========= CUDA-MEMCHECK
========= Leaked 3630 bytes at 0x007d2800
=========
========= Leaked 14740 bytes at 0x008e0700
...
=========
========= LEAK SUMMARY: 11122140 bytes leaked in 39 allocations
========= ERROR SUMMARY: 0 errors
400 bytes at 0x005d2000

我该如何从中获取实际的代码位置？

- Jan Rüegg

Answer 1

提供的泄漏地址并非代码地址，而是数据位置。不幸的是，很难看出这些位置被分配在哪里。

考虑到内存可以在任何地方分配（请记住，指针可以传递、别名等），检查泄漏（即未释放的已分配内存）的唯一方法是程序退出时进行。因此，当您的程序退出时，cuda-memcheck会检查已分配但未释放的内存块，并向您提供内存块的地址，但它无法将其与其分配时间联系起来。

相反，最简单的方法是手动检查代码，确保所有的cudaMalloc()调用都有匹配的cudaFree()调用。然而，这可能是一个相当费力的过程...

- Tom

---

VC++通过为每个分配分配一个序列号来处理此问题，然后在退出时将该序列号与泄漏摘要一起输出。通过使用相同的代码和相同的数据两次运行，可以首先找到泄漏的分配的序列号，然后确定该分配发生的位置。 - Roger Dahl

为了跟进我的上一个评论，您可以创建一个包装器来增加静态声明的序列号，并将序列号和分配的内存地址写入stdout或文件中。程序结束后，在cudaMalloc()包装器打印的内容中查找泄漏的地址，以找到泄漏的分配的序列号。然后修改cudaMalloc()包装器，使其在出现该序列号时中断并重新运行程序。 - Roger Dahl