[SYCL] Add atomic_accessor extension (#2069)

Updates the SYCL_INTEL_extended_atomics extension to describe atomic_accessor, and adds an initial implementation.

Author: Pennycook

sycl/doc/extensions/ExtendedAtomics/SYCL_INTEL_extended_atomics.asciidoc

@@ -331,13 +331,43 @@ The effects of a call to +atomic_fence+ depend on the value of the +order+ param
 - `acq_rel`: Both an acquire fence and a release fence
 - `seq_cst`: A sequentially consistent acquire and release fence
 
+==== Atomic Accessor
+
+The +atomic_accessor+ class replaces accessors created with +access::mode::atomic+.  All operators of +atomic_accessor+ which provide access to an element of the underlying memory do so by wrapping the element in an +atomic_ref+.  In order to guarantee that all accesses to the underlying memory are atomic, an +atomic_accessor+ does not allow direct access to the memory via +get_pointer()+.
+
+All other member functions are as defined in the +accessor+ class.
+
+|===
+| Member Function | Description
+
+| `operator atomic_ref<DataT, DefaultOrder, DefaultScope, access::address_space::global_space>() const`
+| Available only when: `Dimensions == 0`.  Returns an `atomic_ref` associated with the element stored in the underlying `buffer` or work-group local memory.
+
+| `atomic_ref<DataT, DefaultOrder, DefaultScope, access::address_space::global_space> operator[](id<Dimensions> index) const`
+| Available only when: `Dimensions > 0`.  Returns an `atomic_ref` associated with the element stored at the specified _index_ in the underlying `buffer` or work-group local memory.
+
+| `atomic_ref<DataT, DefaultOrder, DefaultScope, access::address_space::global_space> operator[](size_t index) const`
+| Available only when: `Dimensions == 1`.  Returns an `atomic_ref` associated with the element stored at the specified _index_ in the underlying `buffer` or work-group local memory.
+
+| `global_ptr<DataT> get_pointer() const = delete`
+| Direct access to the underlying `buffer` or work-group local memory is not permitted.
+
+|===
+
+To simplify the construction of an +atomic_accessor+, tag objects of type +order_tag_t+ and +scope_tag_t+ may optionally be passed to the constructor.  These tag objects enable the `DefaultOrder` and `DefaultScope` template arguments to be deduced via CTAD, as shown in the example below:
+[source,c++]
+----
+auto acc = atomic_accessor(buf, h, relaxed_order, device_scope);
+----
+
 ==== Sample  Header
 
 [source,c++]
 ----
 namespace cl {
 namespace sycl {
 namespace intel {
+
 enum class memory_order : /* unspecified */ {
   relaxed, acquire, release, acq_rel, seq_cst
 };
@@ -356,6 +386,24 @@ inline constexpr memory_scope memory_scope_work_group = memory_scope::work_group
 inline constexpr memory_scope memory_scope_device = memory_scope::device;
 inline constexpr memory_scope memory_scope_system = memory_scope::system;
 
+template <memory_order> struct order_tag_t {
+  explicit order_tag_t() = default;
+};
+inline constexpr order_tag_t<memory_order::relaxed> relaxed_order{};
+inline constexpr order_tag_t<memory_order::acquire> acquire_order{};
+inline constexpr order_tag_t<memory_order::release> release_order{};
+inline constexpr order_tag_t<memory_order::acq_rel> acq_rel_order{};
+inline constexpr order_tag_t<memory_order::seq_cst> seq_cst_order{};
+
+template <memory_scope> struct scope_tag_t {
+  explicit scope_tag_t() = default;
+};
+inline constexpr scope_tag_t<memory_scope::work_item> work_item_scope{};
+inline constexpr scope_tag_t<memory_scope::sub_group> sub_group_scope{};
+inline constexpr scope_tag_t<memory_scope::work_group> work_group_scope{};
+inline constexpr scope_tag_t<memory_scope::device> device_scope{};
+inline constexpr scope_tag_t<memory_scope::system> system_scope{};
+
 // Exposition only
 template <memory_order ReadModifyWriteOrder>
 struct memory_order_traits;
@@ -578,6 +626,12 @@ class atomic_ref<T*, DefaultOrder, DefaultScope, Space> {
 
 void atomic_fence(memory_order order, memory_scope scope):
 
+template <typename DataT, int Dimensions,
+          memory_order DefaultOrder, memory_scope DefaultScope,
+          access::target AccessTarget = access::target::global_buffer,
+          access::placeholder IsPlaceholder = access::placeholder::false_t>
+class atomic_accessor;
+
 } // namespace intel
 } // namespace sycl
 } // namespace cl
@@ -605,6 +659,7 @@ None.
 |3|2020-04-09|John Pennycook|*Add atomic_fence*
 |4|2020-04-24|John Pennycook|*Add memory scope*
 |5|2020-04-29|John Pennycook|*Fix ambiguous overloads of compare_exchange and typo in fetch_sub*
+|6|2020-07-08|John Pennycook|*Add atomic_accessor*
 |========================================
 
 //************************************************************************

sycl/include/CL/sycl/accessor.hpp

@@ -1359,7 +1359,7 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
 #endif
     public detail::accessor_common<DataT, Dimensions, AccessMode,
                                    access::target::local, IsPlaceholder> {
-
+protected:
   constexpr static int AdjustedDim = Dimensions == 0 ? 1 : Dimensions;
 
   using AccessorCommonT =
@@ -1395,7 +1395,7 @@ class accessor<DataT, Dimensions, AccessMode, access::target::local,
   accessor()
       : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
 
-private:
+protected:
   ConcreteASPtrType getQualifiedPtr() const { return MData; }
 
   ConcreteASPtrType MData;

sycl/include/CL/sycl/intel/atomic.hpp

@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <CL/sycl/intel/atomic_accessor.hpp>
 #include <CL/sycl/intel/atomic_enums.hpp>
 #include <CL/sycl/intel/atomic_fence.hpp>
 #include <CL/sycl/intel/atomic_ref.hpp>

sycl/include/CL/sycl/intel/atomic_accessor.hpp

@@ -0,0 +1,128 @@
+//==--- atomic_accessor.hpp - SYCL_INTEL_extended_atomics atomic_accessor --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl/access/access.hpp>
+#include <CL/sycl/intel/atomic_enums.hpp>
+#include <CL/sycl/intel/atomic_ref.hpp>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace intel {
+
+#if __cplusplus > 201402L
+
+template <memory_order> struct order_tag_t {
+  explicit order_tag_t() = default;
+};
+inline constexpr order_tag_t<memory_order::relaxed> relaxed_order{};
+inline constexpr order_tag_t<memory_order::acquire> acquire_order{};
+inline constexpr order_tag_t<memory_order::release> release_order{};
+inline constexpr order_tag_t<memory_order::acq_rel> acq_rel_order{};
+inline constexpr order_tag_t<memory_order::seq_cst> seq_cst_order{};
+
+template <memory_scope> struct scope_tag_t {
+  explicit scope_tag_t() = default;
+};
+inline constexpr scope_tag_t<memory_scope::work_item> work_item_scope{};
+inline constexpr scope_tag_t<memory_scope::sub_group> sub_group_scope{};
+inline constexpr scope_tag_t<memory_scope::work_group> work_group_scope{};
+inline constexpr scope_tag_t<memory_scope::device> device_scope{};
+inline constexpr scope_tag_t<memory_scope::system> system_scope{};
+
+#endif
+
+template <typename DataT, int Dimensions, memory_order DefaultOrder,
+          memory_scope DefaultScope,
+          access::target AccessTarget = access::target::global_buffer,
+          access::placeholder IsPlaceholder = access::placeholder::false_t>
+class atomic_accessor
+    : public accessor<DataT, Dimensions, access::mode::read_write, AccessTarget,
+                      IsPlaceholder> {
+
+  using AccessorT = accessor<DataT, Dimensions, access::mode::read_write,
+                             AccessTarget, IsPlaceholder>;
+
+private:
+  using AccessorT::getLinearIndex;
+  using AccessorT::getQualifiedPtr;
+
+  // Prevent non-atomic access to atomic accessor
+  multi_ptr<DataT, AccessorT::AS> get_pointer() const = delete;
+
+protected:
+  using AccessorT::AdjustedDim;
+
+public:
+  using value_type = DataT;
+  using reference =
+      atomic_ref<DataT, DefaultOrder, DefaultScope, AccessorT::AS>;
+
+  using AccessorT::AccessorT;
+
+#if __cplusplus > 201402L
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            memory_order Order, memory_scope Scope>
+  atomic_accessor(buffer<T, Dims, AllocatorT> &BufferRef, order_tag_t<Order>,
+                  scope_tag_t<Scope>, const property_list &PropertyList = {})
+      : atomic_accessor(BufferRef, PropertyList) {}
+
+  template <typename T = DataT, int Dims = Dimensions, typename AllocatorT,
+            memory_order Order, memory_scope Scope>
+  atomic_accessor(buffer<T, Dims, AllocatorT> &BufferRef,
+                  handler &CommandGroupHandler, order_tag_t<Order>,
+                  scope_tag_t<Scope>, const property_list &PropertyList = {})
+      : atomic_accessor(BufferRef, CommandGroupHandler, PropertyList) {}
+
+#endif
+
+  // Override subscript operators and conversions to wrap in an atomic_ref
+  template <int Dims = Dimensions>
+  operator typename detail::enable_if_t<Dims == 0, reference>() const {
+    const size_t LinearIndex = getLinearIndex(id<AdjustedDim>());
+    return reference(getQualifiedPtr()[LinearIndex]);
+  }
+
+  template <int Dims = Dimensions>
+  typename detail::enable_if_t<(Dims > 0), reference>
+  operator[](id<Dimensions> Index) const {
+    const size_t LinearIndex = getLinearIndex(Index);
+    return reference(getQualifiedPtr()[LinearIndex]);
+  }
+
+  template <int Dims = Dimensions>
+  typename detail::enable_if_t<Dims == 1, reference>
+  operator[](size_t Index) const {
+    const size_t LinearIndex = getLinearIndex(id<AdjustedDim>(Index));
+    return reference(getQualifiedPtr()[LinearIndex]);
+  }
+};
+
+#if __cplusplus > 201402L
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          memory_order Order, memory_scope Scope>
+atomic_accessor(buffer<DataT, Dimensions, AllocatorT>, order_tag_t<Order>,
+                scope_tag_t<Scope>, property_list = {})
+    ->atomic_accessor<DataT, Dimensions, Order, Scope, target::global_buffer,
+                      access::placeholder::true_t>;
+
+template <typename DataT, int Dimensions, typename AllocatorT,
+          memory_order Order, memory_scope Scope>
+atomic_accessor(buffer<DataT, Dimensions, AllocatorT>, handler,
+                order_tag_t<Order>, scope_tag_t<Scope>, property_list = {})
+    ->atomic_accessor<DataT, Dimensions, Order, Scope, target::global_buffer,
+                      access::placeholder::false_t>;
+
+#endif
+
+} // namespace intel
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/atomic_ref/accessor.cpp

@@ -0,0 +1,110 @@
+// RUN: %clangxx -fsycl -fsycl-unnamed-lambda -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+#include <algorithm>
+#include <cassert>
+#include <iostream>
+#include <numeric>
+#include <vector>
+using namespace sycl;
+using namespace sycl::intel;
+
+// Equivalent to add_test from add.cpp
+// Uses atomic_accessor instead of atomic_ref
+template <typename T> void accessor_test(queue q, size_t N) {
+  T sum = 0;
+  std::vector<T> output(N, 0);
+  {
+    buffer<T> sum_buf(&sum, 1);
+    buffer<T> output_buf(output.data(), output.size());
+
+    q.submit([&](handler &cgh) {
+#if __cplusplus > 201402L
+      static_assert(
+          std::is_same<decltype(atomic_accessor(sum_buf, cgh, relaxed_order,
+                                                device_scope)),
+                       atomic_accessor<T, 1, intel::memory_order::relaxed,
+                                       intel::memory_scope::device>>::value,
+          "atomic_accessor type incorrectly deduced");
+#endif
+      auto sum = atomic_accessor<T, 1, intel::memory_order::relaxed,
+                                 intel::memory_scope::device>(sum_buf, cgh);
+      auto out =
+          output_buf.template get_access<access::mode::discard_write>(cgh);
+      cgh.parallel_for(range<1>(N), [=](item<1> it) {
+        int gid = it.get_id(0);
+        static_assert(
+            std::is_same<
+                decltype(sum[0]),
+                atomic_ref<T, intel::memory_order::relaxed,
+                           intel::memory_scope::device,
+                           access::address_space::global_space>>::value,
+            "atomic_accessor returns incorrect atomic_ref");
+        out[gid] = sum[0].fetch_add(T(1));
+      });
+    });
+  }
+
+  // All work-items increment by 1, so final value should be equal to N
+  assert(sum == N);
+
+  // Intermediate values should be unique
+  std::sort(output.begin(), output.end());
+  assert(std::unique(output.begin(), output.end()) == output.end());
+
+  // Fetch returns original value: will be in [0, N-1]
+  auto min_e = output[0];
+  auto max_e = output[output.size() - 1];
+  assert(min_e == 0 && max_e == N - 1);
+}
+
+// Simplified form of accessor_test for local memory
+template <typename T>
+void local_accessor_test(queue q, size_t N, size_t L = 8) {
+  assert(N % L == 0);
+  std::vector<T> output(N / L, 0);
+  {
+    buffer<T> output_buf(output.data(), output.size());
+    q.submit([&](handler &cgh) {
+      auto sum =
+          atomic_accessor<T, 1, intel::memory_order::relaxed,
+                          intel::memory_scope::device, access::target::local>(
+              1, cgh);
+      auto out = output_buf.template get_access<access::mode::read_write>(cgh);
+      cgh.parallel_for(nd_range<1>(N, L), [=](nd_item<1> it) {
+        int grp = it.get_group(0);
+        sum[0].store(0);
+        it.barrier();
+        static_assert(
+            std::is_same<decltype(sum[0]),
+                         atomic_ref<T, intel::memory_order::relaxed,
+                                    intel::memory_scope::device,
+                                    access::address_space::local_space>>::value,
+            "local atomic_accessor returns incorrect atomic_ref");
+        T result = sum[0].fetch_add(T(1));
+        if (result == it.get_local_range(0) - 1) {
+          out[grp] = result;
+        }
+      });
+    });
+  }
+
+  // All work-items increment by 1, and last in the group writes out old value
+  // All values should be L-1
+  assert(std::all_of(output.begin(), output.end(),
+                     [=](T x) { return x == L - 1; }));
+}
+
+int main() {
+  queue q;
+  constexpr int N = 32;
+  accessor_test<int>(q, N);
+  // TODO: Enable local accessor test for host when barrier is supported
+  if (!q.get_device().is_host()) {
+    local_accessor_test<int>(q, N);
+  }
+  std::cout << "Test passed." << std::endl;
+}