DPC++ Runtime
Runtime libraries for oneAPI DPC++
atomic_accessor.hpp
Go to the documentation of this file.
1 //==-- atomic_accessor.hpp - SYCL_ONEAPI_extended_atomics atomic_accessor --==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 
9 #pragma once
10 
12 #include <CL/sycl/accessor.hpp>
15 
17 namespace sycl {
18 namespace ext {
19 namespace oneapi {
20 
21 #if __cplusplus >= 201703L
22 
23 template <memory_order> struct order_tag_t {
24  explicit order_tag_t() = default;
25 };
26 inline constexpr order_tag_t<memory_order::relaxed> relaxed_order{};
27 inline constexpr order_tag_t<memory_order::acquire> acquire_order{};
28 inline constexpr order_tag_t<memory_order::release> release_order{};
29 inline constexpr order_tag_t<memory_order::acq_rel> acq_rel_order{};
30 inline constexpr order_tag_t<memory_order::seq_cst> seq_cst_order{};
31 
32 template <memory_scope> struct scope_tag_t {
33  explicit scope_tag_t() = default;
34 };
35 inline constexpr scope_tag_t<memory_scope::work_item> work_item_scope{};
36 inline constexpr scope_tag_t<memory_scope::sub_group> sub_group_scope{};
37 inline constexpr scope_tag_t<memory_scope::work_group> work_group_scope{};
38 inline constexpr scope_tag_t<memory_scope::device> device_scope{};
39 inline constexpr scope_tag_t<memory_scope::system> system_scope{};
40 
41 #endif
42 
43 template <typename DataT, int Dimensions, memory_order DefaultOrder,
44  memory_scope DefaultScope,
46  access::placeholder IsPlaceholder = access::placeholder::false_t>
48  : public accessor<DataT, Dimensions, access::mode::read_write, AccessTarget,
49  IsPlaceholder, ext::oneapi::accessor_property_list<>> {
50 
51  using AccessorT =
52  accessor<DataT, Dimensions, access::mode::read_write, AccessTarget,
54 
55 private:
57  using AccessorT::getQualifiedPtr;
58 
59  // Prevent non-atomic access to atomic accessor
60  multi_ptr<DataT, AccessorT::AS> get_pointer() const = delete;
61 
62 protected:
63  using AccessorT::AdjustedDim;
64 
65 public:
66  using value_type = DataT;
67  using reference =
69 
70  using AccessorT::AccessorT;
71 
72 #if __cplusplus >= 201703L
73 
74  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
75  memory_order Order, memory_scope Scope>
76  atomic_accessor(buffer<T, Dims, AllocatorT> &BufferRef, order_tag_t<Order>,
77  scope_tag_t<Scope>, const property_list &PropertyList = {})
78  : atomic_accessor(BufferRef, PropertyList) {}
79 
80  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
81  memory_order Order, memory_scope Scope>
82  atomic_accessor(buffer<T, Dims, AllocatorT> &BufferRef,
83  handler &CommandGroupHandler, order_tag_t<Order>,
84  scope_tag_t<Scope>, const property_list &PropertyList = {})
85  : atomic_accessor(BufferRef, CommandGroupHandler, PropertyList) {}
86 
87 #endif
88 
89  // Override subscript operators and conversions to wrap in an atomic_ref
90  template <int Dims = Dimensions>
91  operator typename detail::enable_if_t<Dims == 0, reference>() const {
92  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
93  return reference(getQualifiedPtr()[LinearIndex]);
94  }
95 
96  template <int Dims = Dimensions>
97  typename detail::enable_if_t<(Dims > 0), reference>
98  operator[](id<Dimensions> Index) const {
99  const size_t LinearIndex = getLinearIndex(Index);
100  return reference(getQualifiedPtr()[LinearIndex]);
101  }
102 
103  template <int Dims = Dimensions>
105  operator[](size_t Index) const {
106  const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
107  return reference(getQualifiedPtr()[LinearIndex]);
108  }
109 };
110 
111 #if __cplusplus >= 201703L
112 
113 template <typename DataT, int Dimensions, typename AllocatorT,
114  memory_order Order, memory_scope Scope>
115 atomic_accessor(buffer<DataT, Dimensions, AllocatorT>, order_tag_t<Order>,
116  scope_tag_t<Scope>, property_list = {})
117  -> atomic_accessor<DataT, Dimensions, Order, Scope, target::device,
118  access::placeholder::true_t>;
119 
120 template <typename DataT, int Dimensions, typename AllocatorT,
121  memory_order Order, memory_scope Scope>
123  order_tag_t<Order>, scope_tag_t<Scope>, property_list = {})
124  -> atomic_accessor<DataT, Dimensions, Order, Scope, target::device,
125  access::placeholder::false_t>;
126 
127 #endif
128 
129 } // namespace oneapi
130 } // namespace ext
131 
132 } // namespace sycl
133 } // __SYCL_INLINE_NAMESPACE(cl)
cl::sycl::detail::getLinearIndex
size_t getLinearIndex(const T< Dims > &Index, const U< Dims > &Range)
Definition: common.hpp:336
cl::sycl::memory_order
memory_order
Definition: memory_enums.hpp:16
cl::sycl::access::placeholder
placeholder
Definition: access.hpp:43
cl::sycl::ext::oneapi::atomic_accessor::operator[]
detail::enable_if_t< Dims==1, reference > operator[](size_t Index) const
Definition: atomic_accessor.hpp:105
atomic_ref.hpp
cl::sycl::id
A unique identifier of an item in an index space.
Definition: array.hpp:17
cl::sycl::ext::oneapi::atomic_accessor
Definition: atomic_accessor.hpp:47
sycl
Definition: invoke_simd.hpp:68
cl::sycl::property_list
Objects of the property_list class are containers for the SYCL properties.
Definition: property_list.hpp:26
cl::sycl::multi_ptr
Provides constructors for address space qualified and non address space qualified pointers to allow i...
Definition: atomic.hpp:33
access.hpp
cl::sycl::buffer
Defines a shared array that can be used by kernels in queues.
Definition: buffer.hpp:78
cl::sycl::memory_scope
memory_scope
Definition: memory_enums.hpp:26
cl::sycl::ext::oneapi::atomic_accessor::value_type
DataT value_type
Definition: atomic_accessor.hpp:66
atomic_enums.hpp
cl::sycl::accessor
Buffer accessor.
Definition: accessor.hpp:224
cl::sycl::access::target
target
Definition: access.hpp:17
cl
We provide new interfaces for matrix muliply in this patch:
Definition: access.hpp:13
cl::sycl::atomic_ref
Definition: atomic_ref.hpp:661
cl::sycl::handler
Command group handler class.
Definition: handler.hpp:362
cl::sycl::ext::oneapi::accessor_property_list
Objects of the accessor_property_list class are containers for the SYCL properties.
Definition: property_list.hpp:19
accessor.hpp
cl::sycl::info::device
device
Definition: info_desc.hpp:53
cl::sycl::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
cl::sycl::Dimensions
Dimensions
Definition: backend.hpp:142
__SYCL_INLINE_NAMESPACE
#define __SYCL_INLINE_NAMESPACE(X)
Definition: defines_elementary.hpp:12