Fix shfl_sync for CUDA8.0

[chengduoZH]

fix #10323

CUDA 8:

T __shfl(T var, int srcLane, int width=warpSize);
T __shfl_down(T var, unsigned int delta, int width=warpSize);   

CUDA 9:

T __shfl_sync(unsigned mask, T var, int srcLane, int width=warpSize);
T __shfl_down_sync(unsigned mask, T var, unsigned detla,  int width=warpSize);

[dzhwinter]

LGTM

[dzhwinter]

I forget to change the wrapper in old paddle, but how it passes the CI test? So weird.