diff --git a/examples/alpaka/asyncblur/asyncblur.cpp b/examples/alpaka/asyncblur/asyncblur.cpp index 622b0e1aea..fc810ed6b2 100644 --- a/examples/alpaka/asyncblur/asyncblur.cpp +++ b/examples/alpaka/asyncblur/asyncblur.cpp @@ -127,7 +127,7 @@ struct BlurKernel LLAMA_INDEPENDENT_DATA for (auto x = start[1]; x < end[1]; ++x) { - auto sum = llama::allocVirtualDatumStack(); + llama::One sum; sum = 0; using ItType = long int; diff --git a/examples/alpaka/nbody/nbody.cpp b/examples/alpaka/nbody/nbody.cpp index ea9e7ccdbe..e9c913c532 100644 --- a/examples/alpaka/nbody/nbody.cpp +++ b/examples/alpaka/nbody/nbody.cpp @@ -217,15 +217,15 @@ int main() LLAMA_INDEPENDENT_DATA for (std::size_t i = 0; i < PROBLEM_SIZE; ++i) { - auto temp = llama::allocVirtualDatumStack(); - temp(tag::Pos(), tag::X()) = distribution(generator); - temp(tag::Pos(), tag::Y()) = distribution(generator); - temp(tag::Pos(), tag::Z()) = distribution(generator); - temp(tag::Vel(), tag::X()) = distribution(generator) / FP(10); - temp(tag::Vel(), tag::Y()) = distribution(generator) / FP(10); - temp(tag::Vel(), tag::Z()) = distribution(generator) / FP(10); - temp(tag::Mass()) = distribution(generator) / FP(100); - hostView(i) = temp; + llama::One p; + p(tag::Pos(), tag::X()) = distribution(generator); + p(tag::Pos(), tag::Y()) = distribution(generator); + p(tag::Pos(), tag::Z()) = distribution(generator); + p(tag::Vel(), tag::X()) = distribution(generator) / FP(10); + p(tag::Vel(), tag::Y()) = distribution(generator) / FP(10); + p(tag::Vel(), tag::Z()) = distribution(generator) / FP(10); + p(tag::Mass()) = distribution(generator) / FP(100); + hostView(i) = p; } chrono.printAndReset("Init"); diff --git a/examples/cuda/nbody/nbody.cu b/examples/cuda/nbody/nbody.cu index 05dab1ba02..ee5d6b50ae 100644 --- a/examples/cuda/nbody/nbody.cu +++ b/examples/cuda/nbody/nbody.cu @@ -103,7 +103,7 @@ __global__ void updateSM(View particles) const auto ti = threadIdx.x + blockIdx.x * blockDim.x; const auto tbi = blockIdx.x; - auto pi = llama::allocVirtualDatumStack(); + llama::One pi; if constexpr (UseAccumulator) pi = particles(ti); for (std::size_t blockOffset = 0; blockOffset < ProblemSize; blockOffset += BlockSize) @@ -132,7 +132,7 @@ __global__ void update(View particles) { const auto ti = threadIdx.x + blockIdx.x * blockDim.x; - auto pi = llama::allocVirtualDatumStack(); + llama::One pi; if constexpr (UseAccumulator) pi = particles(ti); LLAMA_INDEPENDENT_DATA @@ -218,15 +218,15 @@ try std::normal_distribution distribution(FP(0), FP(1)); for (std::size_t i = 0; i < PROBLEM_SIZE; ++i) { - auto temp = llama::allocVirtualDatumStack(); - temp(tag::Pos(), tag::X()) = distribution(generator); - temp(tag::Pos(), tag::Y()) = distribution(generator); - temp(tag::Pos(), tag::Z()) = distribution(generator); - temp(tag::Vel(), tag::X()) = distribution(generator) / FP(10); - temp(tag::Vel(), tag::Y()) = distribution(generator) / FP(10); - temp(tag::Vel(), tag::Z()) = distribution(generator) / FP(10); - temp(tag::Mass()) = distribution(generator) / FP(100); - hostView(i) = temp; + llama::One p; + p(tag::Pos(), tag::X()) = distribution(generator); + p(tag::Pos(), tag::Y()) = distribution(generator); + p(tag::Pos(), tag::Z()) = distribution(generator); + p(tag::Vel(), tag::X()) = distribution(generator) / FP(10); + p(tag::Vel(), tag::Y()) = distribution(generator) / FP(10); + p(tag::Vel(), tag::Z()) = distribution(generator) / FP(10); + p(tag::Mass()) = distribution(generator) / FP(100); + hostView(i) = p; } watch.printAndReset("init"); diff --git a/examples/nbody/nbody.cpp b/examples/nbody/nbody.cpp index ee8dd41079..e4fb515e1f 100644 --- a/examples/nbody/nbody.cpp +++ b/examples/nbody/nbody.cpp @@ -81,7 +81,7 @@ namespace usellama LLAMA_INDEPENDENT_DATA for (std::size_t i = 0; i < PROBLEM_SIZE; i++) { - auto pi = llama::allocVirtualDatumStack(); + llama::One pi; if constexpr (UseAccumulator) pi = particles(i); LLAMA_INDEPENDENT_DATA diff --git a/include/llama/View.hpp b/include/llama/View.hpp index 8afe48ccae..0837511456 100644 --- a/include/llama/View.hpp +++ b/include/llama/View.hpp @@ -54,12 +54,6 @@ namespace llama return allocView(Mapping{}, llama::allocator::Stack>{}); } - template - inline constexpr auto IsView = false; - - template - inline constexpr auto IsView> = true; - template , bool OwnView = false> struct VirtualDatum; @@ -69,20 +63,16 @@ namespace llama template inline constexpr auto is_VirtualDatum> = true; - /// Creates a single \ref VirtualDatum owning a view with stack memory. + /// A \ref VirtualDatum that owns and holds a single value. template - LLAMA_FN_HOST_ACC_INLINE auto allocVirtualDatumStack() - -> VirtualDatum()), DatumCoord<>, true> - { - return {ArrayDomain<1>{}, llama::allocViewStack<1, DatumDomain>()}; - } + using One = VirtualDatum()), DatumCoord<>, true>; /// Creates a single \ref VirtualDatum owning a view with stack memory and /// copies all values from an existing \ref VirtualDatum. template LLAMA_FN_HOST_ACC_INLINE auto copyVirtualDatumStack(const VirtualDatum& vd) -> decltype(auto) { - auto temp = allocVirtualDatumStack(); + One temp; temp = vd; return temp; } @@ -259,6 +249,14 @@ namespace llama /// AccessibleDatumDomain is the same as `Mapping::DatumDomain`. using AccessibleDatumDomain = GetType; + LLAMA_FN_HOST_ACC_INLINE VirtualDatum() + /* requires(OwnView) */ + : userDomainPos({}) + , view{allocViewStack<1, DatumDomain>()} + { + static_assert(OwnView, "The default constructor of VirtualDatum is only available if the "); + } + LLAMA_FN_HOST_ACC_INLINE VirtualDatum(ArrayDomain userDomainPos, std::conditional_t view) : userDomainPos(userDomainPos) @@ -619,6 +617,12 @@ namespace llama } }; + template + inline constexpr auto IsView = false; + + template + inline constexpr auto IsView> = true; + /// Acts like a \ref View, but shows only a smaller and/or shifted part of /// another view it references, the parent view. template diff --git a/tests/view.cpp b/tests/view.cpp index 2b3ccecc62..94de79a1db 100644 --- a/tests/view.cpp +++ b/tests/view.cpp @@ -182,7 +182,7 @@ TEST_CASE("view.assign-one-datum") Mapping mapping{arrayDomain}; auto view = allocView(mapping); - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; datum(tag::Pos{}, tag::X{}) = 14.0f; datum(tag::Pos{}, tag::Y{}) = 15.0f; datum(tag::Pos{}, tag::Z{}) = 16.0f; diff --git a/tests/virtualdatum.cpp b/tests/virtualdatum.cpp index 9241e8b3cf..4100429eff 100644 --- a/tests/virtualdatum.cpp +++ b/tests/virtualdatum.cpp @@ -29,7 +29,7 @@ using Name = llama::DS< TEST_CASE("VirtualDatum.operator=") { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; // scalar to multiple elements in virtual datum datum(tag::Pos{}) = 1; @@ -84,7 +84,7 @@ namespace { auto allocVc() { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; datum(tag::Pos{}, tag::A{}) = 1; datum(tag::Pos{}, tag::Y{}) = 2; datum(tag::Vel{}, tag::X{}) = 3; @@ -267,7 +267,7 @@ using Name2 = llama::DS< TEST_CASE("VirtualDatum.operator=.propagation") { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; datum(tag::Part1{}) = 1; datum(tag::Part2{}) = 2; @@ -295,8 +295,8 @@ TEST_CASE("VirtualDatum.operator=.propagation") TEST_CASE("VirtualDatum.operator=.multiview") { - auto datum1 = llama::allocVirtualDatumStack(); - auto datum2 = llama::allocVirtualDatumStack(); + llama::One datum1; + llama::One datum2; datum2 = 1; datum1 = datum2; @@ -318,7 +318,7 @@ TEST_CASE("VirtualDatum.operator=.multiview") TEST_CASE("VirtualDatum.operator==") { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; datum = 1; @@ -342,7 +342,7 @@ TEST_CASE("VirtualDatum.operator==") TEST_CASE("VirtualDatum.operator<") { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; datum = 1; @@ -380,14 +380,14 @@ TEST_CASE("VirtualDatum.operator<") TEST_CASE("VirtualDatum.asTuple.types") { { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; std::tuple pos = datum(tag::Pos{}).asTuple(); std::tuple vel = datum(tag::Vel{}).asTuple(); std::tuple name = datum.asTuple(); } { - const auto datum = llama::allocVirtualDatumStack(); + const llama::One datum; std::tuple pos = datum(tag::Pos{}).asTuple(); std::tuple vel = datum(tag::Vel{}).asTuple(); @@ -397,7 +397,7 @@ TEST_CASE("VirtualDatum.asTuple.types") TEST_CASE("VirtualDatum.asTuple.assign") { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; datum(tag::Pos{}).asTuple() = std::tuple{1, 1}; CHECK(datum(tag::Pos{}, tag::A{}) == 1); @@ -426,7 +426,7 @@ TEST_CASE("VirtualDatum.asTuple.assign") TEST_CASE("VirtualDatum.asTuple.structuredBindings") { - auto datum = llama::allocVirtualDatumStack(); + llama::One datum; { auto [a, y] = datum(tag::Pos{}).asTuple();