Define SYCL macros, implement async_wg_copy

CMakeLists.txt

@@ -131,6 +131,7 @@ add_library(simsycl
     include/simsycl/sycl/type_traits.hh
     include/simsycl/sycl/usm.hh
     include/simsycl/sycl/vec.hh
+    include/simsycl/macros.hh
     include/simsycl/system.hh
     ${CONFIG_PATH}
     src/simsycl/check.cc

include/CL/sycl.hpp

@@ -1,6 +1,9 @@
 #pragma once
 
+// IWYU pragma: begin_keep
+#include "../simsycl/macros.hh"
 #include "../simsycl/sycl.hh"
+// IWYU pragma: end_keep
 
 namespace cl {
 namespace sycl = simsycl::sycl;

include/simsycl/macros.hh

@@ -0,0 +1,7 @@
+#pragma once
+
+#define SYCL_LANGUAGE_VERSION 202008
+#define SYCL_DEVICE_COPYABLE 1
+#define __SYCL_SINGLE_SOURCE__ 1
+#define SYCL_FEATURE_SET_FULL 1
+#define SYCL_EXTERNAL /* empty */

include/simsycl/sycl/group.hh

@@ -5,6 +5,7 @@
 #include "h_item.hh"
 #include "id.hh"
 #include "item.hh"
+#include "multi_ptr.hh"
 #include "range.hh"
 #include "type_traits.hh"
 
@@ -50,6 +51,11 @@ class hierarchical_group_size_setter {
 
 namespace simsycl::sycl {
 
+class device_event {
+  public:
+    void wait() noexcept {}
+};
+
 template<int Dimensions>
 class group {
   public:
@@ -61,8 +67,10 @@ class group {
 
     group() = delete;
 
+    // TODO not in the spec, remove
     [[deprecated("use sycl::group::get_group_id")]] id_type get_id() const { return get_group_id(); }
 
+    // TODO not in the spec, remove
     [[deprecated("use sycl::group::get_group_id")]] size_t get_id(int dimension) const {
         return get_group_id(dimension);
     }
@@ -199,6 +207,7 @@ class group {
         }
     }
 
+    // TODO not in the spec, remove
     template<access_mode AccessMode = access_mode::read_write>
     void mem_fence(typename std::enable_if_t<AccessMode == access_mode::read || AccessMode == access_mode::write
                            || AccessMode == access_mode::read_write,
@@ -215,6 +224,58 @@ class group {
         // wait_for is a no-op in SimSYCL
     }
 
+    template<typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements) const {
+        std::copy_n(src.get(), num_elements, dest.get());
+    }
+
+    template<typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements) const {
+        std::copy_n(src.get(), num_elements, dest.get());
+    }
+
+    template<typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements, size_t src_stride) const {
+        for(size_t i = 0; i < num_elements; ++i) { dest[i] = src[i * src_stride]; }
+    }
+
+    template<typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements, size_t dest_stride) const {
+        for(size_t i = 0; i < num_elements; ++i) { dest[i * dest_stride] = src[i]; }
+    }
+
+    template<typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(
+        decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src, size_t num_elements) const {
+        std::copy_n(src.get(), num_elements, dest.get());
+    }
+
+    template<typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(
+        decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src, size_t num_elements) const {
+        std::copy_n(src.get(), num_elements, dest.get());
+    }
+
+    template<typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src,
+        size_t num_elements, size_t src_stride) const {
+        for(size_t i = 0; i < num_elements; ++i) { dest[i] = src[i * src_stride]; }
+    }
+
+    template<typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src,
+        size_t num_elements, size_t dest_stride) const {
+        for(size_t i = 0; i < num_elements; ++i) { dest[i * dest_stride] = src[i]; }
+    }
+
     friend bool operator==(const group<Dimensions> &lhs, const group<Dimensions> &rhs) {
         return lhs.m_local_item == rhs.m_local_item && lhs.m_global_item == rhs.m_global_item
             && lhs.m_group_item == rhs.m_group_item && lhs.m_concurrent_group == rhs.m_concurrent_group;

include/simsycl/sycl/nd_item.hh

@@ -4,7 +4,6 @@
 
 #include "group.hh"
 #include "id.hh"
-#include "multi_ptr.hh"
 #include "range.hh"
 #include "sub_group.hh"
 
@@ -35,11 +34,6 @@ sycl::nd_item<Dimensions> make_nd_item(const sycl::item<Dimensions, true> &globa
 
 namespace simsycl::sycl {
 
-class device_event {
-  public:
-    void wait() noexcept {}
-};
-
 template<int Dimensions>
 class nd_item {
   public:
@@ -107,41 +101,57 @@ class nd_item {
 
     template<typename DataT>
     [[deprecated]] device_event async_work_group_copy(
-        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements) const;
+        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements) const {
+        m_group.async_work_group_copy(dest, src, num_elements);
+    }
 
     template<typename DataT>
     [[deprecated]] device_event async_work_group_copy(
-        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements) const;
+        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements) const {
+        m_group.async_work_group_copy(dest, src, num_elements);
+    }
 
     template<typename DataT>
     [[deprecated]] device_event async_work_group_copy(
-        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements, size_t src_stride) const;
+        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements, size_t src_stride) const {
+        m_group.async_work_group_copy(dest, src, num_elements, src_stride);
+    }
 
     template<typename DataT>
     [[deprecated]] device_event async_work_group_copy(
-        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements, size_t dest_stride) const;
+        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements, size_t dest_stride) const {
+        m_group.async_work_group_copy(dest, src, num_elements, dest_stride);
+    }
 
     SIMSYCL_STOP_IGNORING_DEPRECATIONS
 
     template<typename DestDataT, typename SrcDataT>
         requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
     device_event async_work_group_copy(
-        decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src, size_t num_elements) const;
+        decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src, size_t num_elements) const {
+        return m_group.async_work_group_copy(dest, src, num_elements);
+    }
 
     template<typename DestDataT, typename SrcDataT>
         requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
     device_event async_work_group_copy(
-        decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src, size_t num_elements) const;
+        decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src, size_t num_elements) const {
+        return m_group.async_work_group_copy(dest, src, num_elements);
+    }
 
     template<typename DestDataT, typename SrcDataT>
         requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
     device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src,
-        size_t num_elements, size_t src_stride) const;
+        size_t num_elements, size_t src_stride) const {
+        return m_group.async_work_group_copy(dest, src, num_elements, src_stride);
+    }
 
     template<typename DestDataT, typename SrcDataT>
         requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
     device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src,
-        size_t num_elements, size_t dest_stride) const;
+        size_t num_elements, size_t dest_stride) const {
+        return m_group.async_work_group_copy(dest, src, num_elements, dest_stride);
+    }
 
     template<typename... Events>
     void wait_for(Events... events) const {

include/simsycl/sycl/type_traits.hh

@@ -65,8 +65,9 @@ namespace simsycl::sycl {
 // TODO consider moving this to a different header.
 using std::bit_cast;
 
+// approximation. must inherit to allow specialization
 template<typename T>
-using is_device_copyable = std::is_nothrow_copy_constructible<T>; // approximation
+struct is_device_copyable : std::is_nothrow_copy_constructible<T> {};
 
 template<typename T>
 inline constexpr bool is_device_copyable_v = is_device_copyable<T>::value;

include/sycl/sycl.hpp

@@ -1,5 +1,8 @@
 #pragma once
 
-#include "../simsycl/sycl.hh" // IWYU pragma: keep
+// IWYU pragma: begin_keep
+#include "../simsycl/macros.hh"
+#include "../simsycl/sycl.hh"
+// IWYU pragma: end_keep
 
 namespace sycl = simsycl::sycl;