From c0c49c1c970347858b423e128a85887820f39a74 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 | 2 +- examples/bufferguard/bufferguard.cpp | 9 +- examples/cuda/nbody/nbody.cu | 20 +- examples/cuda/pitch/pitch.cu | 4 +- examples/nbody_benchmark/nbody.cpp | 2 +- examples/simpletest/simpletest.cpp | 14 +- examples/vectoradd/vectoradd.cpp | 2 +- include/llama/ArrayExtents.hpp | 208 +++++++++++++----- include/llama/ArrayIndexRange.hpp | 2 +- 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 | 19 +- 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 | 171 +++++++++----- tests/arrayindexrange.cpp | 4 +- tests/bloballocators.cpp | 2 +- tests/common.hpp | 3 + tests/computedprop.cpp | 6 +- tests/heatmap_trace.cpp | 12 +- tests/iterator.cpp | 4 +- tests/mapping.BitPackedFloatSoA.cpp | 36 +-- tests/mapping.BitPackedIntSoA.cpp | 14 +- tests/mapping.Bytesplit.cpp | 2 +- tests/mapping.ChangeType.cpp | 19 +- tests/mapping.Null.cpp | 2 +- tests/mapping.Split.cpp | 8 +- tests/mapping.Tree.cpp | 4 +- tests/mapping.cpp | 98 +++++---- tests/proofs.cpp | 4 +- tests/proxyrefopmixin.cpp | 13 +- tests/recorddimension.cpp | 8 +- tests/vector.cpp | 2 +- tests/view.cpp | 14 +- tests/virtualrecord.cpp | 4 +- tests/virtualview.cpp | 10 +- 58 files changed, 666 insertions(+), 444 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 0264f1a7cc..40e79615d7 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..f5e036b85f 100644 --- a/examples/bitpackint/bitpackint.cpp +++ b/examples/bitpackint/bitpackint.cpp @@ -49,7 +49,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..62a548dd57 100644 --- a/examples/bufferguard/bufferguard.cpp +++ b/examples/bufferguard/bufferguard.cpp @@ -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; @@ -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/cuda/nbody/nbody.cu b/examples/cuda/nbody/nbody.cu index 7410f9a8bd..7823ef87a1 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<1, int>; 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..4c3fb7ddd2 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}; diff --git a/examples/nbody_benchmark/nbody.cpp b/examples/nbody_benchmark/nbody.cpp index 71d983c475..027feebf1d 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<1>; const auto extents = ArrayExtents{PROBLEM_SIZE}; if constexpr(Mapping == 0) return llama::mapping::AoS{extents}; diff --git a/examples/simpletest/simpletest.cpp b/examples/simpletest/simpletest.cpp index 6611ba0680..7817841a75 100644 --- a/examples/simpletest/simpletest.cpp +++ b/examples/simpletest/simpletest.cpp @@ -127,7 +127,7 @@ auto main() -> int try { // Defining two array dimensions - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic<2, int>; // Setting the run time size of the array dimensions to 8192 * 8192 auto extents = ArrayExtents{8192, 8192}; @@ -177,11 +177,11 @@ try << '\n'; // iterating over the array dimensions at run time to do some stuff with the allocated data - for(size_t x = 0; x < extents[0]; ++x) + for(int x = 0; x < extents[0]; ++x) // telling the compiler that all data in the following loop is independent to each other and thus can be // vectorized LLAMA_INDEPENDENT_DATA - for(size_t y = 0; y < extents[1]; ++y) + for(int y = 0; y < extents[1]; ++y) { // Defining a functor for a given virtual record SetZeroFunctor szf{view(x, y)}; @@ -194,9 +194,9 @@ try view({x, y}) = double(x + y) / double(extents[0] + extents[1]); } - for(size_t x = 0; x < extents[0]; ++x) + for(int x = 0; x < extents[0]; ++x) LLAMA_INDEPENDENT_DATA - for(size_t y = 0; y < extents[1]; ++y) + for(int y = 0; y < extents[1]; ++y) { // Showing different options of access data with llama. Internally all do the same data- and mappingwise auto record = view(x, y); @@ -209,9 +209,9 @@ try } double sum = 0.0; LLAMA_INDEPENDENT_DATA - for(size_t x = 0; x < extents[0]; ++x) + for(int x = 0; x < extents[0]; ++x) LLAMA_INDEPENDENT_DATA - for(size_t y = 0; y < extents[1]; ++y) + for(int y = 0; y < extents[1]; ++y) sum += view(x, y)(llama::RecordCoord<1, 0>{}); std::cout << "Sum: " << sum << '\n'; diff --git a/examples/vectoradd/vectoradd.cpp b/examples/vectoradd/vectoradd.cpp index 527d4ac132..5e6b2063e0 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<1>; 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..47fa7d6d6f 100644 --- a/include/llama/ArrayExtents.hpp +++ b/include/llama/ArrayExtents.hpp @@ -13,67 +13,138 @@ 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>; } // 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; }; 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)> + 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 +168,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<0, T>; + using value_type = T; LLAMA_FN_HOST_ACC_INLINE constexpr auto toArray() const -> Index { @@ -119,71 +190,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(dyn)>; - 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..09ae41325c 100644 --- a/include/llama/ArrayIndexRange.hpp +++ b/include/llama/ArrayIndexRange.hpp @@ -75,7 +75,7 @@ namespace llama current[rank - 1]--; for(auto i = static_cast(rank) - 2; i >= 0; i--) { - if(current[i + 1] != std::numeric_limits::max()) + if(current[i + 1] != std::numeric_limits::max()) 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..e3db9d04fb 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..ea632188f7 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 1d7736c1df..bc0a704ee5 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{}); } @@ -329,6 +329,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>, @@ -358,6 +359,12 @@ namespace llama return static_cast(*this); } + template + auto operator()(llama::ArrayIndex) const + { + static_assert(!sizeof(SizeType), "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) { @@ -389,28 +396,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)...}); } @@ -429,14 +434,20 @@ namespace llama return (*this)(ai); } + template + auto operator[](llama::ArrayIndex) const + { + static_assert(!sizeof(SizeType), "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); @@ -504,6 +515,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) @@ -545,8 +558,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}); @@ -559,8 +572,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..95ba89072d 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}; } }; 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..3a8faa2fb3 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..1ec4543911 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,171 @@ 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<3>, + llama::ArrayExtents>); + STATIC_REQUIRE(std::is_same_v< + llama::ArrayExtentsDynamic<3, int>, + llama::ArrayExtents>); +} + +TEST_CASE("ArrayExtentsStatic") +{ + STATIC_REQUIRE(std::is_same_v, llama::ArrayExtents>); + STATIC_REQUIRE( + std::is_same_v, llama::ArrayExtents>); + STATIC_REQUIRE(std::is_same_v, llama::ArrayExtents>); +} + +TEST_CASE("ArrayExtents.ctor_value_init") +{ + [[maybe_unused]] const llama::ArrayExtentsDynamic<1> 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<1, TestType> 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<2, TestType> 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<2, TestType>>{{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<3, TestType> 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..c26eb24f33 100644 --- a/tests/arrayindexrange.cpp +++ b/tests/arrayindexrange.cpp @@ -14,7 +14,7 @@ TEST_CASE("ArrayIndexIterator.concepts") TEST_CASE("ArrayIndexIterator") { - llama::ArrayIndexRange r{llama::ArrayExtentsDynamic<2>{3, 3}}; + llama::ArrayIndexRange r{llama::ArrayExtentsDynamic<2, int>{3, 3}}; llama::ArrayIndexIterator it = std::begin(r); CHECK(*it == llama::ArrayIndex{0, 0}); @@ -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<2, int>{3, 3}}); b &= *it == llama::ArrayIndex{0, 0}; it++; b &= *it == llama::ArrayIndex{0, 1}; 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..3783920ea6 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)); @@ -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/heatmap_trace.cpp b/tests/heatmap_trace.cpp index 905285a0a1..cf0b70dbaa 100644 --- a/tests/heatmap_trace.cpp +++ b/tests/heatmap_trace.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/iterator.cpp b/tests/iterator.cpp index b766278c66..4c57dfb363 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; @@ -256,7 +256,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.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..6e87d2fb32 100644 --- a/tests/mapping.BitPackedIntSoA.cpp +++ b/tests/mapping.BitPackedIntSoA.cpp @@ -99,15 +99,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 +115,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 @@ -189,7 +189,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..4a231da155 100644 --- a/tests/mapping.Bytesplit.cpp +++ b/tests/mapping.Bytesplit.cpp @@ -57,7 +57,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<1, int>, Vec3I, llama::mapping::BindSplit< llama::RecordCoord<1>, 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.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.Split.cpp b/tests/mapping.Split.cpp index 7e131ee680..69abe1d8e5 100644 --- a/tests/mapping.Split.cpp +++ b/tests/mapping.Split.cpp @@ -50,7 +50,7 @@ TEST_CASE("Split.partitionRecordDim.Particle.List") TEST_CASE("Split.SoA_SingleBlob.AoS_Packed.1Buffer") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic<2, int>; auto extents = ArrayExtents{16, 16}; // we layout Pos as SoA, the rest as AoS @@ -77,7 +77,7 @@ TEST_CASE("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::ArrayExtentsDynamic<1, int>; auto extents = ArrayExtents{32}; auto mapping = llama::mapping::Split< ArrayExtents, @@ -123,7 +123,7 @@ TEST_CASE("Split.AoSoA8.AoS_Packed.One.SoA_SingleBlob.4Buffer") TEST_CASE("Split.Multilist.SoA.One") { // split out Pos and Vel into SoA, the rest into One - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic<1, int>; auto extents = ArrayExtents{32}; auto mapping = llama::mapping::Split< ArrayExtents, @@ -163,7 +163,7 @@ TEST_CASE("Split.Multilist.SoA.One") TEST_CASE("Split.BitPacked") { // split out Pos and Vel into SoA, the rest into One - using ArrayExtents = llama::ArrayExtentsDynamic<1>; + using ArrayExtents = llama::ArrayExtentsDynamic<1, int>; auto extents = ArrayExtents{32}; auto mapping = llama::mapping::Split< ArrayExtents, diff --git a/tests/mapping.Tree.cpp b/tests/mapping.Tree.cpp index 2edb7045d5..36ee77cc87 100644 --- a/tests/mapping.Tree.cpp +++ b/tests/mapping.Tree.cpp @@ -757,7 +757,7 @@ TEST_CASE("treemapping.getNode") TEST_CASE("treemapping") { - 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("treemapping") 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 eb30c18d1a..8191ce4bd1 100644 --- a/tests/mapping.cpp +++ b/tests/mapping.cpp @@ -9,11 +9,13 @@ TEST_CASE("mapping.concepts") 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>; STATIC_REQUIRE(llama::PhysicalMapping>); @@ -171,7 +173,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 +204,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 +235,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); @@ -306,9 +308,9 @@ TEST_CASE("address.AoS.Packed") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.AoS.Packed.fortran") @@ -366,9 +368,9 @@ TEST_CASE("address.AoS.Packed.fortran") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.AoS.Packed.morton") @@ -426,9 +428,9 @@ TEST_CASE("address.AoS.Packed.morton") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.AoS.Aligned") @@ -485,9 +487,9 @@ TEST_CASE("address.AoS.Aligned") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.AoS.aligned_min") @@ -544,9 +546,9 @@ TEST_CASE("address.AoS.aligned_min") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.SoA.SingleBlob") @@ -603,9 +605,9 @@ TEST_CASE("address.SoA.SingleBlob") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.SoA.SingleBlob.fortran") @@ -663,9 +665,9 @@ TEST_CASE("address.SoA.SingleBlob.fortran") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.SoA.SingleBlob.morton") @@ -727,9 +729,9 @@ TEST_CASE("address.SoA.SingleBlob.morton") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.SoA.MultiBlob") @@ -785,10 +787,10 @@ TEST_CASE("address.SoA.MultiBlob") 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<2, int>{16, 16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.AoSoA.4") @@ -845,9 +847,9 @@ TEST_CASE("address.AoSoA.4") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.PackedOne") @@ -875,9 +877,9 @@ TEST_CASE("address.PackedOne") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.AlignedOne") @@ -905,9 +907,9 @@ TEST_CASE("address.AlignedOne") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } TEST_CASE("address.MinAlignedOne") @@ -935,9 +937,9 @@ TEST_CASE("address.MinAlignedOne") } }; test(llama::ArrayExtentsDynamic<2>{16, 16}); - test(llama::ArrayExtents<16, llama::dyn>{16}); - test(llama::ArrayExtents{16}); - test(llama::ArrayExtents<16, 16>{}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{16}); + test(llama::ArrayExtents{}); } diff --git a/tests/proofs.cpp b/tests/proofs.cpp index 2be43388ab..a3fbebdefc 100644 --- a/tests/proofs.cpp +++ b/tests/proofs.cpp @@ -49,7 +49,7 @@ namespace template constexpr auto blobNrAndOffset(ArrayIndex, llama::RecordCoord = {}) const - -> llama::NrAndOffset + -> llama::NrAndOffset { return {0, 0}; } @@ -97,7 +97,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) 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..08dabf5cca 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..901dae20f2 100644 --- a/tests/vector.cpp +++ b/tests/vector.cpp @@ -1,7 +1,7 @@ #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") diff --git a/tests/view.cpp b/tests/view.cpp index 8d2b25110c..bc1975af53 100644 --- a/tests/view.cpp +++ b/tests/view.cpp @@ -84,7 +84,7 @@ TEST_CASE("view.non-memory-owning") TEST_CASE("view.access") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic<2, int>; 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<2, int>; 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<2, int>; 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,8 +216,8 @@ 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); } diff --git a/tests/virtualrecord.cpp b/tests/virtualrecord.cpp index bff7fa1610..fcbdc54b3e 100644 --- a/tests/virtualrecord.cpp +++ b/tests/virtualrecord.cpp @@ -954,13 +954,13 @@ TEST_CASE("VirtualRecord.size") [[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.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..7b84ae4edc 100644 --- a/tests/virtualview.cpp +++ b/tests/virtualview.cpp @@ -38,14 +38,14 @@ TEST_CASE("VirtualView.CTAD") TEST_CASE("VirtualView.fast") { - using ArrayExtents = llama::ArrayExtentsDynamic<2>; + using ArrayExtents = llama::ArrayExtentsDynamic<2, int>; 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}}; @@ -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++)