From 3ff629c5164cb8e2ae8bd91c655cccb769a73980 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 1 Dec 2022 18:12:41 +0100 Subject: [PATCH] Support LLAMA frames in IO --- .../picongpu/plugins/openPMD/WriteSpecies.hpp | 15 +-- .../LoadParticleAttributesFromOpenPMD.hpp | 13 ++- .../plugins/openPMD/restart/LoadSpecies.hpp | 11 +-- .../openPMD/writer/ParticleAttribute.hpp | 11 ++- .../plugins/output/WriteSpeciesCommon.hpp | 95 +++++++------------ .../memory/buffers/ParticlesBuffer.hpp | 6 +- .../pmacc/particles/memory/frames/Frame.hpp | 74 +++++++++------ thirdParty/llama | 2 +- 8 files changed, 105 insertions(+), 122 deletions(-) diff --git a/include/picongpu/plugins/openPMD/WriteSpecies.hpp b/include/picongpu/plugins/openPMD/WriteSpecies.hpp index 80e565cc75e..56e8e34ccb9 100644 --- a/include/picongpu/plugins/openPMD/WriteSpecies.hpp +++ b/include/picongpu/plugins/openPMD/WriteSpecies.hpp @@ -113,15 +113,13 @@ namespace picongpu { /* malloc host memory */ log("openPMD: (begin) malloc host memory: %1%") % name; - meta::ForEach> mallocMem; - mallocMem(hostFrame, myNumParticles); + mallocFrameMemory(hostFrame); log("openPMD: ( end ) malloc host memory: %1%") % name; } void free(openPMDFrameType& hostFrame) override { - meta::ForEach> freeMem; - freeMem(hostFrame); + freeFrameMemory(hostFrame); } @@ -182,16 +180,13 @@ namespace picongpu void malloc(std::string name, openPMDFrameType& mappedFrame, uint64_cu const myNumParticles) override { log("openPMD: (begin) malloc mapped memory: %1%") % name; - /*malloc mapped memory*/ - meta::ForEach> mallocMem; - mallocMem(mappedFrame, myNumParticles); + mallocMappedFrameMemory(mappedFrame); log("openPMD: ( end ) malloc mapped memory: %1%") % name; } void free(openPMDFrameType& mappedFrame) override { - meta::ForEach> freeMem; - freeMem(mappedFrame); + freeMappedFrameMemory(mappedFrame); } void prepare(std::string name, openPMDFrameType& mappedFrame, RunParameters rp) override @@ -247,7 +242,7 @@ namespace picongpu using NewParticleDescription = typename ReplaceValueTypeSeq::type; - using openPMDFrameType = Frame; + using openPMDFrameType = Frame; void setParticleAttributes( ::openPMD::ParticleSpecies& record, diff --git a/include/picongpu/plugins/openPMD/restart/LoadParticleAttributesFromOpenPMD.hpp b/include/picongpu/plugins/openPMD/restart/LoadParticleAttributesFromOpenPMD.hpp index 65bb878f606..44d8c12a274 100644 --- a/include/picongpu/plugins/openPMD/restart/LoadParticleAttributesFromOpenPMD.hpp +++ b/include/picongpu/plugins/openPMD/restart/LoadParticleAttributesFromOpenPMD.hpp @@ -68,7 +68,7 @@ namespace picongpu { using Identifier = T_Identifier; using ValueType = typename pmacc::traits::Resolve::type::type; - const uint32_t components = GetNComponents::value; + constexpr uint32_t components = GetNComponents::value; using ComponentType = typename GetComponentsType::type; OpenPMDName openPMDName; @@ -76,6 +76,7 @@ namespace picongpu const std::string name_lookup[] = {"x", "y", "z"}; + // TODO(bgruber): make this a std::shared_ptr with openPMD 0.15 std::shared_ptr loadBfr; if(elements > 0) { @@ -90,7 +91,6 @@ namespace picongpu ::openPMD::RecordComponent rc = components > 1 ? record[name_lookup[n]] : record[::openPMD::RecordComponent::SCALAR]; - ValueType* dataPtr = frame.getIdentifier(Identifier()).getPointer(); if(elements > 0) { @@ -119,12 +119,15 @@ namespace picongpu "%3%") % elements % globalNumElements % openPMDName(); -/* copy component from temporary array to array of structs */ + /* copy component from temporary array to array of structs */ #pragma omp parallel for simd for(size_t i = 0; i < elements; ++i) { - ComponentType* ref = &reinterpret_cast(dataPtr)[i * components + n]; - *ref = loadBfr.get()[i]; + auto& attrib = frame[i][Identifier{}]; + if constexpr(components == 1) + attrib = loadBfr.get()[i]; + else + reinterpret_cast(&attrib)[n] = loadBfr.get()[i]; } } diff --git a/include/picongpu/plugins/openPMD/restart/LoadSpecies.hpp b/include/picongpu/plugins/openPMD/restart/LoadSpecies.hpp index 27070036095..965d4f409e4 100644 --- a/include/picongpu/plugins/openPMD/restart/LoadSpecies.hpp +++ b/include/picongpu/plugins/openPMD/restart/LoadSpecies.hpp @@ -74,7 +74,7 @@ namespace picongpu using NewParticleDescription = typename ReplaceValueTypeSeq::type; - using openPMDFrameType = Frame; + using openPMDFrameType = Frame; /** Load species from openPMD checkpoint storage * @@ -138,9 +138,8 @@ namespace picongpu // memory is visible on host and device openPMDFrameType mappedFrame; log("openPMD: malloc mapped memory: %1%") % speciesName; - /*malloc mapped memory*/ - meta::ForEach> mallocMem; - mallocMem(mappedFrame, totalNumParticles); + + mallocMappedFrameMemory(mappedFrame); meta::ForEach> loadAttributes; @@ -158,9 +157,7 @@ namespace picongpu *(params->cellDescription), picLog::INPUT_OUTPUT()); - /*free host memory*/ - meta::ForEach> freeMem; - freeMem(mappedFrame); + freeMappedFrameMemory(mappedFrame); } log("openPMD: ( end ) load species: %1%") % speciesName; } diff --git a/include/picongpu/plugins/openPMD/writer/ParticleAttribute.hpp b/include/picongpu/plugins/openPMD/writer/ParticleAttribute.hpp index ad9053442f9..85acbe3527b 100644 --- a/include/picongpu/plugins/openPMD/writer/ParticleAttribute.hpp +++ b/include/picongpu/plugins/openPMD/writer/ParticleAttribute.hpp @@ -64,7 +64,7 @@ namespace picongpu { using Identifier = T_Identifier; using ValueType = typename pmacc::traits::Resolve::type::type; - const uint32_t components = GetNComponents::value; + constexpr uint32_t components = GetNComponents::value; using ComponentType = typename GetComponentsType::type; OpenPMDName openPMDName; @@ -93,7 +93,7 @@ namespace picongpu ::openPMD::RecordComponent recordComponent = components > 1 ? record[name_lookup[d]] : record[::openPMD::MeshRecordComponent::SCALAR]; - std::string datasetName = components > 1 ? baseName + "/" + name_lookup[d] : baseName; + const std::string datasetName = components > 1 ? baseName + "/" + name_lookup[d] : baseName; params->initDataset(recordComponent, openPMDType, {globalElements}, datasetName); if(unit.size() >= (d + 1)) @@ -107,7 +107,6 @@ namespace picongpu continue; } - ValueType* dataPtr = frame.getIdentifier(Identifier()).getPointer(); // can be moved up? // ask openPMD to create a buffer for us // in some backends (ADIOS2), this allows avoiding memcopies auto span = storeChunkSpan( @@ -132,7 +131,11 @@ namespace picongpu #pragma omp parallel for simd for(size_t i = 0; i < elements; ++i) { - span[i] = reinterpret_cast(dataPtr)[d + i * components]; + const auto attrib = frame[i][Identifier{}]; + if constexpr(components == 1) + span[i] = attrib; + else + span[i] = reinterpret_cast(&attrib)[d]; } flushSeries(*params->openPMDSeries, PreferredFlushTarget::Disk); diff --git a/include/picongpu/plugins/output/WriteSpeciesCommon.hpp b/include/picongpu/plugins/output/WriteSpeciesCommon.hpp index 81d049f2b43..d8cb9e71ffe 100644 --- a/include/picongpu/plugins/output/WriteSpeciesCommon.hpp +++ b/include/picongpu/plugins/output/WriteSpeciesCommon.hpp @@ -31,51 +31,48 @@ namespace picongpu { using namespace pmacc; - - template - struct MallocMappedMemory + template + void mallocMappedFrameMemory(Frame& frame) { - template - HINLINE void operator()(ValueType& v1, const size_t size) const - { - using type = typename pmacc::traits::Resolve::type::type; - - bool isMappedMemorySupported = alpaka::hasMappedBufSupport<::alpaka::Pltf>; - - PMACC_VERIFY_MSG(isMappedMemorySupported, "Device must support mapped memory!"); + constexpr bool isMappedMemorySupported = alpaka::hasMappedBufSupport<::alpaka::Pltf>; + PMACC_VERIFY_MSG(isMappedMemorySupported, "Device must support mapped memory!"); - type* ptr = nullptr; + int i = 0; + for(std::byte*& ptr : frame.blobs()) + { + const auto size = frame.blobSize(i); if(size != 0) { // Memory is automatically mapped to the device if supported. - CUDA_CHECK(cuplaMallocHost((void**) &ptr, size * sizeof(type))); + CUDA_CHECK(cuplaMallocHost((void**) &ptr, size)); } - v1.getIdentifier(T_Type()) = VectorDataBox(ptr); + else + ptr = nullptr; + i++; } - }; + } /** allocate memory on host * * This functor use `new[]` to allocate memory */ - template - struct MallocHostMemory + template + void mallocFrameMemory(Frame& frame) { - template - HINLINE void operator()(ValueType& v1, const size_t size) const - { - using Attribute = T_Attribute; - using type = typename pmacc::traits::Resolve::type::type; + constexpr bool isMappedMemorySupported = alpaka::hasMappedBufSupport<::alpaka::Pltf>; + PMACC_VERIFY_MSG(isMappedMemorySupported, "Device must support mapped memory!"); - type* ptr = nullptr; + int i = 0; + for(std::byte*& ptr : frame.blobs()) + { + const auto size = frame.blobSize(i); if(size != 0) - { - ptr = new type[size]; - } - v1.getIdentifier(Attribute()) = VectorDataBox(ptr); + ptr = new std::byte[size]; + else + ptr = nullptr; + i++; } - }; - + } /** copy species to host memory * @@ -94,42 +91,18 @@ namespace picongpu } }; - template - struct FreeMappedMemory + template + void freeMappedFrameMemory(Frame& frame) { - template - HINLINE void operator()(ValueType& value) const - { - auto* ptr = value.getIdentifier(T_Type()).getPointer(); + for(auto* ptr : frame.blobs()) if(ptr != nullptr) - { CUDA_CHECK(cuplaFreeHost(ptr)); - } - } - }; + } - //! Free memory - template - struct FreeHostMemory + template + void freeFrameMemory(Frame& frame) { - template - HINLINE void operator()(ValueType& value) const - { - using Attribute = T_Attribute; - - auto* ptr = value.getIdentifier(Attribute()).getPointer(); + for(auto* ptr : frame.blobs()) delete[] ptr; - } - }; - - /*functor to create a pair for a MapTuple map*/ - struct OperatorCreateVectorBox - { - template - struct apply - { - typedef bmpl::pair::type::type>> type; - }; - }; - + } } // namespace picongpu diff --git a/include/pmacc/particles/memory/buffers/ParticlesBuffer.hpp b/include/pmacc/particles/memory/buffers/ParticlesBuffer.hpp index c765f1ab06f..e459fd815a1 100644 --- a/include/pmacc/particles/memory/buffers/ParticlesBuffer.hpp +++ b/include/pmacc/particles/memory/buffers/ParticlesBuffer.hpp @@ -103,9 +103,7 @@ namespace pmacc * * a group of particles is stored as frame */ - using FrameType = Frame< - OperatorCreatePairStaticArray::type::value>, - FrameDescription>; + using FrameType = Frame::type::value, FrameDescription>; using FrameDescriptionBorder = typename ReplaceValueTypeSeq::type; @@ -115,7 +113,7 @@ namespace pmacc * - each frame contains only one particle * - local administration attributes of a particle are removed */ - using FrameTypeBorder = Frame, FrameDescriptionBorder>; + using FrameTypeBorder = Frame<1, FrameDescriptionBorder>; using SuperCellType = SuperCell; diff --git a/include/pmacc/particles/memory/frames/Frame.hpp b/include/pmacc/particles/memory/frames/Frame.hpp index 5633301ac9b..a2346afbf8d 100644 --- a/include/pmacc/particles/memory/frames/Frame.hpp +++ b/include/pmacc/particles/memory/frames/Frame.hpp @@ -61,32 +61,48 @@ namespace pmacc using RecordDimFromValueTypeSeq = boost::mp11:: mp_rename>, llama::Record>; - template - inline constexpr auto nonTypeArgOf = nullptr; // = delete; - - template typename T, auto I> - inline constexpr auto nonTypeArgOf> = I; - - template + template struct ViewHolder { - static constexpr std::size_t particlesPerFrame - = nonTypeArgOf; // T_ParticleDescription::SuperCellSize + using IndexType = int; // TODO(bgruber): where do I get this type from? + static constexpr IndexType particlesPerFrame + = (T_Size == llama::dyn) ? static_cast(llama::dyn) : static_cast(T_Size); + using FrameMapping = typename T_ParticleDescription::FrameMapping; using RawRecordDim = RecordDimFromValueTypeSeq; using SplitRecordDim = llama::TransformLeaves; using RecordDim = std::conditional_t; - using IndexType = int; // TODO(bgruber): where do I get this type from? using ArrayExtents = llama::ArrayExtents; - using Mapping = boost::mp11::mp_if_c< + using Mapping = std::conditional_t< particlesPerFrame == 1, llama::mapping::One, typename FrameMapping::template fn>; - static_assert(Mapping::blobCount == 1, "Only frame mappings with a single blob are supported"); - using LlamaViewType = llama::View>; + static_assert( + particlesPerFrame == llama::dyn || Mapping::blobCount == 1, + "For statically sizes frames, only mappings with a single blob are supported"); + using BlobType = std::conditional_t< + particlesPerFrame == llama::dyn, + std::byte*, + llama::Array>; + using LlamaViewType = llama::View; LlamaViewType view; + + auto& blobs() + { + return view.storageBlobs; + } + + auto& blobPtr(int i) + { + return view.storageBlobs[i]; + } + + auto blobSize(int i) + { + return view.mapping().blobSize(i); + } }; template @@ -125,25 +141,23 @@ namespace pmacc /** Frame is a storage for arbitrary number >0 of Particles with attributes * - * @tparam T_CreatePairOperator unary template operator to create a boost pair - * from single type ( pair ) - * @see MapTupel + * @tparam T_Size Static number of particles this frame stores, or llama::dyn for dynamic size * @tparam T_ValueTypeSeq sequence with value_identifier * @tparam T_MethodsList sequence of classes with particle methods * (e.g. calculate mass, gamma, ...) * @tparam T_Flags sequence with identifiers to add flags on a frame * (e.g. useSolverXY, calcRadiation, ...) */ - template + template struct Frame; - template + template struct Frame : public InheritLinearly - , public detail::ViewHolder + , public detail::ViewHolder , public InheritLinearly>>::type> + bmpl::apply1>>::type> { using ParticleDescription = T_ParticleDescription; using Name = typename ParticleDescription::Name; @@ -152,7 +166,7 @@ namespace pmacc using MethodsList = typename ParticleDescription::MethodsList; using FlagList = typename ParticleDescription::FlagsList; using FrameExtensionList = typename ParticleDescription::FrameExtensionList; - using ThisType = Frame; + using ThisType = Frame; /* type of a single particle*/ using ParticleType = pmacc::Particle; @@ -206,11 +220,11 @@ namespace pmacc namespace traits { - template - struct HasIdentifier, T_IdentifierName> + template + struct HasIdentifier, T_IdentifierName> { private: - using FrameType = pmacc::Frame; + using FrameType = pmacc::Frame; public: using ValueTypeSeq = typename FrameType::ValueTypeSeq; @@ -222,11 +236,11 @@ namespace pmacc using type = bmpl::contains; }; - template - struct HasFlag, T_IdentifierName> + template + struct HasFlag, T_IdentifierName> { private: - using FrameType = pmacc::Frame; + using FrameType = pmacc::Frame; using SolvedAliasName = typename GetFlagType::type; using FlagList = typename FrameType::FlagList; @@ -234,11 +248,11 @@ namespace pmacc using type = bmpl::contains; }; - template - struct GetFlagType, T_IdentifierName> + template + struct GetFlagType, T_IdentifierName> { private: - using FrameType = pmacc::Frame; + using FrameType = pmacc::Frame; using FlagList = typename FrameType::FlagList; public: diff --git a/thirdParty/llama b/thirdParty/llama index 5aa40486d56..f87dd5dba9b 160000 --- a/thirdParty/llama +++ b/thirdParty/llama @@ -1 +1 @@ -Subproject commit 5aa40486d56ec4462839e9ec7db187f8f92f5262 +Subproject commit f87dd5dba9b7ba148c6a9cd8895688243e847deb