adds dpcpp kernel

Author: MarcelKoch

dpcpp/matrix/dense_kernels.dp.cpp

@@ -56,6 +56,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include "dpcpp/base/onemkl_bindings.hpp"
 #include "dpcpp/components/cooperative_groups.dp.hpp"
 #include "dpcpp/components/reduction.dp.hpp"
+#include "dpcpp/components/searching.dp.hpp"
 #include "dpcpp/components/thread_ids.dp.hpp"
 #include "dpcpp/components/uninitialized_array.hpp"
 #include "dpcpp/synthesizer/implementation_selection.hpp"
@@ -606,7 +607,45 @@ template <typename ValueType, typename OutputType, typename IndexType>
 void row_scatter(std::shared_ptr<const DefaultExecutor> exec,
                  const index_set<IndexType>* row_idxs,
                  const matrix::Dense<ValueType>* orig,
-                 matrix::Dense<OutputType>* target) GKO_NOT_IMPLEMENTED;
+                 matrix::Dense<OutputType>* target)
+{
+    const auto num_sets = row_idxs->get_num_subsets();
+    const auto num_rows = row_idxs->get_num_elems();
+    const auto num_cols = orig->get_size()[1];
+
+    const auto* row_set_begins = row_idxs->get_subsets_begin();
+    const auto* row_set_offsets = row_idxs->get_superset_indices();
+
+    const auto orig_stride = orig->get_stride();
+    const auto* orig_values = orig->get_const_values();
+
+    const auto target_stride = target->get_stride();
+    auto* target_values = target->get_values();
+
+    exec->get_queue()->submit([&](sycl::handler& cgh) {
+        cgh.parallel_for(
+            static_cast<size_type>(num_rows * num_cols),
+            [=](sycl::item<1> item) {
+                const auto row = static_cast<size_type>(item[0]) / num_cols;
+                const auto col = static_cast<size_type>(item[0]) % num_cols;
+
+                if (row >= num_rows) {
+                    return;
+                }
+
+                auto set_id =
+                    binary_search<size_type>(
+                        0, num_sets + 1,
+                        [=](auto i) { return row < row_set_offsets[i]; }) -
+                    1;
+                auto set_local_row = row - row_set_offsets[set_id];
+                auto target_row = set_local_row + row_set_begins[set_id];
+
+                target_values[target_row * target_stride + col] =
+                    orig_values[row * orig_stride + col];
+            });
+    });
+}
 
 GKO_INSTANTIATE_FOR_EACH_MIXED_VALUE_AND_INDEX_TYPE_2(
     GKO_DECLARE_DENSE_ROW_SCATTER_INDEX_SET_KERNEL);