Instruction 'shfl' without '.sync' is deprecated

[andrewcorrigan]

When compiling Thrust algorithms with Clang, there are tons of warnings generated. These warnings are reproducible across multiple versions of Clang, but I am currently on:

% clang++ --version
clang version 11.0.0

Minimal reproducer:

#include <thrust/sort.h>
#include <thrust/fill.h>
#include <thrust/device_vector.h>

int main(int argc, char** argv)
{
    thrust::device_vector<float> x(1000);
    thrust::fill(x.begin(), x.end(), 1.0f);
    thrust::sort(x.begin(), x.end());

    return 0;
}

Output:

% clang++ -x cuda --cuda-path=/usr/local/install/cuda-10.1 --cuda-gpu-arch=sm_61 thrust_shfl_without_sync.cpp -I/thrust/ -std=c++14 -L/usr/local/install/cuda-10.1/lib64 -lcudart
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 4090; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 12272; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 12323; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 12355; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 12467; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 12499; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 15067; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 15918; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 15950; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 17176; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 18016; warning : Instruction 'vote' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 19878; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 19929; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 19961; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 20082; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 20114; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 20918; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 20950; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 22209; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 22260; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 22292; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 22413; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/thrust_shfl_without_sync-df0f0e.s, line 22445; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version

[alliepiper]

Does applying NVIDIA/cub#170 to CUB fix it? It sounds like the same issue, though there may be additional places in Thrust that need similar patches.

[andrewcorrigan]

From the description, it must. I'll try it out in a second.

I wasn't sure where to report this (Thrust/CUB/Clang)... Good to see the Google/Clang teams was already on this.

[andrewcorrigan]

That fixed it! I'll close this out and track the CUB PR.

Just to confirm, the warnings are completely gone. My code has pretty wide coverage of Thrust, so I doubt there's anything specific to Thrust.