Various fixes related to the SYCL back-end

  - add the missing specialization of CreateViewPlainPtr for SYCL devices
  - improve the comments on the ALPAKA_FN_INLINE macro
  - remove unnecessary ALPAKA_FN_HOST attributes
  - rename QueueGenericSyclBase::m_impl to m_spQueueImpl, to align with the other back-ends

include/alpaka/acc/AccFpgaSyclIntel.hpp

@@ -1,4 +1,4 @@
-/* Copyright 2022 Jan Stephan
+/* Copyright 2023 Jan Stephan, Aurora Perego
  * SPDX-License-Identifier: MPL-2.0
  */
 
@@ -45,7 +45,7 @@ namespace alpaka::trait
     template<typename TDim, typename TIdx>
     struct GetAccName<AccFpgaSyclIntel<TDim, TIdx>>
     {
-        ALPAKA_FN_HOST static auto getAccName() -> std::string
+        static auto getAccName() -> std::string
         {
             return "AccFpgaSyclIntel<" + std::to_string(TDim::value) + "," + core::demangled<TIdx> + ">";
         }

include/alpaka/core/Common.hpp

@@ -76,13 +76,12 @@
 //! Macro defining the inline function attribute.
 #if BOOST_LANG_CUDA || BOOST_LANG_HIP
 #    define ALPAKA_FN_INLINE __forceinline__
-#else
-#    if BOOST_COMP_MSVC || defined(BOOST_COMP_MSVC_EMULATED)
+#elif BOOST_COMP_MSVC || defined(BOOST_COMP_MSVC_EMULATED)
 // TODO: With C++20 [[msvc::forceinline]] can be used.
-#        define ALPAKA_FN_INLINE __forceinline
-#    else
-#        define ALPAKA_FN_INLINE [[gnu::always_inline]] inline
-#    endif
+#    define ALPAKA_FN_INLINE __forceinline
+#else
+// For gcc, clang, and clang-based compilers like Intel icpx
+#    define ALPAKA_FN_INLINE [[gnu::always_inline]] inline
 #endif
 
 //! This macro defines a variable lying in global accelerator device memory.

include/alpaka/event/EventGenericSycl.hpp

@@ -1,4 +1,4 @@
-/* Copyright 2022 Jan Stephan, Antonio Di Pilato
+/* Copyright 2023 Jan Stephan, Antonio Di Pilato, Aurora Perego
  * SPDX-License-Identifier: MPL-2.0
  */
 
@@ -87,7 +87,7 @@ namespace alpaka::trait
     {
         static auto enqueue(QueueGenericSyclNonBlocking<TDev>& queue, EventGenericSycl<TDev>& event)
         {
-            event.setEvent(queue.m_impl->get_last_event());
+            event.setEvent(queue.m_spQueueImpl->get_last_event());
         }
     };
 
@@ -97,7 +97,7 @@ namespace alpaka::trait
     {
         static auto enqueue(QueueGenericSyclBlocking<TDev>& queue, EventGenericSycl<TDev>& event)
         {
-            event.setEvent(queue.m_impl->get_last_event());
+            event.setEvent(queue.m_spQueueImpl->get_last_event());
         }
     };
 
@@ -120,7 +120,7 @@ namespace alpaka::trait
     {
         static auto waiterWaitFor(QueueGenericSyclNonBlocking<TDev>& queue, EventGenericSycl<TDev> const& event)
         {
-            queue.m_impl->register_dependency(event.getNativeHandle());
+            queue.m_spQueueImpl->register_dependency(event.getNativeHandle());
         }
     };
 
@@ -130,7 +130,7 @@ namespace alpaka::trait
     {
         static auto waiterWaitFor(QueueGenericSyclBlocking<TDev>& queue, EventGenericSycl<TDev> const& event)
         {
-            queue.m_impl->register_dependency(event.getNativeHandle());
+            queue.m_spQueueImpl->register_dependency(event.getNativeHandle());
         }
     };
 

include/alpaka/mem/buf/sycl/Copy.hpp

@@ -59,7 +59,7 @@ namespace alpaka::detail
         }
 
 #    if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
-        ALPAKA_FN_HOST auto printDebug() const -> void
+        auto printDebug() const -> void
         {
             std::cout << __func__ << " e: " << m_extent << " ewb: " << this->m_extentWidthBytes
                       << " de: " << m_dstExtent << " dptr: " << reinterpret_cast<void*>(m_dstMemNative)
@@ -93,8 +93,7 @@ namespace alpaka::detail
 
         using TaskCopySyclBase<TDim, TViewDst, TViewSrc, TExtent>::TaskCopySyclBase;
 
-        ALPAKA_FN_HOST auto operator()(sycl::queue& queue, std::vector<sycl::event> const& requirements) const
-            -> sycl::event
+        auto operator()(sycl::queue& queue, std::vector<sycl::event> const& requirements) const -> sycl::event
         {
             ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
 
@@ -145,8 +144,7 @@ namespace alpaka::detail
         using TaskCopySyclBase<DimInt<1u>, TViewDst, TViewSrc, TExtent>::TaskCopySyclBase;
         using Elem = alpaka::Elem<TViewSrc>;
 
-        ALPAKA_FN_HOST auto operator()(sycl::queue& queue, std::vector<sycl::event> const& requirements) const
-            -> sycl::event
+        auto operator()(sycl::queue& queue, std::vector<sycl::event> const& requirements) const -> sycl::event
         {
             ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
 
@@ -179,10 +177,7 @@ namespace alpaka::detail
         using Elem = alpaka::Elem<TViewSrc>;
 
         template<typename TViewDstFwd>
-        ALPAKA_FN_HOST TaskCopySycl(
-            TViewDstFwd&& viewDst,
-            TViewSrc const& viewSrc,
-            [[maybe_unused]] TExtent const& extent)
+        TaskCopySycl(TViewDstFwd&& viewDst, TViewSrc const& viewSrc, [[maybe_unused]] TExtent const& extent)
             : m_dstMemNative(reinterpret_cast<void*>(getPtrNative(viewDst)))
             , m_srcMemNative(reinterpret_cast<void const*>(getPtrNative(viewSrc)))
         {
@@ -211,10 +206,7 @@ namespace alpaka::trait
     struct CreateTaskMemcpy<TDim, DevGenericSycl<TPltf>, DevCpu>
     {
         template<typename TExtent, typename TViewSrc, typename TViewDstFwd>
-        ALPAKA_FN_HOST static auto createTaskMemcpy(
-            TViewDstFwd&& viewDst,
-            TViewSrc const& viewSrc,
-            TExtent const& extent)
+        static auto createTaskMemcpy(TViewDstFwd&& viewDst, TViewSrc const& viewSrc, TExtent const& extent)
             -> alpaka::detail::TaskCopySycl<TDim, std::remove_reference_t<TViewDstFwd>, TViewSrc, TExtent>
         {
             ALPAKA_DEBUG_FULL_LOG_SCOPE;
@@ -228,10 +220,7 @@ namespace alpaka::trait
     struct CreateTaskMemcpy<TDim, DevCpu, DevGenericSycl<TPltf>>
     {
         template<typename TExtent, typename TViewSrc, typename TViewDstFwd>
-        ALPAKA_FN_HOST static auto createTaskMemcpy(
-            TViewDstFwd&& viewDst,
-            TViewSrc const& viewSrc,
-            TExtent const& extent)
+        static auto createTaskMemcpy(TViewDstFwd&& viewDst, TViewSrc const& viewSrc, TExtent const& extent)
             -> alpaka::detail::TaskCopySycl<TDim, std::remove_reference_t<TViewDstFwd>, TViewSrc, TExtent>
         {
             ALPAKA_DEBUG_FULL_LOG_SCOPE;
@@ -245,10 +234,7 @@ namespace alpaka::trait
     struct CreateTaskMemcpy<TDim, DevGenericSycl<TPltfDst>, DevGenericSycl<TPltfSrc>>
     {
         template<typename TExtent, typename TViewSrc, typename TViewDstFwd>
-        ALPAKA_FN_HOST static auto createTaskMemcpy(
-            TViewDstFwd&& viewDst,
-            TViewSrc const& viewSrc,
-            TExtent const& extent)
+        static auto createTaskMemcpy(TViewDstFwd&& viewDst, TViewSrc const& viewSrc, TExtent const& extent)
             -> alpaka::detail::TaskCopySycl<TDim, std::remove_reference_t<TViewDstFwd>, TViewSrc, TExtent>
         {
             ALPAKA_DEBUG_FULL_LOG_SCOPE;

include/alpaka/mem/buf/sycl/Set.hpp

@@ -54,7 +54,7 @@ namespace alpaka
             }
 
 #    if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
-            ALPAKA_FN_HOST auto printDebug() const -> void
+            auto printDebug() const -> void
             {
                 std::cout << __func__ << " e: " << this->m_extent << " ewb: " << this->m_extentWidthBytes
                           << " de: " << this->m_dstExtent << " dptr: " << reinterpret_cast<void*>(this->m_dstMemNative)
@@ -170,7 +170,7 @@ namespace alpaka
             }
 
 #    if ALPAKA_DEBUG >= ALPAKA_DEBUG_FULL
-            ALPAKA_FN_HOST auto printDebug() const -> void
+            auto printDebug() const -> void
             {
                 std::cout << __func__ << " e: " << Scalar() << " ewb: " << sizeof(Elem) << " de: " << Scalar()
                           << " dptr: " << reinterpret_cast<void*>(m_dstMemNative) << " dpitchb: " << Scalar()
@@ -203,7 +203,7 @@ namespace alpaka
         struct CreateTaskMemset<TDim, DevGenericSycl<TPltf>>
         {
             template<typename TExtent, typename TView>
-            ALPAKA_FN_HOST static auto createTaskMemset(TView& view, std::uint8_t const& byte, TExtent const& extent)
+            static auto createTaskMemset(TView& view, std::uint8_t const& byte, TExtent const& extent)
                 -> detail::TaskSetSycl<TDim, TView, TExtent>
             {
                 return detail::TaskSetSycl<TDim, TView, TExtent>(view, byte, extent);

include/alpaka/mem/view/ViewPlainPtr.hpp

@@ -1,5 +1,5 @@
 /* Copyright 2023 Benjamin Worpitz, Matthias Werner, René Widera, Sergei Bastrakov, Bernhard Manfred Gruber,
- *                Jan Stephan, Andrea Bocci
+ *                Jan Stephan, Andrea Bocci, Aurora Perego
  * SPDX-License-Identifier: MPL-2.0
  */
 
@@ -229,6 +229,26 @@ namespace alpaka
         };
 #endif
 
+#if defined(ALPAKA_ACC_SYCL_ENABLED)
+        //! The SYCL device CreateViewPlainPtr trait specialization.
+        template<typename TPltf>
+        struct CreateViewPlainPtr<DevGenericSycl<TPltf>>
+        {
+            template<typename TElem, typename TExtent, typename TPitch>
+            static auto createViewPlainPtr(
+                DevGenericSycl<TPltf> const& dev,
+                TElem* pMem,
+                TExtent const& extent,
+                TPitch const& pitch)
+            {
+                return alpaka::ViewPlainPtr<DevGenericSycl<TPltf>, TElem, alpaka::Dim<TExtent>, alpaka::Idx<TExtent>>(
+                    pMem,
+                    dev,
+                    extent,
+                    pitch);
+            }
+        };
+#endif
         //! The ViewPlainPtr offset get trait specialization.
         template<typename TIdxIntegralConst, typename TDev, typename TElem, typename TDim, typename TIdx>
         struct GetOffset<TIdxIntegralConst, ViewPlainPtr<TDev, TElem, TDim, TIdx>>

include/alpaka/queue/sycl/QueueGenericSyclBase.hpp

@@ -1,4 +1,4 @@
-/* Copyright 2023 Jan Stephan, Antonio Di Pilato, Luca Ferragina, Andrea Bocci
+/* Copyright 2023 Jan Stephan, Antonio Di Pilato, Luca Ferragina, Andrea Bocci, Aurora Perego
  * SPDX-License-Identifier: MPL-2.0
  */
 
@@ -172,16 +172,16 @@ namespace alpaka::detail
     public:
         QueueGenericSyclBase(TDev const& dev)
             : m_dev{dev}
-            , m_impl{std::make_shared<detail::QueueGenericSyclImpl>(
+            , m_spQueueImpl{std::make_shared<detail::QueueGenericSyclImpl>(
                   dev.getNativeHandle().second,
                   dev.getNativeHandle().first)}
         {
-            m_dev.m_impl->register_queue(m_impl);
+            m_dev.m_impl->register_queue(m_spQueueImpl);
         }
 
         friend auto operator==(QueueGenericSyclBase const& lhs, QueueGenericSyclBase const& rhs) -> bool
         {
-            return (lhs.m_dev == rhs.m_dev) && (lhs.m_impl == rhs.m_impl);
+            return (lhs.m_dev == rhs.m_dev) && (lhs.m_spQueueImpl == rhs.m_spQueueImpl);
         }
 
         friend auto operator!=(QueueGenericSyclBase const& lhs, QueueGenericSyclBase const& rhs) -> bool
@@ -191,11 +191,11 @@ namespace alpaka::detail
 
         [[nodiscard]] auto getNativeHandle() const noexcept
         {
-            return m_impl->getNativeHandle();
+            return m_spQueueImpl->getNativeHandle();
         }
 
         TDev m_dev;
-        std::shared_ptr<detail::QueueGenericSyclImpl> m_impl;
+        std::shared_ptr<detail::QueueGenericSyclImpl> m_spQueueImpl;
     };
 } // namespace alpaka::detail
 
@@ -239,7 +239,7 @@ namespace alpaka::trait
         static auto enqueue(detail::QueueGenericSyclBase<TDev, TBlocking>& queue, TTask const& task) -> void
         {
             ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
-            queue.m_impl->template enqueue<TBlocking>(task);
+            queue.m_spQueueImpl->template enqueue<TBlocking>(task);
         }
     };
 
@@ -250,7 +250,7 @@ namespace alpaka::trait
         static auto empty(detail::QueueGenericSyclBase<TDev, TBlocking> const& queue) -> bool
         {
             ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
-            return queue.m_impl->empty();
+            return queue.m_spQueueImpl->empty();
         }
     };
 
@@ -264,7 +264,7 @@ namespace alpaka::trait
         static auto currentThreadWaitFor(detail::QueueGenericSyclBase<TDev, TBlocking> const& queue) -> void
         {
             ALPAKA_DEBUG_MINIMAL_LOG_SCOPE;
-            queue.m_impl->wait();
+            queue.m_spQueueImpl->wait();
         }
     };