Invalid __global__ write of size 4

  • 5 replies
  • 541 views
Invalid __global__ write of size 4
« 于: 十月 23, 2022, 08:57:44 am »
请问一下大家,利用cuda-memcheck检查代码,出现这样的错误怎么回事呀
Invalid __global__ write of size 4
=========     at 0x000001f0 in cudiv::KerCalBeginEndCube(int, int*, int2*)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7ff320c7b000 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/lib64/libcuda.so.1 [0x2ea4a1]
=========     Host Frame:./sphgpu [0x24bab]
=========     Host Frame:./sphgpu [0x7b1e8]
=========     Host Frame:./sphgpu [0xff54]
=========     Host Frame:./sphgpu [0xf5bc]
=========     Host Frame:./sphgpu [0xf5f4]
=========     Host Frame:./sphgpu [0xd308]
=========     Host Frame:./sphgpu [0x8783]
=========     Host Frame:./sphgpu [0xa494]
=========     Host Frame:./sphgpu [0x5d44]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22555]
=========     Host Frame:./sphgpu [0x57ee]
=========
CUDA Error:
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
    File:       ConfigDomain.cu
    Line:       196
    Error code: 719
=========     Saved host backtrace up to driver entry point at error
    Error text: unspecified launch failure
=========     Host Frame:/lib64/libcuda.so.1 [0x424d26]
=========     Host Frame:./sphgpu [0x565a7]
=========     Host Frame:./sphgpu [0xd387]
=========     Host Frame:./sphgpu [0x8783]
=========     Host Frame:./sphgpu [0xa494]
=========     Host Frame:./sphgpu [0x5d44]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22555]
=========     Host Frame:./sphgpu [0x57ee]
=========
========= ERROR SUMMARY: 2 errors

Re: Invalid __global__ write of size 4
« 回复 #1 于: 十月 31, 2022, 04:47:04 pm »
请问一下大家,利用cuda-memcheck检查代码,出现这样的错误怎么回事呀
Invalid __global__ write of size 4
=========     at 0x000001f0 in cudiv::KerCalBeginEndCube(int, int*, int2*)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7ff320c7b000 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/lib64/libcuda.so.1 [0x2ea4a1]
=========     Host Frame:./sphgpu [0x24bab]
=========     Host Frame:./sphgpu [0x7b1e8]
=========     Host Frame:./sphgpu [0xff54]
=========     Host Frame:./sphgpu [0xf5bc]
=========     Host Frame:./sphgpu [0xf5f4]
=========     Host Frame:./sphgpu [0xd308]
=========     Host Frame:./sphgpu [0x8783]
=========     Host Frame:./sphgpu [0xa494]
=========     Host Frame:./sphgpu [0x5d44]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22555]
=========     Host Frame:./sphgpu [0x57ee]
=========
CUDA Error:
========= Program hit cudaErrorLaunchFailure (error 719) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize.
    File:       ConfigDomain.cu
    Line:       196
    Error code: 719
=========     Saved host backtrace up to driver entry point at error
    Error text: unspecified launch failure
=========     Host Frame:/lib64/libcuda.so.1 [0x424d26]
=========     Host Frame:./sphgpu [0x565a7]
=========     Host Frame:./sphgpu [0xd387]
=========     Host Frame:./sphgpu [0x8783]
=========     Host Frame:./sphgpu [0xa494]
=========     Host Frame:./sphgpu [0x5d44]
=========     Host Frame:/lib64/libc.so.6 (__libc_start_main + 0xf5) [0x22555]
=========     Host Frame:./sphgpu [0x57ee]
=========
========= ERROR SUMMARY: 2 errors

你这里是访存越界/无效指针/下标,导致kernel写入缓冲区的时候,挂了。不过你使用cuda-memcheck没有看到具体的kernel行。

可以用nvcc编译的时候,加上-G -g参数来编译,然后重新执行你的cuda-memcheck过程,这种情况下,cuda-memcheck会告诉你具体某个kernel出错的行号。然后仔细观察对应的源代码行,是否有上一段说明的情况。

Re: Invalid __global__ write of size 4
« 回复 #2 于: 十一月 02, 2022, 11:17:58 am »
/data/apps/cuda/11.4/bin/nvcc -G  --device-c system.cu -o system.o    请问是类似于这样对吗

Re: Invalid __global__ write of size 4
« 回复 #3 于: 十一月 02, 2022, 12:40:16 pm »
/data/apps/cuda/11.4/bin/nvcc -G  --device-c system.cu -o system.o    请问是类似于这样对吗

如果你需要分步编译的话,则是这样的。需要-G和-dc的。没错。(或者你至少使用-dc和-lineinfo也可以,如果-G
编译出来的代码,运行速度极度缓慢,而在cuda-memcheck下半天没结果的话)。

Re: Invalid __global__ write of size 4
« 回复 #4 于: 十一月 07, 2022, 04:40:26 pm »
不好意思,再次打扰了,如果使用makefile 编译的话我这样source=main.cu

RUN:=./sphgpu

CC=/data/apps/cuda/11.4/bin/nvcc

$(RUN):$(source) sph.o system.o SPHGPU.o ConfigDomain.o
   $(CC) -lineinfo $(source) sph.o system.o SPHGPU.o ConfigDomain.o -o $(RUN)
   
sph.o:sph.cu
   $(CC) -lineinfo --device-c sph.cu -o sph.o
   
system.o:system.cu
   $(CC) -lineinfo --device-c system.cu -o system.o
   
SPHGPU.o:SPHGPU.cu
   $(CC) -lineinfo --device-c SPHGPU.cu -o SPHGPU.o
   
ConfigDomain.o:ConfigDomain.cu
   $(CC) -lineinfo --device-c ConfigDomain.cu -o ConfigDomain.o
   
.PHONY:clean
clean:
   -rm -rf $(RUN)
   -rm -rf *.o
或者将 -lineinfo改成 -G之后使用cuda-memcheck还是没有出现错误的行号,麻烦您指正一下错误,我尝试在网上搜相关资料没有搜到,麻烦您了

Re: Invalid __global__ write of size 4
« 回复 #5 于: 十一月 18, 2022, 10:24:55 am »
不好意思,再次打扰了,如果使用makefile 编译的话我这样source=main.cu

RUN:=./sphgpu

CC=/data/apps/cuda/11.4/bin/nvcc

$(RUN):$(source) sph.o system.o SPHGPU.o ConfigDomain.o
   $(CC) -lineinfo $(source) sph.o system.o SPHGPU.o ConfigDomain.o -o $(RUN)
   
sph.o:sph.cu
   $(CC) -lineinfo --device-c sph.cu -o sph.o
   
system.o:system.cu
   $(CC) -lineinfo --device-c system.cu -o system.o
   
SPHGPU.o:SPHGPU.cu
   $(CC) -lineinfo --device-c SPHGPU.cu -o SPHGPU.o
   
ConfigDomain.o:ConfigDomain.cu
   $(CC) -lineinfo --device-c ConfigDomain.cu -o ConfigDomain.o
   
.PHONY:clean
clean:
   -rm -rf $(RUN)
   -rm -rf *.o
或者将 -lineinfo改成 -G之后使用cuda-memcheck还是没有出现错误的行号,麻烦您指正一下错误,我尝试在网上搜相关资料没有搜到,麻烦您了

可以考虑用cuda-gdb来确定行号,有两种方案(精确和非精确的):
(1)直接用cuda-gdb --args 你的原本程序的命令行的方式,来运行你的程序。
然后在cuda-gdb的提示符下,输入:
set cuda memcheck on
然后输入r回车继续运行你的程序。

这种情况下,运行速度会变得较慢(有的时候会非常慢),但如果有kernel访存错误,会自动停留在这一行,并显示这一行的源代码(需要-g -G)。

(2)如果方法(1)慢到了无法忍受(例如经过了很久都没有停住),可以先全速运行,然后等待kernel出错的时候,暂停住进程(而不是返回出错代码或者挂掉):
export CUDA_DEVICE_WAITS_ON_EXCEPTION=1
你原本的程序命令行 回车
(耐心等待程序挂掉)
然后当屏幕上出现类似你的kernel访存错误,请使用调试器来调试进程XXXX(一个进程编号)的时候,继续输入:
sudo cuda-gdb --pid=XXXX
然后回车几次后,会停留在出错的源代码行附近(注意不是精确的)。
如果你看到的不是源代码,而是某处含有intrinsic字样的汇编代码之类的,继续输入bt回溯,一般再上一个就是出错的你的代码的位置(的附近)。

注意第二种方法不是精确的(因为流水线之类的原因),实际位置往往会偏后一点点。不过这种方法几乎万能,稍微往上看看就能找到。