diff --git a/src/libs/blueprint/conduit_blueprint_mesh.cpp b/src/libs/blueprint/conduit_blueprint_mesh.cpp index a45bf163a..30f6bd81d 100644 --- a/src/libs/blueprint/conduit_blueprint_mesh.cpp +++ b/src/libs/blueprint/conduit_blueprint_mesh.cpp @@ -8324,7 +8324,7 @@ void polyhedral_face_centers_normals(const IndexAccessor subelements_connectivit Vector *allFaceNormalsPtr = allFaceNormals.data(); // Compute face centers and normals. - conduit::execution::for_all(0, totalNumFaces, [=](conduit::index_t f) { + conduit::execution::forall(0, totalNumFaces, [=](conduit::index_t f) { const int NUM_VERTS = 4; const auto size = subelements_sizes[f]; const auto offset = subelements_offsets[f]; @@ -8441,7 +8441,7 @@ void polyhedral_elem_centers(const IndexAccessor elements_connectivity, allElemCenters.resize(totalNumElems); Vector *allElemCentersPtr = allElemCenters.data(); const Vector *allFaceCentersPtr = allFaceCenters.data(); - conduit::execution::for_all(0, totalNumElems, [=](conduit::index_t i) { + conduit::execution::forall(0, totalNumElems, [=](conduit::index_t i) { const auto size = elements_sizes[i]; const auto offset = elements_offsets[i]; Vector center {}; @@ -8564,7 +8564,7 @@ static void polyhedral_to_hexes(const conduit::Node &n_topo, conduit::Node &n_ou const Vector *allFaceCentersPtr = allFaceCenters.data(); const Vector *allFaceNormalsPtr = allFaceNormals.data(); const Vector *allElemCentersPtr = allElemCenters.data(); - conduit::execution::for_all(0, nElem, [=](conduit::index_t i) { + conduit::execution::forall(0, nElem, [=](conduit::index_t i) { constexpr int FORWARD = 1; constexpr int BACKWARD = -1; // Determine face orientations with respect to this element. diff --git a/src/libs/blueprint/conduit_blueprint_mesh_topology_metadata.cpp b/src/libs/blueprint/conduit_blueprint_mesh_topology_metadata.cpp index dd8baff39..56c2e5a5f 100644 --- a/src/libs/blueprint/conduit_blueprint_mesh_topology_metadata.cpp +++ b/src/libs/blueprint/conduit_blueprint_mesh_topology_metadata.cpp @@ -60,14 +60,6 @@ namespace mesh namespace utils { -// We may tag certain algorithms as ParallelExec if it is safe to do so. -using SerialExec = conduit::execution::SerialExec; -#if defined(CONDUIT_USE_OPENMP) -using ParallelExec = conduit::execution::OpenMPExec; -#else -using ParallelExec = conduit::execution::SerialExec; -#endif - //--------------------------------------------------------------------------- void yaml_print(std::ostream &os, const conduit::Node &node) @@ -925,7 +917,8 @@ class TopologyMetadata::Implementation : public TopologyMetadataBase CONDUIT_ANNOTATE_MARK_BEGIN("Labeling"); std::vector> faceid_to_ef(nelem_faces); - conduit::execution::for_all(0, nelem, [&](index_t elem) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, nelem, [&](index_t elem) { // Get the element faces, storing them all in face_pts. index_t elemstart = elem * points_per_elem; @@ -963,7 +956,7 @@ class TopologyMetadata::Implementation : public TopologyMetadataBase // general faceid, then by their elemface "ef", which should keep the // elements in order. CONDUIT_ANNOTATE_MARK_BEGIN("Sort labels"); - conduit::execution::sort(faceid_to_ef.begin(), faceid_to_ef.end()); + conduit::execution::sort(policy, faceid_to_ef.begin(), faceid_to_ef.end()); CONDUIT_ANNOTATE_MARK_END("Sort labels"); #ifdef DEBUG_PRINT std::cout << "faceid_to_ef.sorted = " << faceid_to_ef << std::endl; @@ -982,8 +975,8 @@ class TopologyMetadata::Implementation : public TopologyMetadataBase #endif // Sort on ef to get back to a ef->unique mapping. - conduit::execution::sort( - ef_to_unique.begin(), ef_to_unique.end(), + conduit::execution::sort( + policy, ef_to_unique.begin(), ef_to_unique.end(), [&](const std::pair &lhs, const std::pair &rhs) { // Only sort using the ef value. @@ -1135,7 +1128,8 @@ class TopologyMetadata::Implementation : public TopologyMetadataBase std::vector> edgeid_to_ee(nelem_edges); std::vector> ee_to_edge(nelem_edges); - conduit::execution::for_all(0, nelem, [&](index_t elem) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, nelem, [&](index_t elem) { constexpr size_t MAX_VERTS = 32; @@ -1173,7 +1167,7 @@ class TopologyMetadata::Implementation : public TopologyMetadataBase // Sort edgeid_to_ee so any like edges will be sorted. CONDUIT_ANNOTATE_MARK_BEGIN("Sort labels"); - conduit::execution::sort(edgeid_to_ee.begin(), edgeid_to_ee.end()); + conduit::execution::sort(policy, edgeid_to_ee.begin(), edgeid_to_ee.end()); CONDUIT_ANNOTATE_MARK_END("Sort labels"); #ifdef DEBUG_PRINT std::cout << "edgeid_to_ee.sorted = " << edgeid_to_ee << std::endl; @@ -1189,8 +1183,8 @@ class TopologyMetadata::Implementation : public TopologyMetadataBase #endif // Sort on ef to get back to a ef->unique mapping. - conduit::execution::sort( - ee_to_unique.begin(), ee_to_unique.end(), + conduit::execution::sort( + policy, ee_to_unique.begin(), ee_to_unique.end(), [&](const std::pair &lhs, const std::pair &rhs) { // Only sort using the ee value. @@ -2167,11 +2161,11 @@ TopologyMetadata::Implementation::build_edge_key_to_id( #ifdef DEBUG_PRINT std::cout << "edges_key_to_id = {" << std::endl; // Because of the printing. - using Exec = SerialExec; + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::serial(); #else - using Exec = ParallelExec; + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); #endif - conduit::execution::for_all(0, nedges, [&](index_t edge_index) + conduit::execution::forall(policy, 0, nedges, [&](index_t edge_index) { // Make a key for this edge. index_t edge[2]; @@ -2193,8 +2187,9 @@ TopologyMetadata::Implementation::build_edge_key_to_id( #endif // Sort the edges by the ids. - conduit::execution::sort( - edge_key_to_id.begin(), edge_key_to_id.end(), + conduit::execution::ExecutionPolicy sort_policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::sort( + sort_policy, edge_key_to_id.begin(), edge_key_to_id.end(), [&](const std::pair &lhs, const std::pair &rhs) { @@ -2288,7 +2283,8 @@ TopologyMetadata::Implementation::build_association_3_1_and_3_0_nonph() // Iterate over the elements, applying the edge template to make unique // edges for the element. We look up the edge in edge_key_to_id to get // its id. - conduit::execution::for_all(0, nelem, [&](index_t ei) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, nelem, [&](index_t ei) { index_t elem_offset = ei * points_per_elem; @@ -2487,7 +2483,8 @@ TopologyMetadata::Implementation::build_child_to_parent_association(int e, int a std::cout << "p2c=" << p2c << std::endl; #endif // Sort p2c by child. - conduit::execution::sort(p2c.begin(), p2c.end(), + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::sort(policy, p2c.begin(), p2c.end(), [&](const std::pair &lhs, const std::pair &rhs) { diff --git a/src/libs/blueprint/conduit_blueprint_mesh_utils.cpp b/src/libs/blueprint/conduit_blueprint_mesh_utils.cpp index 8050a948a..2b00c106b 100644 --- a/src/libs/blueprint/conduit_blueprint_mesh_utils.cpp +++ b/src/libs/blueprint/conduit_blueprint_mesh_utils.cpp @@ -4528,12 +4528,6 @@ PointQuery::acceleratedSearch(int ndims, int *result_ptr = &result[0]; conduit::index_t numCoordsetPts = coords[0]->dtype().number_of_elements(); -#if defined(CONDUIT_USE_OPENMP) - using policy = conduit::execution::OpenMPExec; -#else - using policy = conduit::execution::SerialExec; -#endif - // Special case a few large searches where the types are the same. if(ndims == 3 && sameTypes && @@ -4548,7 +4542,8 @@ PointQuery::acceleratedSearch(int ndims, conduit::blueprint::mesh::utils::kdtree search; search.initialize(typedCoords, numCoordsetPts); search.setPointTolerance(m_pointTolerance); - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { float64 searchPt[3] = {static_cast(input_ptr[i * 3 + 0]), static_cast(input_ptr[i * 3 + 1]), @@ -4571,7 +4566,8 @@ PointQuery::acceleratedSearch(int ndims, conduit::blueprint::mesh::utils::kdtree search; search.initialize(typedCoords, numCoordsetPts); search.setPointTolerance(static_cast(m_pointTolerance)); - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { float32 searchPt[3] = {static_cast(input_ptr[i * 3 + 0]), static_cast(input_ptr[i * 3 + 1]), @@ -4594,7 +4590,8 @@ PointQuery::acceleratedSearch(int ndims, conduit::blueprint::mesh::utils::kdtree search; search.initialize(typedCoords, numCoordsetPts); search.setPointTolerance(m_pointTolerance); - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { float64 searchPt[2] = {static_cast(input_ptr[i * 3 + 0]), static_cast(input_ptr[i * 3 + 1])}; @@ -4615,7 +4612,8 @@ PointQuery::acceleratedSearch(int ndims, conduit::blueprint::mesh::utils::kdtree search; search.initialize(typedCoords, numCoordsetPts); search.setPointTolerance(static_cast(m_pointTolerance)); - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { float32 searchPt[2] = {static_cast(input_ptr[i * 3 + 0]), static_cast(input_ptr[i * 3 + 1])}; @@ -4645,16 +4643,11 @@ PointQuery::normalSearch(int ndims, conduit::index_t numCoordsetPts = coords[0]->dtype().number_of_elements(); double EPS_SQ = m_pointTolerance * m_pointTolerance; -#if defined(CONDUIT_USE_OPENMP) - using policy = conduit::execution::OpenMPExec; -#else - using policy = conduit::execution::SerialExec; -#endif - // Back up to a brute force search if(ndims == 3) { - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { const double *searchPt = &input_ptr[i * 3]; int found = NotFound; @@ -4679,7 +4672,8 @@ PointQuery::normalSearch(int ndims, } else if(ndims == 2) { - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { const double *searchPt = &input_ptr[i * 3]; int found = NotFound; @@ -4702,7 +4696,8 @@ PointQuery::normalSearch(int ndims, } else if(ndims == 1) { - conduit::execution::for_all(0, numInputPts, [&](conduit::index_t i) + conduit::execution::ExecutionPolicy policy = conduit::execution::ExecutionPolicy::host(); + conduit::execution::forall(policy, 0, numInputPts, [&](conduit::index_t i) { const double *searchPt = &input_ptr[i * 3]; int found = NotFound; diff --git a/src/libs/conduit/CMakeLists.txt b/src/libs/conduit/CMakeLists.txt index afb0410bd..6ee480f8c 100644 --- a/src/libs/conduit/CMakeLists.txt +++ b/src/libs/conduit/CMakeLists.txt @@ -50,8 +50,10 @@ set(conduit_headers conduit_data_type.hpp conduit_endianness.hpp conduit_execution.hpp - conduit_execution_omp.hpp - conduit_execution_serial.hpp + conduit_memory_manager.hpp + conduit_data_array.hpp + conduit_data_accessor.hpp + conduit_data_type.hpp conduit_fixed_size_map.hpp conduit_fixed_size_vector.hpp conduit_geometry_vector.hpp @@ -89,6 +91,8 @@ set(conduit_sources conduit_core.cpp conduit_error.cpp conduit_endianness.cpp + conduit_execution.cpp + conduit_memory_manager.cpp conduit_data_type.cpp conduit_data_array.cpp conduit_data_accessor.cpp diff --git a/src/libs/conduit/conduit.hpp b/src/libs/conduit/conduit.hpp index 681008ac1..ad3dc0393 100644 --- a/src/libs/conduit/conduit.hpp +++ b/src/libs/conduit/conduit.hpp @@ -26,6 +26,8 @@ #include "conduit_generator.hpp" #include "conduit_utils.hpp" #include "conduit_data_accessor.hpp" +#include "conduit_execution.hpp" +#include "conduit_memory_manager.hpp" #endif diff --git a/src/libs/conduit/conduit_data_accessor.cpp b/src/libs/conduit/conduit_data_accessor.cpp index 24eec9923..0359f85b0 100644 --- a/src/libs/conduit/conduit_data_accessor.cpp +++ b/src/libs/conduit/conduit_data_accessor.cpp @@ -16,6 +16,13 @@ #include #include +//----------------------------------------------------------------------------- +// -- conduit includes -- +//----------------------------------------------------------------------------- +#include "conduit_memory_manager.hpp" +#include "conduit_node.hpp" +#include "conduit_data_array.hpp" + //----------------------------------------------------------------------------- // -- begin conduit:: -- //----------------------------------------------------------------------------- @@ -31,15 +38,27 @@ namespace conduit //---------------------------------------------------------------------------// template DataAccessor::DataAccessor() -: m_data(NULL), - m_dtype() +: m_data(nullptr), + m_dtype(DataType::empty()), + m_node_ptr(nullptr), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(0), + m_stride(0) {} //---------------------------------------------------------------------------// template DataAccessor::DataAccessor(const DataAccessor &accessor) : m_data(accessor.m_data), - m_dtype(accessor.m_dtype) + m_dtype(accessor.m_dtype), + m_node_ptr(accessor.m_node_ptr), + m_other_ptr(accessor.m_other_ptr), + m_other_dtype(accessor.m_other_dtype), + m_do_i_own_it(accessor.m_do_i_own_it), + m_offset(accessor.m_offset), + m_stride(accessor.m_stride) {} @@ -47,7 +66,13 @@ DataAccessor::DataAccessor(const DataAccessor &accessor) template DataAccessor::DataAccessor(void *data, const DataType &dtype) : m_data(data), - m_dtype(dtype) + m_dtype(dtype), + m_node_ptr(nullptr), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(0), + m_stride(0) {} @@ -55,13 +80,83 @@ DataAccessor::DataAccessor(void *data, const DataType &dtype) template DataAccessor::DataAccessor(const void *data, const DataType &dtype) : m_data(const_cast(data)), - m_dtype(dtype) + m_dtype(dtype), + m_node_ptr(nullptr), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(0), + m_stride(0) +{} + +//---------------------------------------------------------------------------// +template +DataAccessor::DataAccessor(Node &node) +: m_data(node.data_ptr()), + m_dtype(node.dtype()), + m_node_ptr(&node), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node.dtype().offset()), + m_stride(node.dtype().stride()) +{} + +//---------------------------------------------------------------------------// +template +DataAccessor::DataAccessor(const Node &node) +: m_data(const_cast(node.data_ptr())), + m_dtype(node.dtype()), + m_node_ptr(const_cast(&node)), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node.dtype().offset()), + m_stride(node.dtype().stride()) +{} + +//---------------------------------------------------------------------------// +template +DataAccessor::DataAccessor(Node *node) +: m_data(node->data_ptr()), + m_dtype(node->dtype()), + m_node_ptr(node), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node->dtype().offset()), + m_stride(node->dtype().stride()) +{} + +//---------------------------------------------------------------------------// +template +DataAccessor::DataAccessor(const Node *node) +: m_data(const_cast(node->data_ptr())), + m_dtype(node->dtype()), + m_node_ptr(const_cast(node)), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node->dtype().offset()), + m_stride(node->dtype().stride()) {} //---------------------------------------------------------------------------// template DataAccessor::~DataAccessor() -{} // all data is external +{ + if (m_do_i_own_it) + { + if (execution::DeviceMemory::is_device_ptr(m_other_ptr)) + { + execution::DeviceMemory::deallocate(m_other_ptr); + } + else + { + execution::HostMemory::deallocate(m_other_ptr); + } + } +} //---------------------------------------------------------------------------// @@ -163,6 +258,12 @@ DataAccessor::operator=(const DataAccessor &accessor) { m_data = accessor.m_data; m_dtype = accessor.m_dtype; + m_node_ptr = accessor.m_node_ptr; + m_other_ptr = accessor.m_other_ptr; + m_other_dtype = accessor.m_other_dtype; + m_do_i_own_it = accessor.m_do_i_own_it; + m_offset = accessor.m_offset; + m_stride = accessor.m_stride; } return *this; } @@ -172,7 +273,7 @@ template T DataAccessor::element(index_t idx) const { - switch(m_dtype.id()) + switch(dtype().id()) { // ints case DataType::INT8_ID: @@ -201,16 +302,17 @@ DataAccessor::element(index_t idx) const // error CONDUIT_ERROR("DataAccessor does not support dtype: " - << m_dtype.name()); + << dtype().name()); return (T)0; } //---------------------------------------------------------------------------// template -void +template +typename std::enable_if::value, void>::type DataAccessor::set(index_t idx, T value) { - switch(m_dtype.id()) + switch(dtype().id()) { // ints case DataType::INT8_ID: @@ -268,7 +370,105 @@ DataAccessor::set(index_t idx, T value) default: // error CONDUIT_ERROR("DataAccessor does not support dtype: " - << m_dtype.name()); + << dtype().name()); + } +} + +//---------------------------------------------------------------------------// +template +template +typename std::enable_if::value, void>::type +DataAccessor::set(const T* values, index_t num_elements) +{ + switch(dtype().id()) + { + // ints + case DataType::INT8_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::INT16_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::INT32_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::INT64_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + // uints + case DataType::UINT8_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::UINT16_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::UINT32_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::UINT64_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + // floats + case DataType::FLOAT32_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + case DataType::FLOAT64_ID: + { + for(index_t idx=0;idx(values[idx]); + } + break; + } + default: + // error + CONDUIT_ERROR("DataAccessor does not support dtype: " + << dtype().name()); } } @@ -277,13 +477,13 @@ template void DataAccessor::fill(T value) { - switch(m_dtype.id()) + switch(dtype().id()) { // ints case DataType::INT8_ID: { int8 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(int8*)(element_ptr(i))) = v; } @@ -292,7 +492,7 @@ DataAccessor::fill(T value) case DataType::INT16_ID: { int16 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(int16*)(element_ptr(i))) = v; } @@ -301,7 +501,7 @@ DataAccessor::fill(T value) case DataType::INT32_ID: { int32 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(int32*)(element_ptr(i))) = v; } @@ -310,7 +510,7 @@ DataAccessor::fill(T value) case DataType::INT64_ID: { int64 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(int64*)(element_ptr(i))) = v; } @@ -320,7 +520,7 @@ DataAccessor::fill(T value) case DataType::UINT8_ID: { uint8 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(uint8*)(element_ptr(i))) = v; } @@ -329,7 +529,7 @@ DataAccessor::fill(T value) case DataType::UINT16_ID: { uint16 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(uint16*)(element_ptr(i))) = v; } @@ -338,7 +538,7 @@ DataAccessor::fill(T value) case DataType::UINT32_ID: { uint32 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(uint32*)(element_ptr(i))) = v; } @@ -347,7 +547,7 @@ DataAccessor::fill(T value) case DataType::UINT64_ID: { uint64 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(uint64*)(element_ptr(i))) = v; } @@ -357,7 +557,7 @@ DataAccessor::fill(T value) case DataType::FLOAT32_ID: { float32 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(float32*)(element_ptr(i))) = v; } @@ -366,7 +566,7 @@ DataAccessor::fill(T value) case DataType::FLOAT64_ID: { float64 v = static_cast(value); - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { (*(float64*)(element_ptr(i))) = v; } @@ -375,7 +575,543 @@ DataAccessor::fill(T value) default: // error CONDUIT_ERROR("DataAccessor does not support dtype: " - << m_dtype.name()); + << dtype().name()); + } +} + +//---------------------------------------------------------------------------// +template +const DataType & +DataAccessor::dtype() const +{ + if (nullptr != m_node_ptr) + { + return (m_data == m_node_ptr->data_ptr() ? orig_dtype() : other_dtype()); + } + else + { + return m_dtype; + } +} + +//---------------------------------------------------------------------------// +template +const DataType & +DataAccessor::orig_dtype() const +{ + if (nullptr != m_node_ptr) + { + return m_node_ptr->dtype(); + } + else + { + return m_dtype; + } +} + +//---------------------------------------------------------------------------// +template +const DataType & +DataAccessor::other_dtype() const +{ + if (nullptr != m_node_ptr) + { + return m_other_dtype; + } + else + { + return m_dtype; + } +} + + +//---------------------------------------------------------------------------// +template +void +DataAccessor::use_with(conduit::execution::ExecutionPolicy policy) +{ + if (nullptr == m_node_ptr) + { + // TODO error; we can't do anything + return; + } + + // we are being asked to execute on the device + if (policy.is_device_policy()) + { + // data is already on the device + if (execution::DeviceMemory::is_device_ptr(m_data)) + { + // Do nothing + } + else // m_data is on the host + { + // if we started out on the host + if (m_node_ptr->data_ptr() == m_data) + { + CONDUIT_ASSERT(m_other_ptr == nullptr, + "Using execution accessor in this way will result in a memory leak."); + + // allocate new memory and create a new dtype + m_other_ptr = execution::DeviceMemory::allocate( + dtype().element_bytes() * number_of_elements()); + m_do_i_own_it = true; + m_other_dtype = DataType(dtype().id(), + number_of_elements(), + 0, // offset is 0 + DataType::default_bytes(dtype().id()), // stride + dtype().element_bytes(), + dtype().endianness()); + + // copy data + utils::conduit_memcpy_strided_elements(m_other_ptr, + number_of_elements(), + dtype().element_bytes(), + m_other_dtype.stride(), + m_data, + dtype().stride()); + + // change where our data pointer points and update offset and stride + m_data = m_other_ptr; + m_offset = m_other_dtype.offset(); + m_stride = m_other_dtype.stride(); + } + else // we started out on the device + { + CONDUIT_ASSERT(m_data == m_other_ptr, + "Using execution accessor in this way will result in a memory leak."); + + // call sync to bring our copy of the data on the host back to the device + sync(); + + // dealloc the ptr on the host now that we have copied back + execution::HostMemory::deallocate(m_data); + m_do_i_own_it = false; + m_other_dtype = DataType::empty(); + + // set m_data to device data and update offset and stride + m_data = m_node_ptr->data_ptr(); + // the order of operations is important here; changing the pointer + // will change the result of calling dtype(). + m_offset = dtype().offset(); + m_stride = dtype().stride(); + + // reset m_other_ptr + m_other_ptr = nullptr; + } + } + } + else // we are being asked to execute on the host + { + // data is already on the host + if (! execution::DeviceMemory::is_device_ptr(m_data)) + { + // Do nothing + } + else // m_data is on the device + { + // if we started out on the device + if (m_node_ptr->data_ptr() == m_data) + { + CONDUIT_ASSERT(m_other_ptr == nullptr, + "Using execution accessor in this way will result in a memory leak."); + + // allocate new memory and create a new dtype + m_other_ptr = execution::HostMemory::allocate( + dtype().element_bytes() * number_of_elements()); + m_do_i_own_it = true; + m_other_dtype = DataType(dtype().id(), + number_of_elements(), + 0, // offset is 0 + DataType::default_bytes(dtype().id()), // stride + dtype().element_bytes(), + dtype().endianness()); + + // copy data + utils::conduit_memcpy_strided_elements(m_other_ptr, + number_of_elements(), + dtype().element_bytes(), + m_other_dtype.stride(), + m_data, + dtype().stride()); + + // change where our data pointer points and update offset and stride + m_data = m_other_ptr; + m_offset = m_other_dtype.offset(); + m_stride = m_other_dtype.stride(); + } + else // we started out on the host + { + CONDUIT_ASSERT(m_data == m_other_ptr, + "Using execution accessor in this way will result in a memory leak."); + + // call sync to bring our copy of the data on the device back to the host + sync(); + + // dealloc the ptr on the host now that we have copied back + execution::DeviceMemory::deallocate(m_data); + m_do_i_own_it = false; + m_other_dtype = DataType::empty(); + + // set m_data to host data and update offset and stride + m_data = m_node_ptr->data_ptr(); + m_offset = dtype().offset(); + m_stride = dtype().stride(); + + // reset m_other_ptr + m_other_ptr = nullptr; + } + } + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::sync() +{ + if (nullptr == m_node_ptr) + { + // TODO error; we can't do anything + return; + } + + // if the ptrs don't point to the same place + if (m_data != m_node_ptr->data_ptr()) + { + if (!(m_node_ptr->dtype().compatible(dtype()) && + number_of_elements() == m_node_ptr->dtype().number_of_elements())) + { + m_node_ptr->set(dtype()); + } + utils::conduit_memcpy_strided_elements(m_node_ptr->data_ptr(), + number_of_elements(), + m_node_ptr->dtype().element_bytes(), + m_node_ptr->dtype().stride(), + m_data, + m_stride); + } +} + + +//---------------------------------------------------------------------------// +template +void +DataAccessor::assume() +{ + if (nullptr == m_node_ptr) + { + // TODO error; we can't do anything + return; + } + + // if the ptrs don't point to the same place + if (m_data != m_node_ptr->data_ptr()) + { + CONDUIT_ASSERT(m_data == m_other_ptr, + "Using execution accessor in this way will result in a memory leak."); + + // reset will deallocate the data the node points to + m_node_ptr->reset(); + m_node_ptr->schema_ptr()->set(dtype()); + m_node_ptr->set_data_ptr(m_data); + + // we no longer own the data since we have given it to node + m_other_ptr = nullptr; + m_do_i_own_it = false; + m_other_dtype = DataType::empty(); + } +} + + +//---------------------------------------------------------------------------// +template +conduit::execution::ExecutionPolicy +DataAccessor::active_space() +{ + if (execution::DeviceMemory::is_device_ptr(m_data)) + { + return execution::ExecutionPolicy::device(); + } + else + { + return execution::ExecutionPolicy::host(); + } +} + +//---------------------------------------------------------------------------// +//***************************************************************************// +// Set from DataAccessor +//***************************************************************************// +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +// Set from DataAccessor signed integers +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +// Set from DataAccessor unsigned integers +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +// Set from DataAccessor floating point +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataAccessor &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +//***************************************************************************// +// Set from DataArray +//***************************************************************************// +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +// Set from DataArray signed integers +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +// Set from DataArray unsigned integers +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +// Set from DataArray floating point +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); + } +} + +//---------------------------------------------------------------------------// +template +void +DataAccessor::set(const DataArray &values) +{ + index_t num_elems = dtype().number_of_elements(); + for(index_t i=0; i set(i, (T)values[i]); } } diff --git a/src/libs/conduit/conduit_data_accessor.hpp b/src/libs/conduit/conduit_data_accessor.hpp index 91ab4db0a..4954a2b86 100644 --- a/src/libs/conduit/conduit_data_accessor.hpp +++ b/src/libs/conduit/conduit_data_accessor.hpp @@ -18,6 +18,7 @@ #include "conduit_core.hpp" #include "conduit_data_type.hpp" #include "conduit_utils.hpp" +#include "conduit_execution.hpp" //----------------------------------------------------------------------------- @@ -26,6 +27,13 @@ namespace conduit { +//----------------------------------------------------------------------------- +// -- forward declarations required for conduit::DataAccessor -- +//----------------------------------------------------------------------------- +class Node; +template +class DataArray; + //----------------------------------------------------------------------------- // -- begin conduit::DataArray -- //----------------------------------------------------------------------------- @@ -33,7 +41,8 @@ namespace conduit /// class: conduit::DataAccessor /// /// description: -/// Helps consume array data as desired type with on the fly conversion. +/// Helps consume array data as desired type with on the fly conversion and +/// supports memory movement between host and device. /// //----------------------------------------------------------------------------- template @@ -58,6 +67,15 @@ class CONDUIT_API DataAccessor DataAccessor(void *data, const DataType &dtype); /// Access a const pointer to raw data according to dtype description. DataAccessor(const void *data, const DataType &dtype); + /// Access a pointer to node data according to node dtype description. + DataAccessor(Node &node); + // /// Access a const pointer to node data according to node dtype description. + DataAccessor(const Node &node); + /// Access a pointer to node data according to node dtype description. + DataAccessor(Node *node); + /// Access a const pointer to node data according to node dtype description. + DataAccessor(const Node *node); + /// Destructor. ~DataAccessor(); /// @@ -82,23 +100,83 @@ class CONDUIT_API DataAccessor T element(index_t idx) const; - void set(index_t idx, T value); + // Without the SFINAE features, the compiler doesn't know which of the two + // set methods to call. We need to restrict them based on if the type is a + // pointer or not so that it is unambiguous which method should be called. + template + typename std::enable_if::value, void>::type + set(index_t idx, T value); + + template + typename std::enable_if::value, void>::type + set(const T* values, index_t num_elements); + + // void set(const std::vector::type> &values) + // { set(values.data(), values.size()); } + void fill(T value); const void *element_ptr(index_t idx) const { return static_cast(m_data) + - m_dtype.element_index(idx); + dtype().element_index(idx); } index_t number_of_elements() const - {return m_dtype.number_of_elements();} + {return dtype().number_of_elements();} + + const DataType &dtype() const; + + const DataType &orig_dtype() const; + + const DataType &other_dtype() const; + +//----------------------------------------------------------------------------- +// Data movement +//----------------------------------------------------------------------------- + void use_with(conduit::execution::ExecutionPolicy policy); + + void sync(); - const DataType &dtype() const - { return m_dtype;} + void assume(); + conduit::execution::ExecutionPolicy active_space(); +//----------------------------------------------------------------------------- +// Setters +//----------------------------------------------------------------------------- + /// signed integer arrays via DataArray + void set(const DataArray &values); + void set(const DataArray &values); + void set(const DataArray &values); + void set(const DataArray &values); + + /// unsigned integer arrays via DataArray + void set(const DataArray &values); + void set(const DataArray &values); + void set(const DataArray &values); + void set(const DataArray &values); + + /// floating point arrays via DataArray + void set(const DataArray &values); + void set(const DataArray &values); + + /// signed integer arrays via DataAccessor + void set(const DataAccessor &values); + void set(const DataAccessor &values); + void set(const DataAccessor &values); + void set(const DataAccessor &values); + + /// unsigned integer arrays via DataAccessor + void set(const DataAccessor &values); + void set(const DataAccessor &values); + void set(const DataAccessor &values); + void set(const DataAccessor &values); + + /// floating point arrays via DataAccessor + void set(const DataAccessor &values); + void set(const DataAccessor &values); //----------------------------------------------------------------------------- // Transforms @@ -143,6 +221,18 @@ class CONDUIT_API DataAccessor void *m_data; /// holds data description DataType m_dtype; + + Node *m_node_ptr; + + /// holds data + void *m_other_ptr; + /// holds data description + DataType m_other_dtype; + + bool m_do_i_own_it; + + index_t m_offset; + index_t m_stride; }; //----------------------------------------------------------------------------- diff --git a/src/libs/conduit/conduit_data_array.cpp b/src/libs/conduit/conduit_data_array.cpp index df97f8ca8..9e8533e73 100644 --- a/src/libs/conduit/conduit_data_array.cpp +++ b/src/libs/conduit/conduit_data_array.cpp @@ -23,6 +23,7 @@ #include "conduit_node.hpp" #include "conduit_utils.hpp" #include "conduit_log.hpp" +#include "conduit_data_accessor.hpp" // Easier access to the Conduit logging functions using namespace conduit::utils; @@ -43,36 +44,125 @@ namespace conduit //---------------------------------------------------------------------------// template DataArray::DataArray() -: m_data(NULL), - m_dtype(DataType::empty()) -{} +: m_data(nullptr), + m_dtype(DataType::empty()), + m_node_ptr(nullptr), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(0), + m_stride(0) +{} + +//---------------------------------------------------------------------------// +template +DataArray::DataArray(const DataArray &array) +: m_data(array.m_data), + m_dtype(array.m_dtype), + m_node_ptr(array.m_node_ptr), + m_other_ptr(array.m_other_ptr), + m_other_dtype(array.m_other_dtype), + m_do_i_own_it(array.m_do_i_own_it), + m_offset(array.m_offset), + m_stride(array.m_stride) +{} + //---------------------------------------------------------------------------// template -DataArray::DataArray(void *data,const DataType &dtype) +DataArray::DataArray(void *data, const DataType &dtype) : m_data(data), - m_dtype(dtype) + m_dtype(dtype), + m_node_ptr(nullptr), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(0), + m_stride(0) {} + //---------------------------------------------------------------------------// template -DataArray::DataArray(const void *data,const DataType &dtype) +DataArray::DataArray(const void *data, const DataType &dtype) : m_data(const_cast(data)), - m_dtype(dtype) + m_dtype(dtype), + m_node_ptr(nullptr), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(0), + m_stride(0) {} +//---------------------------------------------------------------------------// +template +DataArray::DataArray(Node &node) +: m_data(node.data_ptr()), + m_dtype(node.dtype()), + m_node_ptr(&node), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node.dtype().offset()), + m_stride(node.dtype().stride()) +{} -//---------------------------------------------------------------------------// +//---------------------------------------------------------------------------// template -DataArray::DataArray(const DataArray &array) -: m_data(array.m_data), - m_dtype(array.m_dtype) +DataArray::DataArray(const Node &node) +: m_data(const_cast(node.data_ptr())), + m_dtype(node.dtype()), + m_node_ptr(const_cast(&node)), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node.dtype().offset()), + m_stride(node.dtype().stride()) +{} + +//---------------------------------------------------------------------------// +template +DataArray::DataArray(Node *node) +: m_data(node->data_ptr()), + m_dtype(node->dtype()), + m_node_ptr(node), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node->dtype().offset()), + m_stride(node->dtype().stride()) +{} + +//---------------------------------------------------------------------------// +template +DataArray::DataArray(const Node *node) +: m_data(const_cast(node->data_ptr())), + m_dtype(node->dtype()), + m_node_ptr(const_cast(node)), + m_other_ptr(nullptr), + m_other_dtype(DataType::empty()), + m_do_i_own_it(false), + m_offset(node->dtype().offset()), + m_stride(node->dtype().stride()) {} //---------------------------------------------------------------------------// template DataArray::~DataArray() -{} // all data is external +{ + if (m_do_i_own_it) + { + if (execution::DeviceMemory::is_device_ptr(m_other_ptr)) + { + execution::DeviceMemory::deallocate(m_other_ptr); + } + else + { + execution::HostMemory::deallocate(m_other_ptr); + } + } +} //---------------------------------------------------------------------------// template @@ -81,8 +171,14 @@ DataArray::operator=(const DataArray &array) { if(this != &array) { - m_data = array.m_data; + m_node_ptr = array.m_node_ptr; + m_other_ptr = array.m_other_ptr; + m_other_dtype = array.m_other_dtype; + m_do_i_own_it = array.m_do_i_own_it; + m_data = array.m_data; m_dtype = array.m_dtype; + m_offset = array.m_offset; + m_stride = array.m_stride; } return *this; } @@ -103,6 +199,51 @@ DataArray::element(index_t idx) const return (*(T*)(element_ptr(idx))); } +//---------------------------------------------------------------------------// +template +const DataType & +DataArray::dtype() const +{ + if (nullptr != m_node_ptr) + { + return (m_data == m_node_ptr->data_ptr() ? orig_dtype() : other_dtype()); + } + else + { + return m_dtype; + } +} + +//---------------------------------------------------------------------------// +template +const DataType & +DataArray::orig_dtype() const +{ + if (nullptr != m_node_ptr) + { + return m_node_ptr->dtype(); + } + else + { + return m_dtype; + } +} + +//---------------------------------------------------------------------------// +template +const DataType & +DataArray::other_dtype() const +{ + if (nullptr != m_node_ptr) + { + return m_other_dtype; + } + else + { + return m_dtype; + } +} + //---------------------------------------------------------------------------// template bool @@ -531,6 +672,220 @@ DataArray::count(T val) const return res; } +//---------------------------------------------------------------------------// +template +void +DataArray::use_with(conduit::execution::ExecutionPolicy policy) +{ + if (nullptr == m_node_ptr) + { + // TODO error; we can't do anything + return; + } + + // we are being asked to execute on the device + if (policy.is_device_policy()) + { + // data is already on the device + if (execution::DeviceMemory::is_device_ptr(m_data)) + { + // Do nothing + } + else // m_data is on the host + { + // if we started out on the host + if (m_node_ptr->data_ptr() == m_data) + { + CONDUIT_ASSERT(m_other_ptr == nullptr, + "Using execution array in this way will result in a memory leak."); + + // allocate new memory and create a new dtype + m_other_ptr = execution::DeviceMemory::allocate( + dtype().element_bytes() * number_of_elements()); + m_do_i_own_it = true; + m_other_dtype = DataType(dtype().id(), + number_of_elements(), + 0, // offset is 0 + DataType::default_bytes(dtype().id()), // stride + dtype().element_bytes(), + dtype().endianness()); + + // copy data + utils::conduit_memcpy_strided_elements(m_other_ptr, + number_of_elements(), + dtype().element_bytes(), + m_other_dtype.stride(), + m_data, + dtype().stride()); + + // change where our data pointer points and update offset and stride + m_data = m_other_ptr; + m_offset = m_other_dtype.offset(); + m_stride = m_other_dtype.stride(); + } + else // we started out on the device + { + CONDUIT_ASSERT(m_data == m_other_ptr, + "Using execution array in this way will result in a memory leak."); + + // call sync to bring our copy of the data on the host back to the device + sync(); + + // dealloc the ptr on the host now that we have copied back + execution::HostMemory::deallocate(m_data); + m_do_i_own_it = false; + m_other_dtype = DataType::empty(); + + // set m_data to device data and update offset and stride + m_data = m_node_ptr->data_ptr(); + // the order of operations is important here; changing the pointer + // will change the result of calling dtype(). + m_offset = dtype().offset(); + m_stride = dtype().stride(); + + // reset m_other_ptr + m_other_ptr = nullptr; + } + } + } + else // we are being asked to execute on the host + { + // data is already on the host + if (! execution::DeviceMemory::is_device_ptr(m_data)) + { + // Do nothing + } + else // m_data is on the device + { + // if we started out on the device + if (m_node_ptr->data_ptr() == m_data) + { + CONDUIT_ASSERT(m_other_ptr == nullptr, + "Using execution array in this way will result in a memory leak."); + + // allocate new memory and create a new dtype + m_other_ptr = execution::HostMemory::allocate( + dtype().element_bytes() * number_of_elements()); + m_do_i_own_it = true; + m_other_dtype = DataType(dtype().id(), + number_of_elements(), + 0, // offset is 0 + DataType::default_bytes(dtype().id()), // stride + dtype().element_bytes(), + dtype().endianness()); + + // copy data + utils::conduit_memcpy_strided_elements(m_other_ptr, + number_of_elements(), + dtype().element_bytes(), + m_other_dtype.stride(), + m_data, + dtype().stride()); + + // change where our data pointer points and update offset and stride + m_data = m_other_ptr; + m_offset = m_other_dtype.offset(); + m_stride = m_other_dtype.stride(); + } + else // we started out on the host + { + CONDUIT_ASSERT(m_data == m_other_ptr, + "Using execution array in this way will result in a memory leak."); + + // call sync to bring our copy of the data on the device back to the host + sync(); + + // dealloc the ptr on the host now that we have copied back + execution::DeviceMemory::deallocate(m_data); + m_do_i_own_it = false; + m_other_dtype = DataType::empty(); + + // set m_data to host data and update offset and stride + m_data = m_node_ptr->data_ptr(); + m_offset = dtype().offset(); + m_stride = dtype().stride(); + + // reset m_other_ptr + m_other_ptr = nullptr; + } + } + } +} + +//---------------------------------------------------------------------------// +template +void +DataArray::sync() +{ + if (nullptr == m_node_ptr) + { + // TODO error; we can't do anything + return; + } + + // if the ptrs don't point to the same place + if (m_data != m_node_ptr->data_ptr()) + { + if (!(m_node_ptr->dtype().compatible(dtype()) && + number_of_elements() == m_node_ptr->dtype().number_of_elements())) + { + m_node_ptr->set(dtype()); + } + utils::conduit_memcpy_strided_elements(m_node_ptr->data_ptr(), + number_of_elements(), + m_node_ptr->dtype().element_bytes(), + m_node_ptr->dtype().stride(), + m_data, + m_stride); + } +} + + +//---------------------------------------------------------------------------// +template +void +DataArray::assume() +{ + if (nullptr == m_node_ptr) + { + // TODO error; we can't do anything + return; + } + + // if the ptrs don't point to the same place + if (m_data != m_node_ptr->data_ptr()) + { + CONDUIT_ASSERT(m_data == m_other_ptr, + "Using execution array in this way will result in a memory leak."); + + // reset will deallocate the data the node points to + m_node_ptr->reset(); + m_node_ptr->schema_ptr()->set(dtype()); + m_node_ptr->set_data_ptr(m_data); + + // we no longer own the data since we have given it to node + m_other_ptr = nullptr; + m_do_i_own_it = false; + m_other_dtype = DataType::empty(); + } +} + + +//---------------------------------------------------------------------------// +template +conduit::execution::ExecutionPolicy +DataArray::active_space() +{ + if (execution::DeviceMemory::is_device_ptr(m_data)) + { + return execution::ExecutionPolicy::device(); + } + else + { + return execution::ExecutionPolicy::host(); + } +} + //---------------------------------------------------------------------------// template std::string @@ -608,7 +963,7 @@ DataArray::to_json_stream(std::ostream &os) const { if(!first) os << ", "; - switch(m_dtype.id()) + switch(dtype().id()) { // ints case DataType::INT8_ID: @@ -649,7 +1004,7 @@ DataArray::to_json_stream(std::ostream &os) const default: { CONDUIT_ERROR("Leaf type \"" - << m_dtype.name() + << dtype().name() << "\"" << "is not supported in conduit::DataArray.") } @@ -894,7 +1249,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -911,7 +1266,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -928,7 +1283,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -945,7 +1300,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -962,7 +1317,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -979,7 +1334,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -996,7 +1351,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1013,7 +1368,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1030,7 +1385,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1047,7 +1402,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1069,7 +1424,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1090,7 +1445,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1107,7 +1462,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1132,7 +1487,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1149,7 +1504,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1174,7 +1529,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1191,7 +1546,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1216,7 +1571,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1233,7 +1588,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1259,7 +1614,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1276,7 +1631,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1299,7 +1654,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1322,7 +1677,7 @@ void DataArray::set(const std::initializer_list &values) { index_t idx = 0; - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); // iterate and set up to the number of elements of this array std::initializer_list::const_iterator itr; for( itr = values.begin(); @@ -1352,7 +1707,7 @@ template void DataArray::fill(int8 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1363,7 +1718,7 @@ template void DataArray::fill(int16 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1374,7 +1729,7 @@ template void DataArray::fill(int32 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1385,7 +1740,7 @@ template void DataArray::fill(int64 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1400,7 +1755,7 @@ template void DataArray::fill(uint8 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1411,7 +1766,7 @@ template void DataArray::fill(uint16 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1422,7 +1777,7 @@ template void DataArray::fill(uint32 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1433,7 +1788,7 @@ template void DataArray::fill(uint64 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1448,7 +1803,7 @@ template void DataArray::fill(float32 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1459,7 +1814,7 @@ template void DataArray::fill(float64 value) { - for(index_t i=0;i < m_dtype.number_of_elements(); i++) + for(index_t i=0;i < dtype().number_of_elements(); i++) { this->element(i) = (T)value; } @@ -1763,7 +2118,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1775,7 +2130,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1787,7 +2142,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1799,7 +2154,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1815,7 +2170,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1827,7 +2182,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1839,7 +2194,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1851,7 +2206,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1867,7 +2222,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1879,7 +2234,7 @@ template void DataArray::set(const DataArray &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1901,7 +2256,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1913,7 +2268,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1925,7 +2280,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1937,7 +2292,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1953,7 +2308,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1965,7 +2320,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1977,7 +2332,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -1989,7 +2344,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -2005,7 +2360,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -2017,7 +2372,7 @@ template void DataArray::set(const DataAccessor &values) { - index_t num_elems = m_dtype.number_of_elements(); + index_t num_elems = dtype().number_of_elements(); for(index_t i=0; i element(i) = (T)values[i]; @@ -2102,7 +2457,7 @@ DataArray::to_summary_string_stream(std::ostream &os, if(idx > 0 ) os << ", "; - switch(m_dtype.id()) + switch(dtype().id()) { // ints case DataType::INT8_ID: @@ -2143,7 +2498,7 @@ DataArray::to_summary_string_stream(std::ostream &os, default: { CONDUIT_ERROR("Leaf type \"" - << m_dtype.name() + << dtype().name() << "\"" << "is not supported in conduit::DataArray.") } diff --git a/src/libs/conduit/conduit_data_array.hpp b/src/libs/conduit/conduit_data_array.hpp index a21466fdd..cf9b035fa 100644 --- a/src/libs/conduit/conduit_data_array.hpp +++ b/src/libs/conduit/conduit_data_array.hpp @@ -19,7 +19,7 @@ #include "conduit_core.hpp" #include "conduit_data_type.hpp" #include "conduit_utils.hpp" -#include "conduit_data_accessor.hpp" +#include "conduit_execution.hpp" //----------------------------------------------------------------------------- // -- begin conduit:: -- @@ -27,6 +27,13 @@ namespace conduit { +//----------------------------------------------------------------------------- +// -- forward declarations required for conduit::DataArray -- +//----------------------------------------------------------------------------- +class Node; +template +class DataAccessor; + //----------------------------------------------------------------------------- // -- begin conduit::DataArray -- //----------------------------------------------------------------------------- @@ -34,7 +41,9 @@ namespace conduit /// class: conduit::DataArray /// /// description: -/// Light weight pointer wrapper that handles addressing for ragged arrays. +/// Light weight pointer wrapper that handles addressing for ragged arrays +/// that may be stored in Nodes; also supports memory movement between host +/// and device. /// //----------------------------------------------------------------------------- template @@ -58,6 +67,14 @@ class CONDUIT_API DataArray DataArray(void *data, const DataType &dtype); /// Access a const pointer to raw data according to dtype description. DataArray(const void *data, const DataType &dtype); + /// Access a pointer to node data according to node dtype description. + DataArray(Node &node); + // /// Access a const pointer to node data according to node dtype description. + DataArray(const Node &node); + /// Access a pointer to node data according to node dtype description. + DataArray(Node *node); + /// Access a const pointer to node data according to node dtype description. + DataArray(const Node *node); /// Destructor ~DataArray(); @@ -80,19 +97,23 @@ class CONDUIT_API DataArray void *element_ptr(index_t idx) { return static_cast(m_data) + - m_dtype.element_index(idx); + dtype().element_index(idx); }; const void *element_ptr(index_t idx) const { return static_cast(m_data) + - m_dtype.element_index(idx); + dtype().element_index(idx); }; index_t number_of_elements() const - {return m_dtype.number_of_elements();} - const DataType &dtype() const - { return m_dtype;} + {return dtype().number_of_elements();} + const DataType &dtype() const; + + const DataType &orig_dtype() const; + + const DataType &other_dtype() const; + void *data_ptr() const { return m_data;} @@ -115,6 +136,17 @@ class CONDUIT_API DataArray /// counts number of occurrences of given value index_t count(T value) const; +//----------------------------------------------------------------------------- +// Data movement +//----------------------------------------------------------------------------- + void use_with(conduit::execution::ExecutionPolicy policy); + + void sync(); + + void assume(); + + conduit::execution::ExecutionPolicy active_space(); + //----------------------------------------------------------------------------- // Setters //----------------------------------------------------------------------------- @@ -394,6 +426,18 @@ class CONDUIT_API DataArray void *m_data; /// holds data description DataType m_dtype; + + Node *m_node_ptr; + + /// holds data + void *m_other_ptr; + /// holds data description + DataType m_other_dtype; + + bool m_do_i_own_it; + + index_t m_offset; + index_t m_stride; }; //----------------------------------------------------------------------------- diff --git a/src/libs/conduit/conduit_execution.cpp b/src/libs/conduit/conduit_execution.cpp new file mode 100644 index 000000000..8594c2e25 --- /dev/null +++ b/src/libs/conduit/conduit_execution.cpp @@ -0,0 +1,329 @@ +// Copyright (c) Lawrence Livermore National Security, LLC and other Conduit +// Project developers. See top-level LICENSE AND COPYRIGHT files for dates and +// other details. No copyright assignment is required to contribute to Conduit. + +//----------------------------------------------------------------------------- +/// +/// file: conduit_execution.cpp +/// +//----------------------------------------------------------------------------- +#include "conduit_execution.hpp" + +//----------------------------------------------------------------------------- +// conduit includes +//----------------------------------------------------------------------------- + + +//----------------------------------------------------------------------------- +// -- begin conduit -- +//----------------------------------------------------------------------------- +namespace conduit +{ + +//----------------------------------------------------------------------------- +// -- begin conduit::execution -- +//----------------------------------------------------------------------------- +namespace execution +{ + +//----------------------------------------------------------------------------- +/// policy constructor helpers +//----------------------------------------------------------------------------- +ExecutionPolicy +ExecutionPolicy::empty() +{ + return ExecutionPolicy(PolicyID::EMPTY_ID); +} + +//---------------------------------------------------------------------------// +ExecutionPolicy +ExecutionPolicy::host() +{ +#if defined(CONDUIT_USE_OPENMP) + return ExecutionPolicy(PolicyID::OPENMP_ID); +#else + return ExecutionPolicy(PolicyID::SERIAL_ID); +#endif +} + +//---------------------------------------------------------------------------// +ExecutionPolicy +ExecutionPolicy::serial() +{ + return ExecutionPolicy(PolicyID::SERIAL_ID); +} + +//---------------------------------------------------------------------------// +ExecutionPolicy +ExecutionPolicy::device() +{ +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_CUDA) + return ExecutionPolicy(PolicyID::CUDA_ID); +#elif defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_HIP) + return ExecutionPolicy(PolicyID::HIP_ID); +#else + CONDUIT_ERROR("Conduit was built with neither CUDA nor HIP."); + return ExecutionPolicy(PolicyID::EMPTY_ID); +#endif +} + +//---------------------------------------------------------------------------// +ExecutionPolicy +ExecutionPolicy::cuda() +{ + return ExecutionPolicy(PolicyID::CUDA_ID); +} + +//---------------------------------------------------------------------------// +ExecutionPolicy +ExecutionPolicy::hip() +{ + return ExecutionPolicy(PolicyID::HIP_ID); +} + +//---------------------------------------------------------------------------// +ExecutionPolicy +ExecutionPolicy::openmp() +{ + return ExecutionPolicy(PolicyID::OPENMP_ID); +} + +//============================================================================= +//----------------------------------------------------------------------------- +// +// +// -- begin conduit::execution::ExecutionPolicy public methods -- +// +// +//----------------------------------------------------------------------------- +//============================================================================= + +//----------------------------------------------------------------------------- +// Construction and Destruction +//----------------------------------------------------------------------------- +//---------------------------------------------------------------------------// +ExecutionPolicy::ExecutionPolicy() +: m_policy_id(PolicyID::EMPTY_ID) +{} + +//---------------------------------------------------------------------------// +ExecutionPolicy::ExecutionPolicy(const ExecutionPolicy& exec_policy) +: m_policy_id(exec_policy.m_policy_id) +{} + +//---------------------------------------------------------------------------// +ExecutionPolicy& ExecutionPolicy::operator=(const ExecutionPolicy& exec_policy) +{ + m_policy_id = exec_policy.m_policy_id; + + return *this; +} + +//---------------------------------------------------------------------------// +ExecutionPolicy::ExecutionPolicy(PolicyID policy_id) +: m_policy_id(policy_id) +{} + +//---------------------------------------------------------------------------// +ExecutionPolicy::ExecutionPolicy(const std::string &policy_name) +: m_policy_id(name_to_policy_id(policy_name)) +{} + +//---------------------------------------------------------------------------// +ExecutionPolicy::~ExecutionPolicy() +{} + +//----------------------------------------------------------------------------- +// Getters and info methods. +//----------------------------------------------------------------------------- + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_empty() const +{ + return m_policy_id == PolicyID::EMPTY_ID; +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_serial() const +{ + return m_policy_id == PolicyID::SERIAL_ID; +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_cuda() const +{ + return m_policy_id == PolicyID::CUDA_ID; +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_hip() const +{ + return m_policy_id == PolicyID::HIP_ID; +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_openmp() const +{ + return m_policy_id == PolicyID::OPENMP_ID; +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_host_policy() const +{ + return is_serial() || is_openmp(); +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_device_policy() const +{ + return is_cuda() || is_hip(); +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_serial_enabled() +{ + return true; +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_cuda_enabled() +{ +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_CUDA) + return true; +#else + return false; +#endif +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_hip_enabled() +{ +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_HIP) + return true; +#else + return false; +#endif +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_openmp_enabled() +{ +#if defined(CONDUIT_USE_OPENMP) + return true; +#else + return false; +#endif +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_host_enabled() +{ + return is_serial_enabled() || is_openmp_enabled(); +} + +//---------------------------------------------------------------------------// +bool +ExecutionPolicy::is_device_enabled() +{ + return is_cuda_enabled() || is_hip_enabled(); +} + +//----------------------------------------------------------------------------- +// PolicyID to string and string to PolicyID +//----------------------------------------------------------------------------- + +//---------------------------------------------------------------------------// +ExecutionPolicy::PolicyID +ExecutionPolicy::name_to_policy_id(const std::string &policy_name) +{ + if (policy_name == "empty") return PolicyID::EMPTY_ID; + else if (policy_name == "serial") return PolicyID::SERIAL_ID; + else if (policy_name == "cuda") return PolicyID::CUDA_ID; + else if (policy_name == "hip") return PolicyID::HIP_ID; + else if (policy_name == "openmp") return PolicyID::OPENMP_ID; + return PolicyID::EMPTY_ID; +} + +//---------------------------------------------------------------------------// +std::string +ExecutionPolicy::policy_id_to_name(const PolicyID policy_id) +{ + if (policy_id == PolicyID::EMPTY_ID) return "empty"; + else if (policy_id == PolicyID::SERIAL_ID) return "serial"; + else if (policy_id == PolicyID::CUDA_ID) return "cuda"; + else if (policy_id == PolicyID::HIP_ID) return "hip"; + else if (policy_id == PolicyID::OPENMP_ID) return "openmp"; + return "empty"; +} + +//---------------------------------------------------------------------------// +void +init_device_memory_handlers() +{ +#if defined(CONDUIT_DEVICE_ENABLED) + // we only need to override the mem handlers in the + // presence of cuda or hip + conduit::utils::set_memcpy_handler(MagicMemory::copy); + conduit::utils::set_memset_handler(MagicMemory::set); +#endif +} + +//---------------------------------------------------------------------------// +void +device_error_check(ExecutionPolicy policy, const char *file, const int line) +{ + if (policy.is_hip()) + { +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_HIP) + hipError_t err = hipGetLastError(); + if ( hipSuccess != err ) + { + std::cerr<<"HIP error reported at: "< +#endif + +//----------------------------------------------------------------------------- +// cpp lib includes +//----------------------------------------------------------------------------- +#include + +//----------------------------------------------------------------------------- +// +/// The CONDUIT_DEVICE_ERROR_CHECK macro is TODO +// +//----------------------------------------------------------------------------- +#define CONDUIT_DEVICE_ERROR_CHECK( policy ) conduit::execution::device_error_check(policy, __FILE__, __LINE__); + +#if defined(CONDUIT_USE_RAJA) && (defined(CONDUIT_USE_CUDA) || defined(CONDUIT_USE_HIP)) +#define EXEC_LAMBDA __device__ __host__ +#else +#define EXEC_LAMBDA #endif //----------------------------------------------------------------------------- @@ -33,50 +53,540 @@ namespace conduit namespace execution { -// NOTE: These are a start. +//----------------------------------------------------------------------------- +// -- begin conduit::execution::ExecutionPolicy -- +//----------------------------------------------------------------------------- +/// +/// class: conduit::execution::ExecutionPolicy +/// +/// description: +/// ExecutionPolicy is a runtime policy object. +/// +//----------------------------------------------------------------------------- +class CONDUIT_API ExecutionPolicy +{ +public: +//----------------------------------------------------------------------------- +/// Policy is an Enumeration used to describe the policy cases supported +/// by conduit: +//----------------------------------------------------------------------------- + enum class PolicyID : conduit::index_t + { + EMPTY_ID, + SERIAL_ID, + CUDA_ID, + HIP_ID, + OPENMP_ID + }; -//--------------------------------------------------------------------------- +//----------------------------------------------------------------------------- +// -- begin conduit::execution::ExecutionPolicy Constructor Helpers -- +//----------------------------------------------------------------------------- + static ExecutionPolicy empty(); + static ExecutionPolicy host(); // prefer openmp to serial + static ExecutionPolicy serial(); + static ExecutionPolicy device(); // prefer cuda to hip + static ExecutionPolicy cuda(); + static ExecutionPolicy hip(); + static ExecutionPolicy openmp(); + +//----------------------------------------------------------------------------- +// -- end conduit::execution::ExecutionPolicy Constructor Helpers -- +//----------------------------------------------------------------------------- + +//----------------------------------------------------------------------------- +// +// -- conduit::execution::ExecutionPolicy public methods -- +// +//----------------------------------------------------------------------------- + +//----------------------------------------------------------------------------- +// Construction and Destruction +//----------------------------------------------------------------------------- + /// standard constructor + ExecutionPolicy(); + /// copy constructor + ExecutionPolicy(const ExecutionPolicy& exec_policy); + /// Assignment operator + ExecutionPolicy& operator=(const ExecutionPolicy& exec_policy); + /// Construct from given policy id + ExecutionPolicy(PolicyID policy_id); + /// Construct from policy name + ExecutionPolicy(const std::string &policy_name); + /// destructor + ~ExecutionPolicy(); + +//----------------------------------------------------------------------------- +// Setters +//----------------------------------------------------------------------------- + void set_policy(PolicyID policy_id) + { m_policy_id = policy_id; } + + // TODO consider if we want to allow fallbacks. This could be the way to + // implement. (i.e. if you ask for openmp but there is no openmp it falls + // back to serial. If this isn't set then you get an error in that case.) + // void set_its_ok_to_lie(bool lies); + +//----------------------------------------------------------------------------- +// Getters and info methods. +//----------------------------------------------------------------------------- + PolicyID policy_id() const { return m_policy_id; } + std::string policy_name() const { return policy_id_to_name(m_policy_id); } + + // these methods answer questions about the chosen policy + bool is_empty() const; + bool is_serial() const; + bool is_cuda() const; + bool is_hip() const; + bool is_openmp() const; + + // these methods answer questions about where the policy can execute + bool is_host_policy() const; + bool is_device_policy() const; + + // these methods answer questions about which policies are able to be instantiated + static bool is_serial_enabled(); + static bool is_cuda_enabled(); + static bool is_hip_enabled(); + static bool is_openmp_enabled(); + + // these methods answer questions about which places we can instantiate policies for + static bool is_host_enabled(); + static bool is_device_enabled(); + +//----------------------------------------------------------------------------- +// Helpers to convert PolicyID Enum Values to human readable strings and +// vice versa. +//----------------------------------------------------------------------------- + static PolicyID name_to_policy_id(const std::string &name); + static std::string policy_id_to_name(const PolicyID policy_id); + +private: +//----------------------------------------------------------------------------- +// +// -- conduit::execution::ExecutionPolicy private data members -- +// +//----------------------------------------------------------------------------- + PolicyID m_policy_id; +}; +//----------------------------------------------------------------------------- +// -- end conduit::execution::ExecutionPolicy -- +//----------------------------------------------------------------------------- + +// registers the fancy conduit memory handlers for +// magic memset and memcpy +void init_device_memory_handlers(); + +// helper for the CONDUIT_DEVICE_ERROR_CHECK macro +void device_error_check(ExecutionPolicy policy, const char *file, const int line); + + + +struct EmptyPolicy +{}; + +#if defined(CONDUIT_USE_RAJA) +//---------------------------------------------------------------------------// +// RAJA_ON policies for when raja is on +//---------------------------------------------------------------------------// struct SerialExec { - using for_policy = seq::for_policy; - using sort_policy = seq::sort_policy; + using for_policy = RAJA::seq_exec; +#if defined(CONDUIT_USE_CUDA) + // the cuda/hip policy for reductions can be used + // by other backends, and this should suppress + // erroneous host device warnings + using reduce_policy = RAJA::cuda_reduce; +#elif defined(CONDUIT_USE_HIP) + using reduce_policy = RAJA::hip_reduce; +#else + using reduce_policy = RAJA::seq_reduce; +#endif + using atomic_policy = RAJA::seq_atomic; + using sort_policy = EmptyPolicy; + static std::string memory_space; +}; + +//--------------------------------------------------------------------------- +#if defined(CONDUIT_USE_CUDA) +struct CudaExec +{ + using for_policy = RAJA::cuda_exec; + using reduce_policy = RAJA::cuda_reduce; + using atomic_policy = RAJA::cuda_atomic; + using sort_policy = EmptyPolicy; + static std::string memory_space; +}; +#endif + +#if defined(CONDUIT_USE_HIP) +//--------------------------------------------------------------------------- +struct HipExec +{ + using for_policy = RAJA::hip_exec; + using reduce_policy = RAJA::hip_reduce; + using atomic_policy = RAJA::hip_atomic; + using sort_policy = EmptyPolicy; + static std::string memory_space; }; +#endif #if defined(CONDUIT_USE_OPENMP) //--------------------------------------------------------------------------- struct OpenMPExec { - using for_policy = omp::for_policy; - using sort_policy = omp::sort_policy; + using for_policy = RAJA::omp_parallel_for_exec; +#if defined(CONDUIT_USE_CUDA) + // the cuda policy for reductions can be used + // by other backends, and this should suppress + // erroneous host device warnings + using reduce_policy = RAJA::cuda_reduce; +#elif defined(CONDUIT_USE_HIP) + using reduce_policy = RAJA::hip_reduce; +#else + using reduce_policy = RAJA::omp_reduce; +#endif + using atomic_policy = RAJA::omp_atomic; + using sort_policy = EmptyPolicy; + static std::string memory_space; }; #endif + +#else +//---------------------------------------------------------------------------// +// RAJA_OFF policies for when raja is OFF +//---------------------------------------------------------------------------// + //--------------------------------------------------------------------------- -template +struct SerialExec +{ + using for_policy = EmptyPolicy; + using reduce_policy = EmptyPolicy; + using atomic_policy = EmptyPolicy; + using sort_policy = EmptyPolicy; + static std::string memory_space; +}; + +// TODO do we want this to be an option when RAJA is off? +// TODO decide about if we want to provide this as an option when RAJA is on as well +// if raja is on we don't want to use openmp not through raja +// and we do want it to be an option when raja is off +#if defined(CONDUIT_USE_OPENMP) +//--------------------------------------------------------------------------- +struct OpenMPExec +{ + using for_policy = EmptyPolicy; + using reduce_policy = EmptyPolicy; + using atomic_policy = EmptyPolicy; + using sort_policy = EmptyPolicy; + static std::string memory_space; +}; +#endif + +//---------------------------------------------------------------------------// +// mock up of a raja like forall implementation +//---------------------------------------------------------------------------// +template +inline void +forall_exec(ExecutionPolicy, + const int& begin, + const int& end, + Kernel&& kernel) noexcept +{ + std::cout << typeid(ExecutionPolicy).name() << " START" << std::endl; + for (int i = begin; i < end; i ++) + { + kernel(i); + } + std::cout << typeid(ExecutionPolicy).name() << " END" << std::endl; +} + +#if defined(CONDUIT_USE_OPENMP) +//---------------------------------------------------------------------------// +template +inline void +forall_exec(OpenMPExec, + const int& begin, + const int& end, + Kernel&& kernel) noexcept +{ + // #pragma message("omp::for_policy -> OMP") + #pragma omp parallel for + for (index_t i = begin; i < end; i ++) + { + kernel(i); + } +} +#endif + +//---------------------------------------------------------------------------// +// invoke forall with concrete template tag +//---------------------------------------------------------------------------// +template +inline void +forall(const int& begin, + const int& end, + Kernel&& kernel) noexcept +{ + forall_exec(ExecutionPolicy{}, begin, end, std::forward(kernel)); +} + +//---------------------------------------------------------------------------// +// mock up of a raja like sort implementation +//---------------------------------------------------------------------------// +template +inline void +sort_exec(ExecutionPolicy, + Iterator begin, + Iterator end) noexcept +{ + std::cout << typeid(ExecutionPolicy).name() << " START" << std::endl; + std::sort(begin, end); + std::cout << typeid(ExecutionPolicy).name() << " END" << std::endl; +} + +//---------------------------------------------------------------------------// +// mock up of a raja like sort implementation +//---------------------------------------------------------------------------// +template +inline void +sort_exec(ExecutionPolicy, + Iterator begin, + Iterator end, + Predicate &&predicate) noexcept +{ + std::cout << typeid(ExecutionPolicy).name() << " START" << std::endl; + std::sort(begin, end, predicate); + std::cout << typeid(ExecutionPolicy).name() << " END" << std::endl; +} + +#if defined(CONDUIT_USE_OPENMP) +//---------------------------------------------------------------------------// +template +inline void +sort_exec(OpenMPExec, + Iterator begin, + Iterator end) noexcept +{ + // #pragma message("omp::sort_policy -> serial") + // TODO: implement an OpenMP sort like in RAJA. + std::sort(begin, end); + // This is only allowed in C++14 or later. + //this->operator()(begin, end, [](auto &lhs, auto &rhs) { return lhs < rhs; }); +} + +//---------------------------------------------------------------------------// +template inline void -for_all(size_t begin, size_t end, Func &&func) +sort_exec(OpenMPExec, + Iterator begin, + Iterator end, + Predicate &&predicate) noexcept { - using policy = typename ExecutionPolicy::for_policy; - policy exec; - exec(begin, end, func); + // #pragma message("omp::sort_policy -> serial") + // TODO: implement an OpenMP sort like in RAJA. + std::sort(begin, end); + // This is only allowed in C++14 or later. + //this->operator()(begin, end, [](auto &lhs, auto &rhs) { return lhs < rhs; }); } +#endif +//---------------------------------------------------------------------------// +// invoke sort with concrete template tag +//---------------------------------------------------------------------------// template inline void -sort(Iterator begin, Iterator end) +sort(Iterator begin, + Iterator end) noexcept { - using policy = typename ExecutionPolicy::sort_policy; - policy exec; - exec(begin, end); + sort_exec(ExecutionPolicy{}, begin, end); } +//---------------------------------------------------------------------------// +// invoke sort with concrete template tag +//---------------------------------------------------------------------------// template inline void -sort(Iterator begin, Iterator end, Predicate &&predicate) +sort(Iterator begin, + Iterator end, + Predicate &&predicate) noexcept +{ + sort_exec(ExecutionPolicy{}, begin, end, std::forward(predicate)); +} + +#endif +//---------------------------------------------------------------------------// +// end RAJA_OFF +//---------------------------------------------------------------------------// + +//---------------------------------------------------------------------------// +// invoke functor with concrete template tag +//---------------------------------------------------------------------------// +template +inline void invoke(ExecPolicyTag &exec, Function&& func) noexcept +{ + func(exec); +} + +//---------------------------------------------------------------------------// +// runtime to concrete template tag dispatch of a functor +//---------------------------------------------------------------------------// +template +void +dispatch(ExecutionPolicy policy, Function&& func) +{ + if (policy.is_serial()) + { + SerialExec se; + invoke(se, func); + } + else if (policy.is_cuda()) + { +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_CUDA) + CudaExec ce; + invoke(ce, func); +#else + CONDUIT_ERROR("Conduit was not built with CUDA."); +#endif + } + else if (policy.is_hip()) + { +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_HIP) + HipExec he; + invoke(he, func); +#else + CONDUIT_ERROR("Conduit was not built with HIP."); +#endif + } + else if (policy.is_openmp()) + { +#if defined(CONDUIT_USE_OPENMP) + OpenMPExec ompe; + invoke(ompe, func); +#else + CONDUIT_ERROR("Conduit was not built with OpenMP."); +#endif + } + else // policy.is_empty() + { + CONDUIT_ERROR("Cannot invoke with an empty policy."); + } +} + +//---------------------------------------------------------------------------// +// runtime to concrete template tag dispatch of a forall +//---------------------------------------------------------------------------// +template +inline void +forall(ExecutionPolicy &policy, + const int& begin, + const int& end, + Kernel&& kernel) noexcept +{ + if (policy.is_serial()) + { + forall(begin, end, std::forward(kernel)); + } + else if (policy.is_cuda()) + { +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_CUDA) + forall(begin, end, std::forward(kernel)); +#else + CONDUIT_ERROR("Conduit was not built with CUDA."); +#endif + } + else if (policy.is_hip()) + { +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_USE_HIP) + forall(begin, end, std::forward(kernel)); +#else + CONDUIT_ERROR("Conduit was not built with HIP."); +#endif + } + else if (policy.is_openmp()) + { +#if defined(CONDUIT_USE_OPENMP) + forall(begin, end, std::forward(kernel)); +#else + CONDUIT_ERROR("Conduit was not built with OpenMP."); +#endif + } + else // policy.is_empty() + { + CONDUIT_ERROR("Cannot call forall with an empty policy."); + } +} + +//---------------------------------------------------------------------------// +// runtime to concrete template tag dispatch of a sort +//---------------------------------------------------------------------------// +template +inline void +sort(ExecutionPolicy &policy, + Iterator begin, + Iterator end) noexcept { - using policy = typename ExecutionPolicy::sort_policy; - policy exec; - exec(begin, end, predicate); + if (policy.is_serial()) + { + sort(begin, end); + } + else if (policy.is_cuda()) + { + CONDUIT_ERROR("sort does not exist for CUDA."); + } + else if (policy.is_hip()) + { + CONDUIT_ERROR("sort does not exist for HIP."); + } + else if (policy.is_openmp()) + { +#if defined(CONDUIT_USE_OPENMP) + sort(begin, end); +#else + CONDUIT_ERROR("Conduit was not built with OpenMP."); +#endif + } + else // policy.is_empty() + { + CONDUIT_ERROR("Cannot call sort with an empty policy."); + } +} + +//---------------------------------------------------------------------------// +// runtime to concrete template tag dispatch of a sort +//---------------------------------------------------------------------------// +template +inline void +sort(ExecutionPolicy &policy, + Iterator begin, + Iterator end, + Predicate &&predicate) noexcept +{ + if (policy.is_serial()) + { + sort(begin, end, std::forward(predicate)); + } + else if (policy.is_cuda()) + { + CONDUIT_ERROR("sort does not exist for CUDA."); + } + else if (policy.is_hip()) + { + CONDUIT_ERROR("sort does not exist for HIP."); + } + else if (policy.is_openmp()) + { +#if defined(CONDUIT_USE_OPENMP) + sort(begin, end, std::forward(predicate)); +#else + CONDUIT_ERROR("Conduit was not built with OpenMP."); +#endif + } + else // policy.is_empty() + { + CONDUIT_ERROR("Cannot call sort with an empty policy."); + } } } diff --git a/src/libs/conduit/conduit_execution_omp.hpp b/src/libs/conduit/conduit_execution_omp.hpp deleted file mode 100644 index 85f0bfc63..000000000 --- a/src/libs/conduit/conduit_execution_omp.hpp +++ /dev/null @@ -1,99 +0,0 @@ -// Copyright (c) Lawrence Livermore National Security, LLC and other Conduit -// Project developers. See top-level LICENSE AND COPYRIGHT files for dates and -// other details. No copyright assignment is required to contribute to Conduit. - -//----------------------------------------------------------------------------- -/// -/// file: conduit_execution_omp.hpp -/// -//----------------------------------------------------------------------------- - -#ifndef CONDUIT_EXECUTION_OMP_HPP -#define CONDUIT_EXECUTION_OMP_HPP - -//----------------------------------------------------------------------------- -// conduit lib includes -//----------------------------------------------------------------------------- -#include "conduit.hpp" - -#include "conduit_config.h" - -#if defined(CONDUIT_USE_OPENMP) -#include -#endif - -//----------------------------------------------------------------------------- -// -- begin conduit -- -//----------------------------------------------------------------------------- -namespace conduit -{ - -//----------------------------------------------------------------------------- -// -- begin conduit::execution -- -//----------------------------------------------------------------------------- -namespace execution -{ - -//----------------------------------------------------------------------------- -// -- begin conduit::execution::omp -- -//----------------------------------------------------------------------------- -namespace omp -{ - -//--------------------------------------------------------------------------- -struct for_policy -{ - template - inline void operator()(index_t begin, index_t end, Func &&func) - { -#if defined(CONDUIT_USE_OPENMP) -// #pragma message("omp::for_policy -> OMP") - #pragma omp parallel for - for(index_t i = begin; i < end; i++) - func(i); -#else -// #pragma message("omp::for_policy -> serial") - for(index_t i = begin; i < end; i++) - func(i); -#endif - } -}; - -//--------------------------------------------------------------------------- -struct sort_policy -{ - template - inline void operator()(Iterator begin, Iterator end) - { - // #pragma message("omp::sort_policy -> serial") - // TODO: implement an OpenMP sort like in RAJA. - std::sort(begin, end); - // This is only allowed in C++14 or later. - //this->operator()(begin, end, [](auto &lhs, auto &rhs) { return lhs < rhs; }); - } - - template - inline void operator()(Iterator begin, Iterator end, Predicate &&predicate) - { - // TODO: implement an OpenMP sort like in RAJA. - // #pragma message("omp::sort_policy -> serial") - std::sort(begin, end, predicate); - } -}; - -} -//----------------------------------------------------------------------------- -// -- end conduit::execution::omp -- -//----------------------------------------------------------------------------- - -} -//----------------------------------------------------------------------------- -// -- end conduit::execution -- -//----------------------------------------------------------------------------- - -} -//----------------------------------------------------------------------------- -// -- end conduit:: -- -//----------------------------------------------------------------------------- - -#endif diff --git a/src/libs/conduit/conduit_memory_manager.cpp b/src/libs/conduit/conduit_memory_manager.cpp new file mode 100644 index 000000000..91235c8d7 --- /dev/null +++ b/src/libs/conduit/conduit_memory_manager.cpp @@ -0,0 +1,310 @@ +// Copyright (c) Lawrence Livermore National Security, LLC and other Conduit +// Project developers. See top-level LICENSE AND COPYRIGHT files for dates and +// other details. No copyright assignment is required to contribute to Conduit. + +//----------------------------------------------------------------------------- +/// +/// file: conduit_memory_manager.cpp +/// +//----------------------------------------------------------------------------- + +//----------------------------------------------------------------------------- +// conduit lib includes +//----------------------------------------------------------------------------- +#include "conduit_memory_manager.hpp" +#include "conduit_config.h" +#include "conduit_utils.hpp" + +#if defined(CONDUIT_UMPIRE_ENABLED) +#include +#include +#include +#endif +#include // memcpy + +#if defined(CONDUIT_HIP_ENABLED) +#if HIP_VERSION_MAJOR >= 6 +#define TYPE_ATTR type +#else +#define TYPE_ATTR memoryType +#endif +#endif + +//----------------------------------------------------------------------------- +// -- begin conduit -- +//----------------------------------------------------------------------------- +namespace conduit +{ + +//----------------------------------------------------------------------------- +// -- begin conduit::execution -- +//----------------------------------------------------------------------------- +namespace execution +{ + +/// +/// Interfaces for host and device memory allocation / deallocation. +/// + + +//----------------------------------------------------------------------------- +//----------------------------------------------------------------------------- +// Host Memory +//----------------------------------------------------------------------------- +//----------------------------------------------------------------------------- +size_t HostMemory::m_total_bytes_alloced = 0; +size_t HostMemory::m_alloc_count = 0; +size_t HostMemory::m_free_count = 0; + +//----------------------------------------------------------------------------- +void * +HostMemory::allocate(size_t bytes) +{ + m_total_bytes_alloced += bytes; + m_alloc_count ++; +#if defined(CONDUIT_UMPIRE_ENABLED) + auto &rm = umpire::ResourceManager::getInstance (); + const int allocator_id = AllocationManager::host_allocator_id(); + umpire::Allocator host_allocator = rm.getAllocator (allocator_id); + return host_allocator.allocate(bytes); +#else + return malloc(bytes); +#endif +} + +//----------------------------------------------------------------------------- +void * +HostMemory::allocate(size_t items, size_t item_size) +{ + return allocate(items * item_size); +} + +//----------------------------------------------------------------------------- +void +HostMemory::deallocate(void *data_ptr) +{ + m_free_count ++; +#if defined(CONDUIT_UMPIRE_ENABLED) + auto &rm = umpire::ResourceManager::getInstance (); + const int allocator_id = AllocationManager::host_allocator_id(); + umpire::Allocator host_allocator = rm.getAllocator (allocator_id); + host_allocator.deallocate(data_ptr); +#else + return free(data_ptr); +#endif +} + +//----------------------------------------------------------------------------- +//----------------------------------------------------------------------------- +// Device Memory +//----------------------------------------------------------------------------- +//----------------------------------------------------------------------------- +size_t DeviceMemory::m_total_bytes_alloced = 0; +size_t DeviceMemory::m_alloc_count = 0; +size_t DeviceMemory::m_free_count = 0; + +//----------------------------------------------------------------------------- +void * +DeviceMemory::allocate(size_t bytes) +{ +#if !defined(CONDUIT_UMPIRE_ENABLED) + CONDUIT_ERROR("Conduit was built without Umpire support. " + "Cannot use DeviceMemory::alloc()."); +#endif + +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_UMPIRE_ENABLED) + m_total_bytes_alloced += bytes; + m_alloc_count ++; + auto &rm = umpire::ResourceManager::getInstance (); + const int allocator_id = AllocationManager::device_allocator_id(); + umpire::Allocator device_allocator = rm.getAllocator (allocator_id); + return device_allocator.allocate(bytes); +#else + (void) bytes; // unused + CONDUIT_ERROR("Calling device allocator when no device is present."); + return nullptr; +#endif +} + +//----------------------------------------------------------------------------- +void * +DeviceMemory::allocate(size_t items, size_t item_size) +{ + return allocate(items * item_size); +} + +//----------------------------------------------------------------------------- +void +DeviceMemory::deallocate(void *data_ptr) +{ +#if !defined(CONDUIT_UMPIRE_ENABLED) + CONDUIT_ERROR("Conduit was built without Umpire support. " + "Cannot use DeviceMemory::free()."); +#endif + +#if defined(CONDUIT_USE_RAJA) && defined(CONDUIT_UMPIRE_ENABLED) + m_free_count++; + auto &rm = umpire::ResourceManager::getInstance (); + const int allocator_id = AllocationManager::device_allocator_id(); + umpire::Allocator device_allocator = rm.getAllocator (allocator_id); + device_allocator.deallocate (data_ptr); +#else + (void) data_ptr; + CONDUIT_ERROR("Calling device allocator when no device is present."); +#endif +} + +// HIP and CUDA are mutually exclusive + +//----------------------------------------------------------------------------- +void +DeviceMemory::is_device_ptr(const void *ptr, bool &is_gpu, bool &is_unified) +{ + is_gpu = false; + is_unified = false; +#if defined(CONDUIT_CUDA_ENABLED) + cudaPointerAttributes atts; + const cudaError_t perr = cudaPointerGetAttributes(&atts, ptr); + + is_gpu = false; + is_unified = false; + + // clear last error so other error checking does + // not pick it up + cudaError_t error = cudaGetLastError(); + is_gpu = (perr == cudaSuccess) && + (atts.type == cudaMemoryTypeDevice || + atts.type == cudaMemoryTypeManaged ); + + is_unified = cudaSuccess && atts.type == cudaMemoryTypeDevice; +#elif defined(CONDUIT_HIP_ENABLED) + hipPointerAttribute_t atts; + const hipError_t perr = hipPointerGetAttributes(&atts, ptr); + + is_gpu = false; + is_unified = false; + + // clear last error so other error checking does + // not pick it up + hipError_t error = hipGetLastError(); + is_gpu = (perr == hipSuccess) && + (atts.TYPE_ATTR == hipMemoryTypeDevice || + atts.TYPE_ATTR == hipMemoryTypeUnified ); + // CYRUSH: this doens't look right: + is_unified = (hipSuccess && atts.TYPE_ATTR == hipMemoryTypeDevice); +#else + (void) ptr; +#endif +} + +//----------------------------------------------------------------------------- +// Adapted from: +// https://gitlab.kitware.com/third-party/nvpipe/blob/master/encode.c +bool +DeviceMemory::is_device_ptr(const void *ptr) +{ +#if defined(CONDUIT_CUDA_ENABLED) + cudaPointerAttributes atts; + const cudaError_t perr = cudaPointerGetAttributes(&atts, ptr); + // clear last error so other error checking does + // not pick it up + cudaError_t error = cudaGetLastError(); + return perr == cudaSuccess && + (atts.type == cudaMemoryTypeDevice || + atts.type == cudaMemoryTypeManaged); + +#elif defined(CONDUIT_HIP_ENABLED) + hipPointerAttribute_t atts; + const hipError_t perr = hipPointerGetAttributes(&atts, ptr); + // clear last error so other error checking does + // not pick it up + hipError_t error = hipGetLastError(); + return perr == hipSuccess && + (atts.TYPE_ATTR == hipMemoryTypeDevice || + atts.TYPE_ATTR == hipMemoryTypeUnified); +#else + (void) ptr; + return false; +#endif +} + +//----------------------------------------------------------------------------- +//----------------------------------------------------------------------------- +// Magic Memory +//----------------------------------------------------------------------------- +//----------------------------------------------------------------------------- + +//----------------------------------------------------------------------------- +void +MagicMemory::set(void * ptr, int value, size_t num ) +{ +#if defined(CONDUIT_USE_RAJA) + bool is_device = DeviceMemory::is_device_ptr(ptr); + if (is_device) + { +#if defined(CONDUIT_CUDA_ENABLED) + cudaMemset(ptr,value,num); +#elif defined(CONDUIT_HIP_ENABLED) + hipMemset(ptr,value,num); +#endif + } + else + { + memset(ptr,value,num); + } +#else + memset(ptr,value,num); +#endif +} + +//----------------------------------------------------------------------------- +void +MagicMemory::copy(void * destination, const void * source, size_t num) +{ +#if defined(CONDUIT_USE_RAJA) + bool src_is_gpu = DeviceMemory::is_device_ptr(source); + bool dst_is_gpu = DeviceMemory::is_device_ptr(destination); + if (src_is_gpu && dst_is_gpu) + { +#if defined(CONDUIT_CUDA_ENABLED) + cudaMemcpy(destination, source, num, cudaMemcpyDeviceToDevice); +#elif defined(CONDUIT_HIP_ENABLED) + hipMemcpy(destination, source, num, hipMemcpyDeviceToDevice); +#endif + } + else if (src_is_gpu && !dst_is_gpu) + { +#if defined(CONDUIT_CUDA_ENABLED) + cudaMemcpy(destination, source, num, cudaMemcpyDeviceToHost); +#elif defined(CONDUIT_HIP_ENABLED) + hipMemcpy(destination, source, num, hipMemcpyDeviceToHost); +#endif + } + else if (!src_is_gpu && dst_is_gpu) + { +#if defined(CONDUIT_CUDA_ENABLED) + cudaMemcpy(destination, source, num, cudaMemcpyHostToDevice); +#elif defined(CONDUIT_HIP_ENABLED) + hipMemcpy(destination, source, num, hipMemcpyHostToDevice); +#endif + } + else + { + // we are the default memcpy in conduit so this is the normal + // path + memcpy(destination,source,num); + } +#else + memcpy(destination,source,num); +#endif +} + +} +//----------------------------------------------------------------------------- +// -- end conduit::execution -- +//----------------------------------------------------------------------------- + +} +//----------------------------------------------------------------------------- +// -- end conduit:: -- +//----------------------------------------------------------------------------- diff --git a/src/libs/conduit/conduit_execution_serial.hpp b/src/libs/conduit/conduit_memory_manager.hpp similarity index 54% rename from src/libs/conduit/conduit_execution_serial.hpp rename to src/libs/conduit/conduit_memory_manager.hpp index 5ecc5f604..07b9fab92 100644 --- a/src/libs/conduit/conduit_execution_serial.hpp +++ b/src/libs/conduit/conduit_memory_manager.hpp @@ -4,17 +4,19 @@ //----------------------------------------------------------------------------- /// -/// file: conduit_execution_serial.hpp +/// file: conduit_memory_manager.hpp /// //----------------------------------------------------------------------------- -#ifndef CONDUIT_EXECUTION_SERIAL_HPP -#define CONDUIT_EXECUTION_SERIAL_HPP +#ifndef CONDUIT_MEMORY_MANAGER_HPP +#define CONDUIT_MEMORY_MANAGER_HPP + +#include //----------------------------------------------------------------------------- // conduit lib includes //----------------------------------------------------------------------------- -#include "conduit.hpp" +#include "conduit_config.h" //----------------------------------------------------------------------------- // -- begin conduit -- @@ -28,43 +30,53 @@ namespace conduit namespace execution { +/// +/// Interfaces for host and device memory allocation / deallocation. +/// + //----------------------------------------------------------------------------- -// -- begin conduit::execution::seq -- +/// Host Memory allocation / deallocation interface (singleton) +/// Uses AllocationManager::host_allocator_id() when Umpire is enabled, +/// Uses malloc/free when Umpire is disabled. //----------------------------------------------------------------------------- -namespace seq +struct HostMemory { + static void *allocate(size_t bytes); + static void *allocate(size_t items, size_t item_size); + static void deallocate(void *data_ptr); -//--------------------------------------------------------------------------- -struct for_policy -{ - template - inline void operator()(index_t begin, index_t end, Func &&func) - { - for(index_t i = begin; i < end; i++) - func(i); - } -}; +private: + static size_t m_total_bytes_alloced; + static size_t m_alloc_count; + static size_t m_free_count; -//--------------------------------------------------------------------------- -struct sort_policy +}; +//----------------------------------------------------------------------------- +/// Device Memory allocation / deallocation interface (singleton) +/// Uses AllocationManager::device_allocator_id() when Umpire is enabled. +/// allocate() and deallocate() throw errors when Umpire is disabled. +//----------------------------------------------------------------------------- +struct DeviceMemory { - template - inline void operator()(Iterator begin, Iterator end) - { - std::sort(begin, end); - } - - template - inline void operator()(Iterator begin, Iterator end, Predicate &&predicate) - { - std::sort(begin, end, predicate); - } + static void *allocate(size_t bytes); + static void *allocate(size_t items, size_t item_size); + static void deallocate(void *data_ptr); + static bool is_device_ptr(const void *ptr); + static void is_device_ptr(const void *ptr, bool &is_gpu, bool &is_unified); + +private: + static size_t m_total_bytes_alloced; + static size_t m_alloc_count; + static size_t m_free_count; + }; -} -//----------------------------------------------------------------------------- -// -- end conduit::execution::seq -- //----------------------------------------------------------------------------- +struct MagicMemory +{ + static void set(void *ptr, int value, size_t num); + static void copy(void *destination, const void *source, size_t num); +}; } //----------------------------------------------------------------------------- diff --git a/src/libs/conduit/conduit_node.cpp b/src/libs/conduit/conduit_node.cpp index 5571e52f6..84f1787f0 100644 --- a/src/libs/conduit/conduit_node.cpp +++ b/src/libs/conduit/conduit_node.cpp @@ -12116,7 +12116,6 @@ Node::Value::operator long_double_accessor() const #endif //---------------------------------------------------------------------------// - //----------------------------------------------------------------------------- // -- ConstValue Helper class --- //----------------------------------------------------------------------------- diff --git a/src/libs/conduit/conduit_node.hpp b/src/libs/conduit/conduit_node.hpp index 6bbcb38da..ad26f9250 100644 --- a/src/libs/conduit/conduit_node.hpp +++ b/src/libs/conduit/conduit_node.hpp @@ -86,6 +86,10 @@ class CONDUIT_API Node friend class NodeIterator; friend class NodeConstIterator; friend class Generator; + template + friend class DataArray; + template + friend class DataAccessor; #if defined(CONDUIT_USE_TOTALVIEW) friend int ::TV_ttf_display_type ( const conduit::Node *n ); @@ -3396,7 +3400,6 @@ class CONDUIT_API Node operator long_double_accessor() const; #endif - private: // This is private we only want conduit::Node to create a // conduit::Node::Value instance @@ -3537,7 +3540,6 @@ class CONDUIT_API Node operator long_double_accessor() const; #endif - private: // This is private we only want conduit::Node to create a // conduit::Node::ConstValue instance @@ -4152,7 +4154,6 @@ class CONDUIT_API Node // index type array accessors index_t_accessor as_index_t_accessor() const; - // char8_str cases char *as_char8_str(); const char *as_char8_str() const; @@ -4412,7 +4413,6 @@ class CONDUIT_API Node long_double_accessor as_long_double_accessor() const; #endif - //----------------------------------------------------------------------------- ///@} //----------------------------------------------------------------------------- diff --git a/src/tests/conduit/CMakeLists.txt b/src/tests/conduit/CMakeLists.txt index 8a1d7c3e2..b7eb11a40 100644 --- a/src/tests/conduit/CMakeLists.txt +++ b/src/tests/conduit/CMakeLists.txt @@ -20,7 +20,6 @@ set(BASIC_TESTS t_conduit_smoke t_conduit_node_static_init t_conduit_node_move_and_swap t_conduit_serialize - t_conduit_array t_conduit_list_of t_conduit_node_binary_io t_conduit_node_save_load @@ -33,6 +32,7 @@ set(BASIC_TESTS t_conduit_smoke t_conduit_json_sanitize t_conduit_yaml t_conduit_generator + t_conduit_data_array t_conduit_data_accessor t_conduit_node_update t_conduit_node_compact @@ -43,6 +43,7 @@ set(BASIC_TESTS t_conduit_smoke t_conduit_node_type_dispatch t_conduit_schema t_conduit_error + t_conduit_execution t_conduit_log t_conduit_utils t_conduit_annotations diff --git a/src/tests/conduit/t_conduit_data_accessor.cpp b/src/tests/conduit/t_conduit_data_accessor.cpp index 3e39c0e91..5b4be1e6f 100644 --- a/src/tests/conduit/t_conduit_data_accessor.cpp +++ b/src/tests/conduit/t_conduit_data_accessor.cpp @@ -206,47 +206,93 @@ TEST(conduit_data_accessor, default_construct) TEST(conduit_data_accessor, set) { - Node n; - n.set(DataType::int8(10)); - - int8_accessor i8_acc = n.value(); - int16_accessor i16_acc = n.value(); - int32_accessor i32_acc = n.value(); - int64_accessor i64_acc = n.value(); - - uint8_accessor ui8_acc = n.value(); - uint16_accessor ui16_acc = n.value(); - uint32_accessor ui32_acc = n.value(); - uint64_accessor ui64_acc = n.value(); - - float32_accessor f32_acc = n.value(); - float64_accessor f64_acc = n.value(); - - i8_acc.set(0,-4); - i16_acc.set(1,-8); - i32_acc.set(2,-16); - i64_acc.set(3,-32); - - ui8_acc.set(4, 4); - ui16_acc.set(5,8); - ui32_acc.set(6,16); - ui64_acc.set(7,32); - - f32_acc.set(8,16); - f64_acc.set(9,32); - - EXPECT_EQ(i32_acc[0],-4); - EXPECT_EQ(i32_acc[1],-8); - EXPECT_EQ(i32_acc[2],-16); - EXPECT_EQ(i32_acc[3],-32); - - EXPECT_EQ(i32_acc[4],4); - EXPECT_EQ(i32_acc[5],8); - EXPECT_EQ(i32_acc[6],16); - EXPECT_EQ(i32_acc[7],32); + { + Node n; + n.set(DataType::int8(10)); + + int8_accessor i8_acc(n); + int16_accessor i16_acc(n); + int32_accessor i32_acc(n); + int64_accessor i64_acc(n); + + uint8_accessor ui8_acc(n); + uint16_accessor ui16_acc(n); + uint32_accessor ui32_acc(n); + uint64_accessor ui64_acc(n); + + float32_accessor f32_acc(n); + float64_accessor f64_acc(n); + + i8_acc.set(0,-4); + i16_acc.set(1,-8); + i32_acc.set(2,-16); + i64_acc.set(3,-32); + + ui8_acc.set(4, 4); + ui16_acc.set(5,8); + ui32_acc.set(6,16); + ui64_acc.set(7,32); + + f32_acc.set(8,16); + f64_acc.set(9,32); + + EXPECT_EQ(i32_acc[0],-4); + EXPECT_EQ(i32_acc[1],-8); + EXPECT_EQ(i32_acc[2],-16); + EXPECT_EQ(i32_acc[3],-32); + + EXPECT_EQ(i32_acc[4],4); + EXPECT_EQ(i32_acc[5],8); + EXPECT_EQ(i32_acc[6],16); + EXPECT_EQ(i32_acc[7],32); + + EXPECT_EQ(i32_acc[8],16); + EXPECT_EQ(i32_acc[9],32); + } - EXPECT_EQ(i32_acc[8],16); - EXPECT_EQ(i32_acc[9],32); + { + Node n; + n.set(DataType::int8(10)); + + int8_accessor i8_acc = n.value(); + int16_accessor i16_acc = n.value(); + int32_accessor i32_acc = n.value(); + int64_accessor i64_acc = n.value(); + + uint8_accessor ui8_acc = n.value(); + uint16_accessor ui16_acc = n.value(); + uint32_accessor ui32_acc = n.value(); + uint64_accessor ui64_acc = n.value(); + + float32_accessor f32_acc = n.value(); + float64_accessor f64_acc = n.value(); + + i8_acc.set(0,-4); + i16_acc.set(1,-8); + i32_acc.set(2,-16); + i64_acc.set(3,-32); + + ui8_acc.set(4, 4); + ui16_acc.set(5,8); + ui32_acc.set(6,16); + ui64_acc.set(7,32); + + f32_acc.set(8,16); + f64_acc.set(9,32); + + EXPECT_EQ(i32_acc[0],-4); + EXPECT_EQ(i32_acc[1],-8); + EXPECT_EQ(i32_acc[2],-16); + EXPECT_EQ(i32_acc[3],-32); + + EXPECT_EQ(i32_acc[4],4); + EXPECT_EQ(i32_acc[5],8); + EXPECT_EQ(i32_acc[6],16); + EXPECT_EQ(i32_acc[7],32); + + EXPECT_EQ(i32_acc[8],16); + EXPECT_EQ(i32_acc[9],32); + } } @@ -289,5 +335,255 @@ TEST(conduit_data_accessor, to_string) } +//----------------------------------------------------------------------------- +TEST(conduit_data_accessor, set_using_data_array) +{ + std::vector v_int8(10,-8); + std::vector v_int16(10,-16); + std::vector v_int32(10,-32); + std::vector v_int64(10,-64); + + std::vector v_uint8(10,8); + std::vector v_uint16(10,16); + std::vector v_uint32(10,32); + std::vector v_uint64(10,64); + + std::vector v_float32(10,32.0); + std::vector v_float64(10,64.0); + + int8_array va_int8(&v_int8[0],DataType::int8(10)); + int16_array va_int16(&v_int16[0],DataType::int16(10)); + int32_array va_int32(&v_int32[0],DataType::int32(10)); + int64_array va_int64(&v_int64[0],DataType::int64(10)); + + uint8_array va_uint8(&v_uint8[0],DataType::uint8(10)); + uint16_array va_uint16(&v_uint16[0],DataType::uint16(10)); + uint32_array va_uint32(&v_uint32[0],DataType::uint32(10)); + uint64_array va_uint64(&v_uint64[0],DataType::uint64(10)); + + float32_array va_float32(&v_float32[0],DataType::float32(10)); + float64_array va_float64(&v_float64[0],DataType::float64(10)); + + + Node n; + + // int8_array + n["vint8"].set(DataType::int8(10)); + n["vint8"].as_int8_accessor().set(va_int8); + int8 *n_int8_ptr = n["vint8"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int8_ptr[i],va_int8[i]); + } + // int16_array + n["vint16"].set(DataType::int16(10)); + n["vint16"].as_int16_accessor().set(va_int16); + int16 *n_int16_ptr = n["vint16"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int16_ptr[i],va_int16[i]); + } + + // int32_array + n["vint32"].set(DataType::int32(10)); + n["vint32"].as_int32_accessor().set(va_int32); + int32 *n_int32_ptr = n["vint32"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int32_ptr[i],va_int32[i]); + } + + // int64_array + n["vint64"].set(DataType::int64(10)); + n["vint64"].as_int64_accessor().set(va_int64); + int64 *n_int64_ptr = n["vint64"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int64_ptr[i],va_int64[i]); + } + + // uint8_array + n["vuint8"].set(DataType::uint8(10)); + n["vuint8"].as_uint8_accessor().set(va_uint8); + uint8 *n_uint8_ptr = n["vuint8"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint8_ptr[i],va_uint8[i]); + } + + // uint16_array + n["vuint16"].set(DataType::uint16(10)); + n["vuint16"].as_uint16_accessor().set(va_uint16); + uint16 *n_uint16_ptr = n["vuint16"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint16_ptr[i],va_uint16[i]); + } + + // uint32_array + n["vuint32"].set(DataType::uint32(10)); + n["vuint32"].as_uint32_accessor().set(va_uint32); + uint32 *n_uint32_ptr = n["vuint32"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint32_ptr[i],va_uint32[i]); + } + + // uint64_array + n["vuint64"].set(DataType::uint64(10)); + n["vuint64"].as_uint64_accessor().set(va_uint64); + uint64 *n_uint64_ptr = n["vuint64"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint64_ptr[i],va_uint64[i]); + } + + + // float32_array + n["vfloat32"].set(DataType::float32(10)); + n["vfloat32"].as_float32_accessor().set(va_float32); + float32 *n_float32_ptr = n["vfloat32"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_float32_ptr[i],va_float32[i]); + } + + // float64_array + n["vfloat64"].set(DataType::float64(10)); + n["vfloat64"].as_float64_accessor().set(va_float64); + float64 *n_float64_ptr = n["vfloat64"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_float64_ptr[i],va_float64[i]); + } + +} + + + +//----------------------------------------------------------------------------- +TEST(conduit_data_accessor, set_using_data_accessor) +{ + std::vector v_int8(10,-8); + std::vector v_int16(10,-16); + std::vector v_int32(10,-32); + std::vector v_int64(10,-64); + + std::vector v_uint8(10,8); + std::vector v_uint16(10,16); + std::vector v_uint32(10,32); + std::vector v_uint64(10,64); + + std::vector v_float32(10,32.0); + std::vector v_float64(10,64.0); + + int8_accessor vacc_int8(&v_int8[0],DataType::int8(10)); + int16_accessor vacc_int16(&v_int16[0],DataType::int16(10)); + int32_accessor vacc_int32(&v_int32[0],DataType::int32(10)); + int64_accessor vacc_int64(&v_int64[0],DataType::int64(10)); + + uint8_accessor vacc_uint8(&v_uint8[0],DataType::uint8(10)); + uint16_accessor vacc_uint16(&v_uint16[0],DataType::uint16(10)); + uint32_accessor vacc_uint32(&v_uint32[0],DataType::uint32(10)); + uint64_accessor vacc_uint64(&v_uint64[0],DataType::uint64(10)); + + float32_accessor vacc_float32(&v_float32[0],DataType::float32(10)); + float64_accessor vacc_float64(&v_float64[0],DataType::float64(10)); + + + Node n; + + // int8_array + n["vint8"].set(DataType::int8(10)); + n["vint8"].as_int8_accessor().set(vacc_int8); + int8 *n_int8_ptr = n["vint8"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int8_ptr[i],v_int8[i]); + } + + // int16_array + n["vint16"].set(DataType::int16(10)); + n["vint16"].as_int16_accessor().set(vacc_int16); + int16 *n_int16_ptr = n["vint16"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int16_ptr[i],v_int16[i]); + } + + // int32_array + n["vint32"].set(DataType::int32(10)); + n["vint32"].as_int32_accessor().set(vacc_int32); + int32 *n_int32_ptr = n["vint32"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int32_ptr[i],v_int32[i]); + } + + // int64_array + n["vint64"].set(DataType::int64(10)); + n["vint64"].as_int64_accessor().set(vacc_int64); + int64 *n_int64_ptr = n["vint64"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_int64_ptr[i],v_int64[i]); + } + + // uint8_array + n["vuint8"].set(DataType::uint8(10)); + n["vuint8"].as_uint8_accessor().set(vacc_uint8); + uint8 *n_uint8_ptr = n["vuint8"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint8_ptr[i],v_uint8[i]); + } + + // uint16_array + n["vuint16"].set(DataType::uint16(10)); + n["vuint16"].as_uint16_accessor().set(vacc_uint16); + uint16 *n_uint16_ptr = n["vuint16"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint16_ptr[i],v_uint16[i]); + } + + // uint32_array + n["vuint32"].set(DataType::uint32(10)); + n["vuint32"].as_uint32_accessor().set(vacc_uint32); + uint32 *n_uint32_ptr = n["vuint32"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint32_ptr[i],v_uint32[i]); + } + + // uint64_array + n["vuint64"].set(DataType::uint64(10)); + n["vuint64"].as_uint64_accessor().set(vacc_uint64); + uint64 *n_uint64_ptr = n["vuint64"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_uint64_ptr[i],v_uint64[i]); + } + + + // float32_array + n["vfloat32"].set(DataType::float32(10)); + n["vfloat32"].as_float32_accessor().set(vacc_float32); + float32 *n_float32_ptr = n["vfloat32"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_float32_ptr[i],v_float32[i]); + } + + // float64_array + n["vfloat64"].set(DataType::float64(10)); + n["vfloat64"].as_float64_accessor().set(vacc_float64); + float64 *n_float64_ptr = n["vfloat64"].value(); + for(index_t i=0;i<10;i++) + { + EXPECT_EQ(n_float64_ptr[i],v_float64[i]); + } + +} diff --git a/src/tests/conduit/t_conduit_array.cpp b/src/tests/conduit/t_conduit_data_array.cpp similarity index 98% rename from src/tests/conduit/t_conduit_array.cpp rename to src/tests/conduit/t_conduit_data_array.cpp index 8265d0465..6b2d4329a 100644 --- a/src/tests/conduit/t_conduit_array.cpp +++ b/src/tests/conduit/t_conduit_data_array.cpp @@ -16,7 +16,7 @@ using namespace conduit; //----------------------------------------------------------------------------- -TEST(conduit_array, basic_construction) +TEST(conduit_data_array, basic_construction) { std::vector data1(10,8); std::vector data2(10,-8); @@ -70,7 +70,7 @@ TEST(conduit_array, basic_construction) } //----------------------------------------------------------------------------- -TEST(conduit_array, array_stride_int8) +TEST(conduit_data_array, array_stride_int8) { std::vector data(20,0); @@ -116,7 +116,7 @@ TEST(conduit_array, array_stride_int8) arr[1] = 100; EXPECT_EQ(data[2],100); - std::cout << "Full Data" << std::endl; + std::cout << "Full Data" << std::endl; for(int i=0;i<20;i++) { @@ -144,7 +144,7 @@ TEST(conduit_array, array_stride_int8) } //----------------------------------------------------------------------------- -TEST(conduit_array, array_stride_int8_external) +TEST(conduit_data_array, array_stride_int8_external) { std::vector data(20,0); @@ -187,7 +187,7 @@ TEST(conduit_array, array_stride_int8_external) //----------------------------------------------------------------------------- -TEST(conduit_array, set_using_ptrs) +TEST(conduit_data_array, set_using_ptrs) { //in this case we are using std vectors to init data conveniently // we are actually testing the pointer set cases @@ -304,7 +304,7 @@ TEST(conduit_array, set_using_ptrs) //----------------------------------------------------------------------------- -TEST(conduit_array, set_using_data_array) +TEST(conduit_data_array, set_using_data_array) { std::vector v_int8(10,-8); std::vector v_int16(10,-16); @@ -429,7 +429,7 @@ TEST(conduit_array, set_using_data_array) } //----------------------------------------------------------------------------- -TEST(conduit_array, set_single_element) +TEST(conduit_data_array, set_single_element) { std::vector v_int8(10,-8); std::vector v_int16(10,-16); @@ -499,7 +499,7 @@ TEST(conduit_array, set_single_element) } //----------------------------------------------------------------------------- -TEST(conduit_array, set_using_data_accessor) +TEST(conduit_data_array, set_using_data_accessor) { std::vector v_int8(10,-8); std::vector v_int16(10,-16); @@ -623,7 +623,7 @@ TEST(conduit_array, set_using_data_accessor) } //----------------------------------------------------------------------------- -TEST(conduit_array, set_using_std_vectors) +TEST(conduit_data_array, set_using_std_vectors) { std::vector v_int8(10,-8); std::vector v_int16(10,-16); @@ -735,7 +735,7 @@ TEST(conduit_array, set_using_std_vectors) } //----------------------------------------------------------------------------- -TEST(conduit_array, print_bells_and_whistles) +TEST(conduit_data_array, print_bells_and_whistles) { Node n; @@ -797,7 +797,7 @@ TEST(conduit_array, print_bells_and_whistles) } //----------------------------------------------------------------------------- -TEST(conduit_array, fill) +TEST(conduit_data_array, fill) { int num_ele = 5; @@ -864,12 +864,12 @@ TEST(conduit_array, fill) va_int8.fill((int8) -1); va_int16.fill((int16) -1); va_int32.fill((int32) -1); - va_int64.fill((int32) -1); + va_int64.fill((int64) -1); va_uint8.fill((uint8) 1); va_uint16.fill((uint16) 1); va_uint32.fill((uint32) 1); - va_uint64.fill((uint32) 1); + va_uint64.fill((uint64) 1); va_float32.fill((float32) 1.0); va_float64.fill((float64) 1.0); @@ -895,7 +895,7 @@ TEST(conduit_array, fill) } //----------------------------------------------------------------------------- -TEST(conduit_array, compact_to_bytes) +TEST(conduit_data_array, compact_to_bytes) { std::vector vals(8,0); @@ -923,7 +923,7 @@ TEST(conduit_array, compact_to_bytes) } //----------------------------------------------------------------------------- -TEST(conduit_array, summary_stats) +TEST(conduit_data_array, summary_stats) { std::vector v_int64(3,-64); std::vector v_uint64(3,64); @@ -959,7 +959,7 @@ TEST(conduit_array, summary_stats) //----------------------------------------------------------------------------- -TEST(conduit_array, summary_print) +TEST(conduit_data_array, summary_print) { std::vector v_int64(5,-64); int64_array va_int64(&v_int64[0],DataType::int64(5)); @@ -1031,7 +1031,7 @@ TEST(conduit_array, summary_print) //----------------------------------------------------------------------------- -TEST(conduit_array, cxx_11_init_lists) +TEST(conduit_data_array, cxx_11_init_lists) { std::vector v_int8(3,-8); std::vector v_int16(3,-16); diff --git a/src/tests/conduit/t_conduit_execution.cpp b/src/tests/conduit/t_conduit_execution.cpp new file mode 100644 index 000000000..25775a74e --- /dev/null +++ b/src/tests/conduit/t_conduit_execution.cpp @@ -0,0 +1,488 @@ +// Copyright (c) Lawrence Livermore National Security, LLC and other Conduit +// Project developers. See top-level LICENSE AND COPYRIGHT files for dates and +// other details. No copyright assignment is required to contribute to Conduit. + +//----------------------------------------------------------------------------- +/// +/// file: t_conduit_execution.cpp +/// +//----------------------------------------------------------------------------- + +#include "conduit.hpp" +#include "conduit_execution.hpp" +#include "conduit_memory_manager.hpp" + +#include +#include "gtest/gtest.h" + +using namespace conduit; +using conduit::execution::ExecutionPolicy; + +void *device_alloc(index_t bytes) +{ +#if defined(CONDUIT_USE_RAJA) + return execution::DeviceMemory::allocate(bytes); +#else + return execution::HostMemory::allocate(bytes); +#endif +} + +void device_free(void *ptr) +{ +#if defined(CONDUIT_USE_RAJA) + return execution::DeviceMemory::deallocate(ptr); +#else + return execution::HostMemory::deallocate(ptr); +#endif +} + +void conduit_device_prepare() +{ + execution::init_device_memory_handlers(); +} + +// TODO someday we want allocator to make sense for nodes when we are done with them + +// TODO turn the strawman into tests? + +//---------------------------------------------------------------------------// +// example functor +//---------------------------------------------------------------------------// +struct MyFunctor +{ + int res; + int size; + template + void operator()(ComboPolicyTag &exec) + { + std::cout << typeid(ComboPolicyTag).name() << std::endl; + using thetag = typename ComboPolicyTag::for_policy; + std::cout << typeid(thetag).name() << std::endl; + res = 0; + conduit::execution::forall(0, size, [=] (int i) + { + std::cout << i << std::endl; + res ++; + }); + } +}; + +//---------------------------------------------------------------------------// +// Mock of a class templated on a concrete tag +// (like a RAJA Reduction Object) +//---------------------------------------------------------------------------// +template +class MySpecialClass +{ +public: + using policy = ExecPolicy; + int val; + + MySpecialClass(int _val) + :val(_val) + {} + + void exec(int i) const + { + std::cout << typeid(policy).name() << " exec " << val << " " << i << std::endl; + } +}; + +//---------------------------------------------------------------------------// +// example functor using MySpecialClass +//---------------------------------------------------------------------------// +struct MySpecialFunctor +{ + int res; + int size; + template + void operator()(ComboPolicyTag &exec) + { + // in this case we use an object + // that is templated on a concrete tag + // (like a RAJA Reduction Object) + using thetag = typename ComboPolicyTag::for_policy; + res = 0; + MySpecialClass s(10); + conduit::execution::forall(0, size, [=] (int i) + { + s.exec(i); + res ++; + }); + } +}; + +//----------------------------------------------------------------------------- +TEST(conduit_execution, test_forall) +{ + conduit_device_prepare(); + const index_t size = 10; + + index_t host_vals[size]; + index_t *dev_vals_ptr = static_cast(device_alloc(sizeof(index_t) * size)); + + ExecutionPolicy serial = ExecutionPolicy::serial(); + conduit::execution::forall(serial, 0, size, [=](index_t i) + { + dev_vals_ptr[i] = i; + }); + CONDUIT_DEVICE_ERROR_CHECK(serial); + + conduit::execution::MagicMemory::copy(&host_vals[0], dev_vals_ptr, sizeof(index_t) * size); + + for (index_t i = 0; i < size; i ++) + { + EXPECT_EQ(host_vals[i],i); + } + + device_free(dev_vals_ptr); +} + +// //----------------------------------------------------------------------------- +// TEST(conduit_execution, test_reductions) +// { +// Conduit::execution::ExecPolicy SerialPolicy(conduit::execution::policy::Serial); +// const index_t size = 4; +// index_t host_vals[size] = {0,-10,10, 5}; +// index_t *dev_vals_ptr = static_cast(device_alloc(sizeof(index_t) * size)); +// MagicMemory::copy(dev_vals_ptr, &host_vals[0], sizeof(index_t) * size); + + +// // sum +// // ascent::ReduceSum sum_reducer; +// using reduce_policy = typename conduit::execution::policy::Serial::reduce_policy; +// conduit::execution::forall(0, size, [=](index_t i) +// { +// sum_reducer += dev_vals_ptr[i]; +// }); +// // CONDUIT_DEVICE_ERROR_CHECK(); + +// EXPECT_EQ(sum_reducer.get(),5); + + +// // // min +// // ascent::ReduceMin min_reducer; +// // conduit::execution::forall(0, size, [=](index_t i) +// // { +// // min_reducer.min(dev_vals_ptr[i]); +// // }); +// // // CONDUIT_DEVICE_ERROR_CHECK(); + +// // EXPECT_EQ(min_reducer.get(),-10); + +// // // minloc +// // ascent::ReduceMinLoc minloc_reducer; +// // conduit::execution::forall(0, size, [=](index_t i) +// // { +// // minloc_reducer.minloc(dev_vals_ptr[i],i); +// // }); +// // // CONDUIT_DEVICE_ERROR_CHECK(); + +// // EXPECT_EQ(minloc_reducer.get(),-10); +// // EXPECT_EQ(minloc_reducer.getLoc(),1); + + +// // // max +// // ascent::ReduceMax max_reducer; +// // conduit::execution::forall(0, size, [=](index_t i) +// // { +// // max_reducer.max(dev_vals_ptr[i]); +// // }); +// // // CONDUIT_DEVICE_ERROR_CHECK(); + +// // EXPECT_EQ(max_reducer.get(),10); + +// // // maxloc +// // ascent::ReduceMaxLoc maxloc_reducer; +// // conduit::execution::forall(0, size, [=](index_t i) +// // { +// // maxloc_reducer.maxloc(dev_vals_ptr[i],i); +// // }); +// // // CONDUIT_DEVICE_ERROR_CHECK(); + +// // EXPECT_EQ(maxloc_reducer.get(),10); +// // EXPECT_EQ(maxloc_reducer.getLoc(),2); + +// device_free(dev_vals_ptr); +// } + +//----------------------------------------------------------------------------- +TEST(conduit_execution, for_all_and_dispatch) +{ + std::cout << "forall cases!" << std::endl; + + const int size = 4; + MyFunctor func; + func.size = size; + MySpecialFunctor sfunc; + sfunc.size = 4; + + auto test_exec_policy = [&](ExecutionPolicy policy) + { + conduit::execution::forall(policy, 0, size, [=] (int i) + { + std::cout << i << std::endl; + }); + + std::cout << "functor cases!" << std::endl; + + conduit::execution::dispatch(policy, func); + std::cout << func.res << std::endl; + + conduit::execution::dispatch(policy, sfunc); + std::cout << func.res << std::endl; + + std::cout << "C++ 20" << std::endl; + + int res =0; + /// c++ 20 allows us to double lambda instead of a functor + + // apparently this works just fine with cpp14...? + + conduit::execution::dispatch(policy, [&] (ComboPolicyTag &exec) + { + using thetag = typename ComboPolicyTag::for_policy; + MySpecialClass s(10); + conduit::execution::forall(0, size, [=] (int i) + { + s.exec(i); + }); + res = 10; + }); + }; + + if (ExecutionPolicy::is_serial_enabled()) + { + ExecutionPolicy serial = ExecutionPolicy::serial(); + test_exec_policy(serial); + } + + if (ExecutionPolicy::is_openmp_enabled()) + { + ExecutionPolicy openmp = ExecutionPolicy::openmp(); + test_exec_policy(openmp); + } + + if (ExecutionPolicy::is_device_enabled()) + { + ExecutionPolicy device = ExecutionPolicy::device(); + test_exec_policy(device); + } +} + +//----------------------------------------------------------------------------- +TEST(conduit_execution, strawman) +{ + // // TODO are there other cases in the notes? + // //------------------------------------------------------ + // // forall cases + // //------------------------------------------------------ + + // //------------------------------------------------------ + // // run on device + // //------------------------------------------------------ + // if (ExecutionPolicy::is_device_enabled()) + // { + // Node node; + // std::vector data_src = {0, 1, 2, 3}; + // node["src"].set(data_src); + // std::vector data_des = {0, 0, 0, 0}; + // node["src"].set(data_des); + // ExecutionAccessor acc_src(node["src"]); + // ExecutionAccessor acc_des(node["des"]); + + // ExecutionPolicy policy = ExecutionPolicy::device(); + + // acc_src.use_with(policy); + // acc_des.use_with(policy); + + // index_t size = acc_src.number_of_elements(); + + // forall(policy, 0, size, [=] EXEC_LAMBDA(index_t idx) + // { + // const float64 val = 2.0 * acc_src[idx]; + // acc_des.set(idx,val); + // }); + // CONDUIT_DEVICE_ERROR_CHECK(); + + // // sync values to node["des"] + // // (no op if node["des"] was originally device memory) + // acc_des.sync(); + // } + + // //------------------------------------------------------ + // // run on device, + // // result stays on device and is owned by node["des"], + // // even if not on the device before hand + // //------------------------------------------------------ + // { + // Node node; + // ExecutionAccessor acc_src(node["src"]); + // ExecutionAccessor acc_des(node["des"]); + + // ExecutionPolicy policy = ExecutionPolicy::device(); + + // acc_src.use_with(policy); + // acc_des.use_with(policy); + + // index_t size = acc_src.number_of_elements(); + + // forall(policy, 0, size, [=] EXEC_LAMBDA(index_t idx) + // { + // const float64 val = 2.0 * acc_src[idx]; + // acc_des.set(idx,val); + // }); + // CONDUIT_DEVICE_ERROR_CHECK(); + + // // move results to be owned by node["des"] + // // (no op if node["des"] was originally device memory) + // acc_des.move(node["des"]); + // } + + // //------------------------------------------------------ + // // run where the src data is + // //------------------------------------------------------ + // { + // Node node; + // ExecutionAccessor acc_src(node["src"]); + // ExecutionAccessor acc_des(node["des"]); + + // ExecutionPolicy policy = acc_src.active_space().execution_policy(); + // acc_des.use_with(policy); + // acc_des.use_with(policy); + + // index_t size = acc_src.number_of_elements(); + + // forall(policy, 0, size, [=] EXEC_LAMBDA(index_t idx) + // { + // const float64 val = 2.0 * acc_src[idx]; + // acc_des.set(idx,val); + // }); + // CONDUIT_DEVICE_ERROR_CHECK(); + + // // sync values to node["des"], + // // (no op if node["des"] was originally in + // // same memory space as node["src"] ) + // acc_des.sync(node["des"]); + // } + + // //------------------------------------------------------ + // // more complex cases + // //------------------------------------------------------ + + // //------------------------------------------------------ + // // complex run on device + // // double lambda forwarding concrete template tag + // // for use in lambda + // // + // // ( requires c++ 20 b/c of templated lambda) + // //------------------------------------------------------ + // { + // Node node; + // ExecutionAccessor acc_src(node["src"]); + // ExecutionAccessor acc_des(node["des"]); + + // ExecutionPolicy policy = ExecutionPolicy::device(); + // acc_des.use_with(policy); + // acc_des.use_with(policy); + + // index_t size = acc_src.number_of_elements(); + + // index_t min_loc = -1; + // float64 min_val = 0; + + // dispatch(policy, [&] (Exec &exec) + // { + // float64 identity = std::numeric_limits::max(); + // using for_policy = typename Exec::for_policy; + // using reduce_policy = typename Exec::reduce_policy; + + // ReduceMinLoc reducer(identity,-1); + + // forall(0, size, [=] EXEC_LAMBDA (int i) + // { + // const float64 val = 2.0 * acc_src[idx]; + // reducer.minloc(val,i); + // acc_des.set(idx,val); + // }); + // CONDUIT_DEVICE_ERROR_CHECK(); + + // min_val = reducer.get(); + // min_loc = reducer.getLoc(); + // }); + + // // sync values to node["des"], + // // (no op if node["des"] was originally in + // // same memory space as node["src"] ) + // acc_des.sync(node["des"]); + // } + + // //------------------------------------------------------ + // // complex run on device using functor + // // (functor implementation) + // //------------------------------------------------------ + // struct ExecFunctor + // { + // float64 min_val; + // index_t min_loc; + + // ExecutionAccessor acc_src; + // ExecutionAccessor acc_des; + + // template + // void operator()(Exec &exec) + // { + // float64 identity = std::numeric_limits::max(); + // using for_policy = typename Exec::for_policy; + // using reduce_policy = typename Exec::reduce_policy; + + // ReduceMinLoc reducer(identity, -1); + + // forall(0, size, [=] (int i) + // { + // const float64 val = 2.0 * acc_src[idx]; + // reducer.minloc(val,i); + // acc_des.set(idx,val); + // }); + // CONDUIT_DEVICE_ERROR_CHECK(); + + // min_val = reducer.get(); + // min_loc = reducer.getLoc(); + // } + // }; + + // //------------------------------------------------------ + // // complex run on device using functor + // // (functor dispatch) + // //------------------------------------------------------ + // { + // Node node; + // ExecutionAccessor acc_src(node["src"]); + // ExecutionAccessor acc_des(node["des"]); + + // ExecutionPolicy policy = ExecutionPolicy::device(); + // acc_des.use_with(policy); + // acc_des.use_with(policy); + + // index_t size = acc_src.number_of_elements(); + + // ExecFunctor f(); + + // // init functor + // f.acc_src = acc_src; + // f.acc_des = acc_des; + + // dispatch(policy,f); + + // // get results stored in functor + // float64 min_val = f.min_val; + // index_t min_loc = f.min_loc; + + // // sync values to node["des"], + // // (no op if node["des"] was originally in + // // same memory space as node["src"]) + // acc_des.sync(node["des"]); + // } +} + +