instaGRAAL icon indicating copy to clipboard operation
instaGRAAL copied to clipboard

usage of __shfl_down is deprecated

Open arielzn opened this issue 1 year ago • 2 comments

It looks the cuda kernels being created are using a deprecated function __shfl_down :

https://github.com/koszullab/instaGRAAL/blob/192fcd056c75d5d4d4821df754b3d046237e2d5c/instagraal/kernels/kernel_sparse.cu#L2947

https://github.com/koszullab/instaGRAAL/blob/192fcd056c75d5d4d4821df754b3d046237e2d5c/instagraal/kernels/kernel_sparse_adapt.cu#L6

When trying to run the code on a computer with recent hardware, which defaults to a compute capability sm_86 we get the error for the kernel compilation:

  kernel.cu(3823): error: identifier "__shfl_down" is undefined

If we force to build for an old arch passing the option -arch sm_53, it builds with the message

kernel.cu(3823): warning: function "__shfl_down(double, unsigned int, int)"                                                                                                                                                                                                                 
/opt/local/easybuild/software/CUDA/10.1.243-GCC-8.3.0/bin/../targets/x86_64-linux/include/sm_30_intrinsics.hpp(295): here was declared deprecated ("__shfl_down() is deprecated in favor of __shfl_down_sync() and may be removed in a future release (Use -Wno-deprecated-declarations to s
uppress this warning).")                                                                                                                                                                                                                                                                    

this makes the code cumbersome to install and use on recent hardware with recent CUDA versions.

maybe the info on this question can help

https://stackoverflow.com/questions/50639194/shfl-down-and-shfl-down-sync-give-different-results

could you consider patching this ?

arielzn avatar Jun 22 '23 15:06 arielzn