From da970d8919221861eee51cdad9da93f09f680b10 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 22 Sep 2021 18:52:47 +0200 Subject: [PATCH] support extents with arbitrary value types --- CMakeLists.txt | 2 +- docs/pages/api.rst | 2 +- examples/alpaka/asyncblur/asyncblur.cpp | 8 +- examples/alpaka/nbody/nbody.cpp | 28 +-- examples/alpaka/pic/pic.cpp | 2 +- examples/alpaka/vectoradd/vectoradd.cpp | 2 +- examples/bitpackfloat/bitpackfloat.cpp | 2 +- examples/bitpackint/bitpackint.cpp | 5 +- examples/bufferguard/bufferguard.cpp | 27 +-- examples/bytesplit/bytesplit.cpp | 2 +- examples/cuda/nbody/nbody.cu | 20 +- examples/cuda/pitch/pitch.cu | 6 +- examples/heatequation/heatequation.cpp | 4 +- examples/nbody/nbody.cpp | 2 +- examples/nbody_benchmark/nbody.cpp | 2 +- examples/raycast/raycast.cpp | 2 +- examples/vectoradd/vectoradd.cpp | 2 +- include/llama/ArrayExtents.hpp | 211 ++++++++++++------ include/llama/ArrayIndexRange.hpp | 3 +- include/llama/Concepts.hpp | 4 +- include/llama/Copy.hpp | 10 +- include/llama/Core.hpp | 30 ++- include/llama/DumpMapping.hpp | 25 ++- include/llama/Vector.hpp | 27 +-- include/llama/View.hpp | 43 ++-- include/llama/mapping/AoS.hpp | 15 +- include/llama/mapping/AoSoA.hpp | 21 +- include/llama/mapping/BitPackedFloatSoA.hpp | 24 +- include/llama/mapping/BitPackedIntSoA.hpp | 26 ++- include/llama/mapping/ChangeType.hpp | 2 +- include/llama/mapping/Common.hpp | 44 ++-- include/llama/mapping/Heatmap.hpp | 2 +- include/llama/mapping/Null.hpp | 3 +- include/llama/mapping/One.hpp | 10 +- include/llama/mapping/SoA.hpp | 21 +- include/llama/mapping/Split.hpp | 27 ++- include/llama/mapping/Trace.hpp | 6 +- include/llama/mapping/tree/Mapping.hpp | 14 +- .../llama/mapping/tree/TreeFromDimensions.hpp | 9 +- tests/arrayextents.cpp | 174 ++++++++++----- tests/arrayindexrange.cpp | 49 ++-- tests/bloballocators.cpp | 2 +- tests/common.hpp | 3 + tests/computedprop.cpp | 8 +- tests/copy.cpp | 2 +- tests/dump.cpp | 2 +- tests/iterator.cpp | 4 +- tests/mapping.AoS.cpp | 40 ++-- tests/mapping.AoSoA.cpp | 12 +- tests/mapping.BitPackedFloatSoA.cpp | 36 +-- tests/mapping.BitPackedIntSoA.cpp | 29 ++- tests/mapping.Bytesplit.cpp | 16 +- tests/mapping.ChangeType.cpp | 19 +- tests/mapping.HeatmapTrace.cpp | 12 +- tests/mapping.Null.cpp | 2 +- tests/mapping.One.cpp | 24 +- tests/mapping.SoA.cpp | 99 ++++---- tests/mapping.Split.cpp | 8 +- tests/mapping.Tree.cpp | 28 +-- tests/mapping.cpp | 78 ++++--- tests/proofs.cpp | 31 +-- tests/proxyrefopmixin.cpp | 13 +- tests/recorddimension.cpp | 8 +- tests/vector.cpp | 18 +- tests/view.cpp | 26 +-- tests/virtualrecord.cpp | 16 +- tests/virtualview.cpp | 14 +- 67 files changed, 863 insertions(+), 605 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index ab016bc060..ac0c4270a6 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -41,7 +41,7 @@ if (BUILD_TESTING) target_compile_options(tests PRIVATE /permissive- /constexpr:steps10000000 /diagnostics:caret) else() target_compile_features(tests PRIVATE cxx_std_20) - target_compile_options(tests PRIVATE -Wall -Wextra -Wno-missing-braces) + target_compile_options(tests PRIVATE -Wall -Wextra -Werror=narrowing -Wno-missing-braces) endif() if (CMAKE_CXX_COMPILER_ID STREQUAL "Clang" OR CMAKE_CXX_COMPILER_ID STREQUAL "AppleClang" OR CMAKE_CXX_COMPILER_ID STREQUAL "IntelLLVM") target_compile_options(tests PRIVATE -fconstexpr-steps=10000000) diff --git a/docs/pages/api.rst b/docs/pages/api.rst index dc121ba46c..835cf3e928 100644 --- a/docs/pages/api.rst +++ b/docs/pages/api.rst @@ -45,7 +45,7 @@ Array dimensions .. doxygenstruct:: llama::ArrayExtents .. doxygentypedef:: llama::ArrayExtentsDynamic -.. doxygentypedef:: llama::ArrayExtentsStatic +.. doxygentypedef:: llama::ArrayExtentsNCube .. doxygenstruct:: llama::ArrayIndex .. doxygenstruct:: llama::ArrayIndexIterator diff --git a/examples/alpaka/asyncblur/asyncblur.cpp b/examples/alpaka/asyncblur/asyncblur.cpp index d9dbcb29a4..3855bdcf1b 100644 --- a/examples/alpaka/asyncblur/asyncblur.cpp +++ b/examples/alpaka/asyncblur/asyncblur.cpp @@ -85,8 +85,10 @@ struct BlurKernel { // Using SoA for the shared memory constexpr auto sharedChunkSize = ElemsPerBlock + 2 * KernelSize; - constexpr auto sharedMapping = llama::mapping:: - SoA, typename View::RecordDim, false>{}; + constexpr auto sharedMapping = llama::mapping::SoA< + llama::ArrayExtents, + typename View::RecordDim, + false>{}; auto& sharedMem = alpaka::declareSharedVar(acc); return llama::View(sharedMapping, llama::Array{&sharedMem[0]}); } @@ -211,7 +213,7 @@ try const auto hostMapping = llama::mapping::tree::Mapping{llama::ArrayExtents{buffer_y, buffer_x}, treeOperationList, Pixel{}}; const auto devMapping = llama::mapping::tree::Mapping{ - llama::ArrayExtents{}, + llama::ArrayExtents{}, treeOperationList, PixelOnAcc{}}; using DevMapping = std::decay_t; diff --git a/examples/alpaka/nbody/nbody.cpp b/examples/alpaka/nbody/nbody.cpp index 68a09702ef..75fc24c898 100644 --- a/examples/alpaka/nbody/nbody.cpp +++ b/examples/alpaka/nbody/nbody.cpp @@ -102,7 +102,7 @@ LLAMA_FN_HOST_ACC_INLINE auto store(FP& dst, Vec v) v.store(&dst); } -template +template struct VecType { // TODO(bgruber): we need a vector type that also works on GPUs @@ -116,7 +116,7 @@ struct VecType<1> using type = FP; }; -template +template LLAMA_FN_HOST_ACC_INLINE void pPInteraction(ViewParticleI pi, VirtualParticleJ pj) { using Vec = typename VecType::type; @@ -143,7 +143,7 @@ LLAMA_FN_HOST_ACC_INLINE void pPInteraction(ViewParticleI pi, VirtualParticleJ p store(pi(tag::Vel{}, tag::Z{}), zdistanceSqr * sts + load(pi(tag::Vel{}, tag::Z{}))); } -template +template struct UpdateKernel { template @@ -158,7 +158,7 @@ struct UpdateKernel { constexpr auto sharedMapping = [] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtents; if constexpr(MappingSM == AoS) return llama::mapping::AoS{}; if constexpr(MappingSM == SoA) @@ -181,7 +181,7 @@ struct UpdateKernel auto pi = [&] { constexpr auto mapping - = llama::mapping::SoA, typename View::RecordDim, false>{}; + = llama::mapping::SoA, typename View::RecordDim, false>{}; return llama::allocViewUninitialized(mapping, llama::bloballoc::Stack{}); }(); // TODO(bgruber): vector load @@ -190,26 +190,26 @@ struct UpdateKernel pi(e) = particles(ti * Elems + e); LLAMA_INDEPENDENT_DATA - for(std::size_t blockOffset = 0; blockOffset < ProblemSize; blockOffset += BlockSize) + for(int blockOffset = 0; blockOffset < ProblemSize; blockOffset += BlockSize) { LLAMA_INDEPENDENT_DATA - for(auto j = tbi; j < BlockSize; j += THREADS_PER_BLOCK) + for(int j = tbi; j < BlockSize; j += THREADS_PER_BLOCK) sharedView(j) = particles(blockOffset + j); alpaka::syncBlockThreads(acc); LLAMA_INDEPENDENT_DATA - for(auto j = std::size_t{0}; j < BlockSize; ++j) + for(int j = 0; j < BlockSize; ++j) pPInteraction(pi(0u), sharedView(j)); alpaka::syncBlockThreads(acc); } // TODO(bgruber): vector store LLAMA_INDEPENDENT_DATA - for(auto e = 0u; e < Elems; e++) + for(int e = 0u; e < Elems; e++) particles(ti * Elems + e) = pi(e); } }; -template +template struct MoveKernel { template @@ -235,7 +235,7 @@ template typename AccTemplate, Mapping MappingGM, M void run(std::ostream& plotFile) { using Dim = alpaka::DimInt<1>; - using Size = std::size_t; + using Size = int; using Acc = AccTemplate; using DevHost = alpaka::DevCpu; using DevAcc = alpaka::Dev; @@ -262,7 +262,7 @@ void run(std::ostream& plotFile) auto mapping = [] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtentsDynamic<1, int>; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(MappingGM == AoS) return llama::mapping::AoS{extents}; @@ -290,7 +290,7 @@ void run(std::ostream& plotFile) std::mt19937_64 generator; std::normal_distribution distribution(FP(0), FP(1)); - for(std::size_t i = 0; i < PROBLEM_SIZE; ++i) + for(int i = 0; i < PROBLEM_SIZE; ++i) { llama::One p; p(tag::Pos(), tag::X()) = distribution(generator); @@ -315,7 +315,7 @@ void run(std::ostream& plotFile) double sumUpdate = 0; double sumMove = 0; - for(std::size_t s = 0; s < STEPS; ++s) + for(int s = 0; s < STEPS; ++s) { auto updateKernel = UpdateKernel{}; alpaka::exec(queue, workdiv, updateKernel, accView); diff --git a/examples/alpaka/pic/pic.cpp b/examples/alpaka/pic/pic.cpp index f2b43d3a32..a778427212 100644 --- a/examples/alpaka/pic/pic.cpp +++ b/examples/alpaka/pic/pic.cpp @@ -309,7 +309,7 @@ auto setup(Queue& queue, const Dev& dev, const DevHost& devHost) auto particleMapping = [&] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtentsDynamic<1>; const auto particleExtents = ArrayExtents{numpart}; if constexpr(ParticleMapping == 0) return llama::mapping::AoS{particleExtents}; diff --git a/examples/alpaka/vectoradd/vectoradd.cpp b/examples/alpaka/vectoradd/vectoradd.cpp index 695faae70c..6677ee3bd6 100644 --- a/examples/alpaka/vectoradd/vectoradd.cpp +++ b/examples/alpaka/vectoradd/vectoradd.cpp @@ -82,7 +82,7 @@ try // LLAMA const auto mapping = [&] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtentsDynamic<1>; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(MAPPING == 0) return llama::mapping::AoS{extents}; diff --git a/examples/bitpackfloat/bitpackfloat.cpp b/examples/bitpackfloat/bitpackfloat.cpp index 3ba48ec9ed..9d6ff297f1 100644 --- a/examples/bitpackfloat/bitpackfloat.cpp +++ b/examples/bitpackfloat/bitpackfloat.cpp @@ -21,7 +21,7 @@ auto main() -> int constexpr auto exponentBits = 5; constexpr auto mantissaBits = 13; const auto mapping - = llama::mapping::BitPackedFloatSoA{llama::ArrayExtents{N}, exponentBits, mantissaBits, Vector{}}; + = llama::mapping::BitPackedFloatSoA{llama::ArrayExtents{N}, exponentBits, mantissaBits, Vector{}}; auto view = llama::allocView(mapping); diff --git a/examples/bitpackint/bitpackint.cpp b/examples/bitpackint/bitpackint.cpp index 08f035aa46..a0f34b6795 100644 --- a/examples/bitpackint/bitpackint.cpp +++ b/examples/bitpackint/bitpackint.cpp @@ -22,7 +22,8 @@ auto main() -> int constexpr auto N = 128; constexpr auto bits = 7; const auto mapping - = llama::mapping::BitPackedIntSoA, Vector, llama::Constant>{{N}}; + = llama::mapping::BitPackedIntSoA, Vector, llama::Constant>{ + {N}}; auto view = llama::allocView(mapping); @@ -49,7 +50,7 @@ auto main() -> int // extract into a view of full size integers auto viewExtracted - = llama::allocViewUninitialized(llama::mapping::AoS, Vector>{{N}}); + = llama::allocViewUninitialized(llama::mapping::AoS, Vector>{{N}}); llama::copy(view, viewExtracted); if(!std::equal(view.begin(), view.end(), viewExtracted.begin(), viewExtracted.end())) fmt::print("ERROR: unpacked view is different\n"); diff --git a/examples/bufferguard/bufferguard.cpp b/examples/bufferguard/bufferguard.cpp index cb6ff047a7..d18e4c5333 100644 --- a/examples/bufferguard/bufferguard.cpp +++ b/examples/bufferguard/bufferguard.cpp @@ -21,16 +21,16 @@ using Vector = llama::Record< // clang-format on template typename InnerMapping, typename TRecordDim> -struct GuardMapping2D : llama::ArrayExtentsDynamic<2> +struct GuardMapping2D : llama::ArrayExtentsDynamic { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; - using ArrayIndex = llama::ArrayIndex<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; + using ArrayIndex = llama::ArrayIndex; using RecordDim = TRecordDim; constexpr GuardMapping2D() = default; constexpr explicit GuardMapping2D(ArrayExtents extents, RecordDim = {}) - : llama::ArrayExtentsDynamic<2>(extents) + : llama::ArrayExtentsDynamic(extents) , left({extents[0] - 2}) , right({extents[0] - 2}) , top({extents[1] - 2}) @@ -69,7 +69,7 @@ struct GuardMapping2D : llama::ArrayExtentsDynamic<2> template constexpr auto blobNrAndOffset(ArrayIndex ai, llama::RecordCoord rc = {}) const - -> llama::NrAndOffset + -> llama::NrAndOffset { // [0][0] is at left top const auto [row, col] = ai; @@ -144,7 +144,8 @@ struct GuardMapping2D : llama::ArrayExtentsDynamic<2> } private: - constexpr auto offsetBlobNr(llama::NrAndOffset nao, std::size_t blobNrOffset) const -> llama::NrAndOffset + constexpr auto offsetBlobNr(llama::NrAndOffset nao, std::size_t blobNrOffset) const + -> llama::NrAndOffset { nao.nr += blobNrOffset; return nao; @@ -162,11 +163,11 @@ struct GuardMapping2D : llama::ArrayExtentsDynamic<2> llama::mapping::One, RecordDim> leftBot; llama::mapping::One, RecordDim> rightTop; llama::mapping::One, RecordDim> rightBot; - InnerMapping, RecordDim> left; - InnerMapping, RecordDim> right; - InnerMapping, RecordDim> top; - InnerMapping, RecordDim> bot; - InnerMapping, RecordDim> center; + InnerMapping, RecordDim> left; + InnerMapping, RecordDim> right; + InnerMapping, RecordDim> top; + InnerMapping, RecordDim> bot; + InnerMapping, RecordDim> center; static constexpr auto leftTopOff = std::size_t{0}; static constexpr auto leftBotOff = leftTopOff + decltype(leftTop)::blobCount; @@ -202,8 +203,8 @@ void run(const std::string& mappingName) { std::cout << "\n===== Mapping " << mappingName << " =====\n\n"; - constexpr auto rows = 7; - constexpr auto cols = 5; + constexpr std::size_t rows = 7; + constexpr std::size_t cols = 5; const auto extents = llama::ArrayExtents{rows, cols}; const auto mapping = GuardMapping2D{extents}; std::ofstream{"bufferguard_" + mappingName + ".svg"} << llama::toSvg(mapping); diff --git a/examples/bytesplit/bytesplit.cpp b/examples/bytesplit/bytesplit.cpp index f8a74e925d..e4c91d785e 100644 --- a/examples/bytesplit/bytesplit.cpp +++ b/examples/bytesplit/bytesplit.cpp @@ -26,7 +26,7 @@ using Data = llama::Record< auto main() -> int { constexpr auto N = 128; - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic; const auto mapping = llama::mapping::Bytesplit::fn>{{N}}; auto view = llama::allocView(mapping); diff --git a/examples/cuda/nbody/nbody.cu b/examples/cuda/nbody/nbody.cu index 7410f9a8bd..edf0452b73 100644 --- a/examples/cuda/nbody/nbody.cu +++ b/examples/cuda/nbody/nbody.cu @@ -86,7 +86,7 @@ __global__ void updateSM(View particles) { constexpr auto sharedMapping = [] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtents; if constexpr(MappingSM == 0) return llama::mapping::AoS{}; if constexpr(MappingSM == 1) @@ -111,15 +111,15 @@ __global__ void updateSM(View particles) const auto tbi = blockIdx.x; llama::One pi = particles(ti); - for(std::size_t blockOffset = 0; blockOffset < PROBLEM_SIZE; blockOffset += SHARED_ELEMENTS_PER_BLOCK) + for(int blockOffset = 0; blockOffset < PROBLEM_SIZE; blockOffset += SHARED_ELEMENTS_PER_BLOCK) { LLAMA_INDEPENDENT_DATA - for(auto j = tbi; j < SHARED_ELEMENTS_PER_BLOCK; j += THREADS_PER_BLOCK) + for(int j = tbi; j < SHARED_ELEMENTS_PER_BLOCK; j += THREADS_PER_BLOCK) sharedView(j) = particles(blockOffset + j); __syncthreads(); LLAMA_INDEPENDENT_DATA - for(auto j = std::size_t{0}; j < SHARED_ELEMENTS_PER_BLOCK; ++j) + for(int j = 0; j < SHARED_ELEMENTS_PER_BLOCK; ++j) pPInteraction(pi, sharedView(j)); __syncthreads(); } @@ -133,7 +133,7 @@ __global__ void update(View particles) llama::One pi = particles(ti); LLAMA_INDEPENDENT_DATA - for(auto j = std::size_t{0}; j < PROBLEM_SIZE; ++j) + for(int j = 0; j < PROBLEM_SIZE; ++j) pPInteraction(pi, particles(j)); particles(ti)(tag::Vel{}) = pi(tag::Vel{}); } @@ -178,7 +178,7 @@ try auto mapping = [] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtentsDynamic; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(Mapping == 0) return llama::mapping::AoS{extents}; @@ -224,7 +224,7 @@ try std::default_random_engine engine; std::normal_distribution distribution(FP(0), FP(1)); - for(std::size_t i = 0; i < PROBLEM_SIZE; ++i) + for(int i = 0; i < PROBLEM_SIZE; ++i) { llama::One p; p(tag::Pos(), tag::X()) = distribution(engine); @@ -267,7 +267,7 @@ try double sumUpdate = 0; double sumMove = 0; - for(std::size_t s = 0; s < STEPS; ++s) + for(int s = 0; s < STEPS; ++s) { if constexpr(RUN_UPATE) { @@ -388,7 +388,7 @@ namespace manual std::default_random_engine engine; std::normal_distribution distribution(FP(0), FP(1)); - for(std::size_t i = 0; i < PROBLEM_SIZE; ++i) + for(int i = 0; i < PROBLEM_SIZE; ++i) { hostPositions[i].x = distribution(engine); hostPositions[i].y = distribution(engine); @@ -426,7 +426,7 @@ namespace manual double sumUpdate = 0; double sumMove = 0; - for(std::size_t s = 0; s < STEPS; ++s) + for(int s = 0; s < STEPS; ++s) { if constexpr(RUN_UPATE) { diff --git a/examples/cuda/pitch/pitch.cu b/examples/cuda/pitch/pitch.cu index 3712fdb2e3..fd75c5f2ff 100644 --- a/examples/cuda/pitch/pitch.cu +++ b/examples/cuda/pitch/pitch.cu @@ -107,7 +107,7 @@ namespace llamaex template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset( typename Base::ArrayIndex ai, - RecordCoord = {}) const -> NrAndOffset + RecordCoord = {}) const -> NrAndOffset { constexpr std::size_t flatFieldIndex = #ifdef __NVCC__ @@ -134,7 +134,7 @@ try prop.totalGlobalMem / 1024 / 1024, prop.sharedMemPerBlock / 1024); - const auto extents = llama::ArrayExtents{600, 800}; // height, width + const auto extents = llama::ArrayExtents{600, 800}; // height, width const auto widthBytes = extents[1] * sizeof(RGB); const auto blockDim = dim3{16, 32, 1}; @@ -151,7 +151,7 @@ try checkError(cudaMallocPitch(&mem, &rowPitch, widthBytes, extents[0])); fmt::print("Row pitch: {} B ({} B padding)\n", rowPitch, rowPitch - widthBytes); - auto mapping = llamaex::PitchedAoS, RGB>{extents, rowPitch}; + auto mapping = llamaex::PitchedAoS, RGB>{extents, rowPitch}; assert(mapping.blobSize(0) == rowPitch * extents[0]); auto view = llama::View{mapping, llama::Array{mem}}; diff --git a/examples/heatequation/heatequation.cpp b/examples/heatequation/heatequation.cpp index 6975738c32..352893f94c 100644 --- a/examples/heatequation/heatequation.cpp +++ b/examples/heatequation/heatequation.cpp @@ -137,7 +137,7 @@ try auto run = [&](std::string_view updateName, auto update) { // init - for(uint32_t i = 0; i < extent; i++) + for(int i = 0; i < extent; i++) uCurr[i] = exactSolution(i * dx, 0.0); uNext[0] = 0; uNext[extent - 1] = 0; @@ -154,7 +154,7 @@ try // calculate error double maxError = 0.0; - for(uint32_t i = 0; i < extent; i++) + for(int i = 0; i < extent; i++) { const auto error = std::abs(uNext[i] - exactSolution(i * dx, tMax)); maxError = std::max(maxError, error); diff --git a/examples/nbody/nbody.cpp b/examples/nbody/nbody.cpp index 0f93987a5d..28b9d82ba2 100644 --- a/examples/nbody/nbody.cpp +++ b/examples/nbody/nbody.cpp @@ -127,7 +127,7 @@ namespace usellama Stopwatch watch; auto mapping = [&] { - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(Mapping == 0) return llama::mapping::AoS{extents}; diff --git a/examples/nbody_benchmark/nbody.cpp b/examples/nbody_benchmark/nbody.cpp index 71d983c475..844e0fc2fc 100644 --- a/examples/nbody_benchmark/nbody.cpp +++ b/examples/nbody_benchmark/nbody.cpp @@ -85,7 +85,7 @@ void run(std::ostream& plotFile) const auto mapping = [&] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtentsDynamic; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(Mapping == 0) return llama::mapping::AoS{extents}; diff --git a/examples/raycast/raycast.cpp b/examples/raycast/raycast.cpp index 1ec15c0a47..c182ddb20a 100644 --- a/examples/raycast/raycast.cpp +++ b/examples/raycast/raycast.cpp @@ -51,7 +51,7 @@ namespace llama::Field, llama::Field>; - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic; using Mapping = llama::mapping::AoS; // using Mapping = llama::mapping::SoA; // using Mapping = llama::mapping::AoSoA; diff --git a/examples/vectoradd/vectoradd.cpp b/examples/vectoradd/vectoradd.cpp index 527d4ac132..a01db094be 100644 --- a/examples/vectoradd/vectoradd.cpp +++ b/examples/vectoradd/vectoradd.cpp @@ -48,7 +48,7 @@ namespace usellama const auto mapping = [&] { - using ArrayExtents = llama::ArrayExtents; + using ArrayExtents = llama::ArrayExtentsDynamic; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(MAPPING == 0) return llama::mapping::AoS{extents, Vector{}}; diff --git a/include/llama/ArrayExtents.hpp b/include/llama/ArrayExtents.hpp index 33cfe6216e..f55cb1b5fc 100644 --- a/include/llama/ArrayExtents.hpp +++ b/include/llama/ArrayExtents.hpp @@ -13,67 +13,137 @@ namespace llama // TODO(bgruber): make this an alias in C++20, when we have CTAD for aliases /// Represents a run-time index into the array dimensions. /// \tparam Dim Compile-time number of dimensions. - template - struct ArrayIndex : Array + template + struct ArrayIndex : Array { static constexpr std::size_t rank = Dim; }; + // allow comparing ArrayIndex with different size types: + template + LLAMA_FN_HOST_ACC_INLINE constexpr auto operator==(ArrayIndex a, ArrayIndex b) -> bool + { + for(std::size_t i = 0; i < Dim; ++i) + if(a[i] != b[i]) + return false; + return true; + } + + template + LLAMA_FN_HOST_ACC_INLINE constexpr auto operator!=(ArrayIndex a, ArrayIndex b) -> bool + { + return !(a == b); + } + static_assert( - std::is_trivially_default_constructible_v>); // so ArrayIndex<1>{} will produce a zeroed - // index. Should hold for all dimensions, - // but just checking for <1> here. - static_assert(std::is_trivially_copy_constructible_v>); - static_assert(std::is_trivially_move_constructible_v>); - static_assert(std::is_trivially_copy_assignable_v>); - static_assert(std::is_trivially_move_assignable_v>); + std::is_trivially_default_constructible_v>); // so ArrayIndex<1>{} will produce a zeroed + // index. Should hold for all dimensions, + // but just checking for <1> here. + static_assert(std::is_trivially_copy_constructible_v>); + static_assert(std::is_trivially_move_constructible_v>); + static_assert(std::is_trivially_copy_assignable_v>); + static_assert(std::is_trivially_move_assignable_v>); + + namespace internal + { + template + struct IndexTypeFromArgs + { + using type = Default; + }; + + template + struct IndexTypeFromArgs + { + static_assert(std::conjunction_v...>, "All index types must be the same"); + using type = FirstInt; + }; + } // namespace internal template - ArrayIndex(Args...) -> ArrayIndex; + ArrayIndex(Args...) + -> ArrayIndex::type, sizeof...(Args)>; } // namespace llama -template -struct std::tuple_size> : std::integral_constant +template +struct std::tuple_size> : std::integral_constant { }; -template -struct std::tuple_element> +template +struct std::tuple_element> { - using type = size_t; + using type = V; }; namespace llama { + namespace internal + { + struct Dyn + { + template + constexpr operator T() const + { + return static_cast(-1); + } + + template + friend constexpr auto operator==(T i, Dyn) -> bool + { + return i == static_cast(-1); + } + + template + friend constexpr auto operator==(Dyn d, T i) -> bool + { + return i == d; + } + + template + friend constexpr auto operator!=(T i, Dyn d) -> bool + { + return !(i == d); + } + + template + friend constexpr auto operator!=(Dyn d, T i) -> bool + { + return !(i == d); + } + }; + } // namespace internal + /// Used as a template argument to \ref ArrayExtents to mark a dynamic extent. - inline constexpr std::size_t dyn = std::numeric_limits::max(); + inline constexpr auto dyn = internal::Dyn{}; /// ArrayExtents holding compile and runtime indices. This is conceptually equivalent to the std::extent of - /// std::mdspan. See: https://wg21.link/P0009 - template - struct ArrayExtents : Array::value_type, ((Sizes == dyn) + ... + 0)> + /// std::mdspan (@see: https://wg21.link/P0009) including the changes to make the size_type controllable (@see: + /// https://wg21.link/P2553). + template + struct ArrayExtents : Array { static constexpr std::size_t rank = sizeof...(Sizes); static constexpr auto rank_dynamic = ((Sizes == dyn) + ... + 0); static constexpr auto rank_static = rank - rank_dynamic; - using Index = ArrayIndex; - using value_type = typename Index::value_type; + using Index = ArrayIndex; + using value_type = T; template - LLAMA_FN_HOST_ACC_INLINE constexpr auto get() const + LLAMA_FN_HOST_ACC_INLINE constexpr auto get() const -> value_type { using namespace boost::mp11; - using TypeList = mp_list_c; + using TypeList = mp_list_c; constexpr auto extent = mp_at_c::value; if constexpr(extent != dyn) return extent; else return static_cast&>( - *this)[+mp_count, mp_size_t>::value]; + *this)[+mp_count, std::integral_constant>::value]; } - LLAMA_FN_HOST_ACC_INLINE constexpr auto operator[](std::size_t i) const + LLAMA_FN_HOST_ACC_INLINE constexpr auto operator[](T i) const -> value_type { return boost::mp11::mp_with_index(i, [&](auto ic) { return get(); }); } @@ -97,15 +167,15 @@ namespace llama } }; - template<> - struct ArrayExtents<> + template + struct ArrayExtents { static constexpr std::size_t rank = 0; static constexpr auto rank_dynamic = 0; static constexpr auto rank_static = 0; - using Index = ArrayIndex; - using value_type = typename Index::value_type; + using Index = ArrayIndex; + using value_type = T; LLAMA_FN_HOST_ACC_INLINE constexpr auto toArray() const -> Index { @@ -119,71 +189,86 @@ namespace llama }; template - ArrayExtents(Args... args) -> ArrayExtents<(Args{}, dyn)...>; + ArrayExtents(Args... args) + -> ArrayExtents::type, (Args{}, dyn)...>; - static_assert(std::is_trivially_default_constructible_v>); - static_assert(std::is_trivially_copy_constructible_v>); - static_assert(std::is_trivially_move_constructible_v>); - static_assert(std::is_trivially_copy_assignable_v>); - static_assert(std::is_trivially_move_assignable_v>); - static_assert(std::is_empty_v>); + static_assert(std::is_trivially_default_constructible_v>); + static_assert(std::is_trivially_copy_constructible_v>); + static_assert(std::is_trivially_move_constructible_v>); + static_assert(std::is_trivially_copy_assignable_v>); + static_assert(std::is_trivially_move_assignable_v>); + static_assert(std::is_empty_v>); - template - LLAMA_FN_HOST_ACC_INLINE constexpr auto operator==(ArrayExtents a, ArrayExtents b) -> bool + template + LLAMA_FN_HOST_ACC_INLINE constexpr auto operator==( + ArrayExtents a, + ArrayExtents b) -> bool { return a.toArray() == b.toArray(); } - template - LLAMA_FN_HOST_ACC_INLINE constexpr auto operator!=(ArrayExtents a, ArrayExtents b) -> bool + template + LLAMA_FN_HOST_ACC_INLINE constexpr auto operator!=( + ArrayExtents a, + ArrayExtents b) -> bool { return !(a == b); } - template - LLAMA_FN_HOST_ACC_INLINE constexpr auto product(ArrayExtents e) -> - typename ArrayExtents::value_type + template + LLAMA_FN_HOST_ACC_INLINE constexpr auto product(ArrayExtents e) -> SizeType { return product(e.toArray()); } - /// N-dimensional ArrayExtents where all values are dynamic. - template - using ArrayExtentsDynamic = internal:: - mp_unwrap_values_into, N>, ArrayExtents>; + namespace internal + { + template + constexpr auto makeArrayExtents(std::index_sequence) + { + return ArrayExtents(Is), Extent)...>{}; + } + } // namespace internal + + /// N-dimensional ArrayExtents where all N extents are Extent. + template + using ArrayExtentsNCube = decltype(internal::makeArrayExtents(std::make_index_sequence{})); - /// N-dimensional ArrayExtents where all values are Extent. - template - using ArrayExtentsStatic = internal:: - mp_unwrap_values_into, N>, ArrayExtents>; + /// N-dimensional ArrayExtents where all values are dynamic. + template + using ArrayExtentsDynamic = ArrayExtentsNCube; - template + template LLAMA_FN_HOST_ACC_INLINE void forEachADCoord( - [[maybe_unused]] ArrayIndex adSize, + [[maybe_unused]] ArrayIndex adSize, Func&& func, OuterIndices... outerIndices) { if constexpr(Dim > 0) - for(std::size_t i = 0; i < adSize[0]; i++) - forEachADCoord(ArrayIndex{pop_front(adSize)}, std::forward(func), outerIndices..., i); + for(SizeType i = 0; i < adSize[0]; i++) + forEachADCoord( + ArrayIndex{pop_front(adSize)}, + std::forward(func), + outerIndices..., + i); else - std::forward(func)(ArrayIndex{outerIndices...}); + std::forward(func)(ArrayIndex{outerIndices...}); } - template - LLAMA_FN_HOST_ACC_INLINE void forEachADCoord(ArrayExtents extents, Func&& func) + template + LLAMA_FN_HOST_ACC_INLINE void forEachADCoord(ArrayExtents extents, Func&& func) { forEachADCoord(extents.toArray(), std::forward(func)); } } // namespace llama -template -struct std::tuple_size> : std::integral_constant +template +struct std::tuple_size> : std::integral_constant { }; -template -struct std::tuple_element> +template +struct std::tuple_element> { - using type = typename llama::ArrayExtents::value_type; + using type = SizeType; }; diff --git a/include/llama/ArrayIndexRange.hpp b/include/llama/ArrayIndexRange.hpp index 4a3ea253a3..60cbd8d76c 100644 --- a/include/llama/ArrayIndexRange.hpp +++ b/include/llama/ArrayIndexRange.hpp @@ -75,7 +75,8 @@ namespace llama current[rank - 1]--; for(auto i = static_cast(rank) - 2; i >= 0; i--) { - if(current[i + 1] != std::numeric_limits::max()) + // return if no underflow + if(current[i + 1] != static_cast(-1)) return *this; current[i + 1] = extents[i] - 1; current[i]--; diff --git a/include/llama/Concepts.hpp b/include/llama/Concepts.hpp index b3b53fc203..9dbc7a3195 100644 --- a/include/llama/Concepts.hpp +++ b/include/llama/Concepts.hpp @@ -21,12 +21,12 @@ namespace llama { m.extents() } -> std::same_as; { M::blobCount } -> std::convertible_to; Array{}; // validates constexpr-ness - { m.blobSize(std::size_t{}) } -> std::same_as; + { m.blobSize(typename M::ArrayExtents::value_type{}) } -> std::same_as; }; template concept PhysicalField = requires(M m, typename M::ArrayIndex ai) { - { m.blobNrAndOffset(ai, RC{}) } -> std::same_as; + { m.blobNrAndOffset(ai, RC{}) } -> std::same_as>; }; template diff --git a/include/llama/Copy.hpp b/include/llama/Copy.hpp index ddc801a5c4..3c455999a3 100644 --- a/include/llama/Copy.hpp +++ b/include/llama/Copy.hpp @@ -97,13 +97,17 @@ namespace llama const auto extents = srcView.mapping().extents().toArray(); const auto workPerThread = (extents[0] + threadCount - 1) / threadCount; const auto start = threadId * workPerThread; - const auto end = std::min((threadId + 1) * workPerThread, extents[0]); + const auto end = std::min((threadId + 1) * workPerThread, static_cast(extents[0])); for(auto i = start; i < end; i++) { + using SrcSizeType = typename SrcMapping::ArrayExtents::value_type; if constexpr(dims > 1) - forEachADCoord(ArrayIndex{pop_front(extents)}, copyOne, static_cast(i)); + forEachADCoord( + ArrayIndex{pop_front(extents)}, + copyOne, + static_cast(i)); else - copyOne(ArrayIndex{static_cast(i)}); + copyOne(ArrayIndex{static_cast(i)}); } } diff --git a/include/llama/Core.hpp b/include/llama/Core.hpp index d932e17c22..1cc21812d9 100644 --- a/include/llama/Core.hpp +++ b/include/llama/Core.hpp @@ -43,20 +43,11 @@ namespace llama static_assert(isAllowedFieldType, "This field's type is not allowed"); }; + template struct NrAndOffset { - std::size_t nr; - std::size_t offset; - - friend auto operator==(const NrAndOffset& a, const NrAndOffset& b) -> bool - { - return a.nr == b.nr && a.offset == b.offset; - } - - friend auto operator!=(const NrAndOffset& a, const NrAndOffset& b) -> bool - { - return !(a == b); - } + T nr; + T offset; friend auto operator<<(std::ostream& os, const NrAndOffset& value) -> std::ostream& { @@ -64,6 +55,21 @@ namespace llama } }; + template + NrAndOffset(Int, Int) -> NrAndOffset; + + template + auto operator==(const NrAndOffset& a, const NrAndOffset& b) -> bool + { + return a.nr == b.nr && a.offset == b.offset; + } + + template + auto operator!=(const NrAndOffset& a, const NrAndOffset& b) -> bool + { + return !(a == b); + } + /// Get the tag from a \ref Field. template using GetFieldTag = boost::mp11::mp_first; diff --git a/include/llama/DumpMapping.hpp b/include/llama/DumpMapping.hpp index 587ccd2125..3c3b05a1c7 100644 --- a/include/llama/DumpMapping.hpp +++ b/include/llama/DumpMapping.hpp @@ -38,8 +38,8 @@ namespace llama return c; } - template - auto formatArrayIndex(const ArrayIndex& ai) + template + auto formatArrayIndex(const ArrayIndex& ai) { if constexpr(Dim == 1) return std::to_string(ai[0]); @@ -57,13 +57,13 @@ namespace llama } } - template + template struct FieldBox { - ArrayIndex arrayIndex; + ArrayIndex arrayIndex; std::vector recordCoord; std::string recordTags; - NrAndOffset nrAndOffset; + NrAndOffset nrAndOffset; std::size_t size; }; @@ -80,12 +80,12 @@ namespace llama View& view, typename View::Mapping::ArrayIndex ai, RecordCoord rc, - std::vector>& infos) + std::vector>& infos) { using Mapping = typename View::Mapping; using RecordDim = typename Mapping::RecordDim; - auto emitInfo = [&](NrAndOffset nrAndOffset, std::size_t size) { + auto emitInfo = [&](auto nrAndOffset, std::size_t size) { infos.push_back({ai, internal::toVec(rc), recordCoordTags(rc), nrAndOffset, size}); }; @@ -149,13 +149,13 @@ namespace llama } // if we come here, we could not find out where the value is coming from - emitInfo(NrAndOffset{Mapping::blobCount, 0}, sizeof(Type)); + emitInfo(NrAndOffset{Mapping::blobCount, std::size_t{0}}, sizeof(Type)); } template - auto boxesFromMapping(const Mapping& mapping) -> std::vector> + auto boxesFromMapping(const Mapping& mapping) -> std::vector> { - std::vector> infos; + std::vector> infos; std::optional view; if constexpr(hasAnyComputedField()) @@ -181,8 +181,9 @@ namespace llama return infos; } - template - auto breakBoxes(std::vector> boxes, std::size_t wrapByteCount) -> std::vector> + template + auto breakBoxes(std::vector> boxes, std::size_t wrapByteCount) + -> std::vector> { for(std::size_t i = 0; i < boxes.size(); i++) { diff --git a/include/llama/Vector.hpp b/include/llama/Vector.hpp index 958bd3838c..3108f54772 100644 --- a/include/llama/Vector.hpp +++ b/include/llama/Vector.hpp @@ -26,14 +26,15 @@ namespace llama using iterator = decltype(std::declval().begin()); using value_type = typename iterator::value_type; + using size_type = typename Mapping::ArrayExtents::value_type; Vector() = default; template> - LLAMA_FN_HOST_ACC_INLINE explicit Vector(std::size_t count, const VirtualRecord& value = {}) + LLAMA_FN_HOST_ACC_INLINE explicit Vector(size_type count, const VirtualRecord& value = {}) { reserve(count); - for(std::size_t i = 0; i < count; i++) + for(size_type i = 0; i < count; i++) push_back(value); } @@ -67,7 +68,7 @@ namespace llama // TODO(bgruber): assign - LLAMA_FN_HOST_ACC_INLINE auto at(std::size_t i) -> decltype(auto) + LLAMA_FN_HOST_ACC_INLINE auto at(size_type i) -> decltype(auto) { if(i >= m_size) throw std::out_of_range{ @@ -75,7 +76,7 @@ namespace llama return m_view(i); } - LLAMA_FN_HOST_ACC_INLINE auto at(std::size_t i) const -> decltype(auto) + LLAMA_FN_HOST_ACC_INLINE auto at(size_type i) const -> decltype(auto) { if(i >= m_size) throw std::out_of_range{ @@ -83,12 +84,12 @@ namespace llama return m_view(i); } - LLAMA_FN_HOST_ACC_INLINE auto operator[](std::size_t i) -> decltype(auto) + LLAMA_FN_HOST_ACC_INLINE auto operator[](size_type i) -> decltype(auto) { return m_view(i); } - LLAMA_FN_HOST_ACC_INLINE auto operator[](std::size_t i) const -> decltype(auto) + LLAMA_FN_HOST_ACC_INLINE auto operator[](size_type i) const -> decltype(auto) { return m_view(i); } @@ -158,18 +159,18 @@ namespace llama return m_size == 0; } - LLAMA_FN_HOST_ACC_INLINE auto size() const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE auto size() const -> size_type { return m_size; } - LLAMA_FN_HOST_ACC_INLINE void reserve(std::size_t cap) + LLAMA_FN_HOST_ACC_INLINE void reserve(size_type cap) { if(cap > capacity()) changeCapacity(cap); } - LLAMA_FN_HOST_ACC_INLINE auto capacity() const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE auto capacity() const -> size_type { return m_view.mapping().extents()[0]; } @@ -228,10 +229,10 @@ namespace llama } template> - LLAMA_FN_HOST_ACC_INLINE void resize(std::size_t count, const VirtualRecord& value = {}) + LLAMA_FN_HOST_ACC_INLINE void resize(size_type count, const VirtualRecord& value = {}) { reserve(count); - for(std::size_t i = m_size; i < count; i++) + for(size_type i = m_size; i < count; i++) m_view[i] = value; m_size = count; } @@ -274,7 +275,7 @@ namespace llama } private: - LLAMA_FN_HOST_ACC_INLINE void changeCapacity(std::size_t cap) + LLAMA_FN_HOST_ACC_INLINE void changeCapacity(size_type cap) { auto newView = allocViewUninitialized(Mapping{typename Mapping::ArrayExtents{cap}}); auto b = begin(); @@ -291,7 +292,7 @@ namespace llama } ViewType m_view = {}; - std::size_t m_size = 0; + size_type m_size = 0; }; diff --git a/include/llama/View.hpp b/include/llama/View.hpp index 2ffeac5851..3aa2125f63 100644 --- a/include/llama/View.hpp +++ b/include/llama/View.hpp @@ -121,7 +121,7 @@ namespace llama template LLAMA_FN_HOST_ACC_INLINE auto allocViewStack() -> decltype(auto) { - constexpr auto mapping = mapping::MinAlignedOne, RecordDim>{}; + constexpr auto mapping = mapping::MinAlignedOne, RecordDim>{}; return allocView(mapping, bloballoc::Stack{}); } @@ -346,6 +346,7 @@ namespace llama using RecordDim = typename Mapping::RecordDim; using iterator = Iterator; using const_iterator = Iterator; + using size_type = typename ArrayExtents::value_type; static_assert( std::is_same_v>, @@ -375,6 +376,12 @@ namespace llama return static_cast(*this); } + template + auto operator()(llama::ArrayIndex) const + { + static_assert(!sizeof(V), "Passed ArrayIndex with SizeType different than Mapping::ArrayExtent"); + } + /// Retrieves the \ref VirtualRecord at the given \ref ArrayIndex index. LLAMA_FN_HOST_ACC_INLINE auto operator()(ArrayIndex ai) const -> decltype(auto) { @@ -406,28 +413,26 @@ namespace llama /// Retrieves the \ref VirtualRecord at the \ref ArrayIndex index constructed from the passed component /// indices. - template + template< + typename... Indices, + std::enable_if_t...>, int> = 0> LLAMA_FN_HOST_ACC_INLINE auto operator()(Indices... indices) const -> decltype(auto) { static_assert( sizeof...(Indices) == ArrayIndex::rank, "Please specify as many indices as you have array dimensions"); - static_assert( - std::conjunction_v...>, - "Indices must be convertible to std::size_t"); LLAMA_FORCE_INLINE_RECURSIVE return (*this)(ArrayIndex{static_cast(indices)...}); } - template + template< + typename... Indices, + std::enable_if_t...>, int> = 0> LLAMA_FN_HOST_ACC_INLINE auto operator()(Indices... indices) -> decltype(auto) { static_assert( sizeof...(Indices) == ArrayIndex::rank, "Please specify as many indices as you have array dimensions"); - static_assert( - std::conjunction_v...>, - "Indices must be convertible to std::size_t"); LLAMA_FORCE_INLINE_RECURSIVE return (*this)(ArrayIndex{static_cast(indices)...}); } @@ -446,14 +451,20 @@ namespace llama return (*this)(ai); } + template + auto operator[](llama::ArrayIndex) const + { + static_assert(!sizeof(V), "Passed ArrayIndex with SizeType different than Mapping::ArrayExtent"); + } + /// Retrieves the \ref VirtualRecord at the 1D \ref ArrayIndex index constructed from the passed index. - LLAMA_FN_HOST_ACC_INLINE auto operator[](std::size_t index) const -> decltype(auto) + LLAMA_FN_HOST_ACC_INLINE auto operator[](size_type index) const -> decltype(auto) { LLAMA_FORCE_INLINE_RECURSIVE return (*this)(index); } - LLAMA_FN_HOST_ACC_INLINE auto operator[](std::size_t index) -> decltype(auto) + LLAMA_FN_HOST_ACC_INLINE auto operator[](size_type index) -> decltype(auto) { LLAMA_FORCE_INLINE_RECURSIVE return (*this)(index); @@ -521,6 +532,8 @@ namespace llama using ArrayExtents = typename Mapping::ArrayExtents; ///< array extents of the parent view using ArrayIndex = typename Mapping::ArrayIndex; ///< array index of the parent view + using size_type = typename ArrayExtents::value_type; + /// Creates a VirtualView given a parent \ref View and offset. template LLAMA_FN_HOST_ACC_INLINE VirtualView(StoredParentViewFwd&& parentView, ArrayIndex offset) @@ -562,8 +575,8 @@ namespace llama sizeof...(Indices) == ArrayIndex::rank, "Please specify as many indices as you have array dimensions"); static_assert( - std::conjunction_v...>, - "Indices must be convertible to std::size_t"); + std::conjunction_v...>, + "Indices must be convertible to ArrayExtents::size_type"); LLAMA_FORCE_INLINE_RECURSIVE return parentView( ArrayIndex{ArrayIndex{static_cast(indices)...} + offset}); @@ -576,8 +589,8 @@ namespace llama sizeof...(Indices) == ArrayIndex::rank, "Please specify as many indices as you have array dimensions"); static_assert( - std::conjunction_v...>, - "Indices must be convertible to std::size_t"); + std::conjunction_v...>, + "Indices must be convertible to ArrayExtents::size_type"); LLAMA_FORCE_INLINE_RECURSIVE return parentView( ArrayIndex{ArrayIndex{static_cast(indices)...} + offset}); diff --git a/include/llama/mapping/AoS.hpp b/include/llama/mapping/AoS.hpp index e8c15f4533..cb051eae12 100644 --- a/include/llama/mapping/AoS.hpp +++ b/include/llama/mapping/AoS.hpp @@ -26,6 +26,7 @@ namespace llama::mapping private: using Base = MappingBase; using Flattener = FlattenRecordDim; + using size_type = typename Base::size_type; public: using LinearizeArrayDimsFunctor = TLinearizeArrayDimsFunctor; @@ -33,7 +34,7 @@ namespace llama::mapping using Base::Base; - LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize(std::size_t) const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize(size_type) const -> size_type { return LinearizeArrayDimsFunctor{}.size(Base::extents()) * flatSizeOf; @@ -42,19 +43,17 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset( typename Base::ArrayIndex ai, - RecordCoord = {}) const -> NrAndOffset + RecordCoord = {}) const -> NrAndOffset { constexpr std::size_t flatFieldIndex = #ifdef __NVCC__ *& // mess with nvcc compiler state to workaround bug #endif Flattener::template flatIndex; - const auto offset - = LinearizeArrayDimsFunctor{}(ai, Base::extents()) - * flatSizeOf< - typename Flattener::FlatRecordDim, - AlignAndPad> + flatOffsetOf; - return {0, offset}; + const auto offset = LinearizeArrayDimsFunctor{}(ai, Base::extents()) + * static_cast(flatSizeOf) + + static_cast(flatOffsetOf); + return {size_type{0}, offset}; } }; diff --git a/include/llama/mapping/AoSoA.hpp b/include/llama/mapping/AoSoA.hpp index e9288dd343..cdb74c83ca 100644 --- a/include/llama/mapping/AoSoA.hpp +++ b/include/llama/mapping/AoSoA.hpp @@ -32,7 +32,7 @@ namespace llama::mapping template< typename TArrayExtents, typename TRecordDim, - std::size_t Lanes, + typename TArrayExtents::value_type Lanes, typename TLinearizeArrayDimsFunctor = LinearizeArrayDimsCpp, template typename FlattenRecordDim = FlattenRecordDimInOrder> struct AoSoA : MappingBase @@ -40,6 +40,7 @@ namespace llama::mapping private: using Base = MappingBase; using Flattener = FlattenRecordDim; + using size_type = typename Base::size_type; public: using LinearizeArrayDimsFunctor = TLinearizeArrayDimsFunctor; @@ -47,17 +48,16 @@ namespace llama::mapping using Base::Base; - LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize(std::size_t) const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize(size_type) const -> size_type { - return roundUpToMultiple( - LinearizeArrayDimsFunctor{}.size(Base::extents()) * sizeOf, - Lanes * sizeOf); + const auto rs = static_cast(sizeOf); + return roundUpToMultiple(LinearizeArrayDimsFunctor{}.size(Base::extents()) * rs, Lanes * rs); } template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset( typename Base::ArrayIndex ai, - RecordCoord = {}) const -> NrAndOffset + RecordCoord = {}) const -> NrAndOffset { constexpr std::size_t flatFieldIndex = #ifdef __NVCC__ @@ -67,9 +67,10 @@ namespace llama::mapping const auto flatArrayIndex = LinearizeArrayDimsFunctor{}(ai, Base::extents()); const auto blockIndex = flatArrayIndex / Lanes; const auto laneIndex = flatArrayIndex % Lanes; - const auto offset = (sizeOf * Lanes) * blockIndex - + flatOffsetOf * Lanes - + sizeof(GetType>) * laneIndex; + const auto offset = static_cast(sizeOf * Lanes) * blockIndex + + static_cast(flatOffsetOf) + * Lanes + + static_cast(sizeof(GetType>)) * laneIndex; return {0, offset}; } }; @@ -86,7 +87,7 @@ namespace llama::mapping template inline constexpr bool isAoSoA = false; - template + template inline constexpr bool isAoSoA> = true; } // namespace llama::mapping diff --git a/include/llama/mapping/BitPackedFloatSoA.hpp b/include/llama/mapping/BitPackedFloatSoA.hpp index 6f6e1a8dad..eca21b777b 100644 --- a/include/llama/mapping/BitPackedFloatSoA.hpp +++ b/include/llama/mapping/BitPackedFloatSoA.hpp @@ -99,11 +99,11 @@ namespace llama::mapping /// specified bit offset. /// @tparam Float Floating-point data type which can be loaded and store through this reference. /// @tparam StoredIntegralPointer Pointer to integral type used for storing the bits. - template + template struct LLAMA_DECLSPEC_EMPTY_BASES BitPackedFloatRef : private VHExp , private VHMan - , ProxyRefOpMixin, Float> + , ProxyRefOpMixin, Float> { private: static_assert( @@ -118,7 +118,8 @@ namespace llama::mapping BitPackedIntRef< FloatBits, StoredIntegralPointer, - decltype(integBits(std::declval(), std::declval()))> + decltype(integBits(std::declval(), std::declval())), + SizeType> intref; public: @@ -126,7 +127,7 @@ namespace llama::mapping LLAMA_FN_HOST_ACC_INLINE constexpr BitPackedFloatRef( StoredIntegralPointer p, - std::size_t bitOffset, + SizeType bitOffset, VHExp vhExp, VHMan vhMan #ifndef NDEBUG @@ -206,6 +207,7 @@ namespace llama::mapping using Base = MappingBase; using VHExp = llama::internal::BoxedValue; using VHMan = llama::internal::BoxedValue; + using size_type = typename TArrayExtents::value_type; public: static constexpr std::size_t blobCount = boost::mp11::mp_size>::value; @@ -223,11 +225,11 @@ namespace llama::mapping } LLAMA_FN_HOST_ACC_INLINE - constexpr auto blobSize(std::size_t /*blobIndex*/) const -> std::size_t + constexpr auto blobSize(size_type /*blobIndex*/) const -> size_type { - constexpr auto bitsPerStoredIntegral = sizeof(StoredIntegral) * CHAR_BIT; - const auto bitsNeeded - = LinearizeArrayDimsFunctor{}.size(Base::extents()) * (VHExp::value() + VHMan::value() + 1); + constexpr auto bitsPerStoredIntegral = static_cast(sizeof(StoredIntegral) * CHAR_BIT); + const auto bitsNeeded = LinearizeArrayDimsFunctor{}.size(Base::extents()) + * (static_cast(VHExp::value()) + static_cast(VHMan::value()) + 1); return roundUpToMultiple(bitsNeeded, bitsPerStoredIntegral) / CHAR_BIT; } @@ -244,12 +246,12 @@ namespace llama::mapping Blobs& blobs) const { constexpr auto blob = llama::flatRecordCoord>; - const auto bitOffset - = LinearizeArrayDimsFunctor{}(ai, Base::extents()) * (VHExp::value() + VHMan::value() + 1); + const auto bitOffset = LinearizeArrayDimsFunctor{}(ai, Base::extents()) + * (static_cast(VHExp::value()) + static_cast(VHMan::value()) + 1); using QualifiedStoredIntegral = CopyConst; using DstType = GetType>; - return internal::BitPackedFloatRef{ + return internal::BitPackedFloatRef{ reinterpret_cast(&blobs[blob][0]), bitOffset, static_cast(*this), diff --git a/include/llama/mapping/BitPackedIntSoA.hpp b/include/llama/mapping/BitPackedIntSoA.hpp index de4b25b198..b245f76f0e 100644 --- a/include/llama/mapping/BitPackedIntSoA.hpp +++ b/include/llama/mapping/BitPackedIntSoA.hpp @@ -17,10 +17,10 @@ namespace llama::mapping /// specified bit offset. /// @tparam Integral Integral data type which can be loaded and store through this reference. /// @tparam StoredIntegralPointer Pointer to integral type used for storing the bits. - template + template struct BitPackedIntRef : private VHBits - , ProxyRefOpMixin, Integral> + , ProxyRefOpMixin, Integral> { private: using StoredIntegral = std::remove_const_t>; @@ -33,19 +33,19 @@ namespace llama::mapping "retrieve"); StoredIntegralPointer ptr; - std::size_t bitOffset; + SizeType bitOffset; #ifndef NDEBUG StoredIntegralPointer endPtr; #endif - static constexpr auto bitsPerStoredIntegral = sizeof(StoredIntegral) * CHAR_BIT; + static constexpr auto bitsPerStoredIntegral = static_cast(sizeof(StoredIntegral) * CHAR_BIT); public: using value_type = Integral; LLAMA_FN_HOST_ACC_INLINE constexpr BitPackedIntRef( StoredIntegralPointer ptr, - std::size_t bitOffset, + SizeType bitOffset, VHBits vhBits #ifndef NDEBUG , @@ -179,6 +179,8 @@ namespace llama::mapping { private: using Base = MappingBase; + using VHBits = llama::internal::BoxedValue; + using size_type = typename TArrayExtents::value_type; template using IsAllowedFieldType = boost::mp11::mp_or, std::is_enum>; @@ -187,8 +189,6 @@ namespace llama::mapping boost::mp11::mp_all_of, IsAllowedFieldType>::value, "All record dimension field types must be integral"); - using VHBits = llama::internal::BoxedValue; - public: static constexpr std::size_t blobCount = boost::mp11::mp_size>::value; @@ -201,10 +201,11 @@ namespace llama::mapping } LLAMA_FN_HOST_ACC_INLINE - constexpr auto blobSize(std::size_t /*blobIndex*/) const -> std::size_t + constexpr auto blobSize(size_type /*blobIndex*/) const -> size_type { - constexpr auto bitsPerStoredIntegral = sizeof(StoredIntegral) * CHAR_BIT; - const auto bitsNeeded = LinearizeArrayDimsFunctor{}.size(Base::extents()) * VHBits::value(); + constexpr auto bitsPerStoredIntegral = static_cast(sizeof(StoredIntegral) * CHAR_BIT); + const auto bitsNeeded + = LinearizeArrayDimsFunctor{}.size(Base::extents()) * static_cast(VHBits::value()); return roundUpToMultiple(bitsNeeded, bitsPerStoredIntegral) / CHAR_BIT; } @@ -221,11 +222,12 @@ namespace llama::mapping Blobs& blobs) const { constexpr auto blob = flatRecordCoord>; - const auto bitOffset = LinearizeArrayDimsFunctor{}(ai, Base::extents()) * VHBits::value(); + const auto bitOffset + = LinearizeArrayDimsFunctor{}(ai, Base::extents()) * static_cast(VHBits::value()); using QualifiedStoredIntegral = CopyConst; using DstType = GetType>; - return internal::BitPackedIntRef{ + return internal::BitPackedIntRef{ reinterpret_cast(&blobs[blob][0]), bitOffset, static_cast(*this) diff --git a/include/llama/mapping/ChangeType.hpp b/include/llama/mapping/ChangeType.hpp index 1334bf2493..bbce13059d 100644 --- a/include/llama/mapping/ChangeType.hpp +++ b/include/llama/mapping/ChangeType.hpp @@ -113,7 +113,7 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset(ArrayIndex ai, RecordCoord rc = {}) - const -> NrAndOffset + const -> NrAndOffset { static_assert(!isComputed(rc)); return Inner::blobNrAndOffset(ai, rc); diff --git a/include/llama/mapping/Common.hpp b/include/llama/mapping/Common.hpp index e76edfa57a..54dd4e198b 100644 --- a/include/llama/mapping/Common.hpp +++ b/include/llama/mapping/Common.hpp @@ -15,6 +15,7 @@ namespace llama::mapping using ArrayExtents = TArrayExtents; using ArrayIndex = typename ArrayExtents::Index; using RecordDim = TRecordDim; + using size_type = typename ArrayExtents::value_type; constexpr MappingBase() = default; @@ -35,7 +36,7 @@ namespace llama::mapping struct LinearizeArrayDimsCpp { template - LLAMA_FN_HOST_ACC_INLINE constexpr auto size(const ArrayExtents& extents) -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto size(const ArrayExtents& extents) -> typename ArrayExtents::value_type { return product(extents); } @@ -46,14 +47,14 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto operator()( const typename ArrayExtents::Index& ai, - const ArrayExtents& extents) const -> std::size_t + const ArrayExtents& extents) const -> typename ArrayExtents::value_type { if constexpr(ArrayExtents::rank == 0) return 0; else { - std::size_t address = ai[0]; - for(std::size_t i = 1; i < ArrayExtents::rank; i++) + auto address = ai[0]; + for(int i = 1; i < static_cast(ArrayExtents::rank); i++) { address *= extents[i]; address += ai[i]; @@ -69,7 +70,7 @@ namespace llama::mapping struct LinearizeArrayDimsFortran { template - LLAMA_FN_HOST_ACC_INLINE constexpr auto size(const ArrayExtents& extents) -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto size(const ArrayExtents& extents) -> typename ArrayExtents::value_type { return product(extents); } @@ -80,13 +81,13 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto operator()( const typename ArrayExtents::Index& ai, - const ArrayExtents& extents) const -> std::size_t + const ArrayExtents& extents) const -> typename ArrayExtents::value_type { if constexpr(ArrayExtents::rank == 0) return 0; else { - std::size_t address = ai[ArrayExtents::rank - 1]; + auto address = ai[ArrayExtents::rank - 1]; for(int i = static_cast(ArrayExtents::rank) - 2; i >= 0; i--) { address *= extents[i]; @@ -101,17 +102,18 @@ namespace llama::mapping struct LinearizeArrayDimsMorton { template - LLAMA_FN_HOST_ACC_INLINE constexpr auto size(const ArrayExtents& extents) const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto size(const ArrayExtents& extents) const -> + typename ArrayExtents::value_type { if constexpr(ArrayExtents::rank == 0) return 0; else { - std::size_t longest = extents[0]; - for(std::size_t i = 1; i < ArrayExtents::rank; i++) + auto longest = extents[0]; + for(int i = 1; i < static_cast(ArrayExtents::rank); i++) longest = std::max(longest, extents[i]); const auto longestPO2 = bit_ceil(longest); - return intPow(longestPO2, ArrayExtents::rank); + return intPow(longestPO2, static_cast(ArrayExtents::rank)); } } @@ -121,25 +123,29 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto operator()( const typename ArrayExtents::Index& ai, - [[maybe_unused]] const ArrayExtents& extents) const -> std::size_t + [[maybe_unused]] const ArrayExtents& extents) const -> typename ArrayExtents::value_type { - std::size_t r = 0; - for(std::size_t bit = 0; bit < (sizeof(std::size_t) * CHAR_BIT) / ArrayExtents::rank; bit++) - for(std::size_t i = 0; i < ArrayExtents::rank; i++) - r |= (ai[i] & (std::size_t{1} << bit)) << ((bit + 1) * (ArrayExtents::rank - 1) - i); + using size_type = typename ArrayExtents::value_type; + constexpr auto rank = static_cast(ArrayExtents::rank); + size_type r = 0; + for(size_type bit = 0; bit < (static_cast(sizeof(size_type)) * CHAR_BIT) / rank; bit++) + for(size_type i = 0; i < rank; i++) + r |= (ai[i] & (size_type{1} << bit)) << ((bit + 1) * (rank - 1) - i); return r; } private: - LLAMA_FN_HOST_ACC_INLINE static constexpr auto bit_ceil(std::size_t n) -> std::size_t + template + LLAMA_FN_HOST_ACC_INLINE static constexpr auto bit_ceil(T n) -> T { - std::size_t r = 1; + T r = 1u; while(r < n) r <<= 1u; return r; } - LLAMA_FN_HOST_ACC_INLINE static constexpr auto intPow(std::size_t b, std::size_t e) -> std::size_t + template + LLAMA_FN_HOST_ACC_INLINE static constexpr auto intPow(T b, T e) -> T { e--; auto r = b; diff --git a/include/llama/mapping/Heatmap.hpp b/include/llama/mapping/Heatmap.hpp index 1f87464e37..31439070b4 100644 --- a/include/llama/mapping/Heatmap.hpp +++ b/include/llama/mapping/Heatmap.hpp @@ -34,7 +34,7 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE auto blobNrAndOffset( typename Mapping::ArrayIndex ai, - RecordCoord rc = {}) const -> NrAndOffset + RecordCoord rc = {}) const -> NrAndOffset { const auto nao = Mapping::blobNrAndOffset(ai, rc); using Type = GetType>; diff --git a/include/llama/mapping/Null.hpp b/include/llama/mapping/Null.hpp index 3a4b8028c4..d2ceabc34c 100644 --- a/include/llama/mapping/Null.hpp +++ b/include/llama/mapping/Null.hpp @@ -33,6 +33,7 @@ namespace llama::mapping { private: using Base = MappingBase; + using size_type = typename TArrayExtents::value_type; public: static constexpr std::size_t blobCount = 0; @@ -40,7 +41,7 @@ namespace llama::mapping using Base::Base; LLAMA_FN_HOST_ACC_INLINE - constexpr auto blobSize(std::size_t /*blobIndex*/) const -> std::size_t + constexpr auto blobSize(size_type /*blobIndex*/) const -> size_type { return 0; } diff --git a/include/llama/mapping/One.hpp b/include/llama/mapping/One.hpp index 31c45071a5..51319cb7ba 100644 --- a/include/llama/mapping/One.hpp +++ b/include/llama/mapping/One.hpp @@ -24,13 +24,14 @@ namespace llama::mapping { private: using Base = MappingBase; + using size_type = typename Base::size_type; public: static constexpr std::size_t blobCount = 1; using Base::Base; - LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize(std::size_t) const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize(size_type) const -> size_type { return flatSizeOf; // no tail padding } @@ -38,15 +39,16 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset( typename Base::ArrayIndex, - RecordCoord = {}) const -> NrAndOffset + RecordCoord = {}) const -> NrAndOffset { constexpr std::size_t flatFieldIndex = #ifdef __NVCC__ *& // mess with nvcc compiler state to workaround bug #endif Flattener::template flatIndex; - constexpr auto offset = flatOffsetOf; - return {0, offset}; + constexpr auto offset + = static_cast(flatOffsetOf); + return {size_type{0}, offset}; } private: diff --git a/include/llama/mapping/SoA.hpp b/include/llama/mapping/SoA.hpp index 624ccf7f9f..7e49d20e20 100644 --- a/include/llama/mapping/SoA.hpp +++ b/include/llama/mapping/SoA.hpp @@ -27,6 +27,7 @@ namespace llama::mapping private: using Base = MappingBase; using Flattener = FlattenRecordDimSingleBlob; + using size_type = typename TArrayExtents::value_type; public: using LinearizeArrayDimsFunctor = TLinearizeArrayDimsFunctor; @@ -36,13 +37,13 @@ namespace llama::mapping using Base::Base; LLAMA_FN_HOST_ACC_INLINE - constexpr auto blobSize([[maybe_unused]] std::size_t blobIndex) const -> std::size_t + constexpr auto blobSize([[maybe_unused]] size_type blobIndex) const -> size_type { if constexpr(SeparateBuffers) { - constexpr Array typeSizes = []() constexpr + constexpr auto typeSizes = []() constexpr { - Array r{}; + Array r{}; forEachLeafCoord([&r, i = 0](auto rc) mutable constexpr { r[i++] = sizeof(GetType); }); return r; @@ -52,20 +53,20 @@ namespace llama::mapping } else { - return LinearizeArrayDimsFunctor{}.size(Base::extents()) * sizeOf; + return LinearizeArrayDimsFunctor{}.size(Base::extents()) * static_cast(sizeOf); } } template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset( typename Base::ArrayIndex ad, - RecordCoord = {}) const -> NrAndOffset + RecordCoord = {}) const -> NrAndOffset { if constexpr(SeparateBuffers) { constexpr auto blob = flatRecordCoord>; const auto offset = LinearizeArrayDimsFunctor{}(ad, Base::extents()) - * sizeof(GetType>); + * static_cast(sizeof(GetType>)); return {blob, offset}; } else @@ -76,11 +77,9 @@ namespace llama::mapping #endif Flattener::template flatIndex; const auto offset = LinearizeArrayDimsFunctor{}(ad, Base::extents()) - * sizeof(GetType>) - + flatOffsetOf< - typename Flattener::FlatRecordDim, - flatFieldIndex, - false> * LinearizeArrayDimsFunctor{}.size(Base::extents()); + * static_cast(sizeof(GetType>)) + + static_cast(flatOffsetOf) + * LinearizeArrayDimsFunctor{}.size(Base::extents()); return {0, offset}; } } diff --git a/include/llama/mapping/Split.hpp b/include/llama/mapping/Split.hpp index 12a062bef9..42300a576b 100644 --- a/include/llama/mapping/Split.hpp +++ b/include/llama/mapping/Split.hpp @@ -106,6 +106,11 @@ namespace llama::mapping static_assert(SeparateBlobs || Mapping1::blobCount == 1); static_assert(SeparateBlobs || Mapping2::blobCount == 1); + private: + using size_type = typename ArrayExtents::value_type; + static constexpr size_type m1bc = static_cast(Mapping1::blobCount); + + public: constexpr Split() = default; LLAMA_FN_HOST_ACC_INLINE @@ -132,13 +137,13 @@ namespace llama::mapping return mapping1.extents(); } - LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize([[maybe_unused]] std::size_t i) const -> std::size_t + LLAMA_FN_HOST_ACC_INLINE constexpr auto blobSize([[maybe_unused]] size_type i) const -> size_type { if constexpr(SeparateBlobs) { - if(i < Mapping1::blobCount) + if(i < m1bc) return mapping1.blobSize(i); - return mapping2.blobSize(i - Mapping1::blobCount); + return mapping2.blobSize(i - m1bc); } else return mapping1.blobSize(0) + mapping2.blobSize(0); @@ -146,7 +151,7 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE constexpr auto blobNrAndOffset(ArrayIndex ai, RecordCoord = {}) const - -> NrAndOffset + -> NrAndOffset { using Tags = GetTags>; @@ -156,10 +161,10 @@ namespace llama::mapping { auto nrAndOffset = mapping2.blobNrAndOffset(ai, GetCoordFromTags{}); if constexpr(SeparateBlobs) - nrAndOffset.nr += Mapping1::blobCount; + nrAndOffset.nr += m1bc; else { - for(std::size_t i = 0; i < Mapping1::blobCount; i++) + for(size_type i = 0; i < m1bc; i++) nrAndOffset.offset += mapping1.blobSize(i); } return nrAndOffset; @@ -167,7 +172,7 @@ namespace llama::mapping } template - static constexpr auto isComputed(llama::RecordCoord) -> bool + static constexpr auto isComputed(RecordCoord) -> bool { using Tags = GetTags>; if constexpr(internal::isSelected, RecordCoordForMapping1>) @@ -177,10 +182,8 @@ namespace llama::mapping } template - LLAMA_FN_HOST_ACC_INLINE constexpr auto compute( - ArrayIndex ai, - llama::RecordCoord, - Blobs& blobs) const + LLAMA_FN_HOST_ACC_INLINE constexpr auto compute(ArrayIndex ai, RecordCoord, Blobs& blobs) + const { using Tags = GetTags>; if constexpr(internal::isSelected, RecordCoordForMapping1>) @@ -188,7 +191,7 @@ namespace llama::mapping else { // only pass on blobs for mapping 2, so it can index starting from 0 - auto* blobs2 = &blobs[0] + Mapping1::blobCount; + auto* blobs2 = &blobs[0] + m1bc; return mapping2.compute(ai, GetCoordFromTags{}, blobs2); } } diff --git a/include/llama/mapping/Trace.hpp b/include/llama/mapping/Trace.hpp index 95bb5100cb..d56062e9ff 100644 --- a/include/llama/mapping/Trace.hpp +++ b/include/llama/mapping/Trace.hpp @@ -14,6 +14,10 @@ namespace llama::mapping template struct Trace : Mapping { + private: + using size_type = typename Mapping::ArrayExtents::value_type; + + public: using RecordDim = typename Mapping::RecordDim; constexpr Trace() = default; @@ -57,7 +61,7 @@ namespace llama::mapping template LLAMA_FN_HOST_ACC_INLINE auto blobNrAndOffset( typename Mapping::ArrayIndex ai, - RecordCoord rc = {}) const -> NrAndOffset + RecordCoord rc = {}) const -> NrAndOffset { ++fieldHits[flatRecordCoord>]; return Mapping::blobNrAndOffset(ai, rc); diff --git a/include/llama/mapping/tree/Mapping.hpp b/include/llama/mapping/tree/Mapping.hpp index cc544d85a1..f319da7540 100644 --- a/include/llama/mapping/tree/Mapping.hpp +++ b/include/llama/mapping/tree/Mapping.hpp @@ -173,12 +173,16 @@ namespace llama::mapping::tree using ArrayExtents = TArrayExtents; using ArrayIndex = typename ArrayExtents::Index; using RecordDim = TRecordDim; - using BasicTree = TreeFromDimensions; + // TODO(bgruber): , support more than one blob static constexpr std::size_t blobCount = 1; - using MergedFunctors = internal::MergeFunctors; + private: + using size_type = typename ArrayExtents::value_type; + public: + using BasicTree = TreeFromDimensions; + using MergedFunctors = internal::MergeFunctors; BasicTree basicTree; MergedFunctors mergedFunctors; @@ -202,15 +206,17 @@ namespace llama::mapping::tree } LLAMA_FN_HOST_ACC_INLINE - auto blobSize(std::size_t const) const -> std::size_t + auto blobSize(size_type const) const -> size_type { + // TODO(bgruber): propagate use of size_type return internal::getTreeBlobSize(resultTree); } template LLAMA_FN_HOST_ACC_INLINE auto blobNrAndOffset(ArrayIndex ai, RecordCoord = {}) const - -> NrAndOffset + -> NrAndOffset { + // TODO(bgruber): propagate use of size_type auto const basicTreeCoord = createTreeCoord>(ai); auto const resultTreeCoord = mergedFunctors.basicCoordToResultCoord(basicTreeCoord, basicTree); const auto offset = internal::getTreeBlobByte(resultTree, resultTreeCoord); diff --git a/include/llama/mapping/tree/TreeFromDimensions.hpp b/include/llama/mapping/tree/TreeFromDimensions.hpp index bab8ef6be6..835bd6c20e 100644 --- a/include/llama/mapping/tree/TreeFromDimensions.hpp +++ b/include/llama/mapping/tree/TreeFromDimensions.hpp @@ -123,14 +123,14 @@ namespace llama::mapping::tree using TreeFromDimensions = typename internal::WrapInNNodes, ArrayExtents::rank - 1>::type; - template - LLAMA_FN_HOST_ACC_INLINE auto createTree(const ArrayIndex& size) + template + LLAMA_FN_HOST_ACC_INLINE auto createTree(const ArrayIndex& size) { if constexpr(Pos == N - 1) return TreeFromRecordDim{size[N - 1]}; else { - Tuple inner{createTree(size)}; + Tuple inner{createTree(size)}; return Node{size[Pos], inner}; } }; @@ -148,7 +148,8 @@ namespace llama::mapping::tree RecordCoord) { return Tuple{ - TreeCoordElement<(ADIndices == ArrayIndex::rank - 1 ? FirstRecordCoord : 0)>{ai[ADIndices]}..., + TreeCoordElement<(ADIndices == ArrayIndex::rank - 1 ? FirstRecordCoord : 0)>{ + (std::size_t) ai[ADIndices]}..., // TODO TreeCoordElement>{}..., TreeCoordElement<0, boost::mp11::mp_size_t<0>>{}}; } diff --git a/tests/arrayextents.cpp b/tests/arrayextents.cpp index 349290b758..f24d852d16 100644 --- a/tests/arrayextents.cpp +++ b/tests/arrayextents.cpp @@ -1,16 +1,23 @@ #include "common.hpp" -TEST_CASE("ArrayExtents.CTAD") +TEMPLATE_LIST_TEST_CASE("dyn", "", SizeTypes) { + STATIC_REQUIRE(static_cast(-1) == llama::dyn); +} + +TEMPLATE_LIST_TEST_CASE("ArrayExtents.CTAD", "", SizeTypes) +{ + TestType one = 1; + llama::ArrayExtents ad0{}; - llama::ArrayExtents ad1{1}; - llama::ArrayExtents ad2{1, 1}; - llama::ArrayExtents ad3{1, 1, 1}; + llama::ArrayExtents ad1{one}; + llama::ArrayExtents ad2{one, one}; + llama::ArrayExtents ad3{one, one, one}; - STATIC_REQUIRE(std::is_same_v>); - STATIC_REQUIRE(std::is_same_v>); - STATIC_REQUIRE(std::is_same_v>); - STATIC_REQUIRE(std::is_same_v>); + STATIC_REQUIRE(std::is_same_v>); + STATIC_REQUIRE(std::is_same_v>); + STATIC_REQUIRE(std::is_same_v>); + STATIC_REQUIRE(std::is_same_v>); STATIC_REQUIRE(decltype(ad0)::rank == 0); STATIC_REQUIRE(decltype(ad1)::rank == 1); @@ -20,7 +27,7 @@ TEST_CASE("ArrayExtents.CTAD") TEST_CASE("ArrayExtents.dim0") { - auto mapping = llama::mapping::SoA{llama::ArrayExtents<>{}, Particle{}}; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{}, Particle{}}; auto view = llama::allocView(mapping); double& x1 = view(llama::ArrayIndex{})(tag::Pos{}, tag::X{}); double& x2 = view()(tag::Pos{}, tag::X{}); @@ -28,117 +35,174 @@ TEST_CASE("ArrayExtents.dim0") x2 = 0; } -TEST_CASE("ArrayExtents.dim1.dynamic") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim1.dynamic", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents{16}, Particle{}}; + const TestType n = 16; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{n}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim1.static") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim1.static", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents<16>{}, Particle{}}; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim2.dynamic") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim2.dynamic", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents{16, 16}, Particle{}}; + const TestType n = 16; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{n, n}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1, 1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i, i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim2.static") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim2.static", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents<16, 16>{}, Particle{}}; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1, 1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i, i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim3.dynamic") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim3.dynamic", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents{16, 16, 16}, Particle{}}; + const TestType n = 16; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{n, n, n}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1, 1, 1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i, i, i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim3.static") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim3.static", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents<16, 16, 16>{}, Particle{}}; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1, 1, 1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i, i, i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim10.dynamic") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim10.dynamic", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents{2, 2, 2, 2, 2, 2, 2, 2, 2, 2}, Particle{}}; + const TestType n = 2; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{n, n, n, n, n, n, n, n, n, n}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1, 1, 1, 1, 1, 1, 1, 1, 1, 1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i, i, i, i, i, i, i, i, i, i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.dim10.static") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.dim10.static", "", SizeTypes) { - auto mapping = llama::mapping::SoA{llama::ArrayExtents<2, 2, 2, 2, 2, 2, 2, 2, 2, 2>{}, Particle{}}; + const TestType i = 1; + auto mapping = llama::mapping::SoA{llama::ArrayExtents{}, Particle{}}; auto view = llama::allocView(mapping); - double& x = view(llama::ArrayIndex{1, 1, 1, 1, 1, 1, 1, 1, 1, 1})(tag::Pos{}, tag::X{}); + double& x = view(llama::ArrayIndex{i, i, i, i, i, i, i, i, i, i})(tag::Pos{}, tag::X{}); x = 0; } -TEST_CASE("ArrayExtents.ctor") +TEST_CASE("ArrayExtentsDynamic") { - llama::ArrayExtentsDynamic<1> extents{}; + STATIC_REQUIRE(std::is_same_v< + llama::ArrayExtentsDynamic, + llama::ArrayExtents>); + STATIC_REQUIRE(std::is_same_v< + llama::ArrayExtentsDynamic, + llama::ArrayExtents>); +} + +TEST_CASE("ArrayExtentsStatic") +{ + STATIC_REQUIRE( + std::is_same_v, llama::ArrayExtents>); + STATIC_REQUIRE(std::is_same_v< + llama::ArrayExtentsNCube, + llama::ArrayExtents>); + STATIC_REQUIRE( + std::is_same_v, llama::ArrayExtents>); +} + +TEST_CASE("ArrayExtents.ctor_value_init") +{ + [[maybe_unused]] const llama::ArrayExtentsDynamic extents{}; CHECK(extents[0] == 0); } -TEST_CASE("ArrayExtents.toArray") +TEMPLATE_LIST_TEST_CASE("ArrayExtents.toArray", "", SizeTypes) { - CHECK(llama::ArrayExtents{}.toArray() == llama::Array{}); - CHECK(llama::ArrayExtents{42}.toArray() == llama::Array{42}); - CHECK(llama::ArrayExtents{42, 43}.toArray() == llama::Array{42, 43}); + CHECK(llama::ArrayExtents{}.toArray() == llama::Array{}); + + // dynamic + CHECK(llama::ArrayExtents{42}.toArray() == llama::Array{42}); + CHECK( + llama::ArrayExtents{42, 43}.toArray() == llama::Array{42, 43}); + CHECK( + llama::ArrayExtents{42, 43, 44}.toArray() + == llama::Array{42, 43, 44}); + + // static + CHECK(llama::ArrayExtents{}.toArray() == llama::Array{42}); + CHECK(llama::ArrayExtents{}.toArray() == llama::Array{42, 43}); + CHECK( + llama::ArrayExtents{42, 44}.toArray() + == llama::Array{42, 43, 44}); + + // mixed + CHECK( + llama::ArrayExtents{43, 44}.toArray() + == llama::Array{42, 43, 44}); CHECK( - llama::ArrayExtents{42, 44}.toArray() == llama::Array{42, 43, 44}); + llama::ArrayExtents{42, 44}.toArray() + == llama::Array{42, 43, 44}); + CHECK(llama::ArrayExtents{42}.toArray() == llama::Array{42, 43, 44}); } -TEST_CASE("forEachADCoord_1D") +TEMPLATE_LIST_TEST_CASE("forEachADCoord_1D", "", SizeTypes) { - llama::ArrayExtents extents{3}; + const auto n = TestType{3}; + llama::ArrayExtents extents{n}; - std::vector> indices; - llama::forEachADCoord(extents, [&](llama::ArrayIndex<1> ai) { indices.push_back(ai); }); + std::vector> indices; + llama::forEachADCoord(extents, [&](llama::ArrayIndex ai) { indices.push_back(ai); }); - CHECK(indices == std::vector>{{0}, {1}, {2}}); + CHECK(indices == std::vector>{{0}, {1}, {2}}); } -TEST_CASE("forEachADCoord_2D") +TEMPLATE_LIST_TEST_CASE("forEachADCoord_2D", "", SizeTypes) { - llama::ArrayExtents extents{3, 3}; + const auto n = TestType{3}; + llama::ArrayExtents extents{n, n}; - std::vector> indices; - llama::forEachADCoord(extents, [&](llama::ArrayIndex<2> ai) { indices.push_back(ai); }); + std::vector> indices; + llama::forEachADCoord(extents, [&](llama::ArrayIndex ai) { indices.push_back(ai); }); CHECK( indices - == std::vector>{{0, 0}, {0, 1}, {0, 2}, {1, 0}, {1, 1}, {1, 2}, {2, 0}, {2, 1}, {2, 2}}); + == std::vector< + llama::ArrayIndex>{{0, 0}, {0, 1}, {0, 2}, {1, 0}, {1, 1}, {1, 2}, {2, 0}, {2, 1}, {2, 2}}); } -TEST_CASE("forEachADCoord_3D") +TEMPLATE_LIST_TEST_CASE("forEachADCoord_3D", "", SizeTypes) { - llama::ArrayExtents extents{3, 3, 3}; + const auto n = TestType{3}; + llama::ArrayExtents extents{n, n, n}; - std::vector> indices; - llama::forEachADCoord(extents, [&](llama::ArrayIndex<3> ai) { indices.push_back(ai); }); + std::vector> indices; + llama::forEachADCoord(extents, [&](llama::ArrayIndex ai) { indices.push_back(ai); }); CHECK( indices - == std::vector>{ + == std::vector>{ {0, 0, 0}, {0, 0, 1}, {0, 0, 2}, {0, 1, 0}, {0, 1, 1}, {0, 1, 2}, {0, 2, 0}, {0, 2, 1}, {0, 2, 2}, {1, 0, 0}, {1, 0, 1}, {1, 0, 2}, {1, 1, 0}, {1, 1, 1}, {1, 1, 2}, {1, 2, 0}, {1, 2, 1}, {1, 2, 2}, {2, 0, 0}, {2, 0, 1}, {2, 0, 2}, {2, 1, 0}, {2, 1, 1}, {2, 1, 2}, {2, 2, 0}, {2, 2, 1}, {2, 2, 2}, diff --git a/tests/arrayindexrange.cpp b/tests/arrayindexrange.cpp index 6e9f247d52..8959f79203 100644 --- a/tests/arrayindexrange.cpp +++ b/tests/arrayindexrange.cpp @@ -2,7 +2,7 @@ TEST_CASE("ArrayIndexIterator.concepts") { - using ArrayExtents = llama::ArrayExtentsDynamic<3>; + using ArrayExtents = llama::ArrayExtentsDynamic; STATIC_REQUIRE(std::is_same_v< std::iterator_traits>::iterator_category, std::random_access_iterator_tag>); @@ -14,7 +14,7 @@ TEST_CASE("ArrayIndexIterator.concepts") TEST_CASE("ArrayIndexIterator") { - llama::ArrayIndexRange r{llama::ArrayExtentsDynamic<2>{3, 3}}; + llama::ArrayIndexRange r{llama::ArrayExtentsDynamic{3, 3}}; llama::ArrayIndexIterator it = std::begin(r); CHECK(*it == llama::ArrayIndex{0, 0}); @@ -69,12 +69,12 @@ TEST_CASE("ArrayIndexIterator") TEST_CASE("ArrayIndexIterator.operator+=") { - std::vector> indices; - llama::ArrayIndexRange r{llama::ArrayExtentsDynamic<2>{3, 4}}; + std::vector> indices; + llama::ArrayIndexRange r{llama::ArrayExtentsDynamic{3, 4}}; for(auto it = r.begin(); it != r.end(); it += 2) indices.push_back(*it); - CHECK(indices == std::vector>{{0, 0}, {0, 2}, {1, 0}, {1, 2}, {2, 0}, {2, 2}}); + CHECK(indices == std::vector>{{0, 0}, {0, 2}, {1, 0}, {1, 2}, {2, 0}, {2, 2}}); } TEST_CASE("ArrayIndexIterator.constexpr") @@ -82,7 +82,7 @@ TEST_CASE("ArrayIndexIterator.constexpr") constexpr auto r = [&]() constexpr { bool b = true; - llama::ArrayIndexIterator it = std::begin(llama::ArrayIndexRange{llama::ArrayExtentsDynamic<2>{3, 3}}); + llama::ArrayIndexIterator it = std::begin(llama::ArrayIndexRange{llama::ArrayExtentsDynamic{3, 3}}); b &= *it == llama::ArrayIndex{0, 0}; it++; b &= *it == llama::ArrayIndex{0, 1}; @@ -96,33 +96,34 @@ TEST_CASE("ArrayIndexIterator.constexpr") TEST_CASE("ArrayIndexRange.1D") { - std::vector> indices; - for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic<1>{3}}) + std::vector> indices; + for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic{3}}) indices.push_back(ai); - CHECK(indices == std::vector>{{0}, {1}, {2}}); + CHECK(indices == std::vector>{{0}, {1}, {2}}); } TEST_CASE("ArrayIndexRange.2D") { - std::vector> indices; - for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic<2>{3, 3}}) + std::vector> indices; + for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic{3, 3}}) indices.push_back(ai); CHECK( indices - == std::vector>{{0, 0}, {0, 1}, {0, 2}, {1, 0}, {1, 1}, {1, 2}, {2, 0}, {2, 1}, {2, 2}}); + == std::vector< + llama::ArrayIndex>{{0, 0}, {0, 1}, {0, 2}, {1, 0}, {1, 1}, {1, 2}, {2, 0}, {2, 1}, {2, 2}}); } TEST_CASE("ArrayIndexRange.3D") { - std::vector> indices; - for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic<3>{3, 3, 3}}) + std::vector> indices; + for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic{3, 3, 3}}) indices.push_back(ai); CHECK( indices - == std::vector>{ + == std::vector>{ {0, 0, 0}, {0, 0, 1}, {0, 0, 2}, {0, 1, 0}, {0, 1, 1}, {0, 1, 2}, {0, 2, 0}, {0, 2, 1}, {0, 2, 2}, {1, 0, 0}, {1, 0, 1}, {1, 0, 2}, {1, 1, 0}, {1, 1, 1}, {1, 1, 2}, {1, 2, 0}, {1, 2, 1}, {1, 2, 2}, {2, 0, 0}, {2, 0, 1}, {2, 0, 2}, {2, 1, 0}, {2, 1, 1}, {2, 1, 2}, {2, 2, 0}, {2, 2, 1}, {2, 2, 2}, @@ -134,7 +135,7 @@ TEST_CASE("ArrayIndexRange.3D") TEST_CASE("ArrayIndexRange.concepts") { - using ArrayExtents = llama::ArrayExtentsDynamic<3>; + using ArrayExtents = llama::ArrayExtentsDynamic; STATIC_REQUIRE(std::ranges::range>); STATIC_REQUIRE(std::ranges::random_access_range>); // STATIC_REQUIRE(std::ranges::view>); @@ -142,15 +143,15 @@ TEST_CASE("ArrayIndexRange.concepts") TEST_CASE("ArrayIndexRange.3D.reverse") { - llama::ArrayExtentsDynamic<3> extents{3, 3, 3}; + llama::ArrayExtentsDynamic extents{3, 3, 3}; - std::vector> indices; + std::vector> indices; for(auto ai : llama::ArrayIndexRange{extents} | std::views::reverse) indices.push_back(ai); CHECK( indices - == std::vector>{ + == std::vector>{ {{2, 2, 2}, {2, 2, 1}, {2, 2, 0}, {2, 1, 2}, {2, 1, 1}, {2, 1, 0}, {2, 0, 2}, {2, 0, 1}, {2, 0, 0}, {1, 2, 2}, {1, 2, 1}, {1, 2, 0}, {1, 1, 2}, {1, 1, 1}, {1, 1, 0}, {1, 0, 2}, {1, 0, 1}, {1, 0, 0}, {0, 2, 2}, {0, 2, 1}, {0, 2, 0}, {0, 1, 2}, {0, 1, 1}, {0, 1, 0}, {0, 0, 2}, {0, 0, 1}, {0, 0, 0}}}); @@ -162,13 +163,13 @@ TEST_CASE("ArrayIndexRange.1D.constexpr") constexpr auto r = []() constexpr { int i = 0; - for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic<1>{3}}) + for(auto ai : llama::ArrayIndexRange{llama::ArrayExtentsDynamic{3}}) { - if(i == 0 && ai != llama::ArrayIndex<1>{0}) + if(i == 0 && ai != llama::ArrayIndex{0}) return false; - if(i == 1 && ai != llama::ArrayIndex<1>{1}) + if(i == 1 && ai != llama::ArrayIndex{1}) return false; - if(i == 2 && ai != llama::ArrayIndex<1>{2}) + if(i == 2 && ai != llama::ArrayIndex{2}) return false; i++; } @@ -181,7 +182,7 @@ TEST_CASE("ArrayIndexRange.1D.constexpr") TEST_CASE("ArrayIndexRange.3D.destructering") { - for(auto [x, y, z] : llama::ArrayIndexRange{llama::ArrayExtentsDynamic<3>{1, 1, 1}}) + for(auto [x, y, z] : llama::ArrayIndexRange{llama::ArrayExtentsDynamic{1, 1, 1}}) { CHECK(x == 0); CHECK(y == 0); diff --git a/tests/bloballocators.cpp b/tests/bloballocators.cpp index f176abcf34..5cfcf57bd9 100644 --- a/tests/bloballocators.cpp +++ b/tests/bloballocators.cpp @@ -1,6 +1,6 @@ #include "common.hpp" -using ArrayExtents = llama::ArrayExtents<16>; +using ArrayExtents = llama::ArrayExtents; using Mapping = llama::mapping::AoS; TEST_CASE("bloballocators.Stack") diff --git a/tests/common.hpp b/tests/common.hpp index 15b8ac1304..ba20d96b7e 100644 --- a/tests/common.hpp +++ b/tests/common.hpp @@ -8,6 +8,9 @@ #include #include +using SizeTypes = boost::mp11::mp_list; +static_assert(boost::mp11::mp_contains::value); + // clang-format off namespace tag { diff --git a/tests/computedprop.cpp b/tests/computedprop.cpp index fe1a780905..4c35059e79 100644 --- a/tests/computedprop.cpp +++ b/tests/computedprop.cpp @@ -30,7 +30,7 @@ namespace llama::RecordCoord, llama::Array& storageBlobs) const { - auto fetch = [&](llama::NrAndOffset nrAndOffset) -> double + auto fetch = [&](llama::NrAndOffset nrAndOffset) -> double { return *reinterpret_cast(&storageBlobs[nrAndOffset.nr][nrAndOffset.offset]); }; const auto ax = fetch(Base::template blobNrAndOffset<0, 0>(ai)); @@ -81,7 +81,7 @@ namespace TEST_CASE("computedprop") { - auto extents = llama::ArrayExtentsDynamic<1>{10}; + auto extents = llama::ArrayExtentsDynamic{10}; auto mapping = AoSWithComputedNormal{extents}; STATIC_REQUIRE(mapping.blobCount == 1); @@ -146,7 +146,7 @@ namespace TEST_CASE("fully_computed_mapping") { - auto mapping = ComputedMapping, int>{{}}; + auto mapping = ComputedMapping, int>{{}}; auto view = llama::allocViewUninitialized(mapping); using namespace tag; @@ -242,7 +242,7 @@ namespace TEST_CASE("compressed_bools") { - auto mapping = CompressedBoolMapping, BoolRecord>{{}}; + auto mapping = CompressedBoolMapping, BoolRecord>{{}}; STATIC_REQUIRE(decltype(mapping)::blobCount == 3); CHECK(mapping.blobSize(0) == 8); CHECK(mapping.blobSize(1) == 8); diff --git a/tests/copy.cpp b/tests/copy.cpp index e58b87eeed..704a2522a7 100644 --- a/tests/copy.cpp +++ b/tests/copy.cpp @@ -2,7 +2,7 @@ namespace { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; using RecordDim = Vec3I; template diff --git a/tests/dump.cpp b/tests/dump.cpp index f13fce15c2..f3cf40c94a 100644 --- a/tests/dump.cpp +++ b/tests/dump.cpp @@ -8,7 +8,7 @@ #if !(defined(__APPLE__) && __clang_major__ == 13 && __clang_minor__ == 0) namespace { - llama::ArrayExtentsDynamic<1> extents{32}; + llama::ArrayExtentsDynamic extents{32}; using ArrayExtents = decltype(extents); template diff --git a/tests/iterator.cpp b/tests/iterator.cpp index 08e2d655f2..060a60cff4 100644 --- a/tests/iterator.cpp +++ b/tests/iterator.cpp @@ -8,7 +8,7 @@ using Position = Vec3I; TEST_CASE("iterator.concepts") { - using Mapping = llama::mapping::AoS, Position>; + using Mapping = llama::mapping::AoS, Position>; using View = llama::View; using Iterator = typename View::iterator; @@ -251,7 +251,7 @@ TEST_CASE("ranges") TEST_CASE("iterator.sort") { constexpr auto N = 10; - auto view = llama::allocViewUninitialized(llama::mapping::AoS{llama::ArrayExtents{}, Position{}}); + auto view = llama::allocViewUninitialized(llama::mapping::AoS{llama::ArrayExtents{}, Position{}}); std::default_random_engine e{}; std::uniform_int_distribution d{0, 1000}; diff --git a/tests/mapping.AoS.cpp b/tests/mapping.AoS.cpp index ac5bd141a8..60433483e4 100644 --- a/tests/mapping.AoS.cpp +++ b/tests/mapping.AoS.cpp @@ -53,10 +53,10 @@ TEST_CASE("mapping.AoS.Packed.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 951); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.AoS.Packed.fortran.address") @@ -113,10 +113,10 @@ TEST_CASE("mapping.AoS.Packed.fortran.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 111); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.AoS.Packed.morton.address") @@ -173,10 +173,10 @@ TEST_CASE("mapping.AoS.Packed.morton.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 167); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.AoS.Aligned.address") @@ -232,10 +232,10 @@ TEST_CASE("mapping.AoS.Aligned.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 1083); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.AoS.aligned_min.address") @@ -291,8 +291,8 @@ TEST_CASE("mapping.AoS.aligned_min.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 899); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } diff --git a/tests/mapping.AoSoA.cpp b/tests/mapping.AoSoA.cpp index 4d4868dbae..f51e96a1ca 100644 --- a/tests/mapping.AoSoA.cpp +++ b/tests/mapping.AoSoA.cpp @@ -74,15 +74,15 @@ TEST_CASE("mapping.AoSoA.4.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 1116); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("AoSoA.size_round_up") { - using AoSoA = llama::mapping::AoSoA, Particle, 4>; + using AoSoA = llama::mapping::AoSoA, Particle, 4>; constexpr auto psize = llama::sizeOf; CHECK(AoSoA{{0}}.blobSize(0) == 0 * psize); @@ -99,7 +99,7 @@ TEST_CASE("AoSoA.size_round_up") TEST_CASE("AoSoA.address_within_bounds") { - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic; using AoSoA = llama::mapping::AoSoA; const auto ad = ArrayExtents{3}; diff --git a/tests/mapping.BitPackedFloatSoA.cpp b/tests/mapping.BitPackedFloatSoA.cpp index 9e342c18c2..9c08cca577 100644 --- a/tests/mapping.BitPackedFloatSoA.cpp +++ b/tests/mapping.BitPackedFloatSoA.cpp @@ -132,7 +132,7 @@ TEMPLATE_TEST_CASE("mapping.BitPackedFloatSoA", "", float, double) TEST_CASE("mapping.BitPackedFloatSoA.ReducedPrecisionComputation") { constexpr auto N = 1000; - auto view = llama::allocView(llama::mapping::AoS, Vec3D>{{}}); + auto view = llama::allocView(llama::mapping::AoS, Vec3D>{{}}); std::default_random_engine engine; std::uniform_real_distribution dist{0.0f, 1e20f}; for(auto i = 0; i < N; i++) @@ -143,10 +143,12 @@ TEST_CASE("mapping.BitPackedFloatSoA.ReducedPrecisionComputation") } // copy into packed representation - auto packedView = llama::allocView( - llama::mapping:: - BitPackedFloatSoA, Vec3D, llama::Constant<8>, llama::Constant<23>>{}); // basically - // float + auto packedView = llama::allocView(llama::mapping::BitPackedFloatSoA< + llama::ArrayExtents, + Vec3D, + llama::Constant<8>, + llama::Constant<23>>{}); // basically + // float llama::copy(view, packedView); // compute on original representation @@ -169,18 +171,24 @@ TEST_CASE("mapping.BitPackedFloatSoA.ReducedPrecisionComputation") TEST_CASE("mapping.BitPackedFloatSoA.Size") { - STATIC_REQUIRE(std::is_empty_v< - llama::mapping:: - BitPackedFloatSoA, float, llama::Constant<7>, llama::Constant<16>>>); + STATIC_REQUIRE(std::is_empty_v, + float, + llama::Constant<7>, + llama::Constant<16>>>); STATIC_REQUIRE( - sizeof(llama::mapping::BitPackedFloatSoA, float, llama::Constant<7>, unsigned>{ - {}, - {}, - 16}) + sizeof(llama::mapping:: + BitPackedFloatSoA, float, llama::Constant<7>, unsigned>{ + {}, + {}, + 16}) == sizeof(unsigned)); STATIC_REQUIRE( - sizeof(llama::mapping::BitPackedFloatSoA, float, unsigned, llama::Constant<16>>{{}, 7}) + sizeof( + llama::mapping:: + BitPackedFloatSoA, float, unsigned, llama::Constant<16>>{{}, 7}) == sizeof(unsigned)); STATIC_REQUIRE( - sizeof(llama::mapping::BitPackedFloatSoA, float>{{}, 7, 16}) == 2 * sizeof(unsigned)); + sizeof(llama::mapping::BitPackedFloatSoA, float>{{}, 7, 16}) + == 2 * sizeof(unsigned)); } diff --git a/tests/mapping.BitPackedIntSoA.cpp b/tests/mapping.BitPackedIntSoA.cpp index 7e68056372..1240de1a75 100644 --- a/tests/mapping.BitPackedIntSoA.cpp +++ b/tests/mapping.BitPackedIntSoA.cpp @@ -17,7 +17,7 @@ TEST_CASE("mapping.BitPackedIntSoA.Constant.SInts") { // 16 elements * 4 fields = 64 integers, iota produces [0;63], which fits int8_t and into 7 bits auto view = llama::allocView( - llama::mapping::BitPackedIntSoA, SInts, llama::Constant<7>>{{16}}); + llama::mapping::BitPackedIntSoA, SInts, llama::Constant<7>>{{16}}); iotaFillView(view); iotaCheckView(view); } @@ -25,7 +25,8 @@ TEST_CASE("mapping.BitPackedIntSoA.Constant.SInts") TEST_CASE("mapping.BitPackedIntSoA.Value.SInts") { // 16 elements * 4 fields = 64 integers, iota produces [0;63], which fits int8_t and into 7 bits - auto view = llama::allocView(llama::mapping::BitPackedIntSoA, SInts>{{16}, 7}); + auto view = llama::allocView( + llama::mapping::BitPackedIntSoA, SInts>{{16}, 7}); iotaFillView(view); iotaCheckView(view); } @@ -34,7 +35,7 @@ TEST_CASE("mapping.BitPackedIntSoA.Constant.UInts") { // 32 elements * 4 fields = 128 integers, iota produces [0;127], which fits uint8_t and into 7 bits auto view = llama::allocView( - llama::mapping::BitPackedIntSoA, UInts, llama::Constant<7>>{{32}}); + llama::mapping::BitPackedIntSoA, UInts, llama::Constant<7>>{{32}}); iotaFillView(view); iotaCheckView(view); } @@ -42,7 +43,8 @@ TEST_CASE("mapping.BitPackedIntSoA.Constant.UInts") TEST_CASE("mapping.BitPackedIntSoA.Value.UInts") { // 32 elements * 4 fields = 128 integers, iota produces [0;127], which fits uint8_t and into 7 bits - auto view = llama::allocView(llama::mapping::BitPackedIntSoA, UInts>{{32}, 7}); + auto view = llama::allocView( + llama::mapping::BitPackedIntSoA, UInts>{{32}, 7}); iotaFillView(view); iotaCheckView(view); } @@ -99,15 +101,15 @@ TEST_CASE("mapping.BitPackedIntSoA.SInts.Cutoff") TEST_CASE("mapping.BitPackedIntSoA.SInts.Roundtrip") { constexpr auto N = 1000; - auto view = llama::allocView(llama::mapping::AoS, Vec3I>{}); + auto view = llama::allocView(llama::mapping::AoS, Vec3I>{}); std::default_random_engine engine; std::uniform_int_distribution dist{-2000, 2000}; // fits into 12 bits for(auto i = 0; i < N; i++) view(i) = dist(engine); // copy into packed representation - auto packedView - = llama::allocView(llama::mapping::BitPackedIntSoA, Vec3I, llama::Constant<12>>{}); + auto packedView = llama::allocView( + llama::mapping::BitPackedIntSoA, Vec3I, llama::Constant<12>>{}); llama::copy(view, packedView); // compute on packed representation @@ -115,7 +117,7 @@ TEST_CASE("mapping.BitPackedIntSoA.SInts.Roundtrip") packedView(i) = packedView(i) + 1; // copy into normal representation - auto view2 = llama::allocView(llama::mapping::AoS, Vec3I>{}); + auto view2 = llama::allocView(llama::mapping::AoS, Vec3I>{}); llama::copy(packedView, view2); // compute on normal representation @@ -130,7 +132,8 @@ TEST_CASE("mapping.BitPackedIntSoA.bool") { // pack 32 bools into 4 bytes const auto n = 32; - const auto mapping = llama::mapping::BitPackedIntSoA, bool, llama::Constant<1>>{{n}}; + const auto mapping + = llama::mapping::BitPackedIntSoA, bool, llama::Constant<1>>{{n}}; CHECK(mapping.blobSize(0) == n / CHAR_BIT); auto view = llama::allocView(mapping); for(auto i = 0; i < n; i++) @@ -171,7 +174,7 @@ TEMPLATE_TEST_CASE("mapping.BitPackedIntSoA.Enum", "", Grades, GradesClass) STATIC_REQUIRE(std::is_same_v); auto view = llama::allocView( - llama::mapping::BitPackedIntSoA, Enum, llama::Constant<3>>{{6}}); + llama::mapping::BitPackedIntSoA, Enum, llama::Constant<3>>{{6}}); view(0) = Enum::A; view(1) = Enum::B; view(2) = Enum::C; @@ -189,7 +192,9 @@ TEMPLATE_TEST_CASE("mapping.BitPackedIntSoA.Enum", "", Grades, GradesClass) TEST_CASE("mapping.BitPackedIntSoA.Size") { + STATIC_REQUIRE(std::is_empty_v< + llama::mapping::BitPackedIntSoA, SInts, llama::Constant<7>>>); STATIC_REQUIRE( - std::is_empty_v, SInts, llama::Constant<7>>>); - STATIC_REQUIRE(sizeof(llama::mapping::BitPackedIntSoA, SInts>{{}, 7}) == sizeof(unsigned)); + sizeof(llama::mapping::BitPackedIntSoA, SInts>{{}, 7}) + == sizeof(unsigned)); } diff --git a/tests/mapping.Bytesplit.cpp b/tests/mapping.Bytesplit.cpp index c2f6848f3c..148c8c441a 100644 --- a/tests/mapping.Bytesplit.cpp +++ b/tests/mapping.Bytesplit.cpp @@ -5,7 +5,8 @@ TEST_CASE("mapping.ByteSplit.AoS") { auto view = llama::allocView( - llama::mapping::Bytesplit, Vec3I, llama::mapping::BindAoS<>::fn>{{128}}); + llama::mapping::Bytesplit, Vec3I, llama::mapping::BindAoS<>::fn>{ + {128}}); iotaFillView(view); iotaCheckView(view); } @@ -13,7 +14,8 @@ TEST_CASE("mapping.ByteSplit.AoS") TEST_CASE("mapping.ByteSplit.SoA") { auto view = llama::allocView( - llama::mapping::Bytesplit, Vec3I, llama::mapping::BindSoA<>::fn>{{128}}); + llama::mapping::Bytesplit, Vec3I, llama::mapping::BindSoA<>::fn>{ + {128}}); iotaFillView(view); iotaCheckView(view); } @@ -21,7 +23,8 @@ TEST_CASE("mapping.ByteSplit.SoA") TEST_CASE("mapping.ByteSplit.AoSoA") { auto view = llama::allocView( - llama::mapping::Bytesplit, Vec3I, llama::mapping::BindAoSoA<16>::fn>{{128}}); + llama::mapping:: + Bytesplit, Vec3I, llama::mapping::BindAoSoA<16>::fn>{{128}}); iotaFillView(view); iotaCheckView(view); } @@ -29,7 +32,7 @@ TEST_CASE("mapping.ByteSplit.AoSoA") TEST_CASE("mapping.ByteSplit.ChangeType.SoA") { using Mapping = llama::mapping::Bytesplit< - llama::ArrayExtentsDynamic<1>, + llama::ArrayExtentsDynamic, Vec3I, llama::mapping::BindChangeType< llama::mapping::BindSoA<>::fn, @@ -57,7 +60,7 @@ TEST_CASE("mapping.ByteSplit.ChangeType.SoA") TEST_CASE("mapping.ByteSplit.Split.BitPackedIntSoA") { auto view = llama::allocView(llama::mapping::Bytesplit< - llama::ArrayExtentsDynamic<1>, + llama::ArrayExtentsDynamic, Vec3I, llama::mapping::BindSplit< llama::RecordCoord<1>, @@ -71,7 +74,8 @@ TEST_CASE("mapping.ByteSplit.Split.BitPackedIntSoA") TEST_CASE("mapping.ByteSplit.SoA.verify") { - using Mapping = llama::mapping::Bytesplit, Vec3I, llama::mapping::BindSoA<>::fn>; + using Mapping + = llama::mapping::Bytesplit, Vec3I, llama::mapping::BindSoA<>::fn>; auto view = llama::allocView(Mapping{{128}}); for(auto i = 0; i < 128; i++) view(i) = i; diff --git a/tests/mapping.ChangeType.cpp b/tests/mapping.ChangeType.cpp index cb3dfb81bb..e6699593a6 100644 --- a/tests/mapping.ChangeType.cpp +++ b/tests/mapping.ChangeType.cpp @@ -4,8 +4,11 @@ TEST_CASE("mapping.ChangeType.AoS") { - auto mapping = llama::mapping:: - ChangeType, Vec3D, llama::mapping::BindAoS::fn, boost::mp11::mp_list<>>{{}}; + auto mapping = llama::mapping::ChangeType< + llama::ArrayExtents, + Vec3D, + llama::mapping::BindAoS::fn, + boost::mp11::mp_list<>>{{}}; CHECK(mapping.blobSize(0) == 3072); auto view = llama::allocView(mapping); iotaFillView(view); @@ -15,7 +18,7 @@ TEST_CASE("mapping.ChangeType.AoS") TEST_CASE("mapping.ChangeType.AoS.doubleToFloat") { auto mapping = llama::mapping::ChangeType< - llama::ArrayExtents<128>, + llama::ArrayExtents, Vec3D, llama::mapping::BindAoS::fn, boost::mp11::mp_list>>{{}}; @@ -28,7 +31,7 @@ TEST_CASE("mapping.ChangeType.AoS.doubleToFloat") TEST_CASE("mapping.ChangeType.AoS.Coord1ToFloat") { auto mapping = llama::mapping::ChangeType< - llama::ArrayExtents<128>, + llama::ArrayExtents, Vec3D, llama::mapping::BindAoS::fn, boost::mp11::mp_list, float>>>{{}}; @@ -41,7 +44,7 @@ TEST_CASE("mapping.ChangeType.AoS.Coord1ToFloat") TEST_CASE("mapping.ChangeType.AoS.CoordsAndTypes") { auto mapping = llama::mapping::ChangeType< - llama::ArrayExtents<128>, + llama::ArrayExtents, Vec3D, llama::mapping::BindAoS::fn, boost::mp11:: @@ -54,11 +57,11 @@ TEST_CASE("mapping.ChangeType.AoS.CoordsAndTypes") TEST_CASE("mapping.ChangeType.SoA.particle.types") { - auto mappingAoS = llama::mapping::PackedAoS, Particle>{{}}; + auto mappingAoS = llama::mapping::PackedAoS, Particle>{{}}; CHECK(mappingAoS.blobSize(0) == 128 * (6 * sizeof(double) + 1 * sizeof(float) + 4 * sizeof(bool))); auto mapping = llama::mapping::ChangeType< - llama::ArrayExtents<128>, + llama::ArrayExtents, Particle, llama::mapping::BindAoS::fn, boost::mp11::mp_list< @@ -75,7 +78,7 @@ TEST_CASE("mapping.ChangeType.SoA.particle.coords") { using namespace boost::mp11; auto mapping = llama::mapping::ChangeType< - llama::ArrayExtents<128>, + llama::ArrayExtents, Particle, llama::mapping::BindAoS::fn, mp_list< diff --git a/tests/mapping.HeatmapTrace.cpp b/tests/mapping.HeatmapTrace.cpp index 905285a0a1..cf0b70dbaa 100644 --- a/tests/mapping.HeatmapTrace.cpp +++ b/tests/mapping.HeatmapTrace.cpp @@ -49,10 +49,12 @@ TEST_CASE("Heatmap.nbody") std::ofstream{"Heatmap." + name + ".sh"} << script; CHECK(script == expectedScript); }; - run("AlignedAoS", heatmapAlignedAoS, llama::mapping::AlignedAoS, ParticleHeatmap>{}); + run("AlignedAoS", + heatmapAlignedAoS, + llama::mapping::AlignedAoS, ParticleHeatmap>{}); run("SingleBlobSoA", heatmapSingleBlobSoA, - llama::mapping::SingleBlobSoA, ParticleHeatmap>{}); + llama::mapping::SingleBlobSoA, ParticleHeatmap>{}); } TEST_CASE("Trace.nbody") @@ -71,8 +73,8 @@ TEST_CASE("Trace.nbody") CHECK(hits[5] == 400); CHECK(hits[6] == 10300); }; - run(llama::mapping::AlignedAoS, ParticleHeatmap>{}); - run(llama::mapping::SingleBlobSoA, ParticleHeatmap>{}); + run(llama::mapping::AlignedAoS, ParticleHeatmap>{}); + run(llama::mapping::SingleBlobSoA, ParticleHeatmap>{}); } TEST_CASE("Trace.print_dtor") @@ -81,7 +83,7 @@ TEST_CASE("Trace.print_dtor") std::streambuf* old = std::cout.rdbuf(buffer.rdbuf()); { auto particles = llama::allocView( - llama::mapping::Trace{llama::mapping::AlignedAoS, ParticleHeatmap>{}}); + llama::mapping::Trace{llama::mapping::AlignedAoS, ParticleHeatmap>{}}); updateAndMove(particles); } std::cout.rdbuf(old); diff --git a/tests/mapping.Null.cpp b/tests/mapping.Null.cpp index 1068963f35..d14c8f11c1 100644 --- a/tests/mapping.Null.cpp +++ b/tests/mapping.Null.cpp @@ -4,7 +4,7 @@ TEST_CASE("mapping.Null") { - auto mapping = llama::mapping::Null, Particle>{{}}; + auto mapping = llama::mapping::Null, Particle>{{}}; STATIC_REQUIRE(decltype(mapping)::blobCount == 0); auto view = llama::allocView(mapping); diff --git a/tests/mapping.One.cpp b/tests/mapping.One.cpp index 4c5a41ebb5..ad642e92b6 100644 --- a/tests/mapping.One.cpp +++ b/tests/mapping.One.cpp @@ -24,10 +24,10 @@ TEST_CASE("mapping.One.Packed.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 55); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.One.Aligned.address") @@ -54,10 +54,10 @@ TEST_CASE("mapping.One.Aligned.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 59); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.One.MinAligned.address") @@ -84,8 +84,8 @@ TEST_CASE("mapping.One.MinAligned.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 3); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } diff --git a/tests/mapping.SoA.cpp b/tests/mapping.SoA.cpp index a39c8b3c80..3981352100 100644 --- a/tests/mapping.SoA.cpp +++ b/tests/mapping.SoA.cpp @@ -53,10 +53,10 @@ TEST_CASE("mapping.SoA.SingleBlob.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 14096); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.SoA.SingleBlob.fortran.address") @@ -113,10 +113,10 @@ TEST_CASE("mapping.SoA.SingleBlob.fortran.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 14081); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.SoA.SingleBlob.morton.address") @@ -177,10 +177,10 @@ TEST_CASE("mapping.SoA.SingleBlob.morton.address") CHECK(mapping.template blobNrAndOffset<3, 3>(ai).offset == 14082); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("mapping.SoA.MultiBlob.address") @@ -190,54 +190,55 @@ TEST_CASE("mapping.SoA.MultiBlob.address") using Mapping = llama::mapping::MultiBlobSoA; auto mapping = Mapping{arrayExtents}; using ArrayIndex = typename Mapping::ArrayIndex; + using SizeType = typename Mapping::ArrayExtents::value_type; { const auto ai = ArrayIndex{0, 0}; - CHECK(mapping.template blobNrAndOffset<0, 0>(ai) == llama::NrAndOffset{0, 0}); - CHECK(mapping.template blobNrAndOffset<0, 1>(ai) == llama::NrAndOffset{1, 0}); - CHECK(mapping.template blobNrAndOffset<0, 2>(ai) == llama::NrAndOffset{2, 0}); - CHECK(mapping.template blobNrAndOffset<1>(ai) == llama::NrAndOffset{3, 0}); - CHECK(mapping.template blobNrAndOffset<2, 0>(ai) == llama::NrAndOffset{4, 0}); - CHECK(mapping.template blobNrAndOffset<2, 1>(ai) == llama::NrAndOffset{5, 0}); - CHECK(mapping.template blobNrAndOffset<2, 2>(ai) == llama::NrAndOffset{6, 0}); - CHECK(mapping.template blobNrAndOffset<3, 0>(ai) == llama::NrAndOffset{7, 0}); - CHECK(mapping.template blobNrAndOffset<3, 1>(ai) == llama::NrAndOffset{8, 0}); - CHECK(mapping.template blobNrAndOffset<3, 2>(ai) == llama::NrAndOffset{9, 0}); - CHECK(mapping.template blobNrAndOffset<3, 3>(ai) == llama::NrAndOffset{10, 0}); + CHECK(mapping.template blobNrAndOffset<0, 0>(ai) == llama::NrAndOffset{0, 0}); + CHECK(mapping.template blobNrAndOffset<0, 1>(ai) == llama::NrAndOffset{1, 0}); + CHECK(mapping.template blobNrAndOffset<0, 2>(ai) == llama::NrAndOffset{2, 0}); + CHECK(mapping.template blobNrAndOffset<1>(ai) == llama::NrAndOffset{3, 0}); + CHECK(mapping.template blobNrAndOffset<2, 0>(ai) == llama::NrAndOffset{4, 0}); + CHECK(mapping.template blobNrAndOffset<2, 1>(ai) == llama::NrAndOffset{5, 0}); + CHECK(mapping.template blobNrAndOffset<2, 2>(ai) == llama::NrAndOffset{6, 0}); + CHECK(mapping.template blobNrAndOffset<3, 0>(ai) == llama::NrAndOffset{7, 0}); + CHECK(mapping.template blobNrAndOffset<3, 1>(ai) == llama::NrAndOffset{8, 0}); + CHECK(mapping.template blobNrAndOffset<3, 2>(ai) == llama::NrAndOffset{9, 0}); + CHECK(mapping.template blobNrAndOffset<3, 3>(ai) == llama::NrAndOffset{10, 0}); } { const auto ai = ArrayIndex{0, 1}; - CHECK(mapping.template blobNrAndOffset<0, 0>(ai) == llama::NrAndOffset{0, 8}); - CHECK(mapping.template blobNrAndOffset<0, 1>(ai) == llama::NrAndOffset{1, 8}); - CHECK(mapping.template blobNrAndOffset<0, 2>(ai) == llama::NrAndOffset{2, 8}); - CHECK(mapping.template blobNrAndOffset<1>(ai) == llama::NrAndOffset{3, 4}); - CHECK(mapping.template blobNrAndOffset<2, 0>(ai) == llama::NrAndOffset{4, 8}); - CHECK(mapping.template blobNrAndOffset<2, 1>(ai) == llama::NrAndOffset{5, 8}); - CHECK(mapping.template blobNrAndOffset<2, 2>(ai) == llama::NrAndOffset{6, 8}); - CHECK(mapping.template blobNrAndOffset<3, 0>(ai) == llama::NrAndOffset{7, 1}); - CHECK(mapping.template blobNrAndOffset<3, 1>(ai) == llama::NrAndOffset{8, 1}); - CHECK(mapping.template blobNrAndOffset<3, 2>(ai) == llama::NrAndOffset{9, 1}); - CHECK(mapping.template blobNrAndOffset<3, 3>(ai) == llama::NrAndOffset{10, 1}); + CHECK(mapping.template blobNrAndOffset<0, 0>(ai) == llama::NrAndOffset{0, 8}); + CHECK(mapping.template blobNrAndOffset<0, 1>(ai) == llama::NrAndOffset{1, 8}); + CHECK(mapping.template blobNrAndOffset<0, 2>(ai) == llama::NrAndOffset{2, 8}); + CHECK(mapping.template blobNrAndOffset<1>(ai) == llama::NrAndOffset{3, 4}); + CHECK(mapping.template blobNrAndOffset<2, 0>(ai) == llama::NrAndOffset{4, 8}); + CHECK(mapping.template blobNrAndOffset<2, 1>(ai) == llama::NrAndOffset{5, 8}); + CHECK(mapping.template blobNrAndOffset<2, 2>(ai) == llama::NrAndOffset{6, 8}); + CHECK(mapping.template blobNrAndOffset<3, 0>(ai) == llama::NrAndOffset{7, 1}); + CHECK(mapping.template blobNrAndOffset<3, 1>(ai) == llama::NrAndOffset{8, 1}); + CHECK(mapping.template blobNrAndOffset<3, 2>(ai) == llama::NrAndOffset{9, 1}); + CHECK(mapping.template blobNrAndOffset<3, 3>(ai) == llama::NrAndOffset{10, 1}); } { const auto ai = ArrayIndex{1, 0}; - CHECK(mapping.template blobNrAndOffset<0, 0>(ai) == llama::NrAndOffset{0, 128}); - CHECK(mapping.template blobNrAndOffset<0, 1>(ai) == llama::NrAndOffset{1, 128}); - CHECK(mapping.template blobNrAndOffset<0, 2>(ai) == llama::NrAndOffset{2, 128}); - CHECK(mapping.template blobNrAndOffset<1>(ai) == llama::NrAndOffset{3, 64}); - CHECK(mapping.template blobNrAndOffset<2, 0>(ai) == llama::NrAndOffset{4, 128}); - CHECK(mapping.template blobNrAndOffset<2, 1>(ai) == llama::NrAndOffset{5, 128}); - CHECK(mapping.template blobNrAndOffset<2, 2>(ai) == llama::NrAndOffset{6, 128}); - CHECK(mapping.template blobNrAndOffset<3, 0>(ai) == llama::NrAndOffset{7, 16}); - CHECK(mapping.template blobNrAndOffset<3, 1>(ai) == llama::NrAndOffset{8, 16}); - CHECK(mapping.template blobNrAndOffset<3, 2>(ai) == llama::NrAndOffset{9, 16}); - CHECK(mapping.template blobNrAndOffset<3, 3>(ai) == llama::NrAndOffset{10, 16}); + CHECK(mapping.template blobNrAndOffset<0, 0>(ai) == llama::NrAndOffset{0, 128}); + CHECK(mapping.template blobNrAndOffset<0, 1>(ai) == llama::NrAndOffset{1, 128}); + CHECK(mapping.template blobNrAndOffset<0, 2>(ai) == llama::NrAndOffset{2, 128}); + CHECK(mapping.template blobNrAndOffset<1>(ai) == llama::NrAndOffset{3, 64}); + CHECK(mapping.template blobNrAndOffset<2, 0>(ai) == llama::NrAndOffset{4, 128}); + CHECK(mapping.template blobNrAndOffset<2, 1>(ai) == llama::NrAndOffset{5, 128}); + CHECK(mapping.template blobNrAndOffset<2, 2>(ai) == llama::NrAndOffset{6, 128}); + CHECK(mapping.template blobNrAndOffset<3, 0>(ai) == llama::NrAndOffset{7, 16}); + CHECK(mapping.template blobNrAndOffset<3, 1>(ai) == llama::NrAndOffset{8, 16}); + CHECK(mapping.template blobNrAndOffset<3, 2>(ai) == llama::NrAndOffset{9, 16}); + CHECK(mapping.template blobNrAndOffset<3, 3>(ai) == llama::NrAndOffset{10, 16}); } }; - test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtentsDynamic{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } diff --git a/tests/mapping.Split.cpp b/tests/mapping.Split.cpp index c5d5ddb98b..135ca60b7d 100644 --- a/tests/mapping.Split.cpp +++ b/tests/mapping.Split.cpp @@ -50,7 +50,7 @@ TEST_CASE("mapping.Split.partitionRecordDim.Particle.List") TEST_CASE("mapping.Split.SoA_SingleBlob.AoS_Packed.1Buffer") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; auto extents = ArrayExtents{16, 16}; // we layout Pos as SoA, the rest as AoS @@ -77,7 +77,7 @@ TEST_CASE("mapping.Split.AoSoA8.AoS_Packed.One.SoA_SingleBlob.4Buffer") { // split out momentum as AoSoA8, mass into a single value, position into AoS, and the flags into SoA, makes 4 // buffers - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtents; auto extents = ArrayExtents{32}; auto mapping = llama::mapping::Split< ArrayExtents, @@ -123,7 +123,7 @@ TEST_CASE("mapping.Split.AoSoA8.AoS_Packed.One.SoA_SingleBlob.4Buffer") TEST_CASE("mapping.Split.Multilist.SoA.One") { // split out Pos and Vel into SoA, the rest into One - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtents; auto extents = ArrayExtents{32}; auto mapping = llama::mapping::Split< ArrayExtents, @@ -163,7 +163,7 @@ TEST_CASE("mapping.Split.Multilist.SoA.One") TEST_CASE("mapping.Split.BitPacked") { // split out Pos and Vel into SoA, the rest into One - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtents; auto extents = ArrayExtents{32}; auto mapping = llama::mapping::Split< ArrayExtents, diff --git a/tests/mapping.Tree.cpp b/tests/mapping.Tree.cpp index 23402c6dd9..88bf789814 100644 --- a/tests/mapping.Tree.cpp +++ b/tests/mapping.Tree.cpp @@ -36,7 +36,7 @@ namespace tag TEST_CASE("mapping.Tree.empty") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{}; @@ -105,7 +105,7 @@ TEST_CASE("mapping.Tree.empty") TEST_CASE("mapping.Tree.Idem") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::Idem()}; @@ -174,7 +174,7 @@ TEST_CASE("mapping.Tree.Idem") TEST_CASE("mapping.Tree.LeafOnlyRT") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::LeafOnlyRT()}; @@ -243,7 +243,7 @@ TEST_CASE("mapping.Tree.LeafOnlyRT") TEST_CASE("mapping.Tree.MoveRTDown<>") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::MoveRTDown>{4}}; @@ -300,7 +300,7 @@ TEST_CASE("mapping.Tree.MoveRTDown<>") TEST_CASE("mapping.Tree.MoveRTDown<0>") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::MoveRTDown>{4}}; @@ -357,7 +357,7 @@ TEST_CASE("mapping.Tree.MoveRTDown<0>") TEST_CASE("mapping.Tree.MoveRTDown<0,0>") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::MoveRTDown>{4}}; @@ -414,7 +414,7 @@ TEST_CASE("mapping.Tree.MoveRTDown<0,0>") TEST_CASE("mapping.Tree.MoveRTDownFixed<>") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::MoveRTDownFixed, 4>{}}; @@ -471,7 +471,7 @@ TEST_CASE("mapping.Tree.MoveRTDownFixed<>") TEST_CASE("mapping.Tree.MoveRTDownFixed<0>") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::MoveRTDownFixed, 4>{}}; @@ -528,7 +528,7 @@ TEST_CASE("mapping.Tree.MoveRTDownFixed<0>") TEST_CASE("mapping.Tree.MoveRTDownFixed<0,0>") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{tree::functor::MoveRTDownFixed, 4>{}}; @@ -585,7 +585,7 @@ TEST_CASE("mapping.Tree.MoveRTDownFixed<0,0>") TEST_CASE("mapping.Tree.vectorblocks.runtime") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; const auto vectorWidth = 8; @@ -649,7 +649,7 @@ TEST_CASE("mapping.Tree.vectorblocks.runtime") TEST_CASE("mapping.Tree.vectorblocks.compiletime") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; constexpr auto vectorWidth = 8; @@ -713,7 +713,7 @@ TEST_CASE("mapping.Tree.vectorblocks.compiletime") TEST_CASE("mapping.Tree.getNode") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; const ArrayExtents extents{16, 16}; auto treeOperationList = llama::Tuple{}; @@ -757,7 +757,7 @@ TEST_CASE("mapping.Tree.getNode") TEST_CASE("mapping.Tree") { - using ArrayExtents = llama::ArrayExtents<12, 12>; + using ArrayExtents = llama::ArrayExtents; constexpr ArrayExtents extents{}; auto treeOperationList = llama::Tuple{tree::functor::Idem(), tree::functor::LeafOnlyRT{}, tree::functor::Idem{}}; @@ -1051,7 +1051,7 @@ TEST_CASE("mapping.Tree") double sum = 0.0; for(size_t x = 0; x < extents[0]; ++x) for(size_t y = 0; y < extents[1]; ++y) - sum += view({x, y})(llama::RecordCoord<0, 1>{}); + sum += view(x, y)(llama::RecordCoord<0, 1>{}); CHECK(sum == 0); } diff --git a/tests/mapping.cpp b/tests/mapping.cpp index c41ee7e725..54a15775a9 100644 --- a/tests/mapping.cpp +++ b/tests/mapping.cpp @@ -1,54 +1,66 @@ #include "common.hpp" #ifdef __cpp_lib_concepts -TEST_CASE("mapping.concepts") +TEMPLATE_LIST_TEST_CASE("mapping.concepts", "", SizeTypes) { - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle, 8>>); + STATIC_REQUIRE( + llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE( + llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE( + llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE( + llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE( + llama::PhysicalMapping, Particle, 8>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); - STATIC_REQUIRE(llama::PhysicalMapping, Particle, 8>>); + STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE(llama::PhysicalMapping, Particle>>); + STATIC_REQUIRE(llama::PhysicalMapping, Particle, 8>>); - using Inner = llama::mapping::AlignedAoS, Particle>; + using Inner = llama::mapping::AlignedAoS, Particle>; STATIC_REQUIRE(llama::PhysicalMapping>); STATIC_REQUIRE(llama::PhysicalMapping>); - STATIC_REQUIRE(llama::FullyComputedMapping, Particle>>); + STATIC_REQUIRE( + llama::FullyComputedMapping, Particle>>); STATIC_REQUIRE(llama::FullyComputedMapping< - llama::mapping::Bytesplit, Particle, llama::mapping::BindAoS<>::fn>>); - STATIC_REQUIRE(llama::FullyComputedMapping, Vec3I>>); + llama::mapping:: + Bytesplit, Particle, llama::mapping::BindAoS<>::fn>>); STATIC_REQUIRE( - llama::FullyComputedMapping, Vec3D>>); + llama::FullyComputedMapping, Vec3I>>); + STATIC_REQUIRE(llama::FullyComputedMapping< + llama::mapping::BitPackedFloatSoA, Vec3D>>); STATIC_REQUIRE(llama::PartiallyComputedMapping, + llama::ArrayExtentsDynamic, Particle, llama::mapping::BindAoS<>::fn, boost::mp11::mp_list>>>); } #endif -TEST_CASE("mapping.traits") +TEMPLATE_LIST_TEST_CASE("mapping.traits", "", SizeTypes) { - using AAoS = llama::mapping::AlignedAoS, Particle>; - using PAoS = llama::mapping::PackedAoS, Particle>; - using SBSoA = llama::mapping::SingleBlobSoA, Particle>; - using MBSoA = llama::mapping::MultiBlobSoA, Particle>; - using AoAoS = llama::mapping::AoSoA, Particle, 8>; - using One = llama::mapping::One, Particle>; + using AAoS = llama::mapping::AlignedAoS, Particle>; + using PAoS = llama::mapping::PackedAoS, Particle>; + using SBSoA = llama::mapping::SingleBlobSoA, Particle>; + using MBSoA = llama::mapping::MultiBlobSoA, Particle>; + using AoAoS = llama::mapping::AoSoA, Particle, 8>; + using One = llama::mapping::One, Particle>; - using Null = llama::mapping::Null, Particle>; - using BS = llama::mapping::Bytesplit, Particle, llama::mapping::BindAoS<>::fn>; - using CT = llama::mapping:: - ChangeType, Particle, llama::mapping::BindAoS<>::fn, boost::mp11::mp_list<>>; - using BPI = llama::mapping::BitPackedIntSoA, Vec3I>; - using BPF = llama::mapping::BitPackedFloatSoA, Vec3D>; + using Null = llama::mapping::Null, Particle>; + using BS + = llama::mapping::Bytesplit, Particle, llama::mapping::BindAoS<>::fn>; + using CT = llama::mapping::ChangeType< + llama::ArrayExtentsDynamic, + Particle, + llama::mapping::BindAoS<>::fn, + boost::mp11::mp_list<>>; + using BPI = llama::mapping::BitPackedIntSoA, Vec3I>; + using BPF = llama::mapping::BitPackedFloatSoA, Vec3D>; STATIC_REQUIRE(llama::mapping::isAoS); STATIC_REQUIRE(llama::mapping::isAoS); @@ -171,7 +183,7 @@ TEST_CASE("mapping.LinearizeArrayDimsCpp.size") TEST_CASE("mapping.LinearizeArrayDimsCpp") { llama::mapping::LinearizeArrayDimsCpp lin; - const auto extents = llama::ArrayExtents<4, 4>{}; + const auto extents = llama::ArrayExtents{}; CHECK(lin(llama::ArrayIndex{0, 0}, extents) == 0); CHECK(lin(llama::ArrayIndex{0, 1}, extents) == 1); CHECK(lin(llama::ArrayIndex{0, 2}, extents) == 2); @@ -202,7 +214,7 @@ TEST_CASE("mapping.LinearizeArrayDimsFortran.size") TEST_CASE("mapping.LinearizeArrayDimsFortran") { llama::mapping::LinearizeArrayDimsFortran lin; - const auto extents = llama::ArrayExtents<4, 4>{}; + const auto extents = llama::ArrayExtents{}; CHECK(lin(llama::ArrayIndex{0, 0}, extents) == 0); CHECK(lin(llama::ArrayIndex{0, 1}, extents) == 4); CHECK(lin(llama::ArrayIndex{0, 2}, extents) == 8); @@ -233,7 +245,7 @@ TEST_CASE("mapping.LinearizeArrayDimsMorton.size") TEST_CASE("mapping.LinearizeArrayDimsMorton") { llama::mapping::LinearizeArrayDimsMorton lin; - const auto extents = llama::ArrayExtents<4, 4>{}; + const auto extents = llama::ArrayExtents{}; CHECK(lin(llama::ArrayIndex{0, 0}, extents) == 0); CHECK(lin(llama::ArrayIndex{0, 1}, extents) == 1); CHECK(lin(llama::ArrayIndex{0, 2}, extents) == 4); diff --git a/tests/proofs.cpp b/tests/proofs.cpp index 2be43388ab..9c74cec4e3 100644 --- a/tests/proofs.cpp +++ b/tests/proofs.cpp @@ -5,7 +5,7 @@ TEST_CASE("mapsNonOverlappingly.PackedAoS") { #ifdef __cpp_constexpr_dynamic_alloc - constexpr auto mapping = llama::mapping::PackedAoS, Particle>{{32, 32}}; + constexpr auto mapping = llama::mapping::PackedAoS, Particle>{{32, 32}}; STATIC_REQUIRE(llama::mapsNonOverlappingly(mapping)); #else INFO("Test disabled because compiler does not support __cpp_constexpr_dynamic_alloc"); @@ -15,7 +15,8 @@ TEST_CASE("mapsNonOverlappingly.PackedAoS") TEST_CASE("mapsNonOverlappingly.AlignedAoS") { #ifdef __cpp_constexpr_dynamic_alloc - constexpr auto mapping = llama::mapping::AlignedAoS, Particle>{{32, 32}}; + constexpr auto mapping + = llama::mapping::AlignedAoS, Particle>{{32, 32}}; STATIC_REQUIRE(llama::mapsNonOverlappingly(mapping)); #else INFO("Test disabled because compiler does not support __cpp_constexpr_dynamic_alloc"); @@ -49,7 +50,7 @@ namespace template constexpr auto blobNrAndOffset(ArrayIndex, llama::RecordCoord = {}) const - -> llama::NrAndOffset + -> llama::NrAndOffset { return {0, 0}; } @@ -59,10 +60,12 @@ namespace TEST_CASE("mapsNonOverlappingly.MapEverythingToZero") { #ifdef __cpp_constexpr_dynamic_alloc - STATIC_REQUIRE(llama::mapsNonOverlappingly(MapEverythingToZero{llama::ArrayExtentsDynamic<1>{1}, double{}})); - STATIC_REQUIRE(!llama::mapsNonOverlappingly(MapEverythingToZero{llama::ArrayExtentsDynamic<1>{2}, double{}})); STATIC_REQUIRE( - !llama::mapsNonOverlappingly(MapEverythingToZero{llama::ArrayExtentsDynamic<2>{32, 32}, Particle{}})); + llama::mapsNonOverlappingly(MapEverythingToZero{llama::ArrayExtentsDynamic{1}, double{}})); + STATIC_REQUIRE( + !llama::mapsNonOverlappingly(MapEverythingToZero{llama::ArrayExtentsDynamic{2}, double{}})); + STATIC_REQUIRE(!llama::mapsNonOverlappingly( + MapEverythingToZero{llama::ArrayExtentsDynamic{32, 32}, Particle{}})); #else INFO("Test disabled because compiler does not support __cpp_constexpr_dynamic_alloc"); #endif @@ -97,7 +100,7 @@ namespace template constexpr auto blobNrAndOffset(ArrayIndex ai, llama::RecordCoord = {}) const - -> llama::NrAndOffset + -> llama::NrAndOffset { const auto blob = llama::flatRecordCoord>; const auto offset = (llama::mapping::LinearizeArrayDimsCpp{}(ai, extents()) % Modulus) @@ -110,12 +113,12 @@ namespace TEST_CASE("mapsNonOverlappingly.ModulusMapping") { #ifdef __cpp_constexpr_dynamic_alloc - using Modulus10Mapping = ModulusMapping, Particle, 10>; - STATIC_REQUIRE(llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic<1>{1}})); - STATIC_REQUIRE(llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic<1>{9}})); - STATIC_REQUIRE(llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic<1>{10}})); - STATIC_REQUIRE(!llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic<1>{11}})); - STATIC_REQUIRE(!llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic<1>{25}})); + using Modulus10Mapping = ModulusMapping, Particle, 10>; + STATIC_REQUIRE(llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic{1}})); + STATIC_REQUIRE(llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic{9}})); + STATIC_REQUIRE(llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic{10}})); + STATIC_REQUIRE(!llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic{11}})); + STATIC_REQUIRE(!llama::mapsNonOverlappingly(Modulus10Mapping{llama::ArrayExtentsDynamic{25}})); #else INFO("Test disabled because compiler does not support __cpp_constexpr_dynamic_alloc"); #endif @@ -123,7 +126,7 @@ TEST_CASE("mapsNonOverlappingly.ModulusMapping") TEST_CASE("maps.ModulusMapping") { - constexpr auto extents = llama::ArrayExtentsDynamic<1>{128}; + constexpr auto extents = llama::ArrayExtentsDynamic{128}; STATIC_REQUIRE(llama::mapsPiecewiseContiguous<1>(llama::mapping::AoS{extents, Particle{}})); STATIC_REQUIRE(!llama::mapsPiecewiseContiguous<8>(llama::mapping::AoS{extents, Particle{}})); STATIC_REQUIRE(!llama::mapsPiecewiseContiguous<16>(llama::mapping::AoS{extents, Particle{}})); diff --git a/tests/proxyrefopmixin.cpp b/tests/proxyrefopmixin.cpp index 113a667391..3dc3fc8ecc 100644 --- a/tests/proxyrefopmixin.cpp +++ b/tests/proxyrefopmixin.cpp @@ -230,29 +230,30 @@ namespace TEST_CASE("proxyrefopmixin.Bytesplit") { - auto view - = llama::allocView(llama::mapping::Bytesplit, Vec3I, llama::mapping::BindAoS<>::fn>{ - llama::ArrayExtents<4>{}}); + auto view = llama::allocView( + llama::mapping::Bytesplit, Vec3I, llama::mapping::BindAoS<>::fn>{ + llama::ArrayExtents{}}); testProxyRef(view(2)(tag::X{})); } TEST_CASE("proxyrefopmixin.BitPackedIntSoA") { - auto view = llama::allocView(llama::mapping::BitPackedIntSoA, Vec3I>{{}, 12}); + auto view = llama::allocView(llama::mapping::BitPackedIntSoA, Vec3I>{{}, 12}); testProxyRef(view(2)(tag::X{})); } TEST_CASE("proxyrefopmixin.BitPackedFloatSoA") { auto view = llama::allocView( - llama::mapping::BitPackedFloatSoA, Vec3D, llama::Constant<6>, llama::Constant<20>>{}); + llama::mapping:: + BitPackedFloatSoA, Vec3D, llama::Constant<6>, llama::Constant<20>>{}); testProxyRef(view(2)(tag::X{})); } TEST_CASE("proxyrefopmixin.ChangeType") { auto view = llama::allocView(llama::mapping::ChangeType< - llama::ArrayExtents<4>, + llama::ArrayExtents, Vec3D, llama::mapping::BindAoS::fn, boost::mp11::mp_list>>{{}}); diff --git a/tests/recorddimension.cpp b/tests/recorddimension.cpp index 4c4e999952..9072b73150 100644 --- a/tests/recorddimension.cpp +++ b/tests/recorddimension.cpp @@ -223,7 +223,7 @@ TEST_CASE("recorddim.int[3]") using RecordDim = int[3]; auto view - = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); + = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); view(0u)(0_RC) = 42; view(0u)(1_RC) = 43; @@ -237,7 +237,7 @@ TEST_CASE("recorddim.int[200]") using RecordDim = int[200]; auto view - = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); + = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); view(0u)(0_RC) = 42; view(0u)(199_RC) = 43; @@ -249,7 +249,7 @@ TEST_CASE("recorddim.int[3][2]") using RecordDim = int[3][2]; auto view - = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); + = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); view(0u)(0_RC)(0_RC) = 42; view(0u)(0_RC)(1_RC) = 43; @@ -265,7 +265,7 @@ TEST_CASE("recorddim.int[1][1][1][1][1][1][1][1][1][1]") using RecordDim = int[1][1][1][1][1][1][1][1][1][1]; auto view - = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); + = llama::allocView(llama::mapping::AoS, RecordDim>{llama::ArrayExtents{1}}); view(0u)(0_RC)(0_RC)(0_RC)(0_RC)(0_RC)(0_RC)(0_RC)(0_RC)(0_RC)(0_RC) = 42; } diff --git a/tests/vector.cpp b/tests/vector.cpp index 60d05b9778..0856c0972b 100644 --- a/tests/vector.cpp +++ b/tests/vector.cpp @@ -1,14 +1,14 @@ #include "common.hpp" using RecordDim = Vec3D; -using Mapping = llama::mapping::AoS, RecordDim>; +using Mapping = llama::mapping::AoS, RecordDim>; using Vector = llama::Vector; TEST_CASE("Vector.ctor.default") { const Vector v; CHECK(v.empty()); - CHECK(v.size() == 0); + CHECK(v.size() == 0); // NOLINT(readability-container-size-empty) } TEST_CASE("Vector.ctor.count") @@ -29,7 +29,7 @@ TEST_CASE("Vector.ctor.count_and_value") TEST_CASE("Vector.ctor.iterator_pair") { - auto view = llama::allocView(Mapping{llama::ArrayExtents{10}}); + auto view = llama::allocView(Mapping{llama::ArrayExtents{std::size_t{10}}}); for(auto i = 0; i < 10; i++) view[i] = i; @@ -162,10 +162,10 @@ TEST_CASE("vector.reserve") { Vector v; CHECK(v.capacity() == 0); - CHECK(v.size() == 0); + CHECK(v.empty()); v.reserve(100); CHECK(v.capacity() == 100); - CHECK(v.size() == 0); + CHECK(v.empty()); } TEST_CASE("vector.clear_shrink_to_fit") @@ -175,10 +175,10 @@ TEST_CASE("vector.clear_shrink_to_fit") CHECK(v.size() == 10); v.clear(); CHECK(v.capacity() == 10); - CHECK(v.size() == 0); + CHECK(v.empty()); v.shrink_to_fit(); CHECK(v.capacity() == 0); - CHECK(v.size() == 0); + CHECK(v.empty()); } TEST_CASE("vector.insert") @@ -223,7 +223,7 @@ TEST_CASE("vector.pop_back") CHECK(v.size() == 1); CHECK(v[0] == 0); v.pop_back(); - CHECK(v.size() == 0); + CHECK(v.empty()); } TEST_CASE("vector.resize") @@ -252,7 +252,7 @@ TEST_CASE("vector.resize") CHECK(v[0] == 0); CHECK(v[1] == 1); v.resize(0); - CHECK(v.size() == 0); + CHECK(v.empty()); } TEST_CASE("vector.erase") diff --git a/tests/view.cpp b/tests/view.cpp index 8d2b25110c..e834d0f62d 100644 --- a/tests/view.cpp +++ b/tests/view.cpp @@ -13,7 +13,7 @@ using RecordDim = llama::Record< TEST_CASE("view.default-ctor") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; [[maybe_unused]] llama::View, std::byte*> view1{}; [[maybe_unused]] llama::View, std::byte*> view2{}; [[maybe_unused]] llama::View, std::byte*> view3{}; @@ -25,7 +25,7 @@ TEST_CASE("view.default-ctor") TEST_CASE("view.move") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; constexpr ArrayExtents viewSize{16, 16}; using Mapping = llama::mapping::SoA; @@ -39,7 +39,7 @@ TEST_CASE("view.move") TEST_CASE("view.swap") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; constexpr ArrayExtents viewSize{16, 16}; using Mapping = llama::mapping::SoA; @@ -61,7 +61,7 @@ TEST_CASE("view.non-memory-owning") { using byte = typename decltype(typeHolder)::type; - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic; ArrayExtents extents{256}; using Mapping = llama::mapping::SoA; @@ -84,7 +84,7 @@ TEST_CASE("view.non-memory-owning") TEST_CASE("view.access") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; ArrayExtents extents{16, 16}; using Mapping = llama::mapping::SoA; @@ -117,7 +117,7 @@ TEST_CASE("view.assign-one-record") { using namespace llama::literals; - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; ArrayExtents extents{16, 16}; using Mapping = llama::mapping::SoA; @@ -154,7 +154,7 @@ TEST_CASE("view.addresses") { using namespace llama::literals; - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; ArrayExtents extents{16, 16}; using Mapping = llama::mapping::SingleBlobSoA; @@ -199,15 +199,15 @@ struct SetZeroFunctor TEST_CASE("view.iteration-and-access") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; ArrayExtents extents{16, 16}; using Mapping = llama::mapping::SoA; Mapping mapping{extents}; auto view = llama::allocView(mapping); - for(size_t x = 0; x < extents[0]; ++x) - for(size_t y = 0; y < extents[1]; ++y) + for(int x = 0; x < extents[0]; ++x) + for(int y = 0; y < extents[1]; ++y) { SetZeroFunctor szf{view(x, y)}; llama::forEachLeafCoord(szf, llama::RecordCoord<0, 0>{}); @@ -216,15 +216,15 @@ TEST_CASE("view.iteration-and-access") } double sum = 0.0; - for(size_t x = 0; x < extents[0]; ++x) - for(size_t y = 0; y < extents[1]; ++y) + for(int x = 0; x < extents[0]; ++x) + for(int y = 0; y < extents[1]; ++y) sum += view(x, y)(llama::RecordCoord<2, 0>{}); CHECK(sum == 120.0); } TEST_CASE("view.record-access") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; ArrayExtents extents{16, 16}; using Mapping = llama::mapping::SoA; diff --git a/tests/virtualrecord.cpp b/tests/virtualrecord.cpp index bff7fa1610..74814a2fdd 100644 --- a/tests/virtualrecord.cpp +++ b/tests/virtualrecord.cpp @@ -948,19 +948,29 @@ TEST_CASE("VirtualRecord.One_from_scalar") CHECK(p(tag::Mass{}) == 42); } -TEST_CASE("VirtualRecord.size") +TEST_CASE("VirtualRecord.size.int") { auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{5}, ParticleInt{}}); [[maybe_unused]] auto vr = view[0]; STATIC_REQUIRE( sizeof(vr) - == sizeof(llama::ArrayExtents::value_type) + == sizeof(llama::ArrayExtents::value_type) + 4 // padding + + sizeof(&view)); // sizeof array dims and view reference // NOLINT +} + +TEST_CASE("VirtualRecord.size.size_t") +{ + auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{std::size_t{5}}, ParticleInt{}}); + [[maybe_unused]] auto vr = view[0]; + STATIC_REQUIRE( + sizeof(vr) + == sizeof(llama::ArrayExtents::value_type) + sizeof(&view)); // sizeof array dims and view reference // NOLINT } TEST_CASE("VirtualRecord.One.size") { - using Mapping = llama::mapping::MinAlignedOne, Particle>; + using Mapping = llama::mapping::MinAlignedOne, Particle>; STATIC_REQUIRE(Mapping{}.blobSize(0) == 56); STATIC_REQUIRE(std::is_empty_v); diff --git a/tests/virtualview.cpp b/tests/virtualview.cpp index cfd9f71a6b..1bb8a8a278 100644 --- a/tests/virtualview.cpp +++ b/tests/virtualview.cpp @@ -2,7 +2,7 @@ TEST_CASE("VirtualView.CTAD") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; constexpr ArrayExtents viewSize{10, 10}; const auto mapping = llama::mapping::SoA(viewSize); auto view = llama::allocViewUninitialized(mapping); @@ -38,14 +38,14 @@ TEST_CASE("VirtualView.CTAD") TEST_CASE("VirtualView.fast") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; constexpr ArrayExtents viewSize{10, 10}; using Mapping = llama::mapping::SoA; auto view = llama::allocViewUninitialized(Mapping(viewSize)); - for(std::size_t x = 0; x < viewSize[0]; ++x) - for(std::size_t y = 0; y < viewSize[1]; ++y) + for(int x = 0; x < viewSize[0]; ++x) + for(int y = 0; y < viewSize[1]; ++y) view(x, y) = x * y; llama::VirtualView virtualView{view, {2, 4}}; @@ -75,7 +75,7 @@ namespace TEST_CASE("VirtualView") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic; constexpr ArrayExtents viewSize{32, 32}; constexpr ArrayExtents miniSize{8, 8}; using Mapping = llama::mapping::SoA; @@ -143,9 +143,9 @@ TEST_CASE("VirtualView.negative_indices") TEST_CASE("VirtualView.negative_offsets") { - auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{10, 10}, int{}}); + auto view = llama::allocView(llama::mapping::AoS{llama::ArrayExtents{10u, 10u}, int{}}); auto shiftedView = llama::VirtualView{view, {2, 4}}; - auto shiftedView2 = llama::VirtualView{shiftedView, {static_cast(-2), static_cast(-4)}}; + auto shiftedView2 = llama::VirtualView{shiftedView, {static_cast(-2), static_cast(-4)}}; int i = 0; for(int y = 0; y < 10; y++)