Sycl device_type mapping

accessor/sycl_helper.hpp

@@ -0,0 +1,192 @@
+// SPDX-FileCopyrightText: 2017 - 2024 The Ginkgo authors
+//
+// SPDX-License-Identifier: BSD-3-Clause
+
+#ifndef GKO_ACCESSOR_SYCL_HELPER_HPP_
+#define GKO_ACCESSOR_SYCL_HELPER_HPP_
+
+
+#include <complex>
+#include <type_traits>
+
+#include "block_col_major.hpp"
+#include "reduced_row_major.hpp"
+#include "row_major.hpp"
+#include "scaled_reduced_row_major.hpp"
+#include "utils.hpp"
+
+
+namespace sycl {
+inline namespace _V1 {
+
+
+class half;
+
+
+}
+}  // namespace sycl
+
+
+namespace gko {
+
+
+class half;
+
+
+namespace acc {
+namespace detail {
+
+
+template <typename T>
+struct sycl_type {
+    using type = T;
+};
+
+template <>
+struct sycl_type<gko::half> {
+    using type = sycl::half;
+};
+
+// Unpack cv and reference / pointer qualifiers
+template <typename T>
+struct sycl_type<const T> {
+    using type = const typename sycl_type<T>::type;
+};
+
+template <typename T>
+struct sycl_type<volatile T> {
+    using type = volatile typename sycl_type<T>::type;
+};
+
+template <typename T>
+struct sycl_type<T*> {
+    using type = typename sycl_type<T>::type*;
+};
+
+template <typename T>
+struct sycl_type<T&> {
+    using type = typename sycl_type<T>::type&;
+};
+
+template <typename T>
+struct sycl_type<T&&> {
+    using type = typename sycl_type<T>::type&&;
+};
+
+
+// Transform the underlying type of std::complex
+template <typename T>
+struct sycl_type<std::complex<T>> {
+    using type = std::complex<typename sycl_type<T>::type>;
+};
+
+
+}  // namespace detail
+
+
+/**
+ * This is an alias for SYCL's equivalent of `T`.
+ *
+ * @tparam T  a type
+ */
+template <typename T>
+using sycl_type_t = typename detail::sycl_type<T>::type;
+
+
+/**
+ * Reinterprets the passed in value as a SYCL type.
+ *
+ * @param val  the value to reinterpret
+ *
+ * @return `val` reinterpreted to SYCL type
+ */
+template <typename T>
+std::enable_if_t<std::is_pointer<T>::value || std::is_reference<T>::value,
+                 sycl_type_t<T>>
+as_sycl_type(T val)
+{
+    return reinterpret_cast<sycl_type_t<T>>(val);
+}
+
+
+/**
+ * @copydoc as_sycl_type()
+ */
+template <typename T>
+std::enable_if_t<!std::is_pointer<T>::value && !std::is_reference<T>::value,
+                 sycl_type_t<T>>
+as_sycl_type(T val)
+{
+    return *reinterpret_cast<sycl_type_t<T>*>(&val);
+}
+
+
+/**
+ * Changes the types and reinterprets the passed in range pointers as a SYCL
+ * types.
+ *
+ * @param r  the range which pointers need to be reinterpreted
+ *
+ * @return `r` with appropriate types and reinterpreted to SYCL pointers
+ */
+template <std::size_t dim, typename Type1, typename Type2>
+GKO_ACC_INLINE auto as_sycl_range(
+    const range<reduced_row_major<dim, Type1, Type2>>& r)
+{
+    return range<
+        reduced_row_major<dim, sycl_type_t<Type1>, sycl_type_t<Type2>>>(
+        r.get_accessor().get_size(),
+        as_sycl_type(r.get_accessor().get_stored_data()),
+        r.get_accessor().get_stride());
+}
+
+/**
+ * @copydoc as_sycl_range()
+ */
+template <std::size_t dim, typename Type1, typename Type2, std::uint64_t mask>
+GKO_ACC_INLINE auto as_sycl_range(
+    const range<scaled_reduced_row_major<dim, Type1, Type2, mask>>& r)
+{
+    return range<scaled_reduced_row_major<dim, sycl_type_t<Type1>,
+                                          sycl_type_t<Type2>, mask>>(
+        r.get_accessor().get_size(),
+        as_sycl_type(r.get_accessor().get_stored_data()),
+        r.get_accessor().get_storage_stride(),
+        as_sycl_type(r.get_accessor().get_scalar()),
+        r.get_accessor().get_scalar_stride());
+}
+
+/**
+ * @copydoc as_sycl_range()
+ */
+template <typename T, size_type dim>
+GKO_ACC_INLINE auto as_sycl_range(const range<block_col_major<T, dim>>& r)
+{
+    return range<block_col_major<sycl_type_t<T>, dim>>(
+        r.get_accessor().lengths, as_sycl_type(r.get_accessor().data),
+        r.get_accessor().stride);
+}
+
+/**
+ * @copydoc as_sycl_range()
+ */
+template <typename T, size_type dim>
+GKO_ACC_INLINE auto as_sycl_range(const range<row_major<T, dim>>& r)
+{
+    return range<block_col_major<sycl_type_t<T>, dim>>(
+        r.get_accessor().lengths, as_sycl_type(r.get_accessor().data),
+        r.get_accessor().stride);
+}
+
+template <typename AccType>
+GKO_ACC_INLINE auto as_device_range(AccType&& acc)
+{
+    return as_device_range(std::forward<AccType>(acc));
+}
+
+
+}  // namespace acc
+}  // namespace gko
+
+
+#endif  // GKO_ACCESSOR_SYCL_HELPER_HPP_

common/unified/base/kernel_launch.hpp

@@ -74,16 +74,7 @@ GKO_INLINE GKO_ATTRIBUTES constexpr unpack_member_type<T> unpack_member(T value)
 namespace gko {
 namespace kernels {
 namespace dpcpp {
-
-
-template <typename T>
-using device_type = T;
-
-template <typename T>
-device_type<T> as_device_type(T value)
-{
-    return value;
-}
+#include "dpcpp/base/types.hpp"
 
 
 template <typename T>
@@ -95,6 +86,7 @@ GKO_INLINE GKO_ATTRIBUTES constexpr unpack_member_type<T> unpack_member(T value)
     return value;
 }
 
+
 }  // namespace dpcpp
 }  // namespace kernels
 }  // namespace gko

dpcpp/base/device_matrix_data_kernels.dp.cpp

@@ -9,6 +9,7 @@
 #include <ginkgo/core/base/exception_helpers.hpp>
 
 #include "dpcpp/base/onedpl.hpp"
+#include "dpcpp/base/types.hpp"
 
 
 namespace gko {
@@ -22,12 +23,13 @@ void remove_zeros(std::shared_ptr<const DefaultExecutor> exec,
                   array<ValueType>& values, array<IndexType>& row_idxs,
                   array<IndexType>& col_idxs)
 {
-    using nonzero_type = matrix_data_entry<ValueType, IndexType>;
+    using device_value_type = device_type<ValueType>;
     auto size = values.get_size();
     auto policy = onedpl_policy(exec);
-    auto nnz = std::count_if(
-        policy, values.get_const_data(), values.get_const_data() + size,
-        [](ValueType val) { return is_nonzero<ValueType>(val); });
+    auto nnz =
+        std::count_if(policy, as_device_type(values.get_const_data()),
+                      as_device_type(values.get_const_data()) + size,
+                      [](device_value_type val) { return is_nonzero(val); });
     if (nnz < size) {
         // allocate new storage
         array<ValueType> new_values{exec, static_cast<size_type>(nnz)};
@@ -36,10 +38,10 @@ void remove_zeros(std::shared_ptr<const DefaultExecutor> exec,
         // copy nonzeros
         auto input_it = oneapi::dpl::make_zip_iterator(
             row_idxs.get_const_data(), col_idxs.get_const_data(),
-            values.get_const_data());
-        auto output_it = oneapi::dpl::make_zip_iterator(new_row_idxs.get_data(),
-                                                        new_col_idxs.get_data(),
-                                                        new_values.get_data());
+            as_device_type(values.get_const_data()));
+        auto output_it = oneapi::dpl::make_zip_iterator(
+            new_row_idxs.get_data(), new_col_idxs.get_data(),
+            as_device_type(new_values.get_data()));
         std::copy_if(policy, input_it, input_it + size, output_it,
                      [](auto tuple) { return is_nonzero(std::get<2>(tuple)); });
         // swap out storage

dpcpp/base/kernel_launch_reduction.dp.hpp

@@ -239,7 +239,8 @@ void run_kernel_reduction_cached(std::shared_ptr<const DpcppExecutor> exec,
         [&](std::uint32_t cfg) { return cfg == desired_cfg; },
         syn::value_list<bool>(), syn::value_list<int>(),
         syn::value_list<size_type>(), syn::type_list<>(), exec, fn, op,
-        finalize, identity, result, size, tmp, map_to_device(args)...);
+        finalize, as_device_type(identity), as_device_type(result), size, tmp,
+        map_to_device(args)...);
 }
 
 
@@ -261,7 +262,8 @@ void run_kernel_reduction_cached(std::shared_ptr<const DpcppExecutor> exec,
         [&](std::uint32_t cfg) { return cfg == desired_cfg; },
         syn::value_list<bool>(), syn::value_list<int>(),
         syn::value_list<size_type>(), syn::type_list<>(), exec, fn, op,
-        finalize, identity, result, size, tmp, map_to_device(args)...);
+        finalize, as_device_type(identity), as_device_type(result), size, tmp,
+        map_to_device(args)...);
 }
 
 
@@ -658,8 +660,8 @@ void run_kernel_row_reduction_cached(std::shared_ptr<const DpcppExecutor> exec,
         [&](std::uint32_t cfg) { return cfg == desired_cfg; },
         syn::value_list<bool>(), syn::value_list<int>(),
         syn::value_list<size_type>(), syn::type_list<>(), exec, fn, op,
-        finalize, identity, result, result_stride, size, tmp,
-        map_to_device(args)...);
+        finalize, as_device_type(identity), as_device_type(result),
+        result_stride, size, tmp, map_to_device(args)...);
 }
 
 
@@ -681,7 +683,8 @@ void run_kernel_col_reduction_cached(std::shared_ptr<const DpcppExecutor> exec,
         [&](std::uint32_t cfg) { return cfg == desired_cfg; },
         syn::value_list<bool>(), syn::value_list<int>(),
         syn::value_list<size_type>(), syn::type_list<>(), exec, fn, op,
-        finalize, identity, result, size, tmp, map_to_device(args)...);
+        finalize, as_device_type(identity), as_device_type(result), size, tmp,
+        map_to_device(args)...);
 }