diff --git a/include/alpaka/acc/AccGenericSycl.hpp b/include/alpaka/acc/AccGenericSycl.hpp index 03d1e23f90c1..2fb0b081977f 100644 --- a/include/alpaka/acc/AccGenericSycl.hpp +++ b/include/alpaka/acc/AccGenericSycl.hpp @@ -77,9 +77,7 @@ namespace alpaka Vec const& threadElemExtent, sycl::nd_item work_item, sycl::local_accessor dyn_shared_acc, - sycl::local_accessor st_shared_acc, - sycl::accessor global_fence_dummy, - sycl::local_accessor local_fence_dummy) + sycl::local_accessor st_shared_acc) : WorkDivGenericSycl{threadElemExtent, work_item} , gb::IdxGbGenericSycl{work_item} , bt::IdxBtGenericSycl{work_item} @@ -89,7 +87,7 @@ namespace alpaka , BlockSharedMemStGenericSycl{st_shared_acc} , BlockSyncGenericSycl{work_item} , IntrinsicGenericSycl{} - , MemFenceGenericSycl{global_fence_dummy, local_fence_dummy} + , MemFenceGenericSycl{} # ifdef ALPAKA_DISABLE_VENDOR_RNG , rand::RandDefault{} # else diff --git a/include/alpaka/kernel/TaskKernelGenericSycl.hpp b/include/alpaka/kernel/TaskKernelGenericSycl.hpp index 8b9f38b8da57..5527dbbb9b98 100644 --- a/include/alpaka/kernel/TaskKernelGenericSycl.hpp +++ b/include/alpaka/kernel/TaskKernelGenericSycl.hpp @@ -43,8 +43,6 @@ [item_elements, \ dyn_shared_accessor, \ st_shared_accessor, \ - global_fence_dummy, \ - local_fence_dummy, \ k_func, \ k_args](sycl::nd_item work_item) [[intel::reqd_sub_group_size(sub_group_size)]] \ { \ @@ -52,9 +50,7 @@ item_elements, \ work_item, \ dyn_shared_accessor, \ - st_shared_accessor, \ - global_fence_dummy, \ - local_fence_dummy}; \ + st_shared_accessor}; \ core::apply( \ [k_func, &acc](typename std::decay_t const&... args) { k_func(acc, args...); }, \ k_args); \ @@ -66,8 +62,6 @@ [item_elements, \ dyn_shared_accessor, \ st_shared_accessor, \ - global_fence_dummy, \ - local_fence_dummy, \ k_func, \ k_args](sycl::nd_item work_item) \ { \ @@ -75,9 +69,7 @@ item_elements, \ work_item, \ dyn_shared_accessor, \ - st_shared_accessor, \ - global_fence_dummy, \ - local_fence_dummy}; \ + st_shared_accessor}; \ core::apply( \ [k_func, &acc](typename std::decay_t const&... args) { k_func(acc, args...); }, \ k_args); \ @@ -90,8 +82,6 @@ [item_elements, \ dyn_shared_accessor, \ st_shared_accessor, \ - global_fence_dummy, \ - local_fence_dummy, \ k_func, \ k_args](sycl::nd_item work_item) {}); @@ -121,7 +111,7 @@ namespace alpaka { } - auto operator()(sycl::handler& cgh, sycl::buffer& global_fence_buf) const -> void + auto operator()(sycl::handler& cgh) const -> void { auto const work_groups = WorkDivMembers::m_gridBlockExtent; @@ -146,10 +136,6 @@ namespace alpaka constexpr auto st_shared_mem_bytes = std::size_t{ALPAKA_BLOCK_SHARED_DYN_MEMBER_ALLOC_KIB * 1024}; auto st_shared_accessor = sycl::local_accessor{sycl::range<1>{st_shared_mem_bytes}, cgh}; - // register memory fence dummies - auto global_fence_dummy = global_fence_buf.get_access(cgh); // Exists once per queue - auto local_fence_dummy = sycl::local_accessor{sycl::range<1>{1}, cgh}; - // copy-by-value so we don't access 'this' on the device auto k_func = m_kernelFnObj; auto k_args = m_args; diff --git a/include/alpaka/mem/fence/MemFenceGenericSycl.hpp b/include/alpaka/mem/fence/MemFenceGenericSycl.hpp index 0b7559f3d61b..ce2f33719531 100644 --- a/include/alpaka/mem/fence/MemFenceGenericSycl.hpp +++ b/include/alpaka/mem/fence/MemFenceGenericSycl.hpp @@ -23,38 +23,24 @@ namespace alpaka struct SyclFenceProps { static constexpr auto scope = sycl::memory_scope::work_group; - static constexpr auto space = sycl::access::address_space::local_space; }; template<> struct SyclFenceProps { static constexpr auto scope = sycl::memory_scope::device; - static constexpr auto space = sycl::access::address_space::global_space; }; template<> struct SyclFenceProps { static constexpr auto scope = sycl::memory_scope::device; - static constexpr auto space = sycl::access::address_space::global_space; }; } // namespace detail //! The SYCL memory fence. class MemFenceGenericSycl : public concepts::Implements { - public: - MemFenceGenericSycl( - sycl::accessor global_dummy, - sycl::local_accessor local_dummy) - : m_global_dummy{global_dummy} - , m_local_dummy{local_dummy} - { - } - - sycl::accessor m_global_dummy; - sycl::local_accessor m_local_dummy; }; } // namespace alpaka @@ -63,17 +49,10 @@ namespace alpaka::trait template struct MemFence { - static auto mem_fence(MemFenceGenericSycl const& fence, TMemScope const&) + static auto mem_fence(MemFenceGenericSycl const& , TMemScope const&) { static constexpr auto scope = detail::SyclFenceProps::scope; - static constexpr auto space = detail::SyclFenceProps::space; - auto dummy - = (scope == sycl::memory_scope::work_group) - ? sycl::atomic_ref{fence.m_local_dummy[0]} - : sycl::atomic_ref{fence.m_global_dummy[0]}; - auto const dummy_val = dummy.load(); sycl::atomic_fence(sycl::memory_order::acq_rel, scope); - dummy.store(dummy_val); } }; } // namespace alpaka::trait diff --git a/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp b/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp index 9b08f3de55ca..ca62f7250831 100644 --- a/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp +++ b/include/alpaka/queue/sycl/QueueGenericSyclBase.hpp @@ -139,7 +139,7 @@ namespace alpaka::detail cgh.depends_on(m_dependencies); if constexpr(is_sycl_kernel) // Kernel - task(cgh, m_fence_dummy); // Will call cgh.parallel_for internally + task(cgh); // Will call cgh.parallel_for internally else // Host cgh.host_task(task); }); @@ -159,7 +159,6 @@ namespace alpaka::detail std::vector m_dependencies; sycl::event m_last_event; - sycl::buffer m_fence_dummy{sycl::range<1>{1}}; std::shared_mutex mutable m_mutex; private: