From 04bf3bfb7f454d521b20b02c2ea6a4f6fa396cf0 Mon Sep 17 00:00:00 2001 From: maelk3 Date: Fri, 3 Jan 2025 01:58:42 +0100 Subject: [PATCH] Applying t8code's v4 breaking API changes and modifying subgrid example grid size --- examples/compressible_euler/main.cu | 2 +- examples/compressible_euler/solver.cu | 8 +- examples/compressible_euler/solver.h | 2 +- examples/subgrid/kernels_2d.cu | 2 +- examples/subgrid/main_2d.cu | 13 +-- examples/subgrid/main_3d.cu | 9 +- examples/subgrid/solver.h | 2 +- examples/subgrid/solver.inl | 14 +-- examples/subgrid/solver_2d.cu | 2 +- t8gpu/mesh/mesh_manager.h | 31 +++---- t8gpu/mesh/mesh_manager.inl | 75 ++++++++-------- t8gpu/mesh/subgrid_mesh_manager.h | 31 +++---- t8gpu/mesh/subgrid_mesh_manager.inl | 124 ++++++++++++++------------ 13 files changed, 164 insertions(+), 151 deletions(-) diff --git a/examples/compressible_euler/main.cu b/examples/compressible_euler/main.cu index 5c1c2e5..d89292a 100644 --- a/examples/compressible_euler/main.cu +++ b/examples/compressible_euler/main.cu @@ -19,7 +19,7 @@ int main(int argc, char* argv[]) { { size_t dim = 3; sc_MPI_Comm comm = MPI_COMM_WORLD; - t8_scheme_cxx_t* scheme = t8_scheme_new_default_cxx(); + t8_scheme* scheme = t8_scheme_new_default(); t8_cmesh_t cmesh = t8_cmesh_new_prismed_spherical_shell_icosahedron(0.8, 0.2, 2, 1, comm); t8_forest_t forest = t8_forest_new_uniform(cmesh, scheme, 2, true, comm); diff --git a/examples/compressible_euler/solver.cu b/examples/compressible_euler/solver.cu index ca3cbce..0c90646 100644 --- a/examples/compressible_euler/solver.cu +++ b/examples/compressible_euler/solver.cu @@ -6,10 +6,10 @@ using namespace t8gpu; -CompressibleEulerSolver::CompressibleEulerSolver(sc_MPI_Comm comm, - t8_scheme_cxx_t* scheme, - t8_cmesh_t cmesh, - t8_forest_t forest) +CompressibleEulerSolver::CompressibleEulerSolver(sc_MPI_Comm comm, + t8_scheme* scheme, + t8_cmesh_t cmesh, + t8_forest_t forest) : m_comm{comm}, m_mesh_manager{comm, scheme, cmesh, forest}, m_device_face_speed_estimate(m_mesh_manager.get_num_local_faces() + diff --git a/examples/compressible_euler/solver.h b/examples/compressible_euler/solver.h index 6a0fc85..47c47be 100644 --- a/examples/compressible_euler/solver.h +++ b/examples/compressible_euler/solver.h @@ -41,7 +41,7 @@ namespace t8gpu { /// This class takes ownership of the cmesh and forest given and /// the constructor computes the necessary face information and /// allocates space to store all of the variables. - CompressibleEulerSolver(sc_MPI_Comm comm, t8_scheme_cxx_t* scheme, t8_cmesh_t cmesh, t8_forest_t forest); + CompressibleEulerSolver(sc_MPI_Comm comm, t8_scheme* scheme, t8_cmesh_t cmesh, t8_forest_t forest); /// @brief Destructor. /// diff --git a/examples/subgrid/kernels_2d.cu b/examples/subgrid/kernels_2d.cu index 626f48d..e5a608a 100644 --- a/examples/subgrid/kernels_2d.cu +++ b/examples/subgrid/kernels_2d.cu @@ -1,7 +1,7 @@ #define KERNELS_INCLUDE_IMPLEMENTATION #include "kernels.h" -using SubgridType = t8gpu::Subgrid<4, 4>; +using SubgridType = t8gpu::Subgrid<16, 16>; // We explicitely instantiate the kernel functions for the 2D subgrid // in a separate compilation unit for better compilation speed diff --git a/examples/subgrid/main_2d.cu b/examples/subgrid/main_2d.cu index 85538e2..bf3dd1f 100644 --- a/examples/subgrid/main_2d.cu +++ b/examples/subgrid/main_2d.cu @@ -1,3 +1,4 @@ + #include #include @@ -21,18 +22,18 @@ int main(int argc, char* argv[]) { // We create a dummy frame so that destructors of solver class are // called before MPI finalize. { - using SubgridType = Subgrid<4, 4>; + using SubgridType = Subgrid<16, 16>; using float_type = typename SubgridCompressibleEulerSolver::float_type; float_type delta_t = - static_cast(0.1 * pow(0.5, + static_cast(0.05 * pow(0.5, SubgridMeshManager::max_level + t8gpu::meta::log2_v>)); - sc_MPI_Comm comm = MPI_COMM_WORLD; - t8_scheme_cxx_t* scheme = t8_scheme_new_default_cxx(); - t8_cmesh_t cmesh = t8_cmesh_new_periodic(comm, SubgridType::rank); - t8_forest_t forest = t8_forest_new_uniform(cmesh, scheme, 4, true, comm); + sc_MPI_Comm comm = MPI_COMM_WORLD; + t8_scheme* scheme = t8_scheme_new_default(); + t8_cmesh_t cmesh = t8_cmesh_new_periodic(comm, SubgridType::rank); + t8_forest_t forest = t8_forest_new_uniform(cmesh, scheme, 8, true, comm); SubgridCompressibleEulerSolver solver{comm, scheme, cmesh, forest}; diff --git a/examples/subgrid/main_3d.cu b/examples/subgrid/main_3d.cu index 99c52fa..de02270 100644 --- a/examples/subgrid/main_3d.cu +++ b/examples/subgrid/main_3d.cu @@ -29,11 +29,10 @@ int main(int argc, char* argv[]) { SubgridMeshManager::max_level + t8gpu::meta::log2_v>)); - sc_MPI_Comm comm = MPI_COMM_WORLD; - t8_scheme_cxx_t* scheme = t8_scheme_new_default_cxx(); - t8_cmesh_t cmesh = t8_cmesh_new_periodic(comm, SubgridType::rank); - t8_forest_t forest = t8_forest_new_uniform(cmesh, scheme, 4, true, comm); - // t8_forest_t forest = t8_forest_new_uniform(cmesh, scheme, 2, true, comm); + sc_MPI_Comm comm = MPI_COMM_WORLD; + t8_scheme* scheme = t8_scheme_new_default(); + t8_cmesh_t cmesh = t8_cmesh_new_periodic(comm, SubgridType::rank); + t8_forest_t forest = t8_forest_new_uniform(cmesh, scheme, 4, true, comm); SubgridCompressibleEulerSolver solver{comm, scheme, cmesh, forest}; diff --git a/examples/subgrid/solver.h b/examples/subgrid/solver.h index 7babc46..b69a7ab 100644 --- a/examples/subgrid/solver.h +++ b/examples/subgrid/solver.h @@ -41,7 +41,7 @@ class SubgridCompressibleEulerSolver { /// This class takes ownership of the cmesh and forest given and /// the constructor computes the necessary face information and /// allocates space to store all of the variables. - SubgridCompressibleEulerSolver(sc_MPI_Comm comm, t8_scheme_cxx_t* scheme, t8_cmesh_t cmesh, t8_forest_t forest); + SubgridCompressibleEulerSolver(sc_MPI_Comm comm, t8_scheme* scheme, t8_cmesh_t cmesh, t8_forest_t forest); /// @brief Destructor. /// diff --git a/examples/subgrid/solver.inl b/examples/subgrid/solver.inl index fbd03e7..50fb2f1 100644 --- a/examples/subgrid/solver.inl +++ b/examples/subgrid/solver.inl @@ -104,10 +104,10 @@ __global__ std::enable_if_t initialize_variables( } template -SubgridCompressibleEulerSolver::SubgridCompressibleEulerSolver(sc_MPI_Comm comm, - t8_scheme_cxx_t* scheme, - t8_cmesh_t cmesh, - t8_forest_t forest) +SubgridCompressibleEulerSolver::SubgridCompressibleEulerSolver(sc_MPI_Comm comm, + t8_scheme* scheme, + t8_cmesh_t cmesh, + t8_forest_t forest) : m_comm{comm}, m_mesh_manager{comm, scheme, cmesh, forest}, m_device_face_speed_estimate(m_mesh_manager.get_num_local_faces() + @@ -119,8 +119,8 @@ SubgridCompressibleEulerSolver::SubgridCompressibleEulerSolver(sc_M t8_locidx_t element_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_local_trees; tree_idx++) { - t8_eclass_t tree_class{t8_forest_get_tree_class(forest, tree_idx)}; - t8_eclass_scheme_c* eclass_scheme{t8_forest_get_eclass_scheme(forest, tree_class)}; + t8_eclass_t tree_class{t8_forest_get_tree_class(forest, tree_idx)}; + t8_scheme* scheme{t8_forest_get_scheme(forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(forest, tree_idx)}; for (t8_locidx_t tree_element_idx = 0; tree_element_idx < num_elements_in_tree; tree_element_idx++) { @@ -132,7 +132,7 @@ SubgridCompressibleEulerSolver::SubgridCompressibleEulerSolver(sc_M centers[3 * element_idx + 1] = static_cast(center[1]); centers[3 * element_idx + 2] = static_cast(center[2]); - levels[element_idx] = eclass_scheme->t8_element_level(element); + levels[element_idx] = scheme->element_get_level(tree_class, element); element_idx++; } diff --git a/examples/subgrid/solver_2d.cu b/examples/subgrid/solver_2d.cu index 6ab584e..533fe18 100644 --- a/examples/subgrid/solver_2d.cu +++ b/examples/subgrid/solver_2d.cu @@ -1,7 +1,7 @@ #define SOLVER_INCLUDE_IMPLEMENTATION #include "solver.h" -using SubgridType = t8gpu::Subgrid<4, 4>; +using SubgridType = t8gpu::Subgrid<16, 16>; // We explicitely instantiate the solver type for the 2D subgrid in a // separate compilation unit for better compilation speed (instead of diff --git a/t8gpu/mesh/mesh_manager.h b/t8gpu/mesh/mesh_manager.h index 05faa07..8b6d8f0 100644 --- a/t8gpu/mesh/mesh_manager.h +++ b/t8gpu/mesh/mesh_manager.h @@ -250,7 +250,7 @@ namespace t8gpu { /// /// This constructor initializes a MeshManager class taking /// ownership of the coarse mesh cmesh and forest. - MeshManager(sc_MPI_Comm comm, t8_scheme_cxx_t* scheme, t8_cmesh_t cmesh, t8_forest_t forest); + MeshManager(sc_MPI_Comm comm, t8_scheme* scheme, t8_cmesh_t cmesh, t8_forest_t forest); /// @brief Destructor of the MeshManager class. /// @@ -424,12 +424,12 @@ namespace t8gpu { /// We make the base class resize member function private. using MemoryManager::resize; - sc_MPI_Comm m_comm; - int m_rank; - int m_nb_ranks; - t8_scheme_cxx_t* m_scheme; - t8_cmesh_t m_cmesh; - t8_forest_t m_forest; + sc_MPI_Comm m_comm; + int m_rank; + int m_nb_ranks; + t8_scheme* m_scheme; + t8_cmesh_t m_cmesh; + t8_forest_t m_forest; t8_locidx_t m_num_local_elements; t8_locidx_t m_num_ghost_elements; @@ -454,14 +454,15 @@ namespace t8gpu { thrust::host_vector* element_refinement_criteria; }; - static int adapt_callback_iteration(t8_forest_t forest, - t8_forest_t forest_from, - t8_locidx_t which_tree, - t8_locidx_t lelement_id, - t8_eclass_scheme_c* ts, - int const is_family, - int const num_elements, - t8_element_t* elements[]); + static int adapt_callback_iteration(t8_forest_t forest, + t8_forest_t forest_from, + t8_locidx_t which_tree, + t8_eclass_t const tree_class, + t8_locidx_t lelement_id, + t8_scheme_c const* ts, + int const is_family, + int const num_elements, + t8_element_t* elements[]); }; } // namespace t8gpu diff --git a/t8gpu/mesh/mesh_manager.inl b/t8gpu/mesh/mesh_manager.inl index 1850944..bf0c21e 100644 --- a/t8gpu/mesh/mesh_manager.inl +++ b/t8gpu/mesh/mesh_manager.inl @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include @@ -18,10 +18,10 @@ #include template -t8gpu::MeshManager::MeshManager(sc_MPI_Comm comm, - t8_scheme_cxx_t* scheme, - t8_cmesh_t cmesh, - t8_forest_t forest) +t8gpu::MeshManager::MeshManager(sc_MPI_Comm comm, + t8_scheme* scheme, + t8_cmesh_t cmesh, + t8_forest_t forest) : m_comm{comm}, m_scheme{scheme}, m_cmesh{cmesh}, @@ -91,8 +91,7 @@ void t8gpu::MeshManager::initialize_variables(Func t8_locidx_t num_local_trees = t8_forest_get_num_local_trees(m_forest); t8_locidx_t element_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_local_trees; tree_idx++) { - t8_eclass_t tree_class{t8_forest_get_tree_class(m_forest, tree_idx)}; - t8_eclass_scheme_c* eclass_scheme{t8_forest_get_eclass_scheme(m_forest, tree_class)}; + t8_scheme* eclass_scheme{t8_forest_get_scheme(m_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t tree_element_idx = 0; tree_element_idx < num_elements_in_tree; tree_element_idx++) { @@ -122,19 +121,20 @@ void t8gpu::MeshManager::initialize_variables(Func } template -int t8gpu::MeshManager::adapt_callback_iteration(t8_forest_t forest, - t8_forest_t forest_from, - t8_locidx_t which_tree, - t8_locidx_t lelement_id, - t8_eclass_scheme_c* ts, - int const is_family, - int const num_elements, - t8_element_t* elements[]) { +int t8gpu::MeshManager::adapt_callback_iteration(t8_forest_t forest, + t8_forest_t forest_from, + t8_locidx_t which_tree, + t8_eclass_t const tree_class, + t8_locidx_t lelement_id, + t8_scheme_c const* ts, + int const is_family, + int const num_elements, + t8_element_t* elements[]) { t8gpu::MeshManager::UserData* forest_user_data = static_cast::UserData*>(t8_forest_get_user_data(forest_from)); assert(forest_user_data != nullptr); - t8_locidx_t element_level{ts->t8_element_level(elements[0])}; + t8_locidx_t element_level{ts->element_get_level(tree_class, elements[0])}; t8_locidx_t tree_offset = t8_forest_get_tree_element_offset(forest_from, which_tree); @@ -229,28 +229,29 @@ void t8gpu::MeshManager::adapt( t8_locidx_t current_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_old_local_trees; tree_idx++) { - t8_eclass_t old_tree_class{t8_forest_get_tree_class(m_forest, tree_idx)}; - t8_eclass_scheme_c* old_scheme = {t8_forest_get_eclass_scheme(m_forest, old_tree_class)}; + t8_scheme* old_scheme = {t8_forest_get_scheme(m_forest)}; + t8_eclass_t old_tree_class = t8_forest_get_tree_class(m_forest, tree_idx); t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t elem_idx = 0; elem_idx < num_elements_in_tree; elem_idx++) { t8_element_t const* element{t8_forest_get_element_in_tree(m_forest, tree_idx, elem_idx)}; - old_levels[current_idx] = old_scheme->t8_element_level(element); + t8_eclass_t tree_class = t8_forest_get_tree_class(m_forest, tree_idx); + old_levels[current_idx] = old_scheme->element_get_level(old_tree_class, element); current_idx++; } } current_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_new_local_trees; tree_idx++) { - t8_eclass_t new_tree_class{t8_forest_get_tree_class(adapted_forest, tree_idx)}; - t8_eclass_scheme_c* new_scheme = {t8_forest_get_eclass_scheme(adapted_forest, new_tree_class)}; + t8_scheme* new_scheme = {t8_forest_get_scheme(adapted_forest)}; + t8_eclass_t new_tree_class = t8_forest_get_tree_class(adapted_forest, tree_idx); t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(adapted_forest, tree_idx)}; for (t8_locidx_t elem_idx = 0; elem_idx < num_elements_in_tree; elem_idx++) { t8_element_t const* element{t8_forest_get_element_in_tree(adapted_forest, tree_idx, elem_idx)}; - new_levels[current_idx] = new_scheme->t8_element_level(element); + new_levels[current_idx] = new_scheme->element_get_level(new_tree_class, element); current_idx++; } } @@ -368,20 +369,20 @@ void t8gpu::MeshManager::compute_connectivity_infor t8_locidx_t num_local_trees{t8_forest_get_num_local_trees(m_forest)}; t8_locidx_t element_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_local_trees; tree_idx++) { - t8_eclass_t tree_class = t8_forest_get_tree_class(m_forest, tree_idx); - t8_eclass_scheme_c* eclass_scheme{t8_forest_get_eclass_scheme(m_forest, tree_class)}; + t8_eclass_t tree_class = t8_forest_get_tree_class(m_forest, tree_idx); + t8_scheme* scheme{t8_forest_get_scheme(m_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t tree_element_idx = 0; tree_element_idx < num_elements_in_tree; tree_element_idx++) { t8_element_t const* element{t8_forest_get_element_in_tree(m_forest, tree_idx, tree_element_idx)}; - t8_locidx_t num_faces{eclass_scheme->t8_element_num_faces(element)}; + t8_locidx_t num_faces{scheme->element_get_num_faces(tree_class, element)}; for (t8_locidx_t face_idx = 0; face_idx < num_faces; face_idx++) { - int num_neighbors{}; - int* dual_faces{}; - t8_locidx_t* neighbor_ids{}; - t8_element_t** neighbors{}; - t8_eclass_scheme_c* neigh_scheme{}; + int num_neighbors{}; + int* dual_faces{}; + t8_locidx_t* neighbor_ids{}; + t8_element_t** neighbors{}; + t8_eclass_t neigh_eclass{}; t8_forest_leaf_face_neighbors(m_forest, tree_idx, element, @@ -390,7 +391,7 @@ void t8gpu::MeshManager::compute_connectivity_infor &dual_faces, &num_neighbors, &neighbor_ids, - &neigh_scheme, + &neigh_eclass, true); for (int i = 0; i < num_neighbors; i++) { @@ -411,7 +412,7 @@ void t8gpu::MeshManager::compute_connectivity_infor if ((num_neighbors == 1) && (neighbor_ids[0] < m_num_local_elements) && ((neighbor_ids[0] > element_idx) || (neighbor_ids[0] < element_idx && - neigh_scheme[0].t8_element_level(neighbors[0]) < eclass_scheme->t8_element_level(element)))) { + scheme->element_get_level(neigh_eclass, neighbors[0]) < scheme->element_get_level(tree_class, element)))) { face_neighbors.push_back(element_idx); face_neighbors.push_back(neighbor_ids[0]); double face_normal[dim]; @@ -422,11 +423,13 @@ void t8gpu::MeshManager::compute_connectivity_infor face_area.push_back( static_cast(t8_forest_element_face_area(m_forest, tree_idx, element, face_idx))); } - neigh_scheme->t8_element_destroy(num_neighbors, neighbors); - T8_FREE(neighbors); - T8_FREE(dual_faces); - T8_FREE(neighbor_ids); + if (num_neighbors > 0) { + scheme->element_destroy(neigh_eclass, num_neighbors, neighbors); + T8_FREE(neighbors); + T8_FREE(neighbor_ids); + T8_FREE(dual_faces); + } if (num_neighbors == 0) { boundary_face_neighbors.push_back(element_idx); diff --git a/t8gpu/mesh/subgrid_mesh_manager.h b/t8gpu/mesh/subgrid_mesh_manager.h index 36a2d57..8751215 100644 --- a/t8gpu/mesh/subgrid_mesh_manager.h +++ b/t8gpu/mesh/subgrid_mesh_manager.h @@ -285,7 +285,7 @@ namespace t8gpu { /// /// This constructor initializes a MeshManager class taking /// ownership of the coarse mesh cmesh and forest. - SubgridMeshManager(sc_MPI_Comm comm, t8_scheme_cxx_t* scheme, t8_cmesh_t cmesh, t8_forest_t forest); + SubgridMeshManager(sc_MPI_Comm comm, t8_scheme* scheme, t8_cmesh_t cmesh, t8_forest_t forest); /// @brief Destructor of the MeshManager class. /// @@ -465,12 +465,12 @@ namespace t8gpu { /// We make the base class resize member function private. using SubgridMemoryManager::resize; - sc_MPI_Comm m_comm; - int m_rank; - int m_nb_ranks; - t8_scheme_cxx_t* m_scheme; - t8_cmesh_t m_cmesh; - t8_forest_t m_forest; + sc_MPI_Comm m_comm; + int m_rank; + int m_nb_ranks; + t8_scheme* m_scheme; + t8_cmesh_t m_cmesh; + t8_forest_t m_forest; t8_locidx_t m_num_local_elements; t8_locidx_t m_num_ghost_elements; @@ -498,14 +498,15 @@ namespace t8gpu { thrust::host_vector* element_refinement_criteria; }; - static int adapt_callback_iteration(t8_forest_t forest, - t8_forest_t forest_from, - t8_locidx_t which_tree, - t8_locidx_t lelement_id, - t8_eclass_scheme_c* ts, - int const is_family, - int const num_elements, - t8_element_t* elements[]); + static int adapt_callback_iteration(t8_forest_t forest, + t8_forest_t forest_from, + t8_locidx_t which_tree, + t8_eclass_t const tree_class, + t8_locidx_t lelement_id, + t8_scheme_c const* ts, + int const is_family, + int const num_elements, + t8_element_t* elements[]); }; } // namespace t8gpu diff --git a/t8gpu/mesh/subgrid_mesh_manager.inl b/t8gpu/mesh/subgrid_mesh_manager.inl index c0f1d7d..dc6e789 100644 --- a/t8gpu/mesh/subgrid_mesh_manager.inl +++ b/t8gpu/mesh/subgrid_mesh_manager.inl @@ -11,7 +11,7 @@ #include #include #include -#include +#include #include #include #include @@ -19,10 +19,10 @@ #include template -t8gpu::SubgridMeshManager::SubgridMeshManager(sc_MPI_Comm comm, - t8_scheme_cxx_t* scheme, - t8_cmesh_t cmesh, - t8_forest_t forest) +t8gpu::SubgridMeshManager::SubgridMeshManager(sc_MPI_Comm comm, + t8_scheme* scheme, + t8_cmesh_t cmesh, + t8_forest_t forest) : m_comm{comm}, m_scheme{scheme}, m_cmesh{cmesh}, @@ -71,7 +71,7 @@ t8gpu::SubgridMeshManager::SubgridMeshManag t8_locidx_t element_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_local_trees; tree_idx++) { t8_eclass_t tree_class{t8_forest_get_tree_class(m_forest, tree_idx)}; - t8_eclass_scheme_c* eclass_scheme{t8_forest_get_eclass_scheme(m_forest, tree_class)}; + t8_scheme* eclass_scheme{t8_forest_get_scheme(m_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t tree_element_idx = 0; tree_element_idx < num_elements_in_tree; tree_element_idx++) { @@ -159,7 +159,7 @@ void t8gpu::SubgridMeshManager::initialize_ t8_locidx_t element_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_local_trees; tree_idx++) { t8_eclass_t tree_class{t8_forest_get_tree_class(m_forest, tree_idx)}; - t8_eclass_scheme_c* eclass_scheme{t8_forest_get_eclass_scheme(m_forest, tree_class)}; + t8_scheme* eclass_scheme{t8_forest_get_scheme(m_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t tree_element_idx = 0; tree_element_idx < num_elements_in_tree; tree_element_idx++) { @@ -194,20 +194,21 @@ void t8gpu::SubgridMeshManager::initialize_ } template -int t8gpu::SubgridMeshManager::adapt_callback_iteration(t8_forest_t forest, - t8_forest_t forest_from, - t8_locidx_t which_tree, - t8_locidx_t lelement_id, - t8_eclass_scheme_c* ts, - int const is_family, - int const num_elements, - t8_element_t* elements[]) { +int t8gpu::SubgridMeshManager::adapt_callback_iteration(t8_forest_t forest, + t8_forest_t forest_from, + t8_locidx_t which_tree, + t8_eclass_t const tree_class, + t8_locidx_t lelement_id, + t8_scheme_c const* ts, + int const is_family, + int const num_elements, + t8_element_t* elements[]) { t8gpu::SubgridMeshManager::UserData* forest_user_data = static_cast::UserData*>( t8_forest_get_user_data(forest_from)); assert(forest_user_data != nullptr); - t8_locidx_t element_level{ts->t8_element_level(elements[0])}; + t8_locidx_t element_level{ts->element_get_level(tree_class, elements[0])}; t8_locidx_t tree_offset = t8_forest_get_tree_element_offset(forest_from, which_tree); @@ -459,27 +460,27 @@ void t8gpu::SubgridMeshManager::adapt( t8_locidx_t current_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_old_local_trees; tree_idx++) { t8_eclass_t old_tree_class{t8_forest_get_tree_class(m_forest, tree_idx)}; - t8_eclass_scheme_c* old_scheme = {t8_forest_get_eclass_scheme(m_forest, old_tree_class)}; + t8_scheme* old_scheme = {t8_forest_get_scheme(m_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t elem_idx = 0; elem_idx < num_elements_in_tree; elem_idx++) { t8_element_t const* element{t8_forest_get_element_in_tree(m_forest, tree_idx, elem_idx)}; - old_levels[current_idx] = old_scheme->t8_element_level(element); + old_levels[current_idx] = old_scheme->element_get_level(old_tree_class, element); current_idx++; } } current_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_new_local_trees; tree_idx++) { - t8_eclass_t new_tree_class{t8_forest_get_tree_class(adapted_forest, tree_idx)}; - t8_eclass_scheme_c* new_scheme = {t8_forest_get_eclass_scheme(adapted_forest, new_tree_class)}; + t8_eclass_t new_tree_class{t8_forest_get_tree_class(adapted_forest, tree_idx)}; + t8_scheme* new_scheme = {t8_forest_get_scheme(adapted_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(adapted_forest, tree_idx)}; for (t8_locidx_t elem_idx = 0; elem_idx < num_elements_in_tree; elem_idx++) { t8_element_t const* element{t8_forest_get_element_in_tree(adapted_forest, tree_idx, elem_idx)}; - new_levels[current_idx] = new_scheme->t8_element_level(element); + new_levels[current_idx] = new_scheme->element_get_level(new_tree_class, element); current_idx++; } } @@ -567,8 +568,9 @@ std::enable_if_t add_face( t8_element_t const* element, t8_element_t const* neighbor_element, t8_locidx_t neighbor_idx, - t8_eclass_scheme_c* scheme_element, - t8_eclass_scheme_c* scheme_neighbor, + t8_scheme* scheme, + t8_eclass_t eclass_element, + t8_eclass_t eclass_neighbor, thrust::host_vector& face_level_difference, thrust::host_vector& face_neighbor_offset, thrust::host_vector& face_neighbors, @@ -576,8 +578,8 @@ std::enable_if_t add_face( thrust::host_vector::float_type>& face_area) { using float_type = typename t8gpu::variable_traits::float_type; - int level = scheme_element->t8_element_level(element); - int neighbor_level = scheme_neighbor->t8_element_level(neighbor_element); + int level = scheme->element_get_level(eclass_element, element); + int neighbor_level = scheme->element_get_level(eclass_neighbor, neighbor_element); double face_normal[3]; t8_forest_element_face_normal(forest, tree_idx, element, face_idx, face_normal); @@ -613,7 +615,7 @@ std::enable_if_t add_face( face_area.push_back(static_cast(t8_forest_element_face_area(forest, tree_idx, element, face_idx)) / static_cast(num_neighbors)); } else if (neighbor_level < level) { - int child_id = scheme_element->t8_element_child_id(element); + int child_id = scheme->element_get_child_id(eclass_element, element); std::array neighbor_offset{}; @@ -646,7 +648,7 @@ std::enable_if_t add_face( static_cast(num_neighbors)); } else { - int neighbor_child_id = scheme_neighbor->t8_element_child_id(neighbor_element); + int neighbor_child_id = scheme->element_get_child_id(eclass_neighbor, neighbor_element); std::array neighbor_offset{}; @@ -690,8 +692,9 @@ std::enable_if_t add_face( t8_element_t const* element, t8_element_t const* neighbor_element, t8_locidx_t neighbor_idx, - t8_eclass_scheme_c* scheme_element, - t8_eclass_scheme_c* scheme_neighbor, + t8_scheme* scheme, + t8_eclass_t eclass_element, + t8_eclass_t eclass_neighbor, thrust::host_vector& face_level_difference, thrust::host_vector& face_neighbor_offset, thrust::host_vector& face_neighbors, @@ -699,8 +702,8 @@ std::enable_if_t add_face( thrust::host_vector::float_type>& face_area) { using float_type = typename t8gpu::variable_traits::float_type; - int level = scheme_element->t8_element_level(element); - int neighbor_level = scheme_neighbor->t8_element_level(neighbor_element); + int level = scheme->element_get_level(eclass_element, element); + int neighbor_level = scheme->element_get_level(eclass_neighbor, neighbor_element); double face_normal[3]; t8_forest_element_face_normal(forest, tree_idx, element, face_idx, face_normal); @@ -730,7 +733,7 @@ std::enable_if_t add_face( face_area.push_back(static_cast(t8_forest_element_face_area(forest, tree_idx, element, face_idx)) / static_cast(num_neighbors)); } else if (neighbor_level < level) { - int child_id = scheme_element->t8_element_child_id(element); + int child_id = scheme->element_get_child_id(eclass_element, element); std::array neighbor_offset{}; @@ -757,7 +760,7 @@ std::enable_if_t add_face( static_cast(num_neighbors)); } else { - int neighbor_child_id = scheme_neighbor->t8_element_child_id(neighbor_element); + int neighbor_child_id = scheme->element_get_child_id(eclass_neighbor, neighbor_element); std::array neighbor_offset{}; @@ -832,19 +835,19 @@ void t8gpu::SubgridMeshManager::compute_con t8_locidx_t element_idx = 0; for (t8_locidx_t tree_idx = 0; tree_idx < num_local_trees; tree_idx++) { t8_eclass_t tree_class = t8_forest_get_tree_class(m_forest, tree_idx); - t8_eclass_scheme_c* eclass_scheme{t8_forest_get_eclass_scheme(m_forest, tree_class)}; + t8_scheme* scheme{t8_forest_get_scheme(m_forest)}; t8_locidx_t num_elements_in_tree{t8_forest_get_tree_num_elements(m_forest, tree_idx)}; for (t8_locidx_t tree_element_idx = 0; tree_element_idx < num_elements_in_tree; tree_element_idx++) { t8_element_t const* element{t8_forest_get_element_in_tree(m_forest, tree_idx, tree_element_idx)}; - t8_locidx_t num_faces{eclass_scheme->t8_element_num_faces(element)}; + t8_locidx_t num_faces{scheme->element_get_num_faces(tree_class, element)}; for (t8_locidx_t face_idx = 0; face_idx < num_faces; face_idx++) { - int num_neighbors{}; - int* dual_faces{}; - t8_locidx_t* neighbor_ids{}; - t8_element_t** neighbors{}; - t8_eclass_scheme_c* neigh_scheme{}; + int num_neighbors{}; + int* dual_faces{}; + t8_locidx_t* neighbor_ids{}; + t8_element_t** neighbors{}; + t8_eclass_t neigh_eclass{}; t8_forest_leaf_face_neighbors(m_forest, tree_idx, element, @@ -853,7 +856,7 @@ void t8gpu::SubgridMeshManager::compute_con &dual_faces, &num_neighbors, &neighbor_ids, - &neigh_scheme, + &neigh_eclass, true); for (int i = 0; i < num_neighbors; i++) { // we treat the case when the neighboring element is a ghost element. @@ -866,8 +869,9 @@ void t8gpu::SubgridMeshManager::compute_con element, neighbors[i], neighbor_ids[i], - eclass_scheme, - neigh_scheme, + scheme, + neigh_eclass, + tree_class, face_level_difference, face_neighbor_offset, face_neighbors, @@ -880,7 +884,7 @@ void t8gpu::SubgridMeshManager::compute_con if ((num_neighbors == 1) && (neighbor_ids[0] < m_num_local_elements) && ((neighbor_ids[0] > element_idx) || (neighbor_ids[0] < element_idx && - neigh_scheme[0].t8_element_level(neighbors[0]) < eclass_scheme->t8_element_level(element)))) { + scheme->element_get_level(neigh_eclass, neighbors[0]) < scheme->element_get_level(tree_class, element)))) { add_face(face_idx, num_neighbors, m_forest, @@ -889,19 +893,22 @@ void t8gpu::SubgridMeshManager::compute_con element, neighbors[0], neighbor_ids[0], - eclass_scheme, - neigh_scheme, + scheme, + neigh_eclass, + tree_class, face_level_difference, face_neighbor_offset, face_neighbors, face_normals, face_area); } - neigh_scheme->t8_element_destroy(num_neighbors, neighbors); - T8_FREE(neighbors); - T8_FREE(dual_faces); - T8_FREE(neighbor_ids); + if (num_neighbors > 0) { + scheme->element_destroy(neigh_eclass, num_neighbors, neighbors); + T8_FREE(neighbors); + T8_FREE(neighbor_ids); + T8_FREE(dual_faces); + } if (num_neighbors == 0) { boundary_face_neighbors.push_back(element_idx); @@ -1064,14 +1071,15 @@ void t8gpu::SubgridMeshManager::save_variab element_variable = device_element_variable; - auto uniform_adaptation = [](t8_forest_t forest, - t8_forest_t forest_from, - t8_locidx_t which_tree, - t8_locidx_t lelement_id, - t8_eclass_scheme_c* ts, - int const is_family, - int const num_elements, - t8_element_t* elements[]) -> int { return 1; }; + auto uniform_adaptation = [](t8_forest_t forest, + t8_forest_t forest_from, + t8_locidx_t which_tree, + t8_eclass_t tree_class, + t8_locidx_t lelement_id, + t8_scheme_c const* ts, + int const is_family, + int const num_elements, + t8_element_t* elements[]) -> int { return 1; }; constexpr int number_adaptation = t8gpu::meta::log2_v>; std::array> + 1> subgrid_forests;