usage of __shfl_down is deprecated

[arielzn]

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

instaGRAAL/instagraal/kernels/kernel_sparse.cu

Line 2947 in 192fcd0

__device__ inline double __shfl_down(double var, unsigned int srcLane, int width=32) {

instaGRAAL/instagraal/kernels/kernel_sparse_adapt.cu

Line 6 in 192fcd0

__device__ inline double __shfl_down(double var, unsigned int srcLane, int width=32) {

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 ?

[ABignaud]

Hi,
Unfortunately, I don't think we will consider patching this as we have no longer cuda developer in our team. InstaGRAAL is currently working only with CUDA v10, I don't think you will manage to use it with more recent version of CUDA. To install it easily, you can use the docker image we have provided which use CUDA v10 without downgrading the version of CUDA on your computer. Otherwise, you can use the Galaxy Europe server where instaGRAAL has been installed.
Amaury

[arielzn]

The issue in my case is present even using CUDA v10, I tried with cuda v10 installed or also creating a singularity container with cuda v10.

My guess is it depends on the hardware being found, with the A5000 we have available it's triggered to use sm_86 as compute capability automatically, and thus we get the error I mention.
It can be eventually circumvent forcing to use an older capability, but is not quite practical as a workaround.

Without modifications, did you have the chance to run instaGRAAL on GPUs generation 30XX or newer 40XX ?