Handle const and reference scan accumulator types (#736)

We recently updated the default accumulator types for
rocprim::inclusive_scan, deterministic_inclusive_scan,
exclusive_scan, and deterministic_exclusive_scan.

The update set the default type to the return type of the
scan operator. The scan operator can return a type that
is const or a reference. This is problematic because internally,
the scan algorithm need to create instances of the accumulator
type and periodically update their value. It also needs to create
pointers of the accumulator type.

This change fixes the problem by stripping any const or reference
from the accumulator type before it is passed to the internal
scan implementation function.

Author: umfranzw

rocprim/include/rocprim/device/device_scan.hpp

@@ -505,23 +505,27 @@ inline hipError_t inclusive_scan(void*             temporary_storage,
                                  const hipStream_t stream            = 0,
                                  bool              debug_synchronous = false)
 {
+    // AccType may be const or a reference. Get the non-const, non-reference type.
+    // This is necessary because we may need to assign to instances of this type or create pointers to it.
+    using safe_acc_type = typename std::remove_const<typename std::remove_reference<AccType>::type>::type;
+
     // input_type() is a dummy initial value (not used)
     return detail::scan_impl<detail::lookback_scan_determinism::default_determinism,
                              false,
                              Config,
                              InputIterator,
                              OutputIterator,
-                             AccType,
+                             safe_acc_type,
                              BinaryFunction,
-                             AccType>(temporary_storage,
-                                      storage_size,
-                                      input,
-                                      output,
-                                      AccType{},
-                                      size,
-                                      scan_op,
-                                      stream,
-                                      debug_synchronous);
+                             safe_acc_type>(temporary_storage,
+                                            storage_size,
+                                            input,
+                                            output,
+                                            safe_acc_type{},
+                                            size,
+                                            scan_op,
+                                            stream,
+                                            debug_synchronous);
 }
 
 /// \brief Bitwise-reproducible parallel inclusive scan primitive for device level.
@@ -546,22 +550,26 @@ inline hipError_t deterministic_inclusive_scan(void*             temporary_stora
                                                const hipStream_t stream  = 0,
                                                bool              debug_synchronous = false)
 {
+    // AccType may be const or a reference. Get the non-const, non-reference type.
+    // This is necessary because we may need to assign to instances of this type or create pointers to it.
+    using safe_acc_type = typename std::remove_const<typename std::remove_reference<AccType>::type>::type;
+
     return detail::scan_impl<detail::lookback_scan_determinism::deterministic,
                              false,
                              Config,
                              InputIterator,
                              OutputIterator,
-                             AccType,
+                             safe_acc_type,
                              BinaryFunction,
-                             AccType>(temporary_storage,
-                                      storage_size,
-                                      input,
-                                      output,
-                                      AccType{},
-                                      size,
-                                      scan_op,
-                                      stream,
-                                      debug_synchronous);
+                             safe_acc_type>(temporary_storage,
+                                            storage_size,
+                                            input,
+                                            output,
+                                            safe_acc_type{},
+                                            size,
+                                            scan_op,
+                                            stream,
+                                            debug_synchronous);
 }
 
 /// \brief Parallel exclusive scan primitive for device level.
@@ -672,22 +680,26 @@ inline hipError_t exclusive_scan(void*               temporary_storage,
                                  const hipStream_t   stream            = 0,
                                  bool                debug_synchronous = false)
 {
+    // AccType may be const or a reference. Get the non-const, non-reference type.
+    // This is necessary because we may need to assign to instances of this type or create pointers to it.
+    using safe_acc_type = typename std::remove_const<typename std::remove_reference<AccType>::type>::type;
+
     return detail::scan_impl<detail::lookback_scan_determinism::default_determinism,
                              true,
                              Config,
                              InputIterator,
                              OutputIterator,
                              InitValueType,
                              BinaryFunction,
-                             AccType>(temporary_storage,
-                                      storage_size,
-                                      input,
-                                      output,
-                                      initial_value,
-                                      size,
-                                      scan_op,
-                                      stream,
-                                      debug_synchronous);
+                             safe_acc_type>(temporary_storage,
+                                            storage_size,
+                                            input,
+                                            output,
+                                            initial_value,
+                                            size,
+                                            scan_op,
+                                            stream,
+                                            debug_synchronous);
 }
 
 /// \brief Bitwise-reproducible parallel exclusive scan primitive for device level.
@@ -714,22 +726,26 @@ inline hipError_t deterministic_exclusive_scan(void*               temporary_sto
                                                const hipStream_t   stream  = 0,
                                                bool                debug_synchronous = false)
 {
+    // AccType may be const or a reference. Get the non-const, non-reference type.
+    // This is necessary because we may need to assign to instances of this type or create pointers to it.
+    using safe_acc_type = typename std::remove_const<typename std::remove_reference<AccType>::type>::type;
+
     return detail::scan_impl<detail::lookback_scan_determinism::deterministic,
                              true,
                              Config,
                              InputIterator,
                              OutputIterator,
                              InitValueType,
                              BinaryFunction,
-                             AccType>(temporary_storage,
-                                      storage_size,
-                                      input,
-                                      output,
-                                      initial_value,
-                                      size,
-                                      scan_op,
-                                      stream,
-                                      debug_synchronous);
+                             safe_acc_type>(temporary_storage,
+                                            storage_size,
+                                            input,
+                                            output,
+                                            initial_value,
+                                            size,
+                                            scan_op,
+                                            stream,
+                                            debug_synchronous);
 }
 
 /// @}