From b16f5ec454a952b581d4bc33d00d21a1e00c968b Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Wed, 2 Nov 2022 15:53:07 +0100
Subject: [PATCH 1/4] Fix clang-tidy warning

---
 tests/simd.cpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/tests/simd.cpp b/tests/simd.cpp
index d8b7398fd7..5d8ef59a24 100644
--- a/tests/simd.cpp
+++ b/tests/simd.cpp
@@ -364,6 +364,7 @@ TEST_CASE("simd.simdForEachN.stdsimd")
             view,
             [](auto simd)
             {
+                using std::sqrt;
                 simd(tag::Y{}) = sqrt(simd(tag::X{}));
                 return simd; // TODO(bgruber): tag::X{} is redundantly stored
             });
@@ -393,6 +394,7 @@ TEST_CASE("simd.simdForEach.stdsimd")
             view,
             [](auto simd)
             {
+                using std::sqrt;
                 simd(tag::Y{}) = sqrt(simd(tag::X{}));
                 return simd; // TODO(bgruber): tag::X{} is redundantly stored
             });

From 701d13528405a3270eaa756201f773a3758777fd Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Sat, 24 Sep 2022 19:05:25 +0200
Subject: [PATCH 2/4] Add accessors to Views

* Add a new template parameter to View
* Add a few new accessors (Default, ReadOnlyByValue, Const, Restrict,
  Atomic)

Fixes: #523
---
 include/llama/Accessors.hpp |  64 +++++++++++++++++
 include/llama/RecordRef.hpp |   4 +-
 include/llama/View.hpp      | 135 ++++++++++++++++++++----------------
 tests/view.cpp              |  92 +++++++++++++++++++++++-
 4 files changed, 233 insertions(+), 62 deletions(-)
 create mode 100644 include/llama/Accessors.hpp

diff --git a/include/llama/Accessors.hpp b/include/llama/Accessors.hpp
new file mode 100644
index 0000000000..a01af31af8
--- /dev/null
+++ b/include/llama/Accessors.hpp
@@ -0,0 +1,64 @@
+#pragma once
+
+#include "macros.hpp"
+
+#include <atomic>
+
+namespace llama::accessor
+{
+    /// Default accessor. Passes through the given reference.
+    struct Default
+    {
+        template<typename Reference>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(Reference&& r) const -> Reference
+        {
+            return std::forward<Reference>(r);
+        }
+    };
+
+    /// Allows only read access and returns values instead of references to memory.
+    struct ReadOnlyByValue
+    {
+        template<typename Reference>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(Reference&& r) const
+        {
+            using ValueType = std::decay_t<Reference>;
+            if constexpr(isProxyReference<ValueType>)
+                return static_cast<typename ValueType::value_type>(r);
+            else
+                return ValueType{r};
+        }
+    };
+
+    /// Allows only read access by qualifying the references to memory with const. Only works on l-value references.
+    struct Const
+    {
+        template<typename T>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(T& r) const -> const T&
+        {
+            return r;
+        }
+    };
+
+    /// Qualifies references to memory with __restrict. Only works on l-value references.
+    struct Restrict
+    {
+        template<typename T>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(T& r) const -> T& __restrict
+        {
+            return r;
+        }
+    };
+
+#ifdef __cpp_lib_atomic_ref
+    /// Accessor wrapping a reference into a std::atomic_ref. Can only wrap l-value references.
+    struct Atomic
+    {
+        template<typename T>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(T& r) const -> std::atomic_ref<T>
+        {
+            return std::atomic_ref<T>{r};
+        }
+    };
+#endif
+} // namespace llama::accessor
diff --git a/include/llama/RecordRef.hpp b/include/llama/RecordRef.hpp
index 5f52140dce..a3a365c307 100644
--- a/include/llama/RecordRef.hpp
+++ b/include/llama/RecordRef.hpp
@@ -442,7 +442,7 @@ namespace llama
             else
             {
                 LLAMA_FORCE_INLINE_RECURSIVE
-                return this->view.accessor(arrayIndex(), AbsolutCoord{});
+                return this->view.access(arrayIndex(), AbsolutCoord{});
             }
         }
 
@@ -460,7 +460,7 @@ namespace llama
             else
             {
                 LLAMA_FORCE_INLINE_RECURSIVE
-                return this->view.accessor(arrayIndex(), AbsolutCoord{});
+                return this->view.access(arrayIndex(), AbsolutCoord{});
             }
         }
 
diff --git a/include/llama/View.hpp b/include/llama/View.hpp
index 4228447463..109fd3e144 100644
--- a/include/llama/View.hpp
+++ b/include/llama/View.hpp
@@ -3,6 +3,7 @@
 
 #pragma once
 
+#include "Accessors.hpp"
 #include "Array.hpp"
 #include "ArrayIndexRange.hpp"
 #include "BlobAllocators.hpp"
@@ -17,9 +18,9 @@
 namespace llama
 {
 #ifdef __cpp_lib_concepts
-    template<typename TMapping, Blob BlobType>
+    template<typename TMapping, Blob BlobType, typename TAccessor>
 #else
-    template<typename TMapping, typename BlobType>
+    template<typename TMapping, typename BlobType, typename TAccessor>
 #endif
     struct View;
 
@@ -46,12 +47,14 @@ namespace llama
 
     /// Same as \ref allocView but does not run field constructors.
 #ifdef __cpp_lib_concepts
-    template<typename Mapping, BlobAllocator Allocator = bloballoc::Vector>
+    template<typename Mapping, BlobAllocator Allocator = bloballoc::Vector, typename Accessor = accessor::Default>
 #else
-    template<typename Mapping, typename Allocator = bloballoc::Vector>
+    template<typename Mapping, typename Allocator = bloballoc::Vector, typename Accessor = accessor::Default>
 #endif
-    LLAMA_FN_HOST_ACC_INLINE auto allocViewUninitialized(Mapping mapping = {}, const Allocator& alloc = {})
-        -> View<Mapping, internal::AllocatorBlobType<Allocator, typename Mapping::RecordDim>>
+    LLAMA_FN_HOST_ACC_INLINE auto allocViewUninitialized(
+        Mapping mapping = {},
+        const Allocator& alloc = {},
+        Accessor = {}) -> View<Mapping, internal::AllocatorBlobType<Allocator, typename Mapping::RecordDim>, Accessor>
     {
         auto blobs = internal::makeBlobArray(alloc, mapping, std::make_index_sequence<Mapping::blobCount>{});
         return {std::move(mapping), std::move(blobs)};
@@ -108,10 +111,10 @@ namespace llama
 
     /// Runs the constructor of all fields reachable through the given view. Computed fields are constructed if they
     /// return l-value references. If the mapping is a computed
-    template<typename Mapping, typename BlobType>
-    LLAMA_FN_HOST_ACC_INLINE void constructFields(View<Mapping, BlobType>& view)
+    template<typename Mapping, typename BlobType, typename Accessor>
+    LLAMA_FN_HOST_ACC_INLINE void constructFields(View<Mapping, BlobType, Accessor>& view)
     {
-        using View = View<Mapping, BlobType>;
+        using View = View<Mapping, BlobType, Accessor>;
         using RecordDim = typename View::RecordDim;
         forEachADCoord(
             view.mapping().extents(),
@@ -128,13 +131,15 @@ namespace llama
                             using FieldType = GetType<RecordDim, decltype(rc)>;
                             using RefType = decltype(view(ai)(rc));
                             // this handles physical and computed mappings
-                            if constexpr(std::is_lvalue_reference_v<RefType>)
+                            if constexpr(isProxyReference<RefType>)
                             {
-                                new(&view(ai)(rc)) FieldType;
+                                view(ai)(rc) = FieldType{};
                             }
-                            else if constexpr(isProxyReference<RefType>)
+                            else if constexpr(
+                                std::is_lvalue_reference_v<
+                                    RefType> && !std::is_const_v<std::remove_reference_t<RefType>>)
                             {
-                                view(ai)(rc) = FieldType{};
+                                new(&view(ai)(rc)) FieldType;
                             }
                         }
 #endif
@@ -144,13 +149,14 @@ namespace llama
                 {
                     // this handles physical and computed mappings
                     using RefType = decltype(view(ai));
-                    if constexpr(std::is_lvalue_reference_v<RefType>)
+                    if constexpr(isProxyReference<RefType>)
                     {
-                        new(&view(ai)) RecordDim;
+                        view(ai) = RecordDim{};
                     }
-                    else if constexpr(isProxyReference<RefType>)
+                    else if constexpr(
+                        std::is_lvalue_reference_v<RefType> && !std::is_const_v<std::remove_reference_t<RefType>>)
                     {
-                        view(ai) = RecordDim{};
+                        new(&view(ai)) RecordDim;
                     }
                 }
             });
@@ -162,14 +168,14 @@ namespace llama
     /// The constructors are run for all fields by calling \ref constructFields. This function is the preferred way to
     /// create a \ref View. See also \ref allocViewUninitialized.
 #ifdef __cpp_lib_concepts
-    template<typename Mapping, BlobAllocator Allocator = bloballoc::Vector>
+    template<typename Mapping, BlobAllocator Allocator = bloballoc::Vector, typename Accessor = accessor::Default>
 #else
-    template<typename Mapping, typename Allocator = bloballoc::Vector>
+    template<typename Mapping, typename Allocator = bloballoc::Vector, typename Accessor = accessor::Default>
 #endif
-    LLAMA_FN_HOST_ACC_INLINE auto allocView(Mapping mapping = {}, const Allocator& alloc = {})
-        -> View<Mapping, internal::AllocatorBlobType<Allocator, typename Mapping::RecordDim>>
+    LLAMA_FN_HOST_ACC_INLINE auto allocView(Mapping mapping = {}, const Allocator& alloc = {}, Accessor accessor = {})
+        -> View<Mapping, internal::AllocatorBlobType<Allocator, typename Mapping::RecordDim>, Accessor>
     {
-        auto view = allocViewUninitialized(std::move(mapping), alloc);
+        auto view = allocViewUninitialized(std::move(mapping), alloc, accessor);
         constructFields(view);
         return view;
     }
@@ -387,22 +393,26 @@ namespace llama
     /// view should be created using \ref allocView.
     /// \tparam TMapping The mapping used by the view to map accesses into memory.
     /// \tparam BlobType The storage type used by the view holding memory.
+    /// \tparam TAccessor The accessor to use when an access is made through this view.
 #ifdef __cpp_lib_concepts
-    template<typename TMapping, Blob BlobType>
+    template<typename TMapping, Blob BlobType, typename TAccessor = accessor::Default>
 #else
-    template<typename TMapping, typename BlobType>
+    template<typename TMapping, typename BlobType, typename TAccessor = accessor::Default>
 #endif
     struct LLAMA_DECLSPEC_EMPTY_BASES View
         : private TMapping
+        , private TAccessor
 #if CAN_USE_RANGES
         , std::ranges::view_base
 #endif
     {
         static_assert(!std::is_const_v<TMapping>);
+        static_assert(std::is_empty_v<TAccessor>, "Stateful accessors are not implemented");
         using Mapping = TMapping;
         using ArrayExtents = typename Mapping::ArrayExtents;
         using ArrayIndex = typename Mapping::ArrayIndex;
         using RecordDim = typename Mapping::RecordDim;
+        using Accessor = TAccessor;
         using iterator = Iterator<View>;
         using const_iterator = Iterator<const View>;
         using size_type = typename ArrayExtents::value_type;
@@ -437,6 +447,16 @@ namespace llama
             return static_cast<const Mapping&>(*this);
         }
 
+        LLAMA_FN_HOST_ACC_INLINE auto accessor() -> Accessor&
+        {
+            return static_cast<Accessor&>(*this);
+        }
+
+        LLAMA_FN_HOST_ACC_INLINE auto accessor() const -> const Accessor&
+        {
+            return static_cast<const Accessor&>(*this);
+        }
+
 #if !(defined(_MSC_VER) && defined(__NVCC__))
         template<typename V>
         auto operator()(llama::ArrayIndex<V, ArrayIndex::rank>) const
@@ -456,7 +476,7 @@ namespace llama
             else
             {
                 LLAMA_FORCE_INLINE_RECURSIVE
-                return accessor(ai, RecordCoord<>{});
+                return access(ai, RecordCoord<>{});
             }
         }
 
@@ -470,7 +490,7 @@ namespace llama
             else
             {
                 LLAMA_FORCE_INLINE_RECURSIVE
-                return accessor(ai, RecordCoord<>{});
+                return access(ai, RecordCoord<>{});
             }
         }
 
@@ -566,23 +586,23 @@ namespace llama
         friend struct RecordRef;
 
         template<std::size_t... Coords>
-        LLAMA_FN_HOST_ACC_INLINE auto accessor(ArrayIndex ai, RecordCoord<Coords...> rc = {}) const -> decltype(auto)
+        LLAMA_FN_HOST_ACC_INLINE auto access(ArrayIndex ai, RecordCoord<Coords...> rc = {}) const -> decltype(auto)
         {
-            return mapToMemory(mapping(), ai, rc, storageBlobs);
+            return accessor()(mapToMemory(mapping(), ai, rc, storageBlobs));
         }
 
         template<std::size_t... Coords>
-        LLAMA_FN_HOST_ACC_INLINE auto accessor(ArrayIndex ai, RecordCoord<Coords...> rc = {}) -> decltype(auto)
+        LLAMA_FN_HOST_ACC_INLINE auto access(ArrayIndex ai, RecordCoord<Coords...> rc = {}) -> decltype(auto)
         {
-            return mapToMemory(mapping(), ai, rc, storageBlobs);
+            return accessor()(mapToMemory(mapping(), ai, rc, storageBlobs));
         }
     };
 
     namespace internal
     {
-        template<typename Mapping, typename BlobType, typename TransformBlobFunc, std::size_t... Is>
+        template<typename Mapping, typename BlobType, typename Accessor, typename TransformBlobFunc, std::size_t... Is>
         LLAMA_FN_HOST_ACC_INLINE auto makeTransformedBlobArray(
-            View<Mapping, BlobType>& view,
+            View<Mapping, BlobType, Accessor>& view,
             const TransformBlobFunc& transformBlob,
             std::integer_sequence<std::size_t, Is...>)
         {
@@ -592,10 +612,12 @@ namespace llama
 
     /// Applies the given transformation to the blobs of a view and creates a new view with the transformed blobs and
     /// the same mapping as the old view.
-    template<typename Mapping, typename BlobType, typename TransformBlobFunc>
-    LLAMA_FN_HOST_ACC_INLINE auto transformBlobs(View<Mapping, BlobType>& view, const TransformBlobFunc& transformBlob)
+    template<typename Mapping, typename BlobType, typename Accessor, typename TransformBlobFunc>
+    LLAMA_FN_HOST_ACC_INLINE auto transformBlobs(
+        View<Mapping, BlobType, Accessor>& view,
+        const TransformBlobFunc& transformBlob)
     {
-        constexpr auto blobCount = View<Mapping, BlobType>::Mapping::blobCount;
+        constexpr auto blobCount = View<Mapping, BlobType, Accessor>::Mapping::blobCount;
         return View{
             view.mapping(),
             internal::makeTransformedBlobArray(view, transformBlob, std::make_index_sequence<blobCount>{})};
@@ -604,8 +626,8 @@ namespace llama
     /// Creates a shallow copy of a view. This copy must not outlive the view, since it references its blob array.
     /// \tparam NewBlobType The blob type of the shallow copy. Must be a non owning pointer like type.
     /// \return A new view with the same mapping as view, where each blob refers to the blob in view.
-    template<typename Mapping, typename BlobType, typename NewBlobType = std::byte*>
-    LLAMA_FN_HOST_ACC_INLINE auto shallowCopy(View<Mapping, BlobType>& view)
+    template<typename Mapping, typename BlobType, typename Accessor, typename NewBlobType = std::byte*>
+    LLAMA_FN_HOST_ACC_INLINE auto shallowCopy(View<Mapping, BlobType, Accessor>& view)
     {
         return transformBlobs(
             view,
@@ -617,11 +639,20 @@ namespace llama
             });
     }
 
+    // Creates a new view from an existing view with the given accessor.
+    // \param view A view which's mapping and blobs are copied into a new view with the different accessor. If you no
+    // longer need the old view, consider moving it into the argument of this function.
+    template<typename NewAccessor, typename Mapping, typename BlobType, typename OldAccessor>
+    LLAMA_FN_HOST_ACC_INLINE auto withAccessor(View<Mapping, BlobType, OldAccessor> view)
+    {
+        return View<Mapping, BlobType, NewAccessor>{std::move(view.mapping()), std::move(view.storageBlobs)};
+    }
+
     template<typename View>
     inline constexpr auto isView = false;
 
-    template<typename Mapping, typename BlobType>
-    inline constexpr auto isView<View<Mapping, BlobType>> = true;
+    template<typename Mapping, typename BlobType, typename Accessor>
+    inline constexpr auto isView<View<Mapping, BlobType, Accessor>> = true;
 
     /// Like a \ref View, but array indices are shifted.
     /// @tparam TStoredParentView Type of the underlying view. May be cv qualified and/or a reference type.
@@ -644,18 +675,6 @@ namespace llama
         {
         }
 
-        template<std::size_t... Coords>
-        LLAMA_FN_HOST_ACC_INLINE auto accessor(ArrayIndex ai) const -> const auto&
-        {
-            return parentView.template accessor<Coords...>(ArrayIndex{ai + offset});
-        }
-
-        template<std::size_t... Coords>
-        LLAMA_FN_HOST_ACC_INLINE auto accessor(ArrayIndex ai) -> auto&
-        {
-            return parentView.template accessor<Coords...>(ArrayIndex{ai + offset});
-        }
-
         /// Same as \ref View::operator()(ArrayIndex), but shifted by the offset of this \ref VirtualView.
         LLAMA_FN_HOST_ACC_INLINE auto operator()(ArrayIndex ai) const -> decltype(auto)
         {
@@ -698,18 +717,16 @@ namespace llama
                 ArrayIndex{ArrayIndex{static_cast<typename ArrayIndex::value_type>(indices)...} + offset});
         }
 
-        template<std::size_t... Coord>
-        LLAMA_FN_HOST_ACC_INLINE auto operator()(RecordCoord<Coord...> = {}) const -> decltype(auto)
+        template<std::size_t... Coords>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(RecordCoord<Coords...> rc = {}) const -> decltype(auto)
         {
-            LLAMA_FORCE_INLINE_RECURSIVE
-            return accessor<Coord...>(ArrayIndex{});
+            return parentView(ArrayIndex{} + offset, rc);
         }
 
-        template<std::size_t... Coord>
-        LLAMA_FN_HOST_ACC_INLINE auto operator()(RecordCoord<Coord...> = {}) -> decltype(auto)
+        template<std::size_t... Coords>
+        LLAMA_FN_HOST_ACC_INLINE auto operator()(RecordCoord<Coords...> rc = {}) -> decltype(auto)
         {
-            LLAMA_FORCE_INLINE_RECURSIVE
-            return accessor<Coord...>(ArrayIndex{});
+            return parentView(ArrayIndex{} + offset, rc);
         }
 
         StoredParentView parentView;
diff --git a/tests/view.cpp b/tests/view.cpp
index 4974a49603..ad1e5deae1 100644
--- a/tests/view.cpp
+++ b/tests/view.cpp
@@ -1,5 +1,6 @@
 #include "common.hpp"
 
+#include <atomic>
 #include <deque>
 
 // clang-format off
@@ -85,7 +86,7 @@ TEST_CASE("view.non-memory-owning")
     test(boost::mp11::mp_identity<const std::byte>{});
 }
 
-TEST_CASE("view.access")
+TEST_CASE("view.subscript")
 {
     using ArrayExtents = llama::ArrayExtentsDynamic<int, 2>;
     ArrayExtents extents{16, 16};
@@ -334,3 +335,92 @@ TEST_CASE("view.allocViewStack")
     auto v4 = llama::allocViewStack<4, Vec3I>();
     v4(llama::ArrayIndex{0, 0, 0, 0})(tag::X{}) = 42;
 }
+
+TEST_CASE("view.allocView.Default")
+{
+    auto mapping = llama::mapping::AoS{llama::ArrayExtents{3, 4}, Particle{}};
+    auto view = llama::allocView(mapping, llama::bloballoc::Vector{}, llama::accessor::Default{});
+    iotaFillView(view);
+    iotaCheckView(view);
+}
+
+TEST_CASE("view.allocView.ReadOnlyByValue")
+{
+    auto mapping = llama::mapping::AoS{llama::ArrayExtents{3, 4}, Vec3I{}};
+    auto view = llama::allocView(mapping, llama::bloballoc::Vector{}, llama::accessor::ReadOnlyByValue{});
+    STATIC_REQUIRE(std::is_same_v<decltype(view(1, 2)(tag::X{})), int>);
+}
+
+TEST_CASE("view.allocView.Const")
+{
+    auto mapping = llama::mapping::AoS{llama::ArrayExtents{3, 4}, Vec3I{}};
+    auto view = llama::allocView(mapping, llama::bloballoc::Vector{}, llama::accessor::Const{});
+    STATIC_REQUIRE(std::is_same_v<decltype(view(1, 2)(tag::X{})), const int&>);
+}
+
+#ifdef __cpp_lib_atomic_ref
+TEST_CASE("view.allocView.Atomic")
+{
+    auto mapping = llama::mapping::AoS{llama::ArrayExtents{3, 4}, Vec3I{}};
+    auto view = llama::allocView(mapping, llama::bloballoc::Vector{}, llama::accessor::Atomic{});
+    STATIC_REQUIRE(std::is_same_v<decltype(view(1, 2)(tag::X{})), std::atomic_ref<int>>);
+    iotaFillView(view);
+    iotaCheckView(view);
+}
+#endif
+
+TEST_CASE("view.withAccessor.Default.Vector")
+{
+    auto view
+        = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{16, 16}, Vec3I{}}, llama::bloballoc::Vector{});
+    auto* addr = &view(1, 2)(tag::X{});
+    auto view2 = llama::withAccessor<llama::accessor::Default>(view); // copies
+    auto* addr2 = &view2(1, 2)(tag::X{});
+    CHECK(addr != addr2);
+}
+
+TEST_CASE("view.withAccessor.Default.Vector.move")
+{
+    auto view
+        = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{16, 16}, Vec3I{}}, llama::bloballoc::Vector{});
+    auto* addr = &view(1, 2)(tag::X{});
+    auto view2 = llama::withAccessor<llama::accessor::Default>(std::move(view));
+    auto* addr2 = &view2(1, 2)(tag::X{});
+    CHECK(addr == addr2);
+}
+
+TEST_CASE("view.withAccessor.Default.SharedPtr")
+{
+    auto view
+        = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{16, 16}, Vec3I{}}, llama::bloballoc::SharedPtr{});
+    auto* addr = &view(1, 2)(tag::X{});
+    auto view2 = llama::withAccessor<llama::accessor::Default>(view); // copies shared pointers, but not memory chunks
+    auto* addr2 = &view2(1, 2)(tag::X{});
+    CHECK(addr == addr2);
+}
+
+TEMPLATE_TEST_CASE("view.withAccessor.shallowCopy.Default", "", llama::bloballoc::Vector, llama::bloballoc::SharedPtr)
+{
+    auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{16, 16}, Particle{}}, TestType{});
+    auto view2 = llama::withAccessor<llama::accessor::Default>(llama::shallowCopy(view));
+    iotaFillView(view2);
+    iotaCheckView(view);
+}
+
+#ifdef __cpp_lib_atomic_ref
+TEMPLATE_TEST_CASE("view.withAccessor.shallowCopy.Atomic", "", llama::bloballoc::Vector, llama::bloballoc::SharedPtr)
+{
+    auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{16, 16}, Particle{}}, TestType{});
+    auto view2 = llama::withAccessor<llama::accessor::Atomic>(llama::shallowCopy(view));
+    iotaFillView(view2);
+    iotaCheckView(view);
+}
+#endif
+
+TEMPLATE_TEST_CASE("view.withAccessor.shallowCopy.Restrict", "", llama::bloballoc::Vector, llama::bloballoc::SharedPtr)
+{
+    auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{16, 16}, Particle{}}, TestType{});
+    auto view2 = llama::withAccessor<llama::accessor::Restrict>(llama::shallowCopy(view));
+    iotaFillView(view2);
+    iotaCheckView(view);
+}
\ No newline at end of file

From c2b0b97706acd39eedd8f99de907c393a2722d92 Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 7 Nov 2022 11:27:06 +0100
Subject: [PATCH 3/4] Disable some example for nvcc 11.3

---
 examples/alpaka/asyncblur/asyncblur.cpp |  7 +++++++
 examples/alpaka/nbody/nbody.cpp         |  1 -
 examples/alpaka/vectoradd/vectoradd.cpp |  8 ++++++++
 examples/cuda/pitch/pitch.cu            | 11 +++++++++--
 4 files changed, 24 insertions(+), 3 deletions(-)

diff --git a/examples/alpaka/asyncblur/asyncblur.cpp b/examples/alpaka/asyncblur/asyncblur.cpp
index a34627e516..5301098d44 100644
--- a/examples/alpaka/asyncblur/asyncblur.cpp
+++ b/examples/alpaka/asyncblur/asyncblur.cpp
@@ -143,6 +143,12 @@ struct BlurKernel
 auto main(int argc, char** argv) -> int
 try
 {
+#if defined(__NVCC__) && __CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ >= 3 && __CUDACC_VER_MINOR__ < 4
+// nvcc 11.3 fails to generate the template signature for llama::View, if it has a forward declaration with a default
+// argument (which we need for the default accessor)
+#    warning "alpaka nbody example disabled for nvcc 11.3, because it generates invalid C++ code for the host compiler"
+    return -1;
+#else
     // ALPAKA
     using Dim = alpaka::DimInt<2>;
 
@@ -390,6 +396,7 @@ try
     }
 
     return 0;
+#endif
 }
 catch(const std::exception& e)
 {
diff --git a/examples/alpaka/nbody/nbody.cpp b/examples/alpaka/nbody/nbody.cpp
index f83d5282fe..8bfc510781 100644
--- a/examples/alpaka/nbody/nbody.cpp
+++ b/examples/alpaka/nbody/nbody.cpp
@@ -317,7 +317,6 @@ try
 #    warning "alpaka nbody example disabled for nvcc <= 11.5, because the compiler segfaults"
     return -1;
 #else
-
     std::cout << problemSize / 1000 << "k particles (" << problemSize * llama::sizeOf<Particle> / 1024 << "kiB)\n"
               << "Caching " << threadsPerBlock << " particles (" << threadsPerBlock * llama::sizeOf<Particle> / 1024
               << " kiB) in shared memory\n"
diff --git a/examples/alpaka/vectoradd/vectoradd.cpp b/examples/alpaka/vectoradd/vectoradd.cpp
index 56123fb055..7334f16fa8 100644
--- a/examples/alpaka/vectoradd/vectoradd.cpp
+++ b/examples/alpaka/vectoradd/vectoradd.cpp
@@ -61,6 +61,13 @@ struct AddKernel
 auto main() -> int
 try
 {
+#if defined(__NVCC__) && __CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ >= 3 && __CUDACC_VER_MINOR__ < 4
+// nvcc 11.3 fails to generate the template signature for llama::View, if it has a forward declaration with a default
+// argument (which we need for the default accessor)
+#    warning "alpaka nbody example disabled for nvcc 11.3, because it generates invalid C++ code for the host compiler"
+    return -1;
+#else
+
     // ALPAKA
     using Dim = alpaka::DimInt<1>;
     using Size = std::size_t;
@@ -160,6 +167,7 @@ try
     chrono.printAndReset("Copy D->H");
 
     return 0;
+#endif
 }
 catch(const std::exception& e)
 {
diff --git a/examples/cuda/pitch/pitch.cu b/examples/cuda/pitch/pitch.cu
index 680c9bd783..7c9e4bd0df 100644
--- a/examples/cuda/pitch/pitch.cu
+++ b/examples/cuda/pitch/pitch.cu
@@ -124,6 +124,12 @@ namespace llamaex
 auto main() -> int
 try
 {
+#if defined(__NVCC__) && __CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ >= 3 && __CUDACC_VER_MINOR__ < 4
+// nvcc 11.3 fails to generate the template signature for llama::View, if it has a forward declaration with a default
+// argument (which we need for the default accessor)
+#    warning "alpaka nbody example disabled for nvcc 11.3, because it generates invalid C++ code for the host compiler"
+    return -1;
+#else
     int device = 0;
     checkError(cudaGetDevice(&device));
     cudaDeviceProp prop{};
@@ -164,7 +170,7 @@ try
     }
 
     // nvcc 11.3 fails to compile the AoS mapping here
-#if !(defined(__NVCC__) && __CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ == 3)
+#    if !(defined(__NVCC__) && __CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ == 3)
     {
         std::byte* mem = nullptr;
         checkError(cudaMalloc(&mem, widthBytes * extents[0]));
@@ -179,12 +185,13 @@ try
 
         stbi_write_png("pitch2.png", extents[1], extents[0], 3, host2.data(), 0);
     }
-#endif
+#    endif
 
     if(host1 != host2)
         fmt::print("ERROR: produced two different images");
 
     return 0;
+#endif
 }
 catch(const std::exception& e)
 {

From 64962f1c1009a42d96c5bf5ee390a84a6e1c65ec Mon Sep 17 00:00:00 2001
From: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
Date: Mon, 7 Nov 2022 15:59:07 +0100
Subject: [PATCH 4/4] Ensure clang-format action can push back to correct
 branch

---
 .github/workflows/ci.yaml | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/.github/workflows/ci.yaml b/.github/workflows/ci.yaml
index c1908806c8..44260497e3 100644
--- a/.github/workflows/ci.yaml
+++ b/.github/workflows/ci.yaml
@@ -19,6 +19,9 @@ jobs:
     runs-on: ubuntu-22.04
     steps:
     - uses: actions/checkout@v3
+      with:
+        repository: ${{ github.event.pull_request.head.repo.full_name }}
+        ref: ${{ github.event.pull_request.head.ref }}
     - uses: DoozyX/clang-format-lint-action@v0.14
       with:
         exclude: './thirdparty'