NVBug 3262468: Bogus data produced in CUDA 11(.1) by thrust::sort_by_key with custom comparison operator over two (zipped) key arrays

[Oblomov]

Hello,

while testing GPUSPH on CUDA 11.1 I discovered that the thrust::sort_by_key procedure that we use to sort the particle indices is now producing bogus data. The sort uses a custom comparator that fetches data from two different arrays (using zip_iterator etc). The sort results in one of the (sorted) key arrays being clobbered with invalid data.

I am currently in the process of finding a minimal test case, but in the mean time the bug can be observed in action on the cuda11-thrust-sort-bug of GPUSPH, by running

make DamBreak3D && ./DamBreak3D --maxiter 1 | grep '64656 255 '

from the git working directly. This should produce no results (as it does on CUDA 10) if the sort is correct, but it results in numerous hits (all clobbered entries with bogus values) in CUDA 11.

Thanks for looking into this,

[alliepiper]

That's troubling...let me know when you've got a pure Thrust reproducer and I'll take a look.

[Oblomov]

I'm still working on trying to get a minimal reproducible, which is turning out to be harder than expected —quite surprising, given that the sorting kernel is only the second kernel to be invoked, preceded by a trivial initialization kernel.

Is there an easy way to print each CUDA API call that is being made by a program, in the order they are being made?

OK, I've found how to produce an API trace with nvprof and apparently the standalone case produces a thrust kernel that consumes one less register than the one in GPUSPH, so it's obviously not exactly equivalent, despite the code being the same. Down the rabbithole we go ...

[Oblomov]

Interesting that by updating the NVIDIA driver now the clobbered entries have a different value (64672 instead of 64656 or 0xFCA0 instead of 0xFC90).

[Oblomov]

OK I've made some significant progress in producing the minimal test case AND found a workaround from within GPUSPH while the bug gets fixed: apparently the issue is somehow connected to the comparator functor being defined inside the scope of a template class. Moving it to the top-level scope (it's actually independent from the class, so this can and probably should be done anyway) seems to fix the issue. So there is something in the class template definition that somehow interferes with the comparator functor and its usage in thrust.

[Oblomov]

I've pushed a minimal reproducible on the thrust-cuda11-bug branch of my titanxstall project (sorry for recycling an older repo —it's still relevant since that was for another issue we had with thrust ;-)).

git clone https://github.com/Oblomov/titanxstall -b thrust-cuda11-bug
cd titanxstall
make test

On CUDA 10, All OK!, but with CUDA 11.1 you'll get something like:

AFTER: 1: 64880 255 64880 255 0 22477
terminate called after throwing an instance of 'std::runtime_error'
  what():  clobbered info

It's possible to change CUDA version by setting NVCC=/path/to/nvcc in Makefile.local (create the file if not present).

The issue can be solved simply by pulling the ptype_has_compare functor definition out of the CUDAEngine class template. Note also that if the class is not a class template the problem doesn't seem to arise, and if it's a class template that depends on some integer value type (e.g. an enum) it doesn't arise either.

Hopefully this is enough information to find the issue!

[alliepiper]

Thanks, I'll try to take a look at this in the next few weeks.

[alliepiper]

I found some time to look at this, but got different results (CUDA 11.2 and g++ 7.5.0):

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  merge_sort: failed to synchronize: cudaErrorIllegalAddress: an illegal memory access was encountered
Aborted 

The source file refers to a data.idx file that doesn't seem to exist in the repo, that may be why.

[Oblomov]

Hello, I've added the data file to the repository. Hopefully this helps.

[alliepiper]

Thanks, I can repro the issue now. I'll try to debug this a bit more soon.