Got inclusive scan working

LLNL · Dec 17, 2024 · 5fe4a54 · 5fe4a54
1 parent e4801c9
commit 5fe4a54
Showing 1 changed file with 87 additions and 32 deletions.
diff --git a/include/RAJA/policy/sycl/scan.hpp b/include/RAJA/policy/sycl/scan.hpp
@@ -18,6 +18,7 @@
 #ifndef RAJA_scan_sycl_HPP
 #define RAJA_scan_sycl_HPP
 
+#include <cstddef>
 #include "RAJA/config.hpp"
 #include "camp/resource/sycl.hpp"
 
@@ -60,17 +61,47 @@ inclusive_inplace(
      // Calculate the size of the input range
     size_t size = std::distance(begin, end);
 
-    ::sycl::buffer<valueT, 1> buf(begin, ::sycl::range<1>(size));
-
-    // Submit the kernel to the SYCL queue
-    sycl_queue->submit([&](::sycl::handler& cgh) {
-        // ::sycl::accessor<valueT, 1, ::sycl::access::mode::read_write
-        cgh.parallel_for(::sycl::range<1>(size), [=](::sycl::id<1> idx) {
-            if (idx[0] != 0) {
-                *(begin + idx[0]) = binary_op(*(begin + idx[0] - 1), *(begin + idx[0]));
-            }
+    ::sycl::buffer<valueT, 1> inBuff(begin, end);
+    ::sycl::buffer<valueT, 1> outBuff(begin, ::sycl::range<1>(size));
+
+    int iterations = 0;
+    for (size_t ii = size >> 1; ii > 0; ii >>= 1) {
+        iterations++;
+    }
+    if ((size & (size - 1)) != 0) {
+        iterations++;
+    }
+
+    auto inPtr = &inBuff;
+    auto outPtr = &outBuff;
+
+    if (iterations % 2 == 0) { 
+        outPtr = &inBuff;
+        inPtr = &outBuff;
+    } 
+
+    int ii = 1;
+    do {
+        // Submit the kernel to the SYCL queue  
+        sycl_queue->submit([&](::sycl::handler& cgh) {
+            auto inAccessor = inPtr->get_access(cgh);
+            auto outAccessor = outPtr->get_access(cgh);
+            // outBuffAccessor = outBuff->get_access<::sycl::access::mode::read(cgh);
+
+            cgh.parallel_for(::sycl::range<1>(size), [=](::sycl::item<1> idx) {
+                size_t td = 1 << (ii - 1);
+                size_t thisID = idx.get_id(0);
+                if (thisID < size and thisID >= td) { 
+                    outAccessor[thisID] = binary_op(inAccessor[thisID - td], inAccessor[thisID]);
+                } else {
+                    outAccessor[thisID] = inAccessor[thisID];
+                }
+            });
         });
-    });
+
+        std::swap(inPtr, outPtr);
+        ii++;
+    } while (  ii <= iterations);
 
   sycl_res.wait();
   return camp::resources::EventProxy<camp::resources::Sycl>(sycl_res);
@@ -91,28 +122,52 @@ exclusive_inplace(
     Function binary_op,
     TT initVal)
 {
-
-    ::sycl::queue* sycl_queue = sycl_res.get_queue();
-
-     // Calculate the size of the input range
-    size_t size = std::distance(begin, end);
-    using valueT = typename std::remove_reference<decltype(*begin)>::type;
-    valueT agg = *begin;
-    // Submit the kernel to the SYCL queue
-    sycl_queue->submit([&](::sycl::handler& cgh) {
-
-        cgh.parallel_for(::sycl::range<1>(size), [=](::sycl::id<1> idx) {
-            if (idx[0] == 0) {
-                *begin = binary_op(*begin, initVal);
-            } else {
-                // agg = binary_op(agg, begin + idx[0]);
-                *(begin + idx[0])  = binary_op(*(begin + idx[0] - 1), *(begin + idx[0]));
-            }
-        });
-    });
-
-    sycl_res.wait();
-    return camp::resources::EventProxy<camp::resources::Sycl>(sycl_res);
+//     ::sycl::queue* sycl_queue = sycl_res.get_queue();
+
+//     using valueT = typename std::remove_reference<decltype(*begin)>::type;
+
+//      // Calculate the size of the input range
+//     size_t size = std::distance(begin, end);
+
+//     ::sycl::buffer<valueT, 1> outBuff(begin, ::sycl::range<1>(size));
+
+//     int iterations = 0;
+//     for (size_t ii = size >> 1; ii > 0; ii >>= 1) {
+//         iterations++;
+//     }
+//     if ((size & (size - 1)) != 0) {
+//         iterations++;
+//     }
+
+//     auto inPtr = begin;
+//     auto outPtr = &outBuff;
+
+//     if (iterations % 2 != 0) { 
+//         outPtr = begin;
+//         inPtr = &outBuff;
+//     } 
+
+//     // Submit the kernel to the SYCL queue
+//     sycl_queue->submit([&](::sycl::handler& cgh) {
+//         // outBuffAccessor = outBuff->get_access<::sycl::access::mode::read(cgh);
+
+//         // ::sycl::accessor<valueT, 1, ::sycl::access::mode::read_write
+//         cgh.parallel_for(::sycl::range<1>(size), [=](::sycl::nd_item<1> idx) {
+//             size_t thisID = idx.get_global_id(0);
+//             if (thisID > 0 ) { 
+//                 outPtr[thisID] = inPtr[thisID - 1];
+//             }
+//             else { 
+//                 outPtr[thisID] = initVal;
+//             }
+//             // if (idx[0] != 0) {
+//             //     *(begin + idx[0]) = binary_op(*(begin + idx[0] - 1), *(begin + idx[0]));
+//             // }
+//         });
+//     });
+
+//   sycl_res.wait();
+  return camp::resources::EventProxy<camp::resources::Sycl>(sycl_res);
 }
 
 template <size_t BLOCK_SIZE,