[SYCLomatic] Block Store headers core

clang/runtime/dpct-rt/include/dpct/group_utils.hpp

@@ -433,11 +433,10 @@ class radix_sort {
 
 /// Load linear segment items into block format across threads
 /// Helper for Block Load
-enum load_algorithm {
+enum class load_algorithm {
 
   BLOCK_LOAD_DIRECT,
   BLOCK_LOAD_STRIPED,
-  // To-do: BLOCK_LOAD_WARP_TRANSPOSE
 
 };
 
@@ -517,16 +516,115 @@ class workgroup_load {
   __dpct_inline__ void load(const Item &item, InputIteratorT block_itr,
                             InputT (&items)[ITEMS_PER_WORK_ITEM]) {
 
-    if constexpr (ALGORITHM == BLOCK_LOAD_DIRECT) {
+    if constexpr (ALGORITHM == load_algorithm::BLOCK_LOAD_DIRECT) {
       load_blocked<ITEMS_PER_WORK_ITEM>(item, block_itr, items);
-    } else if constexpr (ALGORITHM == BLOCK_LOAD_STRIPED) {
+    } else if constexpr (ALGORITHM == load_algorithm::BLOCK_LOAD_STRIPED) {
       load_striped<ITEMS_PER_WORK_ITEM>(item, block_itr, items);
     }
   }
 
 private:
   uint8_t *_local_memory;
 };
+
+/// Store blocked/warped or striped work items into linear segment of items.
+/// Helper for Block Store
+enum class store_algorithm {
+
+  BLOCK_STORE_DIRECT,
+  BLOCK_STORE_STRIPED,
+
+};
+
+/// Stores a blocked arrangement of work items linear segment of items.
+template <size_t ITEMS_PER_WORK_ITEM, typename InputT, typename OutputIteratorT,
+          typename Item>
+__dpct_inline__ void store_blocked(const Item &item, OutputIteratorT block_itr,
+                                   InputT (&items)[ITEMS_PER_WORK_ITEM]) {
+
+  // This implementation does not take in account range storage across
+  // workgroup items To-do: Decide whether range storage is required for group
+  // storage
+  size_t linear_tid = item.get_local_linear_id();
+  OutputIteratorT workitem_itr = block_itr + (linear_tid * ITEMS_PER_WORK_ITEM);
+#pragma unroll
+  for (uint32_t idx = 0; idx < ITEMS_PER_WORK_ITEM; idx++) {
+    workitem_itr[idx] = items[idx];
+  }
+}
+
+/// Stores a striped arrangement of work items linear segment of items.
+template <size_t ITEMS_PER_WORK_ITEM, typename InputT, typename OutputIteratorT,
+          typename Item>
+__dpct_inline__ void store_striped(const Item &item, OutputIteratorT block_itr,
+                                   InputT (&items)[ITEMS_PER_WORK_ITEM]) {
+
+  // This implementation does not take in account range storage across
+  // workgroup items To-do: Decide whether range storage is required for group
+  // storage
+  size_t linear_tid = item.get_local_linear_id();
+  OutputIteratorT workitem_itr = block_itr + linear_tid;
+  size_t GROUP_WORK_ITEMS = item.get_local_range().size();
+#pragma unroll
+  for (uint32_t idx = 0; idx < ITEMS_PER_WORK_ITEM; idx++) {
+    workitem_itr[(idx * GROUP_WORK_ITEMS)] = items[idx];
+  }
+}
+
+/// Stores a warp-striped arrangement of work items linear segment of items.
+// Created as free function until exchange mechanism is
+// implemented.
+// To-do: inline this function with BLOCK_STORE_WARP_TRANSPOSE mechanism
+template <size_t ITEMS_PER_WORK_ITEM, typename InputT, typename OutputIteratorT,
+          typename Item>
+__dpct_inline__ void
+store_subgroup_striped(const Item &item, OutputIteratorT block_itr,
+                       InputT (&items)[ITEMS_PER_WORK_ITEM]) {
+
+  // This implementation does not take in account range loading across
+  // workgroup items To-do: Decide whether range loading is required for group
+  // loading
+  // This implementation loads linear segments into warp striped arrangement.
+  uint32_t subgroup_offset = item.get_sub_group().get_local_linear_id();
+  uint32_t subgroup_size = item.get_sub_group().get_local_linear_range();
+  uint32_t subgroup_idx = item.get_sub_group().get_group_linear_id();
+  uint32_t initial_offset =
+      (subgroup_idx * ITEMS_PER_WORK_ITEM * subgroup_size) + subgroup_offset;
+  OutputIteratorT workitem_itr = block_itr + initial_offset;
+#pragma unroll
+  for (uint32_t idx = 0; idx < ITEMS_PER_WORK_ITEM; idx++) {
+    workitem_itr[(idx * subgroup_size)] = items[idx];
+  }
+}
+
+// template parameters :
+// ITEMS_PER_WORK_ITEM: size_t variable controlling the number of items per
+// thread/work_item
+// ALGORITHM: store_algorithm variable controlling the type of store operation.
+// InputT: type for input sequence.
+// OutputIteratorT:  output iterator type
+// Item : typename parameter resembling sycl::nd_item<3> .
+template <size_t ITEMS_PER_WORK_ITEM, store_algorithm ALGORITHM,
+          typename InputT, typename OutputIteratorT, typename Item>
+class workgroup_store {
+public:
+  static size_t get_local_memory_size(size_t group_work_items) { return 0; }
+  workgroup_store(uint8_t *local_memory) : _local_memory(local_memory) {}
+
+  __dpct_inline__ void store(const Item &item, OutputIteratorT block_itr,
+                             InputT (&items)[ITEMS_PER_WORK_ITEM]) {
+
+    if constexpr (ALGORITHM == store_algorithm::BLOCK_STORE_DIRECT) {
+      store_blocked<ITEMS_PER_WORK_ITEM>(item, block_itr, items);
+    } else if constexpr (ALGORITHM == store_algorithm::BLOCK_STORE_STRIPED) {
+      store_striped<ITEMS_PER_WORK_ITEM>(item, block_itr, items);
+    }
+  }
+
+private:
+  uint8_t *_local_memory;
+};
+
 } // namespace group
 } // namespace dpct