Merge pull request NVIDIA#6 from egaburov/cub-fixes

http://nvbugs/1904217 Former-commit-id: f5c46e7272b9e33b3c6bbdf9145a0e43dc47ea39
GaryShen2008 · May 20, 2017 · aa6ea7c · aa6ea7c
2 parents 3a22fac + 6d3d084
commit aa6ea7c
Show file tree

Hide file tree

Showing 3 changed files with 6 additions and 10 deletions.
diff --git a/cub/thread/thread_operators.cuh b/cub/thread/thread_operators.cuh
@@ -224,7 +224,10 @@ public:
     __host__ __device__ __forceinline__
     T operator()(const T &a, const T &b)
     {
-        return scan_op(b, a);
+      T _a(a);
+      T _b(b);
+
+      return scan_op(_b, _a);
     }
 };
 

diff --git a/cub/warp/specializations/warp_reduce_shfl.cuh b/cub/warp/specializations/warp_reduce_shfl.cuh
@@ -414,10 +414,7 @@ struct WarpReduceShfl
         int             offset,                 ///< [in] Up-offset to pull from
         Int2Type<true>  /*is_small_unsigned*/)  ///< [in] Marker type indicating whether T is a small unsigned integer
     {
-        // Recast as uint32 to take advantage of any specializations
-        unsigned int temp = reinterpret_cast<unsigned int &>(input);
-        temp = ReduceStep(temp, reduction_op, last_lane, offset);
-        return reinterpret_cast<_T&>(temp);
+        return ReduceStep(input, reduction_op, last_lane, offset);
     }
 
 

diff --git a/cub/warp/specializations/warp_scan_shfl.cuh b/cub/warp/specializations/warp_scan_shfl.cuh
@@ -411,11 +411,7 @@ struct WarpScanShfl
         int             offset,             ///< [in] Up-offset to pull from
         Int2Type<true>  /*is_small_unsigned*/)  ///< [in] Marker type indicating whether T is a small integer
     {
-        unsigned int temp = reinterpret_cast<unsigned int &>(input);
-
-        temp = InclusiveScanStep(temp, scan_op, first_lane, offset);
-
-        return reinterpret_cast<_T&>(temp);
+        return InclusiveScanStep(input, scan_op, first_lane, offset);
     }