template<> __device__ __forceinline__ int saturate_cast<int>(double v) 
 {
 #if __CUDA_ARCH__ >= 130 
     return __double2int_rn(v);
 #else
     return saturate_cast<int>((float)v);
 #endif
 }
 template<> static __device__ __forceinline__ schar saturate_cast<schar>(double v)
 {             
 #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 130
     int iv = __double2int_rn(v); return saturate_cast<schar>(iv);
 #else
     return saturate_cast<schar>((float)v);
 #endif
 }
 template<> __device__ __forceinline__ uchar saturate_cast<uchar>(double v)
 {
 #if __CUDA_ARCH__ >= 130
     int iv = __double2int_rn(v); 
     return saturate_cast<uchar>(iv);
 #else
     return saturate_cast<uchar>((float)v);
 #endif
 }