Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4

[blueardour]

Hi,

Thanks for sharing the great example for the effcient computing.

When compling the code, I met the following error. Could I ask if you have any clue about it?

nvcc -std=c++11 -O3 -w  -arch=sm_75 -maxrregcount=64 -rdc=true  -o imagenet_vgg imagenet_vgg.cu data.cpp -ljpeg
ptxas /tmp/tmpxft_000009e5_00000000-5_imagenet_vgg.ptx, line 459; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_000009e5_00000000-5_imagenet_vgg.ptx, line 515; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_000009e5_00000000-5_imagenet_vgg.ptx, line 541; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_000009e5_00000000-5_imagenet_vgg.ptx, line 566; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_000009e5_00000000-5_imagenet_vgg.ptx, line 591; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_000009e5_00000000-5_imagenet_vgg.ptx, line 689; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
.....

I am compiling on a machine with 2080ti gpu/ cuda10.1/ gpu driver 418.87.00/ ubuntu 18.04 installed.

[uuudown]

Yes, it seems the latest ptxas strictly requires the "sync" suffix for the voting instructions (e.g., ballot, shuffle), let me update my code.

[uuudown]

Please let me know if the compile error still occurs. Thanks!

[blueardour]

Hi, @uuudown

Thank you for your attention and effort on the problem.

I update the code and seems the issue is still there.

> make imagenet_vgg
nvcc -std=c++11 -O3 -w  -arch=sm_75 -maxrregcount=64 -rdc=true  -o imagenet_vgg imagenet_vgg.cu data.cpp -ljpeg
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 459; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 515; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 541; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 566; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 591; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 689; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 741; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 780; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 819; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 858; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 955; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 983; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 1012; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 1053; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 1075; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 1097; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas /tmp/tmpxft_00002c3a_00000000-5_imagenet_vgg.ptx, line 1118; error   : Instruction 'vote' without '.sync' is not supported on .target sm_70 and higher from PTX ISA version 6.4
ptxas fatal   : Ptx assembly aborted due to errors
Makefile:42: recipe for target 'imagenet_vgg' failed
make: *** [imagenet_vgg] Error 255

[uuudown]

Yes, you're right. There are other "__ballot()" functions in param.h and they should be changed to "__ballot_sync()". I just tested using cuda/10.2, it should work now.

[blueardour]

Awesome! It works. Thanks very much.

[MiaoDX]

Just for the record, I built our existing codebase with clang, triggered similar errors.

Maybe https://bugs.llvm.org/show_bug.cgi?id=43505 can be helpful.