Fix begin_bit == end_bit == 0 for device-wide and segmented sort

cub/device/dispatch/dispatch_radix_sort.cuh

@@ -1694,7 +1694,7 @@ struct DispatchRadixSort :
         #endif
         cudaError_t error = cudaSuccess;
         error = cudaMemcpyAsync(d_keys.Alternate(), d_keys.Current(), num_items * sizeof(KeyT),
-                                cudaMemcpyDeviceToDevice, stream);
+                                cudaMemcpyDefault, stream);
         if (CubDebug(error))
         {
             return error;
@@ -1713,7 +1713,7 @@ struct DispatchRadixSort :
                     (long long)num_items, (long long)stream);
             #endif
             error = cudaMemcpyAsync(d_values.Alternate(), d_values.Current(),
-                                    num_items * sizeof(ValueT), cudaMemcpyDeviceToDevice, stream);
+                                    num_items * sizeof(ValueT), cudaMemcpyDefault, stream);
             if (CubDebug(error))
             {
                 return error;
@@ -1747,7 +1747,10 @@ struct DispatchRadixSort :
         }
 
         // Check if simple copy suffices (is_overwrite_okay == false at this point)
-        if (begin_bit == end_bit)
+        cudaError_t error = cudaSuccess;
+        bool has_uva = false;
+        if ((error = HasUVA(has_uva)) != cudaSuccess) return error;
+        if (begin_bit == end_bit & has_uva)
         {
             return InvokeCopy();
         }

cub/util_device.cuh

@@ -125,6 +125,23 @@ CUB_RUNTIME_FUNCTION inline int CurrentDevice()
     return device;
 }
 
+/** \brief Gets whether the current device supports unified addressing */
+CUB_RUNTIME_FUNCTION cudaError_t HasUVA(bool& has_uva)
+{
+    has_uva = false;
+    cudaError_t error = cudaSuccess;
+    int device = -1;
+    if (CubDebug(error = cudaGetDevice(&device)) != cudaSuccess) return error;
+    int uva = 0;
+    if (CubDebug(error = cudaDeviceGetAttribute(&uva, cudaDevAttrUnifiedAddressing, device))
+        != cudaSuccess)
+    {
+        return error;
+    }
+    has_uva = uva == 1;
+    return error;
+}
+
 /**
  * \brief RAII helper which saves the current device and switches to the
  *        specified device on construction and switches to the saved device on