bm3d-gpu icon indicating copy to clipboard operation
bm3d-gpu copied to clipboard

Published 20 hours ago • verified publisher icon DawyD

Illegal memory access on some images

Open trougnouf opened this issue 6 years ago • 1 comments

I get the following error:

Sigma = 10
Number of Steps: 2
Color denoising: yes
Noise variance for individual channels (YCrCb if color): 32 29 33 
width: 4231 height: 3051
Shared memory : 46KB/48KB
Number of warps: 20
Processing 0%terminate called after throwing an instance of 'thrust::system::system_error'
  what():  bm3d.hpp(788):: an illegal memory access was encountered
test_bm3d.sh: line 6:  5767 Aborted                 (core dumped) bm3d ${NOISYTESTDIR}/${TESTSET}/${TESTIMG} results/test2/bm3d-${SIGMA}/${TESTIMG} ${SIGMA} color twostep

when trying to denoise any of the images in that set https://commons.wikimedia.org/wiki/File:NIND_MuseeL-skull-C500D_ISO200.jpg

trougnouf avatar Feb 23 '19 09:02 trougnouf

Same problem here with a different image. Cropping the image a bit solves the problem. cuda-memcheck reports:

========= CUDA-MEMCHECK
Sigma = 15
Number of Steps: 2
Color denoising: yes
Noise variance for individual channels (YCrCb if color): 74 66 74 
width: 3846 height: 2885
bm3d object created
bm3d parameters set
Shared memory : 46KB/48KB
Number of warps: 20
Processing 0%========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,7,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,4,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,8,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,6,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,5,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,3,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,2,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,1,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Invalid __shared__ read of size 4
=========     at 0x00001228 in /home/dword/Desktop/media_process/denoise/bm3d-gpu/blockmatching.cu:179:block_matching(unsigned char const *, unsigned short*, unsigned int*, uint2, uint2, Params, uint2)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0xfffffff8 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 (cuLaunchKernel + 0x2c5) [0x27cf05]
=========     Host Frame:./bm3d [0x2fb212]
=========     Host Frame:./bm3d [0x2fb407]
=========     Host Frame:./bm3d [0x32f7c5]
=========     Host Frame:./bm3d [0x1a486]
=========     Host Frame:./bm3d [0x345198]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaDeviceSynchronize. 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0x38c0d3]
=========     Host Frame:./bm3d [0x319106]
=========     Host Frame:./bm3d [0x3451b2]
=========     Host Frame:./bm3d [0x34cde2]
=========     Host Frame:./bm3d [0x107ee]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
========= Program hit cudaErrorLaunchFailure (error 4) due to "unspecified launch failure" on CUDA API call to cudaFree. 
=========     Saved host backtrace up to driver entry point at error
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1 [0x38c0d3]
=========     Host Frame:./bm3d [0x326ff6]
=========     Host Frame:./bm3d [0x3470f6]
=========     Host Frame:./bm3d [0x10047]
=========     Host Frame:/lib/x86_64-linux-gnu/libc.so.6 (__libc_start_main + 0xeb) [0x2409b]
=========     Host Frame:./bm3d [0x10d2a]
=========
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  bm3d.hpp(789):: unspecified launch failure
========= Error: process didn't terminate successfully
========= No CUDA-MEMCHECK results found

The error occurred near

			//Sum column distances to obtain patch distance
			uint diff = 0;
			for (uint i = 0; i < params.k; ++i) 
				diff += s_diff[inner_p_x + i]; // <-- This line

dword1511 avatar May 27 '19 14:05 dword1511