sporadic error: "after cub_::DeviceRadixSort::SortPairs(1): invalid argument"

[schiller-manuel]

In my program the following error ocurrs only sporadically:

after cub_::DeviceRadixSort::SortPairs(1): invalid argument

I only found a reference to that error message in stable_radix_sort.inl, so probably the thrust::sort_by_key in my program is causing this.
Since this error only occurs very infrequently, it is very hard to gain any more insights.
Do you have any clue what could be the cause of this problem? Any idea on how to debug this?

---

latest thrust from github
CUDA 6.0
Ubuntu 14.04 (x64)
gcc 4.8.2
NVIDIA Driver 343.22
NVIDIA GTX 680

[jaredhoberock]

To debug it, first find a data set, which, when sorted, reproduces the bug in isolation. Once you find this data set, write a small program which sorts the data repeatedly in a loop. (In a loop since the bug occurs only sporadically.) Once you have a reproducer program, post it here and I'll take a look.

[jaredhoberock]

Also, since you said you're using CUDA 6, I recommend updating to the latest version of CUDA, which as of today is CUDA 6.5. Thrust does not get tested against old nvcc.

[jvalenciag]

Same problem here when I try to sort a vector of size > 1.5GB
Cuda 7.0
Ubuntu 14.04 (x64)
gcc 4.8.4
Tesla K40c

[renbozqin]

Same problem.
Cuda 7.5
Centos7.0 (x64)
gcc 4.8.4
gtx Titan X

[aliyekta]

I saw it on two different GPUs with the same code, tried going after this bug, after a lengthy period of debug, i saw that the following line defines a num_block size which makes the consequent kernel call , to have num_block greater than max block size supported by the card which in my kepler card is 65536.

thrust/system/cuda/detail/detail/stable_merge_sort.inl:397
Size num_blocks = thrust::detail::util::divide_ri(n, work_per_block);

[3gx]

What is the compilation line. Make sure you pass at least -arch=sm_30. Kepler+ support block sizes > 65535

[aliyekta]

Thanks, I did not have arch in the compile line, this solved my issue with the call and it no longer gives an error.
Thanks

[schiller-manuel]

@egaburov does this mean thrust only supports sm_30 and up?

[Dredok]

same problem here.
I have tested this on a GTX 1070 and a K2200, in windows 10 vs2013 x64.
I use CUDA 8.0, compile support for 30, 50 and 61.

I have reproduced this bug in a fairly complex code, with lots of libraries. Everything runs smoothly until I make a call to a library which uses constant memory, after calling this library I can not call thrust::sort_by_key because it will produce this error.
I will try to isolate the code and fill a bug

[jbedorf]

I saw this error as well, but it turned out to be caused in a completely different part of my own code.
In case anyone runs into this, this happened for me:
An invalid value was set using cudaDeviceSetLimit. The return value of this function was not checked and hence the error got unnoticed. Follow up function and kernel calls did have their error codes checked, but progressed fine.
The error only surfaced inside the Thrust/cub code, because it uses the 'cudaPeekAtLastError' function which made the 'invalid argument' error surface again.

[mcfatealan]

Met this bug on Tesla K40M. The weird thing is that using "-arch=sm_35" does not solve this issue. After struggling with all the solutions I can come out (including switching to CUDPP library, later found that it used thrust too). In the morning it suddenly stroke me that I should test "-arch=sm_30". The code no longer crashed, instead it can become extremely slow periodically.

[mcfatealan]

reboot machine, now works

[brycelelbach]

Is this still an issue?