Add accessors to Views

.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'

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)
 {

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"

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)
 {

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)
 {

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

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{});
             }
         }