fix nvcc errors

alpaka-group · Jul 2, 2021 · 708dda4 · 708dda4
1 parent d08e4d2
commit 708dda4
Show file tree

Hide file tree

Showing 4 changed files with 18 additions and 17 deletions.
diff --git a/examples/alpaka/asyncblur/asyncblur.cpp b/examples/alpaka/asyncblur/asyncblur.cpp
@@ -210,11 +210,12 @@ try
     using ArrayDims = llama::ArrayDims<2>;
 
     auto treeOperationList = llama::Tuple{llama::mapping::tree::functor::LeafOnlyRT()};
-    const auto hostMapping = llama::mapping::tree::Mapping{ArrayDims{buffer_y, buffer_x}, treeOperationList, Pixel{}};
-    const auto devMapping = llama::mapping::tree::Mapping{
-        ArrayDims{CHUNK_SIZE + 2 * KERNEL_SIZE, CHUNK_SIZE + 2 * KERNEL_SIZE},
-        treeOperationList,
-        PixelOnAcc{}};
+    const auto hostMapping = llama::mapping::tree::Mapping<ArrayDims, Pixel, decltype(treeOperationList)>{
+        {buffer_y, buffer_x},
+        treeOperationList};
+    const auto devMapping = llama::mapping::tree::Mapping<ArrayDims, PixelOnAcc, decltype(treeOperationList)>{
+        {CHUNK_SIZE + 2 * KERNEL_SIZE, CHUNK_SIZE + 2 * KERNEL_SIZE},
+        treeOperationList};
 
     const auto hostBufferSize = hostMapping.blobSize(0);
     const auto devBufferSize = devMapping.blobSize(0);

diff --git a/examples/alpaka/nbody/nbody.cpp b/examples/alpaka/nbody/nbody.cpp
@@ -158,7 +158,7 @@ struct UpdateKernel
             {
                 constexpr auto sharedMapping = []
                 {
-                    constexpr auto arrayDims = llama::ArrayDims{BlockSize};
+                    constexpr auto arrayDims = llama::ArrayDims<1>{BlockSize};
                     if constexpr (MappingSM == AoS)
                         return llama::mapping::AoS{arrayDims, Particle{}};
                     if constexpr (MappingSM == SoA)
@@ -180,7 +180,7 @@ struct UpdateKernel
         // TODO: we could optimize here, because only velocity is ever updated
         auto pi = [&]
         {
-            constexpr auto arrayDims = llama::ArrayDims{Elems};
+            constexpr auto arrayDims = llama::ArrayDims<1>{Elems};
             constexpr auto mapping
                 = llama::mapping::SoA<typename View::ArrayDims, typename View::RecordDim, false>{arrayDims};
             constexpr auto blobAlloc = llama::bloballoc::Stack<llama::sizeOf<typename View::RecordDim> * Elems>{};
@@ -266,7 +266,7 @@ void run(std::ostream& plotFile)
     {
         const auto arrayDims = llama::ArrayDims{PROBLEM_SIZE};
         if constexpr (MappingGM == AoS)
-            return llama::mapping::AoS{arrayDims, Particle{}};
+            return llama::mapping::AoS<decltype(arrayDims), Particle>{arrayDims};
         if constexpr (MappingGM == SoA)
             return llama::mapping::SoA<decltype(arrayDims), Particle, false>{arrayDims};
         // if constexpr (MappingGM == 2)
@@ -284,8 +284,8 @@ void run(std::ostream& plotFile)
 
     watch.printAndReset("alloc");
 
-    auto hostView = llama::View{mapping, llama::Array{alpaka::getPtrNative(hostBuffer)}};
-    auto accView = llama::View{mapping, llama::Array{alpaka::getPtrNative(accBuffer)}};
+    auto hostView = llama::View(mapping, llama::Array{alpaka::getPtrNative(hostBuffer)});
+    auto accView = llama::View(mapping, llama::Array{alpaka::getPtrNative(accBuffer)});
 
     watch.printAndReset("views");
 

diff --git a/examples/alpaka/vectoradd/vectoradd.cpp b/examples/alpaka/vectoradd/vectoradd.cpp
@@ -84,9 +84,9 @@ try
     {
         const auto arrayDims = llama::ArrayDims{PROBLEM_SIZE};
         if constexpr (MAPPING == 0)
-            return llama::mapping::AoS{arrayDims, Vector{}};
+            return llama::mapping::AoS<decltype(arrayDims), Vector>{arrayDims};
         if constexpr (MAPPING == 1)
-            return llama::mapping::SoA<decltype(arrayDims), Vector, false>{arrayDims, Vector{}};
+            return llama::mapping::SoA<decltype(arrayDims), Vector, false>{arrayDims};
         if constexpr (MAPPING == 2)
             return llama::mapping::SoA<decltype(arrayDims), Vector, true>{arrayDims};
         if constexpr (MAPPING == 3)

diff --git a/examples/cuda/nbody/nbody.cu b/examples/cuda/nbody/nbody.cu
@@ -88,9 +88,9 @@ __global__ void updateSM(View particles)
         {
             constexpr auto arrayDims = llama::ArrayDims{SHARED_ELEMENTS_PER_BLOCK};
             if constexpr (MappingSM == 0)
-                return llama::mapping::AoS{arrayDims, SharedMemoryParticle{}};
+                return llama::mapping::AoS<decltype(arrayDims), SharedMemoryParticle>{arrayDims};
             if constexpr (MappingSM == 1)
-                return llama::mapping::SoA{arrayDims, SharedMemoryParticle{}};
+                return llama::mapping::SoA<decltype(arrayDims), SharedMemoryParticle, false>{arrayDims};
             if constexpr (MappingSM == 2)
                 return llama::mapping::SoA<decltype(arrayDims), SharedMemoryParticle, true>{arrayDims};
             if constexpr (MappingSM == 3)
@@ -180,11 +180,11 @@ try
     {
         const auto arrayDims = llama::ArrayDims{PROBLEM_SIZE};
         if constexpr (Mapping == 0)
-            return llama::mapping::AoS{arrayDims, Particle{}};
+            return llama::mapping::AoS<decltype(arrayDims), Particle>{arrayDims};
         if constexpr (Mapping == 1)
-            return llama::mapping::SoA{arrayDims, Particle{}};
+            return llama::mapping::SoA<decltype(arrayDims), Particle, false>{arrayDims};
         if constexpr (Mapping == 2)
-            return llama::mapping::SoA<decltype(arrayDims), Particle, true>{arrayDims, Particle{}};
+            return llama::mapping::SoA<decltype(arrayDims), Particle, true>{arrayDims};
         if constexpr (Mapping == 3)
             return llama::mapping::AoSoA<decltype(arrayDims), Particle, AOSOA_LANES>{arrayDims};
         if constexpr (Mapping == 4)