关于error.cuh中的CHECK宏报错信息的疑问。

您好，最近看了您的作品收获很多，并使用了您书中提到的一些技巧。在使用过程中我遇到了如下问题。我在PyTorch框架中编写了一些自定义CUDA算子，其中使用了error.cuh中的CHECK宏进行返回值检测，有如下几个现象。

int *pts_mask = NULL;
cudaMalloc(&pts_mask, boxes_num * pts_num * sizeof(int));  // (N, M)
CHECK(cudaMemset(pts_mask, -1, boxes_num * pts_num * sizeof(int)));

在程序运行过程中（已经稳定运行了很久，cudaMemset函数已经被执行了很多次），被CHECK的这行代码报如下错误：

CUDA Error:
File:       /torch_custom_ext/torchex/src/sparse_roi_voxelization/sparse_roiaware_pool3d_kernel.cu
Line:       264
Error code: 1
Error text: invalid argument

在某些时候（随机发生），错误会发生在其他自定义算子中，并且提示out of memory错误。

如果将代码改为

int *pts_mask = NULL;
CHECK(cudaMalloc(&pts_mask, boxes_num * pts_num * sizeof(int)));  // (N, M)
cudaMemset(pts_mask, -1, boxes_num * pts_num * sizeof(int));

错误码会变成: out of memory

根据以上现象我初步判断真正的错误原因就是显存不足（我的运行模型确实也比较大），我的问题是为什么在某些情况下会报出上面提到的invalid argument的错误。是否是因为分配内存失败过后将NULL指针传给了cudaMenset?

感谢提问。我不知道确切的原因，但建议尝试在运行你的程序之前，试试在命令行设置如下环境变量： export CUDA_LAUNCH_BLOCKING=1 这样会让所有核函数调用变成同步的，也许能帮助找到错误的真正发生的地方。该做法只适合在找错时用，因为它会降低程序的执行速度。

哲勇

On Fri, Sep 17, 2021 at 9:22 PM Lue Fan @.***> wrote:

您好，最近看了您的作品收获很多，并使用了您书中提到的一些技巧。在使用过程中我遇到了如下问题。我在PyTorch框架中编写了一些自定义CUDA算子，其中使用了error.cuh中的CHECK宏进行返回值检测，有如下几个现象。

1.

int pts_mask = NULL; cudaMalloc(&pts_mask, boxes_num pts_num sizeof(int)); // (N, M) CHECK(cudaMemset(pts_mask, -1, boxes_num pts_num * sizeof(int)));

在程序运行过程中（已经稳定运行了很久，cudaMemset函数已经被执行了很多次），被CHECK的这行代码报如下错误：

CUDA Error: File: /torch_custom_ext/torchex/src/sparse_roi_voxelization/sparse_roiaware_pool3d_kernel.cu Line: 264 Error code: 1 Error text: invalid argument

在某些时候（随机发生），错误会发生在其他自定义算子中，并且提示out of memory错误。

1.

如果将代码改为

int pts_mask = NULL; CHECK(cudaMalloc(&pts_mask, boxes_num pts_num sizeof(int))); // (N, M) cudaMemset(pts_mask, -1, boxes_num pts_num * sizeof(int));

错误码会变成: out of memory

根据以上现象我初步判断真正的错误原因就是显存不足（我的运行模型确实也比较大），我的问题是为什么在某些情况下会报出上面提到的invalid argument的错误。是否是因为分配内存失败过后将NULL指针传给了cudaMenset?

— You are receiving this because you are subscribed to this thread. Reply to this email directly, view it on GitHub https://github.com/brucefan1983/CUDA-Programming/issues/12, or unsubscribe https://github.com/notifications/unsubscribe-auth/AF546OKP56ZEUXIIMFVHX53UCM6JZANCNFSM5EHDO33Q . Triage notifications on the go with GitHub Mobile for iOS https://apps.apple.com/app/apple-store/id1477376905?ct=notification-email&mt=8&pt=524675 or Android https://play.google.com/store/apps/details?id=com.github.android&referrer=utm_campaign%3Dnotification-email%26utm_medium%3Demail%26utm_source%3Dgithub.

感谢回复，我在运行上述代码的时候已经使用了CUDA_LAUNCH_BLOCKING=1。我会进行进一步check，看cudaMemset函数在遇到NULL指针时是否会提示argument invalid。

另外我有一个话题之外的疑问：不使用CUDA_LAUNCH_BLOCKING=1的时候，报错信息可能并不准确。我个人理解是由于主机和设备或者不同核函数之间的不同步导致的，比如当某个核函数报错的时候，主机代码运行到了另外的地方，所以此时的报错信息就错误的指向了主机代码正在运行的地方。

请问您我的这一理解是否准确？我查找过书中关于CUDA stream的章节，可惜并没有找到相关知识。

如果没有使用 CUDA_LAUNCH_BLOCKING=1，也没有在调用核函数后用某种方式同步主机与设备，那么主机在调用一个核函数后将不会等待核函数执行。此时，核函数的错误可能不能被捕捉到，而是被后面其它某个API的错误覆盖了。此时，CHECK宏捕捉到的错误就可能不是最早的错误。

On Fri, Sep 17, 2021 at 10:09 PM Lue Fan @.***> wrote:

感谢回复，我在运行上述代码的时候已经使用了CUDA_LAUNCH_BLOCKING=1。我会进行进一步check，看cudaMemset函数在遇到NULL指针时是否会提示argument invalid。

另外我有一个话题之外的疑问：不使用CUDA_LAUNCH_BLOCKING=1的时候，报错信息可能并不准确。我个人理解是由于主机和设备或者不同核函数之间的不同步导致的，比如当某个核函数报错的时候，主机代码运行到了另外的地方，所以此时的报错信息就错误的指向了主机代码正在运行的地方。

请问您我的这一理解是否准确？我查找过书中关于CUDA stream的章节，可惜并没有找到相关知识。

— You are receiving this because you commented. Reply to this email directly, view it on GitHub https://github.com/brucefan1983/CUDA-Programming/issues/12#issuecomment-921827098, or unsubscribe https://github.com/notifications/unsubscribe-auth/AF546OIWTNDTJGC4RW3PQZTUCNDYJANCNFSM5EHDO33Q . Triage notifications on the go with GitHub Mobile for iOS https://apps.apple.com/app/apple-store/id1477376905?ct=notification-email&mt=8&pt=524675 or Android https://play.google.com/store/apps/details?id=com.github.android&referrer=utm_campaign%3Dnotification-email%26utm_medium%3Demail%26utm_source%3Dgithub.

另外，如果要捕捉核函数的错误，还需要在每个核函数调用后加上如下代码： // a kernel here CHECK(cudaGetLastError()); //必须有 CHECK(cudaDeviceSynchronize()); // 如果设置了 CUDA_LAUNCH_BLOCKING=1 就可以去掉该同步

On Fri, Sep 17, 2021 at 10:16 PM Bruce Fan @.***> wrote:

如果没有使用 CUDA_LAUNCH_BLOCKING=1，也没有在调用核函数后用某种方式同步主机与设备，那么主机在调用一个核函数后将不会等待核函数执行。此时，核函数的错误可能不能被捕捉到，而是被后面其它某个API的错误覆盖了。此时，CHECK宏捕捉到的错误就可能不是最早的错误。

On Fri, Sep 17, 2021 at 10:09 PM Lue Fan @.***> wrote:

感谢回复，我在运行上述代码的时候已经使用了CUDA_LAUNCH_BLOCKING=1。我会进行进一步check，看cudaMemset函数在遇到NULL指针时是否会提示argument invalid。

另外我有一个话题之外的疑问：不使用CUDA_LAUNCH_BLOCKING=1的时候，报错信息可能并不准确。我个人理解是由于主机和设备或者不同核函数之间的不同步导致的，比如当某个核函数报错的时候，主机代码运行到了另外的地方，所以此时的报错信息就错误的指向了主机代码正在运行的地方。

请问您我的这一理解是否准确？我查找过书中关于CUDA stream的章节，可惜并没有找到相关知识。

— You are receiving this because you commented. Reply to this email directly, view it on GitHub https://github.com/brucefan1983/CUDA-Programming/issues/12#issuecomment-921827098, or unsubscribe https://github.com/notifications/unsubscribe-auth/AF546OIWTNDTJGC4RW3PQZTUCNDYJANCNFSM5EHDO33Q . Triage notifications on the go with GitHub Mobile for iOS https://apps.apple.com/app/apple-store/id1477376905?ct=notification-email&mt=8&pt=524675 or Android https://play.google.com/store/apps/details?id=com.github.android&referrer=utm_campaign%3Dnotification-email%26utm_medium%3Demail%26utm_source%3Dgithub.

brucefan1983 / CUDA-Programming

关于error.cuh中的CHECK宏报错信息的疑问。 #12