cudaErrorIllegalAddress error in DeviceSelect::If() about cub HOT 5 CLOSED

Launching my app under cuda-memcheck I got following message from it:
Invoking scan_init_kernel<<<306, 128, 0, 139992313554640>>>()

========= Invalid __global__ write of size 4
=========     at 0x00000128 in void cub::DeviceCompactInitKernel<cub::ScanTileState<int, bool=1>, unsigned int*>(int, int, bool=1)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7f52921b9268 is out of bounds

from cub.

After I compiled the code with -lineinfo option I got the row number:
at 0x00000128 in /home/daktfi/polymatica/cub/cub/device/dispatch/dispatch_scan.cuh:88:void cub::DeviceCompactInitKernel<cub::ScanTileState<int, bool=1>, unsigned int*>(int, int, bool=1)
Hope this helps.
The -G option did not worked out - my code is part of larger project and at the moment I'm unable to compile it separately, as according to this NVIDIA/thrust#864 it requires extra nvcc linking pass.

from cub.

I finally got the working standalone project (attached) which reproduces error and get it commented. There's quite complicated logic behind the the issue, so if You have question - I'll gladly explain.

devselect.zip

from cub.

Hi, I think the "an illegal memory access was encountered" issue is probably related to cub::DeviceSelect:If() expecting a pointer to device memory for where to output the number of items selected. (You are simply referencing the auto-var length on the host thread's program stack, which isn't visible to the GPU threads.) You'll want to create a device-allocation for that.

Let me know if that helps!

from cub.

Only one word: "oops". :-(
The most stupid bugs are hardest to find...
Worked like a charm! Thanks a lot!!!

from cub.