Change AccType in inclusive

LLNL · Apr 27, 2024 · 6ba0ffa · 6ba0ffa
1 parent 82cb3b6
commit 6ba0ffa
Showing 1 changed file with 27 additions and 14 deletions.
diff --git a/include/RAJA/policy/hip/scan.hpp b/include/RAJA/policy/hip/scan.hpp
@@ -220,18 +220,27 @@ inclusive(
 {
   hipStream_t stream = hip_res.get_stream();
 
+#if defined(__HIPCC__)
+  using Config = rocprof::default_config;
+  using AccType = typename std::iterator_traits<OutputIter>::value_type;
+#endif
+
   int len = std::distance(begin, end);
   // Determine temporary device storage requirements
   void* d_temp_storage = nullptr;
   size_t temp_storage_bytes = 0;
 #if defined(__HIPCC__)
-  hipErrchk(::rocprim::inclusive_scan(d_temp_storage,
-                                      temp_storage_bytes,
-                                      begin,
-                                      out,
-                                      len,
-                                      binary_op,
-                                      stream));
+  hipErrchk(::rocprim::inclusive_scan<Config,
+                                      InputIter,
+                                      OutputIter,
+                                      Function,
+                                      AccType>(d_temp_storage,
+                                               temp_storage_bytes,
+                                               begin,
+                                               out,
+                                               len,
+                                               binary_op,
+                                               stream));
 #elif defined(__CUDACC__)
   hipErrchk(::cub::DeviceScan::InclusiveScan(d_temp_storage,
                                              temp_storage_bytes,
@@ -247,13 +256,17 @@ inclusive(
           temp_storage_bytes);
   // Run
 #if defined(__HIPCC__)
-  hipErrchk(::rocprim::inclusive_scan(d_temp_storage,
-                                      temp_storage_bytes,
-                                      begin,
-                                      out,
-                                      len,
-                                      binary_op,
-                                      stream));
+  hipErrchk(::rocprim::inclusive_scan<Config,
+                                      InputIter,
+                                      OutputIter,
+                                      Function,
+                                      AccType>(d_temp_storage,
+                                               temp_storage_bytes,
+                                               begin,
+                                               out,
+                                               len,
+                                               binary_op,
+                                               stream));
 #elif defined(__CUDACC__)
   hipErrchk(::cub::DeviceScan::InclusiveScan(d_temp_storage,
                                              temp_storage_bytes,