From 5d2290e8d2d74f8baef50ce82c71bfc2f9d5708a Mon Sep 17 00:00:00 2001 From: Max Yang Date: Fri, 28 Jun 2024 15:36:53 -0700 Subject: [PATCH 1/8] Array: use libc++ resize logic to avoid excessive allocations --- src/axom/core/Array.hpp | 7 ++++++- src/axom/core/tests/core_array.hpp | 4 +++- 2 files changed, 9 insertions(+), 2 deletions(-) diff --git a/src/axom/core/Array.hpp b/src/axom/core/Array.hpp index a7d5d2d21f..b8b9e9330f 100644 --- a/src/axom/core/Array.hpp +++ b/src/axom/core/Array.hpp @@ -1685,7 +1685,12 @@ template inline void Array::dynamicRealloc(IndexType new_num_elements) { assert(m_resize_ratio >= 1.0); - IndexType new_capacity = new_num_elements * m_resize_ratio + 0.5; + + // Using resize strategy from LLVM libc++ (vector::__recommend()): + // new_capacity = max(capacity() * resize_ratio, new_num_elements) + IndexType new_capacity = + axom::utilities::max(this->capacity() * m_resize_ratio + 0.5, + new_num_elements); const IndexType block_size = this->blockSize(); const IndexType remainder = new_capacity % block_size; if(remainder != 0) diff --git a/src/axom/core/tests/core_array.hpp b/src/axom/core/tests/core_array.hpp index 6dc6eebf91..b7cacc83f5 100644 --- a/src/axom/core/tests/core_array.hpp +++ b/src/axom/core/tests/core_array.hpp @@ -25,7 +25,9 @@ axom::IndexType calc_new_capacity(axom::Array& v, axom::IndexType increase) axom::IndexType new_num_elements = v.size() + increase; if(new_num_elements > v.capacity()) { - return new_num_elements * v.getResizeRatio() + 0.5; + axom::IndexType capacity_expanded = v.capacity() * v.getResizeRatio() + 0.5; + return axom::utilities::max(capacity_expanded, + new_num_elements); } return v.capacity(); From 6ba754e3ee97e5c78ad8bf86ccddd9ed9f2b2a48 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Fri, 28 Jun 2024 17:24:54 -0700 Subject: [PATCH 2/8] Array: separate device push_back into separate functions --- src/axom/core/Array.hpp | 73 ++++++++++++++++++---- src/axom/core/tests/core_array_for_all.hpp | 6 +- 2 files changed, 64 insertions(+), 15 deletions(-) diff --git a/src/axom/core/Array.hpp b/src/axom/core/Array.hpp index b8b9e9330f..6df6668387 100644 --- a/src/axom/core/Array.hpp +++ b/src/axom/core/Array.hpp @@ -625,13 +625,34 @@ class Array : public ArrayBase> * \param [in] value the value to be added to the back. * * \note Reallocation is done if the new size will exceed the capacity. - * \note If used in a device kernel, the number of push_backs must not exceed - * the capacity, since device-side reallocations aren't supported. - * \note Array must be allocated in unified memory if calling on the device. - * + * * \pre DIM == 1 */ - AXOM_HOST_DEVICE void push_back(const T& value); + void push_back(const T& value); + + /*! + * \brief Push a value to the back of the array. + * + * \param [in] value the value to move to the back. + * + * \note Reallocation is done if the new size will exceed the capacity. + * + * \pre DIM == 1 + */ + void push_back(T&& value); + + /*! + * \brief Inserts new element at the end of the Array. + * + * \param [in] args the arguments to forward to constructor of the element. + * + * \note Reallocation is done if the new size will exceed the capacity. + * \note The size increases by 1. + * + * \pre DIM == 1 + */ + template + void emplace_back(Args&&... args); /*! * \brief Push a value to the back of the array. @@ -642,10 +663,13 @@ class Array : public ArrayBase> * \note If used in a device kernel, the number of push_backs must not exceed * the capacity, since device-side reallocations aren't supported. * \note Array must be allocated in unified memory if calling on the device. - * + * * \pre DIM == 1 */ - AXOM_HOST_DEVICE void push_back(T&& value); + /// @{ + AXOM_HOST_DEVICE void pushBackDevice(const T& value); + AXOM_HOST_DEVICE void pushBackDevice(T&& value); + /// @} /*! * \brief Inserts new element at the end of the Array. @@ -657,11 +681,11 @@ class Array : public ArrayBase> * \note If used in a device kernel, the number of push_backs must not exceed * the capacity, since device-side reallocations aren't supported. * \note Array must be allocated in unified memory if calling on the device. - * + * * \pre DIM == 1 */ template - AXOM_HOST_DEVICE void emplace_back(Args&&... args); + AXOM_HOST_DEVICE void emplaceBackDevice(Args&&... args); /// @} @@ -1425,7 +1449,7 @@ inline typename Array::ArrayIterator Array::emplac //------------------------------------------------------------------------------ template -AXOM_HOST_DEVICE inline void Array::push_back(const T& value) +inline void Array::push_back(const T& value) { static_assert(DIM == 1, "push_back is only supported for 1D arrays"); emplace_back(value); @@ -1433,7 +1457,7 @@ AXOM_HOST_DEVICE inline void Array::push_back(const T& value) //------------------------------------------------------------------------------ template -AXOM_HOST_DEVICE inline void Array::push_back(T&& value) +inline void Array::push_back(T&& value) { static_assert(DIM == 1, "push_back is only supported for 1D arrays"); emplace_back(std::move(value)); @@ -1442,8 +1466,33 @@ AXOM_HOST_DEVICE inline void Array::push_back(T&& value) //------------------------------------------------------------------------------ AXOM_SUPPRESS_HD_WARN template +AXOM_HOST_DEVICE inline void Array::pushBackDevice(const T& value) +{ + static_assert(DIM == 1, "pushBackDevice is only supported for 1D arrays"); + emplaceBackDevice(value); +} + +//------------------------------------------------------------------------------ +template +AXOM_HOST_DEVICE inline void Array::pushBackDevice(T&& value) +{ + static_assert(DIM == 1, "pushBackDevice is only supported for 1D arrays"); + emplaceBackDevice(std::move(value)); +} + +//------------------------------------------------------------------------------ +template +template +inline void Array::emplace_back(Args&&... args) +{ + static_assert(DIM == 1, "emplace_back is only supported for 1D arrays"); + emplace(size(), std::forward(args)...); +} + +//------------------------------------------------------------------------------ +template template -AXOM_HOST_DEVICE inline void Array::emplace_back(Args&&... args) +AXOM_HOST_DEVICE inline void Array::emplaceBackDevice(Args&&... args) { static_assert(DIM == 1, "emplace_back is only supported for 1D arrays"); #ifdef AXOM_DEVICE_CODE diff --git a/src/axom/core/tests/core_array_for_all.hpp b/src/axom/core/tests/core_array_for_all.hpp index b5174a3e96..6d278a1281 100644 --- a/src/axom/core/tests/core_array_for_all.hpp +++ b/src/axom/core/tests/core_array_for_all.hpp @@ -1180,15 +1180,15 @@ AXOM_TYPED_TEST(core_array_for_all, device_insert) { #pragma omp critical { - arr_v[0].emplace_back(3 * idx + 5); + arr_v[0].emplaceBackDevice(3 * idx + 5); } } else { - arr_v[0].emplace_back(3 * idx + 5); + arr_v[0].emplaceBackDevice(3 * idx + 5); } #else - arr_v[0].emplace_back(3 * idx + 5); + arr_v[0].emplaceBackDevice(3 * idx + 5); #endif }); From 20faf7db7d16c2582dae11ec3273eef4b248279a Mon Sep 17 00:00:00 2001 From: Max Yang Date: Fri, 28 Jun 2024 17:25:19 -0700 Subject: [PATCH 3/8] Silence some MDMapping host/device warnings --- src/axom/core/MDMapping.hpp | 25 +++++++++++++------------ 1 file changed, 13 insertions(+), 12 deletions(-) diff --git a/src/axom/core/MDMapping.hpp b/src/axom/core/MDMapping.hpp index 5167dcdc81..78492a190e 100644 --- a/src/axom/core/MDMapping.hpp +++ b/src/axom/core/MDMapping.hpp @@ -49,9 +49,9 @@ class MDMapping @param [in] fastestStrideLength Stride in the fastest-changing direction. */ - MDMapping(const axom::StackArray& shape, - axom::ArrayStrideOrder arrayStrideOrder, - int fastestStrideLength = 1) + AXOM_HOST_DEVICE MDMapping(const axom::StackArray& shape, + axom::ArrayStrideOrder arrayStrideOrder, + int fastestStrideLength = 1) { initializeShape(shape, arrayStrideOrder, fastestStrideLength); } @@ -66,9 +66,9 @@ class MDMapping fastest-changing direction. */ template - MDMapping(const axom::StackArray& shape, - const axom::StackArray& slowestDirs, - int fastestStrideLength = 1) + AXOM_HOST_DEVICE MDMapping(const axom::StackArray& shape, + const axom::StackArray& slowestDirs, + int fastestStrideLength = 1) { initializeShape(shape, slowestDirs, fastestStrideLength); } @@ -81,8 +81,8 @@ class MDMapping @param [in] orderSource ArrayIndex to copy stride order from. */ - MDMapping(const axom::StackArray& shape, - const axom::MDMapping& orderSource) + AXOM_HOST_DEVICE MDMapping(const axom::StackArray& shape, + const axom::MDMapping& orderSource) { initializeShape(shape, orderSource); } @@ -98,7 +98,8 @@ class MDMapping clash with the more prevalent usage of constructing from the array's shape. */ - MDMapping(const axom::StackArray& strides) : m_strides(strides) + AXOM_HOST_DEVICE MDMapping(const axom::StackArray& strides) + : m_strides(strides) { initializeStrides(strides); } @@ -117,7 +118,7 @@ class MDMapping shape is to be determined. Initialize the mapping with its own slowestDirs() to preserve stride order as its shape changes. */ - MDMapping(ArrayStrideOrder arrayStrideOrder) + AXOM_HOST_DEVICE MDMapping(ArrayStrideOrder arrayStrideOrder) { axom::StackArray shape; for(int d = 0; d < DIM; ++d) @@ -238,8 +239,8 @@ class MDMapping << "Likely, multi-dim array shape is 1 in some direction.\n" << "Impossible to compute index ordering.\n" << "Please use a different MDMapping initializer.\n"; -#endif utilities::processAbort(); +#endif } // 2nd argument doesn't matter because strides are unique. @@ -275,7 +276,7 @@ class MDMapping { if(m_strides[m_slowestDirs[s]] < m_strides[m_slowestDirs[d]]) { - std::swap(m_slowestDirs[s], m_slowestDirs[d]); + axom::utilities::swap(m_slowestDirs[s], m_slowestDirs[d]); } } } From 6cf94a99a1c82b3c0770c65224cfa04837741338 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Mon, 29 Jul 2024 13:51:37 -0700 Subject: [PATCH 4/8] Mint: fixup a buffer overflow detected by asan --- src/axom/mint/execution/internal/for_all_faces.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/axom/mint/execution/internal/for_all_faces.hpp b/src/axom/mint/execution/internal/for_all_faces.hpp index 3c8a92c093..e688d3c8dc 100644 --- a/src/axom/mint/execution/internal/for_all_faces.hpp +++ b/src/axom/mint/execution/internal/for_all_faces.hpp @@ -993,7 +993,7 @@ inline void for_all_faces_impl(xargs::coords, coords[3 * i + 2] = z[nodeID]; } - numerics::Matrix coordsMatrix(dimension, 4, coords, NO_COPY); + numerics::Matrix coordsMatrix(dimension, numNodes, coords, NO_COPY); kernel(faceID, coordsMatrix, nodeIDs); }); } From 5e7ce862024f0c1ec021cf8841bf64a416f5810f Mon Sep 17 00:00:00 2001 From: Max Yang Date: Thu, 1 Aug 2024 14:59:30 -0700 Subject: [PATCH 5/8] Sidre: use libc++ reserve logic for sidre::Array --- src/axom/sidre/core/Array.hpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/axom/sidre/core/Array.hpp b/src/axom/sidre/core/Array.hpp index 7a1877cdc3..52ba8ec195 100644 --- a/src/axom/sidre/core/Array.hpp +++ b/src/axom/sidre/core/Array.hpp @@ -571,7 +571,9 @@ inline void Array::dynamicRealloc(axom::IndexType new_num_elements) "Resize ratio of " << this->m_resize_ratio << " doesn't support dynamic resizing"); - IndexType new_capacity = new_num_elements * this->m_resize_ratio + 0.5; + IndexType new_capacity = axom::utilities::max( + this->capacity() * this->getResizeRatio() + 0.5, + new_num_elements); const IndexType block_size = this->blockSize(); const IndexType remainder = new_capacity % block_size; if(remainder != 0) From c5ab959f5657b43f435c0a4c4feb4955ea05e3c2 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Thu, 1 Aug 2024 15:18:47 -0700 Subject: [PATCH 6/8] Fixup Mint and Sidre tests to test new resize logic --- .../mint/tests/mint_mesh_connectivity_array.cpp | 13 +++++++++---- src/axom/mint/tests/mint_mesh_unstructured_mesh.cpp | 6 ++++-- src/axom/sidre/tests/sidre_mcarray.cpp | 9 ++++++++- 3 files changed, 21 insertions(+), 7 deletions(-) diff --git a/src/axom/mint/tests/mint_mesh_connectivity_array.cpp b/src/axom/mint/tests/mint_mesh_connectivity_array.cpp index 6d25ce5bfd..78765c7453 100644 --- a/src/axom/mint/tests/mint_mesh_connectivity_array.cpp +++ b/src/axom/mint/tests/mint_mesh_connectivity_array.cpp @@ -41,8 +41,9 @@ IndexType calc_ID_capacity(const ConnectivityArray& connec, if(new_n_IDs > connec.getIDCapacity()) { IndexType stride = connec.getNumberOfValuesForID(); - IndexType newCapacity = - static_cast(new_n_IDs * stride * connec.getResizeRatio() + 0.5); + IndexType newCapacity = axom::utilities::max( + connec.getValueCapacity() * connec.getResizeRatio() + 0.5, + new_n_IDs * stride); IndexType remainder = newCapacity % stride; if(remainder != 0) { @@ -67,7 +68,9 @@ IndexType calc_ID_capacity(const ConnectivityArray& connec, IndexType new_n_IDs = connec.getNumberOfIDs() + increase; if(new_n_IDs > connec.getIDCapacity()) { - return static_cast(new_n_IDs * connec.getResizeRatio() + 0.5); + return axom::utilities::max( + connec.getIDCapacity() * connec.getResizeRatio() + 0.5, + new_n_IDs); } return connec.getIDCapacity(); @@ -101,7 +104,9 @@ IndexType calc_value_capacity(const ConnectivityArray& connec, IndexType new_n_values = connec.getNumberOfValues() + increase; if(new_n_values > connec.getValueCapacity()) { - return static_cast(new_n_values * connec.getResizeRatio() + 0.5); + return axom::utilities::max( + connec.getValueCapacity() * connec.getResizeRatio() + 0.5, + new_n_values); } return connec.getValueCapacity(); diff --git a/src/axom/mint/tests/mint_mesh_unstructured_mesh.cpp b/src/axom/mint/tests/mint_mesh_unstructured_mesh.cpp index bdd1c301f3..7d172dc137 100644 --- a/src/axom/mint/tests/mint_mesh_unstructured_mesh.cpp +++ b/src/axom/mint/tests/mint_mesh_unstructured_mesh.cpp @@ -2200,7 +2200,8 @@ void resize_cells(UnstructuredMesh* mesh) /* Append one more, should trigger a resize. */ append_cell_single(mesh, 1); n_cells++; - cell_capacity = static_cast(n_cells * resize_ratio + 0.5); + cell_capacity = + axom::utilities::max(cell_capacity * resize_ratio + 0.5, n_cells); connec_capacity = mesh->getCellNodesCapacity(); ASSERT_EQ(n_cells, mesh->getNumberOfCells()); ASSERT_EQ(cell_capacity, mesh->getCellCapacity()); @@ -2217,7 +2218,8 @@ void resize_cells(UnstructuredMesh* mesh) mesh->setCellResizeRatio(resize_ratio); append_cell_multiple(mesh, 100); n_cells += 100; - cell_capacity = static_cast(resize_ratio * n_cells + 0.5); + cell_capacity = + axom::utilities::max(cell_capacity * resize_ratio + 0.5, n_cells); ASSERT_EQ(n_cells, mesh->getNumberOfCells()); ASSERT_EQ(cell_capacity, mesh->getCellCapacity()); diff --git a/src/axom/sidre/tests/sidre_mcarray.cpp b/src/axom/sidre/tests/sidre_mcarray.cpp index eb9e714d78..23a39c0c93 100644 --- a/src/axom/sidre/tests/sidre_mcarray.cpp +++ b/src/axom/sidre/tests/sidre_mcarray.cpp @@ -34,7 +34,14 @@ axom::IndexType calc_new_capacity(MCArray& v, axom::IndexType increase) axom::IndexType new_num_tuples = (v.size() / num_components) + increase; if((new_num_tuples * num_components) > v.capacity()) { - return new_num_tuples * v.getResizeRatio() + 0.5; + axom::IndexType new_capacity = v.capacity() * v.getResizeRatio() + 0.5; + axom::IndexType remainder = new_capacity % num_components; + if(remainder > 0) + { + new_capacity += num_components - remainder; + } + return axom::utilities::max(new_capacity / num_components, + new_num_tuples); } return v.capacity() / num_components; From 2b7daa9de7c1f1c1b97ce1ec5281daa6b3be9295 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Thu, 1 Aug 2024 15:23:01 -0700 Subject: [PATCH 7/8] Array: use snake_case for device push_back/emplace_back --- src/axom/core/Array.hpp | 20 ++++++++++---------- src/axom/core/tests/core_array_for_all.hpp | 6 +++--- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/src/axom/core/Array.hpp b/src/axom/core/Array.hpp index 6df6668387..303753b46f 100644 --- a/src/axom/core/Array.hpp +++ b/src/axom/core/Array.hpp @@ -667,8 +667,8 @@ class Array : public ArrayBase> * \pre DIM == 1 */ /// @{ - AXOM_HOST_DEVICE void pushBackDevice(const T& value); - AXOM_HOST_DEVICE void pushBackDevice(T&& value); + AXOM_HOST_DEVICE void push_back_device(const T& value); + AXOM_HOST_DEVICE void push_back_device(T&& value); /// @} /*! @@ -685,7 +685,7 @@ class Array : public ArrayBase> * \pre DIM == 1 */ template - AXOM_HOST_DEVICE void emplaceBackDevice(Args&&... args); + AXOM_HOST_DEVICE void emplace_back_device(Args&&... args); /// @} @@ -1466,18 +1466,18 @@ inline void Array::push_back(T&& value) //------------------------------------------------------------------------------ AXOM_SUPPRESS_HD_WARN template -AXOM_HOST_DEVICE inline void Array::pushBackDevice(const T& value) +AXOM_HOST_DEVICE inline void Array::push_back_device(const T& value) { - static_assert(DIM == 1, "pushBackDevice is only supported for 1D arrays"); - emplaceBackDevice(value); + static_assert(DIM == 1, "push_back_device is only supported for 1D arrays"); + emplace_back_device(value); } //------------------------------------------------------------------------------ template -AXOM_HOST_DEVICE inline void Array::pushBackDevice(T&& value) +AXOM_HOST_DEVICE inline void Array::push_back_device(T&& value) { - static_assert(DIM == 1, "pushBackDevice is only supported for 1D arrays"); - emplaceBackDevice(std::move(value)); + static_assert(DIM == 1, "push_back_device is only supported for 1D arrays"); + emplace_back_device(std::move(value)); } //------------------------------------------------------------------------------ @@ -1492,7 +1492,7 @@ inline void Array::emplace_back(Args&&... args) //------------------------------------------------------------------------------ template template -AXOM_HOST_DEVICE inline void Array::emplaceBackDevice(Args&&... args) +AXOM_HOST_DEVICE inline void Array::emplace_back_device(Args&&... args) { static_assert(DIM == 1, "emplace_back is only supported for 1D arrays"); #ifdef AXOM_DEVICE_CODE diff --git a/src/axom/core/tests/core_array_for_all.hpp b/src/axom/core/tests/core_array_for_all.hpp index 6d278a1281..04ee2459dc 100644 --- a/src/axom/core/tests/core_array_for_all.hpp +++ b/src/axom/core/tests/core_array_for_all.hpp @@ -1180,15 +1180,15 @@ AXOM_TYPED_TEST(core_array_for_all, device_insert) { #pragma omp critical { - arr_v[0].emplaceBackDevice(3 * idx + 5); + arr_v[0].emplace_back_device(3 * idx + 5); } } else { - arr_v[0].emplaceBackDevice(3 * idx + 5); + arr_v[0].emplace_back_device(3 * idx + 5); } #else - arr_v[0].emplaceBackDevice(3 * idx + 5); + arr_v[0].emplace_back_device(3 * idx + 5); #endif }); From 655bd1e7b85e740643a9f839dcc627048a662847 Mon Sep 17 00:00:00 2001 From: Max Yang Date: Fri, 2 Aug 2024 11:19:10 -0700 Subject: [PATCH 8/8] Primal: fix for Polygon::addVertex --- src/axom/primal/geometry/Polygon.hpp | 21 +++++++++++++++++++-- 1 file changed, 19 insertions(+), 2 deletions(-) diff --git a/src/axom/primal/geometry/Polygon.hpp b/src/axom/primal/geometry/Polygon.hpp index 2a54ffbe37..6eee85d4b4 100644 --- a/src/axom/primal/geometry/Polygon.hpp +++ b/src/axom/primal/geometry/Polygon.hpp @@ -193,8 +193,25 @@ class Polygon * \sa axom::StaticArray::push_back() for behavior when array type is static * and the list of vertices is full. */ - AXOM_HOST_DEVICE - void addVertex(const PointType& pt) { m_vertices.push_back(pt); } + /// @{ + template = 0> + AXOM_HOST_DEVICE void addVertex(const PointType& pt) + { + m_vertices.push_back(pt); + } + + template = 0> + AXOM_HOST_DEVICE void addVertex(const PointType& pt) + { +#ifdef AXOM_DEVICE_CODE + m_vertices.push_back_device(pt); +#else + m_vertices.push_back(pt); +#endif + } + /// @} /// Clears the list of vertices (dynamic array specialization). /// Specializations are necessary to remove __host__ __device__ warning for