Merge pull request #611 from kaatish/bug-fix/github/warp-scan-vector-…

…type Fix warp scan exlusive sum for vector types
NVIDIA · Jan 7, 2023 · 8bdb54d · 8bdb54d
2 parents 12332f0 + 074ce95
commit 8bdb54d
Show file tree

Hide file tree

Showing 2 changed files with 8 additions and 6 deletions.
diff --git a/cub/warp/warp_scan.cuh b/cub/warp/warp_scan.cuh
@@ -347,7 +347,7 @@ public:
         T               input,              ///< [in] Calling thread's input item.
         T               &exclusive_output)  ///< [out] Calling thread's output item.  May be aliased with \p input.
     {
-        T initial_value = 0;
+        T initial_value {};
         ExclusiveScan(input, exclusive_output, initial_value, cub::Sum());
     }
 
@@ -393,7 +393,7 @@ public:
         T               &exclusive_output,  ///< [out] Calling thread's output item.  May be aliased with \p input.
         T               &warp_aggregate)    ///< [out] Warp-wide aggregate reduction of input items.
     {
-        T initial_value = 0;
+        T initial_value {};
         ExclusiveScan(input, exclusive_output, initial_value, cub::Sum(), warp_aggregate);
     }
 

diff --git a/test/test_warp_scan.cu b/test/test_warp_scan.cu
@@ -124,15 +124,16 @@ __device__ __forceinline__ void DeviceTest(
 /// Exclusive sum basic
 template <
     typename    WarpScanT,
-    typename    T>
+    typename    T,
+    typename    IsPrimitiveT>
 __device__ __forceinline__ void DeviceTest(
     WarpScanT                       &warp_scan,
     T                               &data,
     T                               &initial_value,
     Sum                             &scan_op,
     T                               &aggregate,
     Int2Type<BASIC>                 test_mode,
-    Int2Type<true>                  is_primitive)
+    IsPrimitiveT                    is_primitive)
 {
     // Test basic warp scan
     warp_scan.ExclusiveSum(data, data);
@@ -142,15 +143,16 @@ __device__ __forceinline__ void DeviceTest(
 /// Exclusive sum aggregate
 template <
     typename    WarpScanT,
-    typename    T>
+    typename    T,
+    typename    IsPrimitiveT>
 __device__ __forceinline__ void DeviceTest(
     WarpScanT                       &warp_scan,
     T                               &data,
     T                               &initial_value,
     Sum                             &scan_op,
     T                               &aggregate,
     Int2Type<AGGREGATE>             test_mode,
-    Int2Type<true>                  is_primitive)
+    IsPrimitiveT                    is_primitive)
 {
     // Test with cumulative aggregate
     warp_scan.ExclusiveSum(data, data, aggregate);