From 02e0cd7b4fdbfbfed0b6d927c380f1bc34575480 Mon Sep 17 00:00:00 2001 From: joshia5 Date: Fri, 5 Nov 2021 23:50:47 -0400 Subject: [PATCH 01/40] builds with egadslite and cuda and meshb Conflicts: src/CMakeLists.txt --- CMakeLists.txt | 2 + src/CMakeLists.txt | 34 +++++ src/Omega_h_egads_lite.cpp | 253 +++++++++++++++++++++++++++++++++++++ src/Omega_h_egads_lite.hpp | 21 +++ src/egads_adapt_test.cpp | 101 +++++++++++++++ 5 files changed, 411 insertions(+) create mode 100644 src/Omega_h_egads_lite.cpp create mode 100644 src/Omega_h_egads_lite.hpp create mode 100644 src/egads_adapt_test.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 94b7287d6..9c172dd7c 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -21,6 +21,7 @@ bob_option(Omega_h_CHECK_BOUNDS "Check array bounds (makes code slow too)" OFF) bob_option(Omega_h_THROW "Errors throw exceptions instead of abort" ${USE_XSDK_DEFAULTS}) bob_input(Omega_h_DATA "" PATH "Path to omega_h-data test files") bob_option(Omega_h_USE_EGADS "Use EGADS from ESP for geometry" OFF) +bob_option(Omega_h_USE_EGADSlite "Use EGADSlite from ESP for geometry" OFF) bob_input(EGADS_PREFIX "" PATH "EGADS (or ESP) installation directory") bob_option(Omega_h_USE_Kokkos "Use Kokkos as a backend" OFF) bob_input(Kokkos_PREFIX "" PATH "Path to Kokkos install") @@ -117,6 +118,7 @@ set(Omega_h_KEY_BOOLS Omega_h_USE_ZLIB Omega_h_USE_libMeshb Omega_h_USE_EGADS + Omega_h_USE_EGADSlite Omega_h_USE_SEACASExodus Omega_h_USE_DOLFIN Omega_h_USE_dwarf diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 87e996d78..b6093a36a 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -152,6 +152,32 @@ if(Omega_h_USE_EGADS) set(Omega_h_SOURCES ${Omega_h_SOURCES} Omega_h_egads.cpp) endif() +if(Omega_h_USE_EGADSlite) + message(STATUS "Omega_h_USE_EGADSlite: ${Omega_h_USE_EGADSlite}") + if(NOT EGADSlite_PREFIX) + message(FATAL_ERROR "Omega_h_USE_EGADSlite=${Omega_h_USE_EGADSlite} but EGADSlite_PREFIX not set!") + endif() + find_path(EGADSlite_INCLUDE_DIR NAMES egads_lite.h + PATHS + "${EGADSlite_PREFIX}/EGADSlite/include" + "${EGADSlite_PREFIX}/include" + NO_DEFAULT_PATH) + if(NOT EGADSlite_INCLUDE_DIR) + message(FATAL_ERROR "could not find path to \"egads_lite.h\"") + endif() + find_library(EGADSlite_LIBRARY NAMES egadslite + PATHS + "${EGADSlite_PREFIX}/EGADSlite/lib" + "${EGADSlite_PREFIX}/lib" + NO_DEFAULT_PATH) + if(NOT EGADSlite_LIBRARY) + message(FATAL_ERROR "could not find EGADSlite_LIBRARY") + endif() + message(STATUS "EGADSlite_INCLUDE_DIR: ${EGADSlite_INCLUDE_DIR}") + message(STATUS "EGADSlite_LIBRARY: ${EGADSlite_LIBRARY}") + set(Omega_h_SOURCES ${Omega_h_SOURCES} Omega_h_egads_lite.cpp) +endif() + if(Omega_h_USE_SEACASExodus) set(Omega_h_SOURCES ${Omega_h_SOURCES} Omega_h_exodus.cpp) endif() @@ -194,6 +220,11 @@ if(Omega_h_USE_EGADS) target_link_libraries(omega_h PUBLIC "${EGADS_LIBRARY}") endif() +if(Omega_h_USE_EGADSlite) + target_include_directories(omega_h PUBLIC "${EGADSlite_INCLUDE_DIR}") + target_link_libraries(omega_h PUBLIC "${EGADSlite_LIBRARY}") +endif() + bob_link_dependency(omega_h PUBLIC SEACASExodus) bob_link_dependency(omega_h PUBLIC ZLIB) @@ -395,6 +426,9 @@ if(BUILD_TESTING) if(Omega_h_USE_EGADS) osh_add_exe(egads_test) endif() + if(Omega_h_USE_EGADSlite) + osh_add_exe(egads_adapt_test) + endif() osh_add_exe(advect2d_test) if(Omega_h_DATA) set(TEST_EXES ${TEST_EXES} advect2d_test) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp new file mode 100644 index 000000000..fda2a33d4 --- /dev/null +++ b/src/Omega_h_egads_lite.cpp @@ -0,0 +1,253 @@ +#include "Omega_h_egads_lite.hpp" +#include "Omega_h_array_ops.hpp" +#include "Omega_h_map.hpp" +#include "Omega_h_mesh.hpp" +#include "Omega_h_timer.hpp" + +#include +#include +#include +#include +#include + +#ifdef __clang__ +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wreserved-id-macro" +#endif + +#include + +enum EgadsObjectClass { + EGADS_CONTXT = CONTXT, + EGADS_TRANSFORM = TRANSFORM, + EGADS_TESSELATION = TESSELLATION, + EGADS_NIL = NIL, + /*EGADS_EMPTY = EMPTY, not doing this one + * because an EGADS error exists by the same name + */ + EGADS_REFERENCE = REFERENCE, + EGADS_PCURVE = PCURVE, + EGADS_CURVE = CURVE, + EGADS_SURFACE = SURFACE, + EGADS_NODE = NODE, + EGADS_EDGE = EDGE, + EGADS_LOOP = LOOP, + EGADS_FACE = FACE, + EGADS_SHELL = SHELL, + EGADS_BODY = BODY, + EGADS_MODEL = MODEL +}; + +#undef CONTXT +#undef TRANSFORM +#undef TESSELLATION +#undef NIL +#undef EMPTY +#undef REFERENCE +#undef PCURVE +#undef CURVE +#undef SURFACE +#undef NODE +#undef EDGE +#undef LOOP +#undef FACE +#undef SHELL +#undef BODY +#undef MODEL + +#ifdef __clang__ +#pragma clang diagnostic pop +#endif + +namespace Omega_h { + +static void call_egads( + int result, char const* code, char const* file, int line) { + if (EGADS_SUCCESS == result) return; + Omega_h_fail( + "EGADS call %s returned %d at %s +%d\n", code, result, file, line); +} + +#define CALL(f) call_egads((f), #f, __FILE__, __LINE__) + +static int const dims2oclass[4] = { + EGADS_NODE, EGADS_EDGE, EGADS_FACE, EGADS_BODY}; + +struct Egads { + ego context; + ego model; + ego body; + int counts[3]; + ego* entities[3]; + std::map, ego> classifier; +}; + +Egads* egads_lite_load(std::string const& filename) { + auto eg = new Egads; + CALL(EGlite_open(&eg->context)); + CALL(EGlite_loadModel(eg->context, 0, filename.c_str(), &eg->model)); + ego model_geom; + int model_oclass; + int model_mtype; + int nbodies; + ego* bodies; + int* body_senses; + CALL(EGlite_getTopology(eg->model, &model_geom, &model_oclass, &model_mtype, + nullptr, &nbodies, &bodies, &body_senses)); + OMEGA_H_CHECK(nbodies == 1); + eg->body = bodies[0]; + for (int i = 0; i < 3; ++i) { + CALL(EGlite_getBodyTopos( + eg->body, nullptr, dims2oclass[i], &eg->counts[i], &eg->entities[i])); + } + // preprocess edge and vertex adjacency to faces + for (int i = 0; i < 2; ++i) { + std::vector> idxs2adj_faces(eg->counts[i]); + for (int j = 0; j < eg->counts[2]; ++j) { + auto face = eg->entities[2][j]; + int nadj_ents; + ego* adj_ents; + CALL(EGlite_getBodyTopos( + eg->body, face, dims2oclass[i], &nadj_ents, &adj_ents)); + for (int k = 0; k < nadj_ents; ++k) { + auto adj_ent = adj_ents[k]; + auto idx = EGlite_indexBodyTopo(eg->body, adj_ent) - 1; + idxs2adj_faces[idx].insert(face); + } + } + for (int j = 0; j < eg->counts[i]; ++j) { + auto adj_faces = idxs2adj_faces[j]; + // HACK!: we have a really insane CAD model with nonsensical topology. + // this essentially manifests as edges that are adjacent to only one + // model face. + // we actually want to just ignore these edges, so we won't create + // classifier entries for them. + if (adj_faces.size() == 1) continue; + eg->classifier[adj_faces] = eg->entities[i][j]; + } + } + return eg; +} + +static int get_dim(ego e) { + ego ref; + int oclass; + int mtype; + int nchild; + ego* children; + int* senses; + CALL(EGlite_getTopology( + e, &ref, &oclass, &mtype, nullptr, &nchild, &children, &senses)); + for (int i = 0; i <= 3; ++i) + if (dims2oclass[i] == oclass) return i; + return -1; +} + +void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], + int* class_dim, int* class_id) { + std::set uniq_adj_faces; + for (int i = 0; i < nadj_faces; ++i) { + auto adj_face = eg->entities[2][adj_face_ids[i] - 1]; + uniq_adj_faces.insert(adj_face); + } + auto it = eg->classifier.find(uniq_adj_faces); + if (it != eg->classifier.end()) { + auto ent = it->second; + *class_dim = get_dim(ent); + *class_id = EGlite_indexBodyTopo(eg->body, ent); + } +} + +void egads_lite_free(Egads* eg) { + for (int i = 0; i < 3; ++i) { + EGlite_free(eg->entities[i]); + } + CALL(EGlite_deleteObject(eg->model)); + CALL(EGlite_close(eg->context)); + delete eg; +} + +void egads_lite_reclassify(Mesh* mesh, Egads* eg) { + OMEGA_H_CHECK(mesh->dim() == 3); + auto face_class_dims = mesh->get_array(FACE, "class_dim"); + auto face_class_ids = mesh->get_array(FACE, "class_id"); + for (Int dim = 0; dim < 2; ++dim) { + auto ents2faces = mesh->ask_up(dim, FACE); + auto adj_class_dims = read(unmap(ents2faces.ab2b, face_class_dims, 1)); + auto keep_edges = each_eq_to(adj_class_dims, I8(2)); + auto ents2eq_faces = filter_graph_edges(ents2faces, keep_edges); + auto adj_eq_face_ids = unmap(ents2eq_faces.ab2b, face_class_ids, 1); + auto host_a2ab = HostRead(ents2eq_faces.a2ab); + auto host_face_ids = HostRead(adj_eq_face_ids); + auto class_dims = mesh->get_array(dim, "class_dim"); + auto class_ids = mesh->get_array(dim, "class_id"); + auto host_class_dims = HostWrite(deep_copy(class_dims)); + auto host_class_ids = HostWrite(deep_copy(class_ids)); + for (LO i = 0; i < mesh->nents(dim); ++i) { + auto b = host_a2ab[i]; + auto e = host_a2ab[i + 1]; + Int class_dim = host_class_dims[i]; + LO class_id = host_class_ids[i]; + egads_classify( + eg, e - b, host_face_ids.data() + b, &class_dim, &class_id); + host_class_dims[i] = I8(class_dim); + host_class_ids[i] = class_id; + } + class_dims = Read(host_class_dims.write()); + class_ids = Read(host_class_ids.write()); + mesh->set_tag(dim, "class_id", class_ids); + mesh->set_tag(dim, "class_dim", class_dims); + } +} + +static Vector<3> get_closest_point(ego g, Vector<3> in) { + double ignored[2] = {}; + Vector<3> out = in; + CALL(EGlite_invEvaluate(g, in.data(), ignored, out.data())); + return out; +} + +Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { + OMEGA_H_CHECK(mesh->dim() == 3); + if (verbose) std::cout << "Querying closest points for surface vertices...\n"; + auto t0 = now(); + auto class_dims = mesh->get_array(VERT, "class_dim"); + auto class_ids = mesh->get_array(VERT, "class_id"); + auto coords = mesh->coords(); + auto host_class_dims = HostRead(class_dims); + auto host_class_ids = HostRead(class_ids); + auto host_coords = HostRead(coords); + auto host_warp = HostWrite(mesh->nverts() * 3); + for (LO i = 0; i < mesh->nverts(); ++i) { + //auto a = get_vector<3>(coords, i); + Vector<3> a; + for (Int j = 0; j < 3; ++j) a[j] = host_coords[i * 3 + j]; + //auto a = get_vector<3>(host_coords, i); + //auto device_a = get_vector<3>(coords, i); + Int class_dim = host_class_dims[i]; + OMEGA_H_CHECK(class_dim >= 0); + OMEGA_H_CHECK(class_dim <= 3); + auto d = vector_3(0, 0, 0); + if (0 < class_dim && class_dim < 3) { + auto index = host_class_ids[i] - 1; + OMEGA_H_CHECK(index >= 0); + OMEGA_H_CHECK(index < eg->counts[class_dim]); + auto g = eg->entities[class_dim][index]; + auto index2 = EGlite_indexBodyTopo(eg->body, g); + OMEGA_H_CHECK(index2 == index + 1); + auto b = get_closest_point(g, a); + d = b - a; + } + //set_vector(host_warp, i, d); + for (Int j = 0; j < 3; ++j) host_warp[i * 3 + j] = d[j]; + } + auto warp = Reals(host_warp.write()); + auto t1 = now(); + if (verbose) { + std::cout << "Querying closest points for surface vertices took " + << (t1 - t0) << " seconds\n"; + } + return warp; +} + +} // namespace Omega_h diff --git a/src/Omega_h_egads_lite.hpp b/src/Omega_h_egads_lite.hpp new file mode 100644 index 000000000..cb5457733 --- /dev/null +++ b/src/Omega_h_egads_lite.hpp @@ -0,0 +1,21 @@ +#ifndef OMEGA_H_EGADS_LITE_HPP +#define OMEGA_H_EGADS_LITE_HPP + +#include +#include +#include "Omega_h_egads.hpp" + +namespace Omega_h { + +class Mesh; + +Egads* egads_lite_load(std::string const& filename); +void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], + int* class_dim, int* class_id); +void egads_lite_free(Egads* eg); +void egads_lite_reclassify(Mesh* mesh, Egads* eg); +Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose); + +} // namespace Omega_h + +#endif diff --git a/src/egads_adapt_test.cpp b/src/egads_adapt_test.cpp new file mode 100644 index 000000000..e1137df64 --- /dev/null +++ b/src/egads_adapt_test.cpp @@ -0,0 +1,101 @@ +#include +#include "Omega_h_egads.hpp" +#include +#include +#include +#include +#include + +#include + +int main(int argc, char** argv) { + auto lib = Omega_h::Library(&argc, &argv); + Omega_h::CmdLine cmdline; + cmdline.add_arg("input.osh"); + cmdline.add_arg("desired-num-elements"); + cmdline.add_arg("output.osh"); + cmdline.add_arg("egads-model"); + auto const world = lib.world(); + if (!cmdline.parse_final(world, &argc, argv)) { + return -1; + } + Omega_h::ScopedTimer scoped_timer("main"); + auto const world_size = world->size(); + auto const world_rank = world->rank(); + auto const inpath = cmdline.get("input.osh"); + auto const desired_nelems = cmdline.get("desired-num-elements"); + auto const outpath = cmdline.get("output.osh"); + auto const egads_path = cmdline.get("egads-model"); + auto eg = Omega_h::egads_load(egads_path); + auto const desired_nelems_per_rank = desired_nelems/world_size; + Omega_h::Mesh mesh(&lib); + for (int shift = 0; (1 << shift) <= world_size; ++shift) { + int const group_size = (1 << shift); + if (world_rank == 0) + std::cout << "going to group size " << group_size << '\n'; + int const group_in_world = world_rank / group_size; + int const group_rank = world_rank % group_size; + auto const group = world->split(group_in_world, group_rank); + OMEGA_H_CHECK(group->rank() == group_rank); + if (group_in_world == 0) { + if (shift == 0) { + if (world_rank == 0) std::cout << "reading mesh...\n"; + mesh = Omega_h::binary::read(inpath, group, true); + mesh.set_curved(-1); + Omega_h::egads_reclassify(&mesh, eg); + } else { + if (world_rank == 0) std::cout << "repartitioning...\n"; + mesh.set_comm(group); + mesh.balance(); + } + if (!mesh.has_tag(1, "global")) { + mesh.add_tag(1, "global", 1, Omega_h::GOs(mesh.nedges(), 0, 1)); + } + if (!mesh.has_tag(0, "global")) { + mesh.add_tag(0, "global", 1, Omega_h::GOs(mesh.nverts(), 0, 1)); + } + if (!mesh.has_tag(2, "global")) { + mesh.add_tag(2, "global", 1, Omega_h::GOs(mesh.nfaces(), 0, 1)); + } + Omega_h::AdaptOpts opts(&mesh); + auto nelems = mesh.nglobal_ents(mesh.dim()); + if (world_rank == 0) + std::cout << "mesh has " << nelems << " total elements\n"; + auto const desired_group_nelems = desired_nelems_per_rank * group_size; + if (double(nelems) >= desired_group_nelems) { + if (world_rank == 0) + std::cout << "element count " << nelems << " >= target " + << desired_group_nelems << ", will not adapt\n"; + } + while (double(nelems) < desired_group_nelems) { + if (world_rank == 0) + std::cout << "element count " << nelems << " < target " + << desired_group_nelems << ", will adapt\n"; + if (!mesh.has_tag(0, "metric")) { + if (world_rank == 0) + std::cout + << "mesh had no metric, adding implied and adapting to it\n"; + opts.egads_model = eg; + Omega_h::add_implied_metric_tag(&mesh); + Omega_h::adapt(&mesh, opts); + nelems = mesh.nglobal_ents(mesh.dim()); + if (world_rank == 0) + std::cout << "mesh now has " << nelems << " total elements\n"; + } + auto metrics = mesh.get_array(0, "metric"); + metrics = Omega_h::multiply_each_by(metrics, 1.2); + auto const metric_ncomps = + Omega_h::divide_no_remainder(metrics.size(), mesh.nverts()); + mesh.add_tag(0, "metric", metric_ncomps, metrics); + if (world_rank == 0) std::cout << "adapting to scaled metric\n"; + Omega_h::adapt(&mesh, opts); + nelems = mesh.nglobal_ents(mesh.dim()); + if (world_rank == 0) + std::cout << "mesh now has " << nelems << " total elements\n"; + } + } + world->barrier(); + } + Omega_h::egads_free(eg); + return 0; +} From 1562a278a41ef3c4d35b45f8ecde9d9f21667688 Mon Sep 17 00:00:00 2001 From: joshia5 Date: Sat, 6 Nov 2021 21:44:02 -0400 Subject: [PATCH 02/40] fix bugs for standalone compilation with lite --- src/Omega_h_adapt.cpp | 44 ++++++++++++++++++++++++++++++++++++++ src/Omega_h_adapt.hpp | 4 ++-- src/Omega_h_egads_lite.cpp | 2 +- src/egads_adapt_test.cpp | 8 +++---- 4 files changed, 51 insertions(+), 7 deletions(-) diff --git a/src/Omega_h_adapt.cpp b/src/Omega_h_adapt.cpp index f5f3e4848..d11da5cbd 100644 --- a/src/Omega_h_adapt.cpp +++ b/src/Omega_h_adapt.cpp @@ -18,6 +18,9 @@ #ifdef OMEGA_H_USE_EGADS #include "Omega_h_egads.hpp" #endif +#ifdef OMEGA_H_USE_EGADSlite +#include "Omega_h_egads_lite.hpp" +#endif namespace Omega_h { @@ -76,6 +79,12 @@ AdaptOpts::AdaptOpts(Int dim) { should_smooth_snap = true; snap_smooth_tolerance = 1e-2; allow_snap_failure = false; +#endif +#ifdef OMEGA_H_USE_EGADSlite + egads_model = nullptr; + should_smooth_snap = true; + snap_smooth_tolerance = 1e-2; + allow_snap_failure = false; #endif should_refine = true; should_coarsen = true; @@ -237,6 +246,38 @@ static void snap_and_satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { } } } else +#endif +#ifdef OMEGA_H_USE_EGADSlite + if (opts.egads_model) { + ScopedTimer snap_timer("snap"); + + mesh->change_all_rcFieldsTorc(); + mesh->set_parting(OMEGA_H_GHOSTED); + mesh->change_all_rcFieldsToMesh(); + + auto warp = egads_lite_get_snap_warp( + mesh, opts.egads_model, opts.verbosity >= EACH_REBUILD); + if (opts.should_smooth_snap) { + if (opts.verbosity >= EACH_REBUILD) { + std::cout << "Solving Laplacian of warp field...\n"; + } + auto t0 = now(); + warp = + solve_laplacian(mesh, warp, mesh->dim(), opts.snap_smooth_tolerance); + auto t1 = now(); + if (opts.verbosity >= EACH_REBUILD) { + std::cout << "Solving Laplacian of warp field took " << (t1 - t0) + << " seconds\n"; + } + } + mesh->add_tag(VERT, "warp", mesh->dim(), warp); + while (warp_to_limit(mesh, opts, opts.allow_snap_failure)) { + if (!satisfy_quality(mesh, opts)) { + mesh->remove_tag(VERT, "warp"); + break; + } + } + } else #endif satisfy_quality(mesh, opts); } @@ -254,6 +295,9 @@ static void post_adapt( if (opts.verbosity > SILENT && !mesh->comm()->rank()) { #ifdef OMEGA_H_USE_EGADS if (opts.egads_model) std::cout << "snapping while "; +#endif +#ifdef OMEGA_H_USE_EGADSlite + if (opts.egads_model) std::cout << "snapping while "; #endif std::cout << "addressing element qualities took " << (t3 - t2); std::cout << " seconds\n"; diff --git a/src/Omega_h_adapt.hpp b/src/Omega_h_adapt.hpp index 32906b35c..4c2dbee78 100644 --- a/src/Omega_h_adapt.hpp +++ b/src/Omega_h_adapt.hpp @@ -43,7 +43,7 @@ struct TransferOpts { enum Verbosity { SILENT, EACH_ADAPT, EACH_REBUILD, EXTRA_STATS }; -#ifdef OMEGA_H_USE_EGADS +#if defined (OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) struct Egads; #endif @@ -62,7 +62,7 @@ struct AdaptOpts { Real length_histogram_max; Int nlength_histogram_bins; Int nquality_histogram_bins; -#ifdef OMEGA_H_USE_EGADS +#if defined (OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) Egads* egads_model; bool should_smooth_snap; Real snap_smooth_tolerance; diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index fda2a33d4..4b6f37e09 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -188,7 +188,7 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { auto e = host_a2ab[i + 1]; Int class_dim = host_class_dims[i]; LO class_id = host_class_ids[i]; - egads_classify( + egads_lite_classify( eg, e - b, host_face_ids.data() + b, &class_dim, &class_id); host_class_dims[i] = I8(class_dim); host_class_ids[i] = class_id; diff --git a/src/egads_adapt_test.cpp b/src/egads_adapt_test.cpp index e1137df64..016a93c7e 100644 --- a/src/egads_adapt_test.cpp +++ b/src/egads_adapt_test.cpp @@ -1,5 +1,5 @@ #include -#include "Omega_h_egads.hpp" +#include "Omega_h_egads_lite.hpp" #include #include #include @@ -26,7 +26,7 @@ int main(int argc, char** argv) { auto const desired_nelems = cmdline.get("desired-num-elements"); auto const outpath = cmdline.get("output.osh"); auto const egads_path = cmdline.get("egads-model"); - auto eg = Omega_h::egads_load(egads_path); + auto eg = Omega_h::egads_lite_load(egads_path); auto const desired_nelems_per_rank = desired_nelems/world_size; Omega_h::Mesh mesh(&lib); for (int shift = 0; (1 << shift) <= world_size; ++shift) { @@ -42,7 +42,7 @@ int main(int argc, char** argv) { if (world_rank == 0) std::cout << "reading mesh...\n"; mesh = Omega_h::binary::read(inpath, group, true); mesh.set_curved(-1); - Omega_h::egads_reclassify(&mesh, eg); + Omega_h::egads_lite_reclassify(&mesh, eg); } else { if (world_rank == 0) std::cout << "repartitioning...\n"; mesh.set_comm(group); @@ -96,6 +96,6 @@ int main(int argc, char** argv) { } world->barrier(); } - Omega_h::egads_free(eg); + Omega_h::egads_lite_free(eg); return 0; } From ae3ac3c2770183c1a0b6098d4e8beaa0cc23f372 Mon Sep 17 00:00:00 2001 From: joshia5 Date: Sun, 7 Nov 2021 16:16:08 -0500 Subject: [PATCH 03/40] working egads_lite_adapt_test from device --- src/Omega_h_adapt.cpp | 15 +++---- src/Omega_h_adapt.hpp | 7 ++- src/Omega_h_egads_lite.cpp | 1 + src/egads_lite_adapt_test.cpp | 82 +++++++++++++++++++++++++++++++++++ 4 files changed, 94 insertions(+), 11 deletions(-) create mode 100644 src/egads_lite_adapt_test.cpp diff --git a/src/Omega_h_adapt.cpp b/src/Omega_h_adapt.cpp index d11da5cbd..88a7c7a62 100644 --- a/src/Omega_h_adapt.cpp +++ b/src/Omega_h_adapt.cpp @@ -74,17 +74,14 @@ AdaptOpts::AdaptOpts(Int dim) { length_histogram_max = 3.0; nlength_histogram_bins = 10; nquality_histogram_bins = 10; -#ifdef OMEGA_H_USE_EGADS - egads_model = nullptr; should_smooth_snap = true; snap_smooth_tolerance = 1e-2; allow_snap_failure = false; +#ifdef OMEGA_H_USE_EGADS + egads_model = nullptr; #endif #ifdef OMEGA_H_USE_EGADSlite - egads_model = nullptr; - should_smooth_snap = true; - snap_smooth_tolerance = 1e-2; - allow_snap_failure = false; + egads_lite_model = nullptr; #endif should_refine = true; should_coarsen = true; @@ -248,7 +245,7 @@ static void snap_and_satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { } else #endif #ifdef OMEGA_H_USE_EGADSlite - if (opts.egads_model) { + if (opts.egads_lite_model) { ScopedTimer snap_timer("snap"); mesh->change_all_rcFieldsTorc(); @@ -256,7 +253,7 @@ static void snap_and_satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { mesh->change_all_rcFieldsToMesh(); auto warp = egads_lite_get_snap_warp( - mesh, opts.egads_model, opts.verbosity >= EACH_REBUILD); + mesh, opts.egads_lite_model, opts.verbosity >= EACH_REBUILD); if (opts.should_smooth_snap) { if (opts.verbosity >= EACH_REBUILD) { std::cout << "Solving Laplacian of warp field...\n"; @@ -297,7 +294,7 @@ static void post_adapt( if (opts.egads_model) std::cout << "snapping while "; #endif #ifdef OMEGA_H_USE_EGADSlite - if (opts.egads_model) std::cout << "snapping while "; + if (opts.egads_lite_model) std::cout << "snapping while "; #endif std::cout << "addressing element qualities took " << (t3 - t2); std::cout << " seconds\n"; diff --git a/src/Omega_h_adapt.hpp b/src/Omega_h_adapt.hpp index 4c2dbee78..d33c9a787 100644 --- a/src/Omega_h_adapt.hpp +++ b/src/Omega_h_adapt.hpp @@ -62,11 +62,14 @@ struct AdaptOpts { Real length_histogram_max; Int nlength_histogram_bins; Int nquality_histogram_bins; -#if defined (OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) - Egads* egads_model; bool should_smooth_snap; Real snap_smooth_tolerance; bool allow_snap_failure; +#ifdef OMEGA_H_USE_EGADS + Egads* egads_model; +#endif +#ifdef OMEGA_H_USE_EGADSLITE + Egads* egads_lite_model; #endif bool should_refine; bool should_coarsen; diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 4b6f37e09..49eb26ae3 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -234,6 +234,7 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { OMEGA_H_CHECK(index < eg->counts[class_dim]); auto g = eg->entities[class_dim][index]; auto index2 = EGlite_indexBodyTopo(eg->body, g); + std::cout << "index, index2 " << index << index2 << "\n"; OMEGA_H_CHECK(index2 == index + 1); auto b = get_closest_point(g, a); d = b - a; diff --git a/src/egads_lite_adapt_test.cpp b/src/egads_lite_adapt_test.cpp new file mode 100644 index 000000000..5ba41369f --- /dev/null +++ b/src/egads_lite_adapt_test.cpp @@ -0,0 +1,82 @@ +#include +#include +#include +#include +#include +#include + +#ifdef OMEGA_H_USE_EGADSLITE +#include +#endif + +#include + +static void compute_implied_metric(Omega_h::Mesh* mesh) { + auto metrics = Omega_h::get_implied_metrics(mesh); + metrics = Omega_h::limit_metric_gradation(mesh, metrics, 1.0); + mesh->add_tag( + Omega_h::VERT, "metric", Omega_h::symm_ncomps(mesh->dim()), metrics); +} + +static void compute_target_metric(Omega_h::Mesh* mesh) { + auto metric = Omega_h::diagonal(Omega_h::metric_eigenvalues_from_lengths( + Omega_h::vector_3(0.1, 0.1, 0.1))); + auto metrics = Omega_h::repeat_symm(mesh->nverts(), metric); + mesh->add_tag(Omega_h::VERT, "target_metric", + Omega_h::symm_ncomps(mesh->dim()), metrics); +} + +int main(int argc, char** argv) { + auto lib = Omega_h::Library(&argc, &argv); + Omega_h::CmdLine cmdline; + cmdline.add_arg("mesh_in.meshb"); + cmdline.add_arg("mesh_out.meshb"); +#ifdef OMEGA_H_USE_EGADSLITE + auto& model_flag = cmdline.add_flag("--model", "optional EGADS model"); + model_flag.add_arg("model.step"); +#endif + auto& viz_flag = cmdline.add_flag("--viz", "optional VTK progress log"); + viz_flag.add_arg("path_vtk"); + if (!cmdline.parse_final(lib.world(), &argc, argv)) return -1; + auto path_in = cmdline.get("mesh_in.meshb"); + auto path_out = cmdline.get("mesh_out.meshb"); + Omega_h::Mesh mesh(&lib); + std::cout << "reading in " << path_in << '\n'; + Omega_h::meshb::read(&mesh, path_in); + std::cout << "computing metric tags\n"; + compute_implied_metric(&mesh); + compute_target_metric(&mesh); + std::cout << "computing minimum quality\n"; + Omega_h::AdaptOpts opts(&mesh); +#ifdef OMEGA_H_USE_EGADSLITE + auto has_model = cmdline.parsed("--model"); + if (has_model) { + auto model_path = cmdline.get("--model", "model.step"); + std::cout << "reading in " << model_path << '\n'; + auto eg = Omega_h::egads_lite_load(model_path); + Omega_h::egads_lite_reclassify(&mesh, eg); + opts.egads_lite_model = eg; + } +#endif + auto has_viz = cmdline.parsed("--viz"); + Omega_h::vtk::Writer writer; + if (has_viz) { + auto viz_path = cmdline.get("--viz", "path_vtk"); + writer = Omega_h::vtk::Writer(viz_path, &mesh); + writer.write(); + } + opts.verbosity = Omega_h::EXTRA_STATS; + opts.max_length_allowed = opts.max_length_desired * 2.0; + while (Omega_h::approach_metric(&mesh, opts)) { + Omega_h::adapt(&mesh, opts); + if (has_viz) writer.write(); + } + std::cout << "writing out " << path_out << '\n'; + mesh.remove_tag(Omega_h::VERT, "metric"); + Omega_h::meshb::write(&mesh, path_out); +#ifdef OMEGA_H_USE_EGADSLITE + if (has_model) { + Omega_h::egads_lite_free(opts.egads_lite_model); + } +#endif +} From 272bd0789c1064d4531850234d244bfabbf95044 Mon Sep 17 00:00:00 2001 From: joshia5 Date: Mon, 14 Feb 2022 10:10:37 -0500 Subject: [PATCH 04/40] last egads changes cherry picked from adityayjoshi/crv @ afa9729 with curved adapt changes removed --- src/Omega_h_egads_lite.cpp | 36 +++++-------- src/egads_adapt_test.cpp | 101 ------------------------------------- 2 files changed, 13 insertions(+), 124 deletions(-) delete mode 100644 src/egads_adapt_test.cpp diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 49eb26ae3..c0023daea 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -3,6 +3,7 @@ #include "Omega_h_map.hpp" #include "Omega_h_mesh.hpp" #include "Omega_h_timer.hpp" +#include "Omega_h_for.hpp" #include #include @@ -61,11 +62,9 @@ enum EgadsObjectClass { namespace Omega_h { -static void call_egads( +OMEGA_H_INLINE void call_egads( int result, char const* code, char const* file, int line) { if (EGADS_SUCCESS == result) return; - Omega_h_fail( - "EGADS call %s returned %d at %s +%d\n", code, result, file, line); } #define CALL(f) call_egads((f), #f, __FILE__, __LINE__) @@ -200,10 +199,10 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { } } -static Vector<3> get_closest_point(ego g, Vector<3> in) { - double ignored[2] = {}; +OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { + Vector<2> ignored; Vector<3> out = in; - CALL(EGlite_invEvaluate(g, in.data(), ignored, out.data())); + CALL(EGlite_invEvaluate(g, in.data(), ignored.data(), out.data())); return out; } @@ -214,35 +213,26 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto class_dims = mesh->get_array(VERT, "class_dim"); auto class_ids = mesh->get_array(VERT, "class_id"); auto coords = mesh->coords(); - auto host_class_dims = HostRead(class_dims); - auto host_class_ids = HostRead(class_ids); - auto host_coords = HostRead(coords); - auto host_warp = HostWrite(mesh->nverts() * 3); - for (LO i = 0; i < mesh->nverts(); ++i) { - //auto a = get_vector<3>(coords, i); - Vector<3> a; - for (Int j = 0; j < 3; ++j) a[j] = host_coords[i * 3 + j]; - //auto a = get_vector<3>(host_coords, i); - //auto device_a = get_vector<3>(coords, i); - Int class_dim = host_class_dims[i]; + auto warp = Write(mesh->nverts() * 3); + auto calc_warp = OMEGA_H_LAMBDA(LO i) { + auto a = get_vector<3>(coords, i); + Int class_dim = class_dims[i]; OMEGA_H_CHECK(class_dim >= 0); OMEGA_H_CHECK(class_dim <= 3); auto d = vector_3(0, 0, 0); if (0 < class_dim && class_dim < 3) { - auto index = host_class_ids[i] - 1; + auto index = class_ids[i] - 1; OMEGA_H_CHECK(index >= 0); OMEGA_H_CHECK(index < eg->counts[class_dim]); auto g = eg->entities[class_dim][index]; auto index2 = EGlite_indexBodyTopo(eg->body, g); - std::cout << "index, index2 " << index << index2 << "\n"; OMEGA_H_CHECK(index2 == index + 1); auto b = get_closest_point(g, a); d = b - a; } - //set_vector(host_warp, i, d); - for (Int j = 0; j < 3; ++j) host_warp[i * 3 + j] = d[j]; - } - auto warp = Reals(host_warp.write()); + set_vector(warp, i, d); + }; + parallel_for(mesh->nverts(), std::move(calc_warp), "calc_warp"); auto t1 = now(); if (verbose) { std::cout << "Querying closest points for surface vertices took " diff --git a/src/egads_adapt_test.cpp b/src/egads_adapt_test.cpp deleted file mode 100644 index 016a93c7e..000000000 --- a/src/egads_adapt_test.cpp +++ /dev/null @@ -1,101 +0,0 @@ -#include -#include "Omega_h_egads_lite.hpp" -#include -#include -#include -#include -#include - -#include - -int main(int argc, char** argv) { - auto lib = Omega_h::Library(&argc, &argv); - Omega_h::CmdLine cmdline; - cmdline.add_arg("input.osh"); - cmdline.add_arg("desired-num-elements"); - cmdline.add_arg("output.osh"); - cmdline.add_arg("egads-model"); - auto const world = lib.world(); - if (!cmdline.parse_final(world, &argc, argv)) { - return -1; - } - Omega_h::ScopedTimer scoped_timer("main"); - auto const world_size = world->size(); - auto const world_rank = world->rank(); - auto const inpath = cmdline.get("input.osh"); - auto const desired_nelems = cmdline.get("desired-num-elements"); - auto const outpath = cmdline.get("output.osh"); - auto const egads_path = cmdline.get("egads-model"); - auto eg = Omega_h::egads_lite_load(egads_path); - auto const desired_nelems_per_rank = desired_nelems/world_size; - Omega_h::Mesh mesh(&lib); - for (int shift = 0; (1 << shift) <= world_size; ++shift) { - int const group_size = (1 << shift); - if (world_rank == 0) - std::cout << "going to group size " << group_size << '\n'; - int const group_in_world = world_rank / group_size; - int const group_rank = world_rank % group_size; - auto const group = world->split(group_in_world, group_rank); - OMEGA_H_CHECK(group->rank() == group_rank); - if (group_in_world == 0) { - if (shift == 0) { - if (world_rank == 0) std::cout << "reading mesh...\n"; - mesh = Omega_h::binary::read(inpath, group, true); - mesh.set_curved(-1); - Omega_h::egads_lite_reclassify(&mesh, eg); - } else { - if (world_rank == 0) std::cout << "repartitioning...\n"; - mesh.set_comm(group); - mesh.balance(); - } - if (!mesh.has_tag(1, "global")) { - mesh.add_tag(1, "global", 1, Omega_h::GOs(mesh.nedges(), 0, 1)); - } - if (!mesh.has_tag(0, "global")) { - mesh.add_tag(0, "global", 1, Omega_h::GOs(mesh.nverts(), 0, 1)); - } - if (!mesh.has_tag(2, "global")) { - mesh.add_tag(2, "global", 1, Omega_h::GOs(mesh.nfaces(), 0, 1)); - } - Omega_h::AdaptOpts opts(&mesh); - auto nelems = mesh.nglobal_ents(mesh.dim()); - if (world_rank == 0) - std::cout << "mesh has " << nelems << " total elements\n"; - auto const desired_group_nelems = desired_nelems_per_rank * group_size; - if (double(nelems) >= desired_group_nelems) { - if (world_rank == 0) - std::cout << "element count " << nelems << " >= target " - << desired_group_nelems << ", will not adapt\n"; - } - while (double(nelems) < desired_group_nelems) { - if (world_rank == 0) - std::cout << "element count " << nelems << " < target " - << desired_group_nelems << ", will adapt\n"; - if (!mesh.has_tag(0, "metric")) { - if (world_rank == 0) - std::cout - << "mesh had no metric, adding implied and adapting to it\n"; - opts.egads_model = eg; - Omega_h::add_implied_metric_tag(&mesh); - Omega_h::adapt(&mesh, opts); - nelems = mesh.nglobal_ents(mesh.dim()); - if (world_rank == 0) - std::cout << "mesh now has " << nelems << " total elements\n"; - } - auto metrics = mesh.get_array(0, "metric"); - metrics = Omega_h::multiply_each_by(metrics, 1.2); - auto const metric_ncomps = - Omega_h::divide_no_remainder(metrics.size(), mesh.nverts()); - mesh.add_tag(0, "metric", metric_ncomps, metrics); - if (world_rank == 0) std::cout << "adapting to scaled metric\n"; - Omega_h::adapt(&mesh, opts); - nelems = mesh.nglobal_ents(mesh.dim()); - if (world_rank == 0) - std::cout << "mesh now has " << nelems << " total elements\n"; - } - } - world->barrier(); - } - Omega_h::egads_lite_free(eg); - return 0; -} From 9daf68167b5c8ffd929a3af77bda503f222a0fd9 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 14 Feb 2022 10:31:09 -0500 Subject: [PATCH 05/40] egadslite: fix test name --- src/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index b6093a36a..b5fde5074 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -427,7 +427,7 @@ if(BUILD_TESTING) osh_add_exe(egads_test) endif() if(Omega_h_USE_EGADSlite) - osh_add_exe(egads_adapt_test) + osh_add_exe(egads_lite_adapt_test) endif() osh_add_exe(advect2d_test) if(Omega_h_DATA) From c95c6be9c4d406557ae9bec6ec3fb009f8017a2d Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 14 Feb 2022 10:31:42 -0500 Subject: [PATCH 06/40] egadslite: add input var --- CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/CMakeLists.txt b/CMakeLists.txt index 9c172dd7c..2c892a58f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -23,6 +23,7 @@ bob_input(Omega_h_DATA "" PATH "Path to omega_h-data test files") bob_option(Omega_h_USE_EGADS "Use EGADS from ESP for geometry" OFF) bob_option(Omega_h_USE_EGADSlite "Use EGADSlite from ESP for geometry" OFF) bob_input(EGADS_PREFIX "" PATH "EGADS (or ESP) installation directory") +bob_input(EGADSlite_PREFIX "" PATH "EGADSlite (or ESP) installation directory") bob_option(Omega_h_USE_Kokkos "Use Kokkos as a backend" OFF) bob_input(Kokkos_PREFIX "" PATH "Path to Kokkos install") bob_option(Omega_h_USE_CUDA_AWARE_MPI "Assume MPI is CUDA-aware, make use of that" OFF) From 5acdf05e090becce7edec7bb158855e3840bfb4b Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 14 Feb 2022 10:32:24 -0500 Subject: [PATCH 07/40] egadslite: using version without renamed header --- src/CMakeLists.txt | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index b5fde5074..e3ddde71a 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -157,15 +157,15 @@ if(Omega_h_USE_EGADSlite) if(NOT EGADSlite_PREFIX) message(FATAL_ERROR "Omega_h_USE_EGADSlite=${Omega_h_USE_EGADSlite} but EGADSlite_PREFIX not set!") endif() - find_path(EGADSlite_INCLUDE_DIR NAMES egads_lite.h + find_path(EGADSlite_INCLUDE_DIR NAMES egads.h PATHS "${EGADSlite_PREFIX}/EGADSlite/include" "${EGADSlite_PREFIX}/include" NO_DEFAULT_PATH) if(NOT EGADSlite_INCLUDE_DIR) - message(FATAL_ERROR "could not find path to \"egads_lite.h\"") + message(FATAL_ERROR "could not find path to \"egads.h\"") endif() - find_library(EGADSlite_LIBRARY NAMES egadslite + find_library(EGADSlite_LIBRARY NAMES egadsliteCuda PATHS "${EGADSlite_PREFIX}/EGADSlite/lib" "${EGADSlite_PREFIX}/lib" From 405ae23f50e205f6aeab221f3398de7fbc443474 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 14 Feb 2022 12:28:10 -0500 Subject: [PATCH 08/40] egadslite: fwd declare egads struct --- src/Omega_h_egads_lite.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/Omega_h_egads_lite.hpp b/src/Omega_h_egads_lite.hpp index cb5457733..64a61c576 100644 --- a/src/Omega_h_egads_lite.hpp +++ b/src/Omega_h_egads_lite.hpp @@ -3,11 +3,11 @@ #include #include -#include "Omega_h_egads.hpp" namespace Omega_h { class Mesh; +struct Egads; Egads* egads_lite_load(std::string const& filename); void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], From 563d580c2706904aa564e4581b30ea16fd86a612 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 14 Feb 2022 12:29:46 -0500 Subject: [PATCH 09/40] egadslite: drop lite from names --- src/Omega_h_egads_lite.cpp | 28 ++++++++++++++-------------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index c0023daea..e90efc025 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -16,7 +16,7 @@ #pragma clang diagnostic ignored "-Wreserved-id-macro" #endif -#include +#include enum EgadsObjectClass { EGADS_CONTXT = CONTXT, @@ -83,20 +83,20 @@ struct Egads { Egads* egads_lite_load(std::string const& filename) { auto eg = new Egads; - CALL(EGlite_open(&eg->context)); - CALL(EGlite_loadModel(eg->context, 0, filename.c_str(), &eg->model)); + CALL(EG_open(&eg->context)); + CALL(EG_loadModel(eg->context, 0, filename.c_str(), &eg->model)); ego model_geom; int model_oclass; int model_mtype; int nbodies; ego* bodies; int* body_senses; - CALL(EGlite_getTopology(eg->model, &model_geom, &model_oclass, &model_mtype, + CALL(EG_getTopology(eg->model, &model_geom, &model_oclass, &model_mtype, nullptr, &nbodies, &bodies, &body_senses)); OMEGA_H_CHECK(nbodies == 1); eg->body = bodies[0]; for (int i = 0; i < 3; ++i) { - CALL(EGlite_getBodyTopos( + CALL(EG_getBodyTopos( eg->body, nullptr, dims2oclass[i], &eg->counts[i], &eg->entities[i])); } // preprocess edge and vertex adjacency to faces @@ -106,11 +106,11 @@ Egads* egads_lite_load(std::string const& filename) { auto face = eg->entities[2][j]; int nadj_ents; ego* adj_ents; - CALL(EGlite_getBodyTopos( + CALL(EG_getBodyTopos( eg->body, face, dims2oclass[i], &nadj_ents, &adj_ents)); for (int k = 0; k < nadj_ents; ++k) { auto adj_ent = adj_ents[k]; - auto idx = EGlite_indexBodyTopo(eg->body, adj_ent) - 1; + auto idx = EG_indexBodyTopo(eg->body, adj_ent) - 1; idxs2adj_faces[idx].insert(face); } } @@ -135,7 +135,7 @@ static int get_dim(ego e) { int nchild; ego* children; int* senses; - CALL(EGlite_getTopology( + CALL(EG_getTopology( e, &ref, &oclass, &mtype, nullptr, &nchild, &children, &senses)); for (int i = 0; i <= 3; ++i) if (dims2oclass[i] == oclass) return i; @@ -153,16 +153,16 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], if (it != eg->classifier.end()) { auto ent = it->second; *class_dim = get_dim(ent); - *class_id = EGlite_indexBodyTopo(eg->body, ent); + *class_id = EG_indexBodyTopo(eg->body, ent); } } void egads_lite_free(Egads* eg) { for (int i = 0; i < 3; ++i) { - EGlite_free(eg->entities[i]); + EG_free(eg->entities[i]); } - CALL(EGlite_deleteObject(eg->model)); - CALL(EGlite_close(eg->context)); + CALL(EG_deleteObject(eg->model)); + CALL(EG_close(eg->context)); delete eg; } @@ -202,7 +202,7 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { Vector<2> ignored; Vector<3> out = in; - CALL(EGlite_invEvaluate(g, in.data(), ignored.data(), out.data())); + CALL(EG_invEvaluate(g, in.data(), ignored.data(), out.data())); return out; } @@ -225,7 +225,7 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { OMEGA_H_CHECK(index >= 0); OMEGA_H_CHECK(index < eg->counts[class_dim]); auto g = eg->entities[class_dim][index]; - auto index2 = EGlite_indexBodyTopo(eg->body, g); + auto index2 = EG_indexBodyTopo(eg->body, g); OMEGA_H_CHECK(index2 == index + 1); auto b = get_closest_point(g, a); d = b - a; From eadcd2033471784bde60dd55e323bddd084b290e Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Tue, 15 Feb 2022 14:08:03 -0500 Subject: [PATCH 10/40] egadslite: use egadslite cmake config file --- src/CMakeLists.txt | 26 ++------------------------ 1 file changed, 2 insertions(+), 24 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index e3ddde71a..2265ad37c 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -153,28 +153,7 @@ if(Omega_h_USE_EGADS) endif() if(Omega_h_USE_EGADSlite) - message(STATUS "Omega_h_USE_EGADSlite: ${Omega_h_USE_EGADSlite}") - if(NOT EGADSlite_PREFIX) - message(FATAL_ERROR "Omega_h_USE_EGADSlite=${Omega_h_USE_EGADSlite} but EGADSlite_PREFIX not set!") - endif() - find_path(EGADSlite_INCLUDE_DIR NAMES egads.h - PATHS - "${EGADSlite_PREFIX}/EGADSlite/include" - "${EGADSlite_PREFIX}/include" - NO_DEFAULT_PATH) - if(NOT EGADSlite_INCLUDE_DIR) - message(FATAL_ERROR "could not find path to \"egads.h\"") - endif() - find_library(EGADSlite_LIBRARY NAMES egadsliteCuda - PATHS - "${EGADSlite_PREFIX}/EGADSlite/lib" - "${EGADSlite_PREFIX}/lib" - NO_DEFAULT_PATH) - if(NOT EGADSlite_LIBRARY) - message(FATAL_ERROR "could not find EGADSlite_LIBRARY") - endif() - message(STATUS "EGADSlite_INCLUDE_DIR: ${EGADSlite_INCLUDE_DIR}") - message(STATUS "EGADSlite_LIBRARY: ${EGADSlite_LIBRARY}") + find_package(egadslite REQUIRED) set(Omega_h_SOURCES ${Omega_h_SOURCES} Omega_h_egads_lite.cpp) endif() @@ -221,8 +200,7 @@ if(Omega_h_USE_EGADS) endif() if(Omega_h_USE_EGADSlite) - target_include_directories(omega_h PUBLIC "${EGADSlite_INCLUDE_DIR}") - target_link_libraries(omega_h PUBLIC "${EGADSlite_LIBRARY}") + target_link_libraries(omega_h PUBLIC egadslite::egadslite) endif() bob_link_dependency(omega_h PUBLIC SEACASExodus) From 1e2f7e556e10651151a6caef642713a57d48430e Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Tue, 15 Feb 2022 14:09:05 -0500 Subject: [PATCH 11/40] egadslite: separable compilation appears to be required i really don't understand how separable compilation is supposed to work --- src/CMakeLists.txt | 1 + 1 file changed, 1 insertion(+) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 2265ad37c..e97590ba1 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -201,6 +201,7 @@ endif() if(Omega_h_USE_EGADSlite) target_link_libraries(omega_h PUBLIC egadslite::egadslite) + set_target_properties(omega_h PROPERTIES CUDA_SEPARABLE_COMPILATION ON) endif() bob_link_dependency(omega_h PUBLIC SEACASExodus) From 0ef432f5bc4d4dc7d51f5f4b5fa0cfaf2fff3321 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Tue, 15 Feb 2022 16:07:21 -0500 Subject: [PATCH 12/40] egadslite: links all exe except egads_lite_adapt_test --- src/CMakeLists.txt | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index e97590ba1..8d51881d3 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -154,7 +154,14 @@ endif() if(Omega_h_USE_EGADSlite) find_package(egadslite REQUIRED) - set(Omega_h_SOURCES ${Omega_h_SOURCES} Omega_h_egads_lite.cpp) + add_library(ohEgadslite Omega_h_egads_lite.cpp) + set_source_files_properties(Omega_h_egads_lite.cpp PROPERTIES LANGUAGE CUDA) + set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite) + target_include_directories(ohEgadslite PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR} + ${CMAKE_CURRENT_BINARY_DIR}) + target_compile_options(ohEgadslite PRIVATE "--expt-extended-lambda") endif() if(Omega_h_USE_SEACASExodus) @@ -188,6 +195,7 @@ if (Omega_h_USE_OpenMP) target_compile_options(omega_h PUBLIC -fopenmp) endif() + bob_link_dependency(omega_h PUBLIC Kokkos) bob_link_dependency(omega_h PUBLIC libMeshb) @@ -200,8 +208,7 @@ if(Omega_h_USE_EGADS) endif() if(Omega_h_USE_EGADSlite) - target_link_libraries(omega_h PUBLIC egadslite::egadslite) - set_target_properties(omega_h PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + bob_link_dependency(omega_h PRIVATE ohEgadslite) endif() bob_link_dependency(omega_h PUBLIC SEACASExodus) From 3aef4ac5ebdac77941d9c34b8c13075b99041a01 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Tue, 15 Feb 2022 16:15:16 -0500 Subject: [PATCH 13/40] egadslite: use tll bob_link_dependency requires Omega_h_USE_ but here depname is not a tpl being enabled --- src/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 8d51881d3..5b103f7ca 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -208,7 +208,7 @@ if(Omega_h_USE_EGADS) endif() if(Omega_h_USE_EGADSlite) - bob_link_dependency(omega_h PRIVATE ohEgadslite) + target_link_libraries(omega_h PUBLIC ohEgadslite) endif() bob_link_dependency(omega_h PUBLIC SEACASExodus) From 1e01733d7e3892bb17339ba50310b98230c54a26 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Tue, 15 Feb 2022 16:53:28 -0500 Subject: [PATCH 14/40] egadslite: a few dirty hacks, links --- cmake/bob.cmake | 3 ++- src/CMakeLists.txt | 2 ++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/cmake/bob.cmake b/cmake/bob.cmake index 22f9e9612..3e06b9cb9 100644 --- a/cmake/bob.cmake +++ b/cmake/bob.cmake @@ -494,7 +494,8 @@ function(bob_get_link_libs tgt var) endif() set(link_libs) foreach(lib IN LISTS sublibs) - if (TARGET ${lib}) + if (${lib} STREQUAL "ohEgadslite") #HACK + elseif (TARGET ${lib}) get_target_property(subtgt_type "${lib}" TYPE) if (subtgt_type MATCHES "STATIC_LIBRARY|SHARED_LIBRARY") get_target_property(sublibtgt_loc "${lib}" LOCATION) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 5b103f7ca..ebac8190d 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -157,11 +157,13 @@ if(Omega_h_USE_EGADSlite) add_library(ohEgadslite Omega_h_egads_lite.cpp) set_source_files_properties(Omega_h_egads_lite.cpp PROPERTIES LANGUAGE CUDA) set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + set_target_properties(ohEgadslite PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite) target_include_directories(ohEgadslite PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_BINARY_DIR}) target_compile_options(ohEgadslite PRIVATE "--expt-extended-lambda") + install(TARGETS ohEgadslite EXPORT omega_h-target ARCHIVE) #HACK endif() if(Omega_h_USE_SEACASExodus) From f68480cecac39d235b7320c0b406005f4044fd5b Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 11:15:23 -0500 Subject: [PATCH 15/40] egadslite: getTopo on device doesn't output needed data --- src/Omega_h_egads_lite.cpp | 61 ++++++++++++++++++++++++++++++-------- 1 file changed, 49 insertions(+), 12 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index e90efc025..e35e2d759 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -85,20 +85,57 @@ Egads* egads_lite_load(std::string const& filename) { auto eg = new Egads; CALL(EG_open(&eg->context)); CALL(EG_loadModel(eg->context, 0, filename.c_str(), &eg->model)); - ego model_geom; - int model_oclass; - int model_mtype; int nbodies; ego* bodies; - int* body_senses; - CALL(EG_getTopology(eg->model, &model_geom, &model_oclass, &model_mtype, - nullptr, &nbodies, &bodies, &body_senses)); - OMEGA_H_CHECK(nbodies == 1); - eg->body = bodies[0]; - for (int i = 0; i < 3; ++i) { - CALL(EG_getBodyTopos( - eg->body, nullptr, dims2oclass[i], &eg->counts[i], &eg->entities[i])); - } + + for (int i = 0; i < 3; ++i) + printf("dims2oclass[%d] %d\n", i, dims2oclass[i]); + const auto egModel = eg->model; + Omega_h::LOs d2oc = {dims2oclass[0], + dims2oclass[1], + dims2oclass[2], + dims2oclass[3]}; + auto getTopo = OMEGA_H_LAMBDA(int i) { + printf("cuda eg_getTopo\n"); + //-not used as output { + //-passing them through the capture-by-value lambda + // requires them to be const... which they are not + ego model_geom; + int model_oclass; + int model_mtype; + int* body_senses; + //} + int nbodies_local; + ego* bodies_local; + printf("eg_getTopo 0.1\n"); + EG_getTopology(egModel, + &model_geom, + &model_oclass, + &model_mtype, + nullptr, + &nbodies_local, + &bodies_local, + &body_senses); + printf("nbodies_local %d\n", nbodies_local); + assert(nbodies_local == 1); + //eg->body = bodies_local[0]; //FIXME skip this for now + auto body = bodies_local[0]; //FIXME + for (int i = 0; i < 3; ++i) { + printf("d2oc[%d] %d\n", i, d2oc[i]); + int counts; + ego* ents; + EG_getBodyTopos( + body, //FIXME + nullptr, + d2oc[i], + &counts, + &ents); + } + printf("eg_getTopo 0.3\n"); + }; + //FIXME need to set eg->body and output counts and ents from eg_getBodyTopos + parallel_for(1, getTopo, "getEgadsTopo"); + assert(cudaSuccess == cudaDeviceSynchronize()); // preprocess edge and vertex adjacency to faces for (int i = 0; i < 2; ++i) { std::vector> idxs2adj_faces(eg->counts[i]); From 7bc5d96dccdaf3e1598071daefbb4e9b09de153a Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 13:22:47 -0500 Subject: [PATCH 16/40] egadslite: store outputs from getTopo --- src/Omega_h_egads_lite.cpp | 30 ++++++++++++++++++++++++++---- 1 file changed, 26 insertions(+), 4 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index e35e2d759..7ad91bad2 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -95,6 +95,10 @@ Egads* egads_lite_load(std::string const& filename) { dims2oclass[1], dims2oclass[2], dims2oclass[3]}; + Omega_h::Write egCounts_d(3); + assert(sizeof(Omega_h::GO) == sizeof(ego)); //HACK + Omega_h::Write egEnts_d(3); + Omega_h::Write egBody_d(1); auto getTopo = OMEGA_H_LAMBDA(int i) { printf("cuda eg_getTopo\n"); //-not used as output { @@ -118,24 +122,42 @@ Egads* egads_lite_load(std::string const& filename) { &body_senses); printf("nbodies_local %d\n", nbodies_local); assert(nbodies_local == 1); - //eg->body = bodies_local[0]; //FIXME skip this for now - auto body = bodies_local[0]; //FIXME + egBody_d[0] = (Omega_h::GO) bodies_local[0]; + printf("device body %p\n", bodies_local[0]); for (int i = 0; i < 3; ++i) { printf("d2oc[%d] %d\n", i, d2oc[i]); int counts; ego* ents; EG_getBodyTopos( - body, //FIXME + bodies_local[0], nullptr, d2oc[i], &counts, &ents); + egCounts_d[i] = counts; + egEnts_d[i] = (Omega_h::GO) ents; + printf("device %d count %d ents %p\n", + i, egCounts_d[i], egEnts_d[i]); } printf("eg_getTopo 0.3\n"); }; - //FIXME need to set eg->body and output counts and ents from eg_getBodyTopos parallel_for(1, getTopo, "getEgadsTopo"); assert(cudaSuccess == cudaDeviceSynchronize()); + const auto egEnts = Omega_h::HostRead(egEnts_d); + const auto egCounts = Omega_h::HostRead(egCounts_d); + printf("created reads\n"); + for (int i = 0; i < 3; ++i) { + eg->counts[i] = egCounts[i]; + eg->entities[i] = (ego*)egEnts[i]; + printf("host %d count %d ents %p\n", + i, eg->counts[i], eg->entities[i]); + } + printf("3.0\n"); + const auto egBody = Omega_h::HostRead(egBody_d); + printf("host body %p\n", egBody[0]); + eg->body = (ego)egBody[0]; + printf("3.1\n"); + // preprocess edge and vertex adjacency to faces for (int i = 0; i < 2; ++i) { std::vector> idxs2adj_faces(eg->counts[i]); From 5ffd58e5bec8f6a480bbe151ff94c8700b4e3569 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 14:10:45 -0500 Subject: [PATCH 17/40] egadslite: white space, helper fns --- src/Omega_h_egads_lite.cpp | 61 +++++++++++++++++++++----------------- 1 file changed, 33 insertions(+), 28 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 7ad91bad2..287b62065 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -81,6 +81,27 @@ struct Egads { std::map, ego> classifier; }; +OMEGA_H_INLINE Omega_h::GO egoToGo(ego obj) { + return (Omega_h::GO) obj; +} + +OMEGA_H_INLINE Omega_h::GO egoPtrToGo(ego* obj) { + return (Omega_h::GO) obj; +} + +OMEGA_H_INLINE ego* goToEgoPtr(Omega_h::GO obj) { + return (ego*) obj; +} + +OMEGA_H_INLINE ego goToEgo(Omega_h::GO obj) { + return (ego) obj; +} + +Omega_h::Write OhWriteEgo(int n) { + assert(sizeof(Omega_h::GO) == sizeof(ego)); + return Omega_h::Write(n); +} + Egads* egads_lite_load(std::string const& filename) { auto eg = new Egads; CALL(EG_open(&eg->context)); @@ -96,46 +117,31 @@ Egads* egads_lite_load(std::string const& filename) { dims2oclass[2], dims2oclass[3]}; Omega_h::Write egCounts_d(3); - assert(sizeof(Omega_h::GO) == sizeof(ego)); //HACK - Omega_h::Write egEnts_d(3); - Omega_h::Write egBody_d(1); + auto egEnts_d = OhWriteEgo(3); + auto egBody_d = OhWriteEgo(1); auto getTopo = OMEGA_H_LAMBDA(int i) { printf("cuda eg_getTopo\n"); - //-not used as output { - //-passing them through the capture-by-value lambda - // requires them to be const... which they are not ego model_geom; int model_oclass; int model_mtype; int* body_senses; - //} + //needed outputs int nbodies_local; ego* bodies_local; printf("eg_getTopo 0.1\n"); - EG_getTopology(egModel, - &model_geom, - &model_oclass, - &model_mtype, - nullptr, - &nbodies_local, - &bodies_local, - &body_senses); + EG_getTopology(egModel, &model_geom, &model_oclass, &model_mtype, + nullptr, &nbodies_local, &bodies_local, &body_senses); printf("nbodies_local %d\n", nbodies_local); assert(nbodies_local == 1); - egBody_d[0] = (Omega_h::GO) bodies_local[0]; + egBody_d[0] = egoToGo(bodies_local[0]); printf("device body %p\n", bodies_local[0]); for (int i = 0; i < 3; ++i) { printf("d2oc[%d] %d\n", i, d2oc[i]); int counts; ego* ents; - EG_getBodyTopos( - bodies_local[0], - nullptr, - d2oc[i], - &counts, - &ents); + EG_getBodyTopos(bodies_local[0], nullptr, d2oc[i], &counts, &ents); egCounts_d[i] = counts; - egEnts_d[i] = (Omega_h::GO) ents; + egEnts_d[i] = egoPtrToGo(ents); printf("device %d count %d ents %p\n", i, egCounts_d[i], egEnts_d[i]); } @@ -148,14 +154,13 @@ Egads* egads_lite_load(std::string const& filename) { printf("created reads\n"); for (int i = 0; i < 3; ++i) { eg->counts[i] = egCounts[i]; - eg->entities[i] = (ego*)egEnts[i]; - printf("host %d count %d ents %p\n", - i, eg->counts[i], eg->entities[i]); + eg->entities[i] = goToEgoPtr(egEnts[i]); + printf("host %d count %d ents %p\n", i, eg->counts[i], eg->entities[i]); } printf("3.0\n"); const auto egBody = Omega_h::HostRead(egBody_d); printf("host body %p\n", egBody[0]); - eg->body = (ego)egBody[0]; + eg->body = goToEgo(egBody[0]); printf("3.1\n"); // preprocess edge and vertex adjacency to faces @@ -181,7 +186,7 @@ Egads* egads_lite_load(std::string const& filename) { // we actually want to just ignore these edges, so we won't create // classifier entries for them. if (adj_faces.size() == 1) continue; - eg->classifier[adj_faces] = eg->entities[i][j]; + eg->classifier[adj_faces] = eg->entities[i][j]; //TODO copy arrays of ego pointers to host } } return eg; From 7cb1f8ec9d41c0079c36887337a745b85979ba9c Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 14:11:34 -0500 Subject: [PATCH 18/40] egadslite: build vector of sets --- src/Omega_h_egads_lite.cpp | 70 ++++++++++++++++++++++++++++++++------ 1 file changed, 60 insertions(+), 10 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 287b62065..89d5f1130 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -165,19 +165,69 @@ Egads* egads_lite_load(std::string const& filename) { // preprocess edge and vertex adjacency to faces for (int i = 0; i < 2; ++i) { + printf("3.11\n"); + Omega_h::Write setSizes_d(eg->counts[i]); + auto egBody = eg->body; + auto egCounts = eg->counts[2]; + auto egEnts = eg->entities[2]; + //count the set sizes + auto countIndexBody = OMEGA_H_LAMBDA(int i) { + for (int j = 0; j < egCounts; ++j) { + auto face = egEnts[j]; + int nadj_ents; + ego* adj_ents; + EG_getBodyTopos(egBody, face, d2oc[i], &nadj_ents, &adj_ents); + for (int k = 0; k < nadj_ents; ++k) { + auto adj_ent = adj_ents[k]; + auto idx = EG_indexBodyTopo(egBody, adj_ent) - 1; + setSizes_d[idx]++; + } + } + }; + parallel_for(1, countIndexBody, "getIndexBody"); + assert(cudaSuccess == cudaDeviceSynchronize()); + printf("3.12\n"); + + const auto setSizes = Omega_h::HostRead(setSizes_d); + int totSize = 0; + for(int j=0; j setCounts_d(eg->counts[i]); + printf("3.2\n"); + //fill the sets + auto getIndexBody = OMEGA_H_LAMBDA(int i) { + for (int j = 0; j < egCounts; ++j) { + auto face = egEnts[j]; + int nadj_ents; + ego* adj_ents; + EG_getBodyTopos(egBody, face, d2oc[i], &nadj_ents, &adj_ents); + for (int k = 0; k < nadj_ents; ++k) { + auto adj_ent = adj_ents[k]; + auto idx = EG_indexBodyTopo(egBody, adj_ent) - 1; + auto ohIdx = setCounts_d[idx]++; + idxs2adj_faces_d[ohIdx] = egoToGo(face); + } + } + }; + parallel_for(1, getIndexBody, "getIndexBody"); + assert(cudaSuccess == cudaDeviceSynchronize()); + + printf("3.3\n"); std::vector> idxs2adj_faces(eg->counts[i]); - for (int j = 0; j < eg->counts[2]; ++j) { - auto face = eg->entities[2][j]; - int nadj_ents; - ego* adj_ents; - CALL(EG_getBodyTopos( - eg->body, face, dims2oclass[i], &nadj_ents, &adj_ents)); - for (int k = 0; k < nadj_ents; ++k) { - auto adj_ent = adj_ents[k]; - auto idx = EG_indexBodyTopo(eg->body, adj_ent) - 1; - idxs2adj_faces[idx].insert(face); + //copy array into vector of sets + const auto idxs2adj_faces_h = Omega_h::HostRead(idxs2adj_faces_d); + int idx = 0; + for(int j = 0; j < setSizes.size(); j++) { + for(int k = 0; k < setSizes[j]; k++) { + auto face = goToEgo(idxs2adj_faces_h[idx++]); + idxs2adj_faces[j].insert(face); } } + for (int j = 0; j < eg->counts[i]; ++j) { auto adj_faces = idxs2adj_faces[j]; // HACK!: we have a really insane CAD model with nonsensical topology. From b02a5d66266421916883ac4451aa707597201f91 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 17:01:20 -0500 Subject: [PATCH 19/40] egadslite: remove lambda index, copy dev ent ptrs to host --- src/Omega_h_egads_lite.cpp | 53 +++++++++++++++++++++++++++++++++----- 1 file changed, 46 insertions(+), 7 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 89d5f1130..d8f61eacb 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -119,7 +119,7 @@ Egads* egads_lite_load(std::string const& filename) { Omega_h::Write egCounts_d(3); auto egEnts_d = OhWriteEgo(3); auto egBody_d = OhWriteEgo(1); - auto getTopo = OMEGA_H_LAMBDA(int i) { + auto getTopo = OMEGA_H_LAMBDA(int) { printf("cuda eg_getTopo\n"); ego model_geom; int model_oclass; @@ -171,7 +171,7 @@ Egads* egads_lite_load(std::string const& filename) { auto egCounts = eg->counts[2]; auto egEnts = eg->entities[2]; //count the set sizes - auto countIndexBody = OMEGA_H_LAMBDA(int i) { + auto countIndexBody = OMEGA_H_LAMBDA(int) { for (int j = 0; j < egCounts; ++j) { auto face = egEnts[j]; int nadj_ents; @@ -199,7 +199,7 @@ Egads* egads_lite_load(std::string const& filename) { Omega_h::Write setCounts_d(eg->counts[i]); printf("3.2\n"); //fill the sets - auto getIndexBody = OMEGA_H_LAMBDA(int i) { + auto getIndexBody = OMEGA_H_LAMBDA(int) { for (int j = 0; j < egCounts; ++j) { auto face = egEnts[j]; int nadj_ents; @@ -228,6 +228,23 @@ Egads* egads_lite_load(std::string const& filename) { } } + printf("3.4\n"); + //copy each device entity pointer to the host + auto entPtrs_d = OhWriteEgo(eg->counts[i]); + auto copyDevPtrs = OMEGA_H_LAMBDA(int j) { + ego* ents = goToEgoPtr(egEnts_d[i]); + auto entPtr = ents[j]; + entPtrs_d[j] = egoToGo(entPtr); + }; + parallel_for(eg->counts[i], copyDevPtrs, "copyDevPtrs"); + assert(cudaSuccess == cudaDeviceSynchronize()); + auto entPtrs_h = Omega_h::HostRead(entPtrs_d); + eg->entities[i] = new ego[eg->counts[i]]; + for(int j=0; j < eg->counts[i]; j++) { + eg->entities[i][j] = goToEgo(entPtrs_h[j]); + } + + printf("3.5\n"); for (int j = 0; j < eg->counts[i]; ++j) { auto adj_faces = idxs2adj_faces[j]; // HACK!: we have a really insane CAD model with nonsensical topology. @@ -236,9 +253,28 @@ Egads* egads_lite_load(std::string const& filename) { // we actually want to just ignore these edges, so we won't create // classifier entries for them. if (adj_faces.size() == 1) continue; - eg->classifier[adj_faces] = eg->entities[i][j]; //TODO copy arrays of ego pointers to host + eg->classifier[adj_faces] = eg->entities[i][j]; } + printf("3.6\n"); + } + + //copy each device face pointer to the host + const int fdim = 2; + auto entPtrs_d = OhWriteEgo(eg->counts[fdim]); + auto copyDevPtrs = OMEGA_H_LAMBDA(int j) { + ego* ents = goToEgoPtr(egEnts_d[fdim]); + auto entPtr = ents[j]; + entPtrs_d[j] = egoToGo(entPtr); + }; + parallel_for(eg->counts[fdim], copyDevPtrs, "copyDevPtrs"); + assert(cudaSuccess == cudaDeviceSynchronize()); + auto entPtrs_h = Omega_h::HostRead(entPtrs_d); + eg->entities[fdim] = new ego[eg->counts[fdim]]; + for(int j=0; j < eg->counts[fdim]; j++) { + eg->entities[fdim][j] = goToEgo(entPtrs_h[j]); } + printf("3.7\n"); + return eg; } @@ -260,20 +296,23 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], int* class_dim, int* class_id) { std::set uniq_adj_faces; for (int i = 0; i < nadj_faces; ++i) { - auto adj_face = eg->entities[2][adj_face_ids[i] - 1]; + const auto adjFaceId = adj_face_ids[i]-1; + auto adj_face = eg->entities[2][adjFaceId]; uniq_adj_faces.insert(adj_face); } auto it = eg->classifier.find(uniq_adj_faces); if (it != eg->classifier.end()) { auto ent = it->second; - *class_dim = get_dim(ent); + //TODO port to GPU { + *class_dim = get_dim(ent); //FIXME failing here *class_id = EG_indexBodyTopo(eg->body, ent); + //} } } void egads_lite_free(Egads* eg) { for (int i = 0; i < 3; ++i) { - EG_free(eg->entities[i]); + EG_free(eg->entities[i]); //TODO will this work with new ego[...] ? } CALL(EG_deleteObject(eg->model)); CALL(EG_close(eg->context)); From 126da8bd6b164381ffc99babc57dca4de3eff10d Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 18:22:06 -0500 Subject: [PATCH 20/40] egadslite: egads_lite_adapt_test executes without segfault --- src/Omega_h_egads_lite.cpp | 35 +++++++++++++++++++++++++++++++++-- 1 file changed, 33 insertions(+), 2 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index d8f61eacb..ef7cf6165 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -304,9 +304,40 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], if (it != eg->classifier.end()) { auto ent = it->second; //TODO port to GPU { - *class_dim = get_dim(ent); //FIXME failing here - *class_id = EG_indexBodyTopo(eg->body, ent); + //*class_dim = get_dim(ent); //FIXME failing here + //*class_id = EG_indexBodyTopo(eg->body, ent); //} + Omega_h::LOs d2oc = {dims2oclass[0], + dims2oclass[1], + dims2oclass[2], + dims2oclass[3]}; + Omega_h::Write classDimAndId_d(2); + auto egBody = eg->body; + auto getEntClass = OMEGA_H_LAMBDA(int) { + //get model entity dimension + ego ref; + int oclass; + int mtype; + int nchild; + ego* children; + int* senses; + EG_getTopology(ent, &ref, &oclass, &mtype, nullptr, &nchild, &children, &senses); + classDimAndId_d[0] = -1; + for (int i = 0; i <= 3; ++i) { + if (d2oc[i] == oclass) { + classDimAndId_d[0] = i; + break; + } + } + //get model entity id + classDimAndId_d[1] = EG_indexBodyTopo(egBody, ent); + }; + parallel_for(1, getEntClass, "getEntClass"); + assert(cudaSuccess == cudaDeviceSynchronize()); + auto classDimAndId = Omega_h::HostRead(classDimAndId_d); + //set host vars + *class_dim = classDimAndId[0]; + *class_id = classDimAndId[1]; } } From d1b65a110a1ffef5eeb26cdadf62b1fbcbae2758 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 18:33:05 -0500 Subject: [PATCH 21/40] egadslite: fix pproc conditionals, remove rc api calls egads_lite_test now segfaults --- src/Omega_h_adapt.cpp | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/src/Omega_h_adapt.cpp b/src/Omega_h_adapt.cpp index 88a7c7a62..1f7bff047 100644 --- a/src/Omega_h_adapt.cpp +++ b/src/Omega_h_adapt.cpp @@ -18,7 +18,7 @@ #ifdef OMEGA_H_USE_EGADS #include "Omega_h_egads.hpp" #endif -#ifdef OMEGA_H_USE_EGADSlite +#ifdef OMEGA_H_USE_EGADSLITE #include "Omega_h_egads_lite.hpp" #endif @@ -244,13 +244,11 @@ static void snap_and_satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { } } else #endif -#ifdef OMEGA_H_USE_EGADSlite +#ifdef OMEGA_H_USE_EGADSLITE if (opts.egads_lite_model) { ScopedTimer snap_timer("snap"); - mesh->change_all_rcFieldsTorc(); mesh->set_parting(OMEGA_H_GHOSTED); - mesh->change_all_rcFieldsToMesh(); auto warp = egads_lite_get_snap_warp( mesh, opts.egads_lite_model, opts.verbosity >= EACH_REBUILD); From 9cf6bc0eea697ea69380bffcce4d198d6203edd9 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Thu, 17 Feb 2022 19:50:55 -0500 Subject: [PATCH 22/40] egadslite: store dev ptrs, dies in EG_invEvaLimits --- src/Omega_h_egads_lite.cpp | 37 +++++++++++++++++++++++++++---------- 1 file changed, 27 insertions(+), 10 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index ef7cf6165..38ec0fdbf 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -79,6 +79,8 @@ struct Egads { int counts[3]; ego* entities[3]; std::map, ego> classifier; + LOs counts_d; + ego* entities_d[3]; //HACK stored as arrays of GOs... see conversion fns below }; OMEGA_H_INLINE Omega_h::GO egoToGo(ego obj) { @@ -155,6 +157,8 @@ Egads* egads_lite_load(std::string const& filename) { for (int i = 0; i < 3; ++i) { eg->counts[i] = egCounts[i]; eg->entities[i] = goToEgoPtr(egEnts[i]); + //store the pointer to the array of egads faces + eg->entities_d[i] = goToEgoPtr(egEnts[i]); printf("host %d count %d ents %p\n", i, eg->counts[i], eg->entities[i]); } printf("3.0\n"); @@ -256,7 +260,7 @@ Egads* egads_lite_load(std::string const& filename) { eg->classifier[adj_faces] = eg->entities[i][j]; } printf("3.6\n"); - } + } //done loop over vertices and edges //copy each device face pointer to the host const int fdim = 2; @@ -275,6 +279,9 @@ Egads* egads_lite_load(std::string const& filename) { } printf("3.7\n"); + //set struct device pointer for entity count + eg->counts_d = LOs{eg->counts[0], eg->counts[1], eg->counts[2]}; + return eg; } @@ -303,10 +310,6 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], auto it = eg->classifier.find(uniq_adj_faces); if (it != eg->classifier.end()) { auto ent = it->second; - //TODO port to GPU { - //*class_dim = get_dim(ent); //FIXME failing here - //*class_id = EG_indexBodyTopo(eg->body, ent); - //} Omega_h::LOs d2oc = {dims2oclass[0], dims2oclass[1], dims2oclass[2], @@ -386,11 +389,12 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { Vector<2> ignored; Vector<3> out = in; - CALL(EG_invEvaluate(g, in.data(), ignored.data(), out.data())); + EG_invEvaluate(g, in.data(), ignored.data(), out.data()); return out; } Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { + fprintf(stderr, "numverts %d\n", mesh->nverts()); OMEGA_H_CHECK(mesh->dim() == 3); if (verbose) std::cout << "Querying closest points for surface vertices...\n"; auto t0 = now(); @@ -398,25 +402,38 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto class_ids = mesh->get_array(VERT, "class_id"); auto coords = mesh->coords(); auto warp = Write(mesh->nverts() * 3); + GOs egEnts_d{egoPtrToGo(eg->entities_d[0]), + egoPtrToGo(eg->entities_d[1]), + egoPtrToGo(eg->entities_d[2])}; + auto egCounts_d = eg->counts_d; + auto egBody_d = eg->body; auto calc_warp = OMEGA_H_LAMBDA(LO i) { auto a = get_vector<3>(coords, i); Int class_dim = class_dims[i]; OMEGA_H_CHECK(class_dim >= 0); OMEGA_H_CHECK(class_dim <= 3); auto d = vector_3(0, 0, 0); - if (0 < class_dim && class_dim < 3) { + if (0 < class_dim && class_dim < 3) { //edges and faces only auto index = class_ids[i] - 1; OMEGA_H_CHECK(index >= 0); - OMEGA_H_CHECK(index < eg->counts[class_dim]); - auto g = eg->entities[class_dim][index]; - auto index2 = EG_indexBodyTopo(eg->body, g); + OMEGA_H_CHECK(index < egCounts_d[class_dim]); + auto ents = goToEgoPtr(egEnts_d[class_dim]); + auto g = ents[index]; + auto index2 = EG_indexBodyTopo(egBody_d, g); OMEGA_H_CHECK(index2 == index + 1); + { + int isEdge = (g->oclass == EGADS_EDGE); + int isFace = (g->oclass == EGADS_FACE); + printf("vtx %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", + i, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); + } auto b = get_closest_point(g, a); d = b - a; } set_vector(warp, i, d); }; parallel_for(mesh->nverts(), std::move(calc_warp), "calc_warp"); + assert(cudaSuccess == cudaDeviceSynchronize()); auto t1 = now(); if (verbose) { std::cout << "Querying closest points for surface vertices took " From 6d0e8f562182ceb27fbc547f12c3078cdc637174 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Sat, 19 Feb 2022 06:55:43 -0500 Subject: [PATCH 23/40] egadslite: wrap egads calls with return check --- src/Omega_h_egads_lite.cpp | 35 ++++++++++++++++++++++------------- 1 file changed, 22 insertions(+), 13 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 38ec0fdbf..603f88778 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -65,6 +65,8 @@ namespace Omega_h { OMEGA_H_INLINE void call_egads( int result, char const* code, char const* file, int line) { if (EGADS_SUCCESS == result) return; + OMEGA_H_CHECK_PRINTF(false, + "EGADS call %s returned %d at %s +%d\n", code, result, file, line); } #define CALL(f) call_egads((f), #f, __FILE__, __LINE__) @@ -131,8 +133,8 @@ Egads* egads_lite_load(std::string const& filename) { int nbodies_local; ego* bodies_local; printf("eg_getTopo 0.1\n"); - EG_getTopology(egModel, &model_geom, &model_oclass, &model_mtype, - nullptr, &nbodies_local, &bodies_local, &body_senses); + CALL(EG_getTopology(egModel, &model_geom, &model_oclass, &model_mtype, + nullptr, &nbodies_local, &bodies_local, &body_senses)); printf("nbodies_local %d\n", nbodies_local); assert(nbodies_local == 1); egBody_d[0] = egoToGo(bodies_local[0]); @@ -141,7 +143,7 @@ Egads* egads_lite_load(std::string const& filename) { printf("d2oc[%d] %d\n", i, d2oc[i]); int counts; ego* ents; - EG_getBodyTopos(bodies_local[0], nullptr, d2oc[i], &counts, &ents); + CALL(EG_getBodyTopos(bodies_local[0], nullptr, d2oc[i], &counts, &ents)); egCounts_d[i] = counts; egEnts_d[i] = egoPtrToGo(ents); printf("device %d count %d ents %p\n", @@ -180,10 +182,12 @@ Egads* egads_lite_load(std::string const& filename) { auto face = egEnts[j]; int nadj_ents; ego* adj_ents; - EG_getBodyTopos(egBody, face, d2oc[i], &nadj_ents, &adj_ents); + CALL(EG_getBodyTopos(egBody, face, d2oc[i], &nadj_ents, &adj_ents)); for (int k = 0; k < nadj_ents; ++k) { auto adj_ent = adj_ents[k]; - auto idx = EG_indexBodyTopo(egBody, adj_ent) - 1; + auto egIdx = EG_indexBodyTopo(egBody, adj_ent); + assert(egIdx > 0); //egads error codes are <=0 + auto idx = egIdx-1; setSizes_d[idx]++; } } @@ -208,10 +212,12 @@ Egads* egads_lite_load(std::string const& filename) { auto face = egEnts[j]; int nadj_ents; ego* adj_ents; - EG_getBodyTopos(egBody, face, d2oc[i], &nadj_ents, &adj_ents); + CALL(EG_getBodyTopos(egBody, face, d2oc[i], &nadj_ents, &adj_ents)); for (int k = 0; k < nadj_ents; ++k) { auto adj_ent = adj_ents[k]; - auto idx = EG_indexBodyTopo(egBody, adj_ent) - 1; + auto egIdx = EG_indexBodyTopo(egBody, adj_ent); + assert(egIdx > 0); + auto idx = egIdx-1; auto ohIdx = setCounts_d[idx]++; idxs2adj_faces_d[ohIdx] = egoToGo(face); } @@ -324,7 +330,7 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], int nchild; ego* children; int* senses; - EG_getTopology(ent, &ref, &oclass, &mtype, nullptr, &nchild, &children, &senses); + CALL(EG_getTopology(ent, &ref, &oclass, &mtype, nullptr, &nchild, &children, &senses)); classDimAndId_d[0] = -1; for (int i = 0; i <= 3; ++i) { if (d2oc[i] == oclass) { @@ -333,7 +339,9 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], } } //get model entity id - classDimAndId_d[1] = EG_indexBodyTopo(egBody, ent); + auto egIdx = EG_indexBodyTopo(egBody, ent); + assert(egIdx > 0); + classDimAndId_d[1] = egIdx; }; parallel_for(1, getEntClass, "getEntClass"); assert(cudaSuccess == cudaDeviceSynchronize()); @@ -389,7 +397,7 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { Vector<2> ignored; Vector<3> out = in; - EG_invEvaluate(g, in.data(), ignored.data(), out.data()); + CALL(EG_invEvaluate(g, in.data(), ignored.data(), out.data())); return out; } @@ -420,12 +428,13 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto ents = goToEgoPtr(egEnts_d[class_dim]); auto g = ents[index]; auto index2 = EG_indexBodyTopo(egBody_d, g); + assert(index2 > 0); OMEGA_H_CHECK(index2 == index + 1); - { + if(i == 22 && index2 == 5 && class_dim == 2) { int isEdge = (g->oclass == EGADS_EDGE); int isFace = (g->oclass == EGADS_FACE); - printf("vtx %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", - i, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); + printf("vtx %d class_id %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", + i, index2, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); } auto b = get_closest_point(g, a); d = b - a; From b402d32bedd1cd99aecceb1ccf31092aee5156b6 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Sat, 19 Feb 2022 06:57:00 -0500 Subject: [PATCH 24/40] debug: write vtk, isolate face vertex --- src/Omega_h_egads_lite.cpp | 20 ++++++++++++++++++-- 1 file changed, 18 insertions(+), 2 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 603f88778..ab013757e 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -4,6 +4,7 @@ #include "Omega_h_mesh.hpp" #include "Omega_h_timer.hpp" #include "Omega_h_for.hpp" +#include //vtk #include #include @@ -394,21 +395,28 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { } } -OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { +OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in, int isDebug=0) { Vector<2> ignored; Vector<3> out = in; + if(isDebug) { + printf("in %.3f %.3f %.3f out %.3f %.3f %.3f\\n", + in[0], in[1], in[2], + out[0], out[1], out[2]); + } CALL(EG_invEvaluate(g, in.data(), ignored.data(), out.data())); return out; } Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { fprintf(stderr, "numverts %d\n", mesh->nverts()); + Omega_h::vtk::write_parallel("preWarp", mesh, mesh->dim()); OMEGA_H_CHECK(mesh->dim() == 3); if (verbose) std::cout << "Querying closest points for surface vertices...\n"; auto t0 = now(); auto class_dims = mesh->get_array(VERT, "class_dim"); auto class_ids = mesh->get_array(VERT, "class_id"); auto coords = mesh->coords(); + auto closePts = Write(mesh->nverts() * 3); auto warp = Write(mesh->nverts() * 3); GOs egEnts_d{egoPtrToGo(eg->entities_d[0]), egoPtrToGo(eg->entities_d[1]), @@ -421,6 +429,7 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { OMEGA_H_CHECK(class_dim >= 0); OMEGA_H_CHECK(class_dim <= 3); auto d = vector_3(0, 0, 0); + auto clPt = vector_3(0, 0, 0); if (0 < class_dim && class_dim < 3) { //edges and faces only auto index = class_ids[i] - 1; OMEGA_H_CHECK(index >= 0); @@ -430,20 +439,27 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto index2 = EG_indexBodyTopo(egBody_d, g); assert(index2 > 0); OMEGA_H_CHECK(index2 == index + 1); + int debug = 0; if(i == 22 && index2 == 5 && class_dim == 2) { + debug=1; int isEdge = (g->oclass == EGADS_EDGE); int isFace = (g->oclass == EGADS_FACE); printf("vtx %d class_id %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", i, index2, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); } - auto b = get_closest_point(g, a); + auto b = get_closest_point(g, a, debug); + clPt = b; d = b - a; } set_vector(warp, i, d); + set_vector(closePts, i, clPt); }; parallel_for(mesh->nverts(), std::move(calc_warp), "calc_warp"); assert(cudaSuccess == cudaDeviceSynchronize()); auto t1 = now(); + mesh->add_tag(0, "warpVec", 3, read(warp)); + mesh->add_tag(0, "closePts", 3, read(closePts)); + Omega_h::vtk::write_parallel("warpVec", mesh, mesh->dim()); if (verbose) { std::cout << "Querying closest points for surface vertices took " << (t1 - t0) << " seconds\n"; From a13b4df6ac20ba85fc34b614d1b4f7f779a6d671 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 21 Feb 2022 20:49:24 -0500 Subject: [PATCH 25/40] debug: target a single known vertex --- src/Omega_h_egads_lite.cpp | 17 +++++++---------- 1 file changed, 7 insertions(+), 10 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index ab013757e..0bf767a44 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -395,14 +395,9 @@ void egads_lite_reclassify(Mesh* mesh, Egads* eg) { } } -OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in, int isDebug=0) { +OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { Vector<2> ignored; Vector<3> out = in; - if(isDebug) { - printf("in %.3f %.3f %.3f out %.3f %.3f %.3f\\n", - in[0], in[1], in[2], - out[0], out[1], out[2]); - } CALL(EG_invEvaluate(g, in.data(), ignored.data(), out.data())); return out; } @@ -424,6 +419,7 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto egCounts_d = eg->counts_d; auto egBody_d = eg->body; auto calc_warp = OMEGA_H_LAMBDA(LO i) { + i=22; auto a = get_vector<3>(coords, i); Int class_dim = class_dims[i]; OMEGA_H_CHECK(class_dim >= 0); @@ -446,15 +442,16 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { int isFace = (g->oclass == EGADS_FACE); printf("vtx %d class_id %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", i, index2, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); + auto b = get_closest_point(g, a); + printf("clPt %.3f %.3f %.3f\n", b[0], b[1], b[2]); + clPt = b; + d = b - a; } - auto b = get_closest_point(g, a, debug); - clPt = b; - d = b - a; } set_vector(warp, i, d); set_vector(closePts, i, clPt); }; - parallel_for(mesh->nverts(), std::move(calc_warp), "calc_warp"); + parallel_for(1, std::move(calc_warp), "calc_warp"); assert(cudaSuccess == cudaDeviceSynchronize()); auto t1 = now(); mesh->add_tag(0, "warpVec", 3, read(warp)); From 92f3b5ccb5735a11ec7a54f2281db1c9c173a222 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 21 Feb 2022 20:50:40 -0500 Subject: [PATCH 26/40] egadslite: increase the cuda stack size is this a felony? without this change the call from EG_invEvaluate -> ... -> EG_invEvaGeomLimits fails --- src/CMakeLists.txt | 1 + src/egads_lite_adapt_test.cpp | 29 +++++++++++++++++++++++++++++ 2 files changed, 30 insertions(+) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index ebac8190d..d790fdaad 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -211,6 +211,7 @@ endif() if(Omega_h_USE_EGADSlite) target_link_libraries(omega_h PUBLIC ohEgadslite) + target_link_libraries(omega_h PUBLIC cuda) endif() bob_link_dependency(omega_h PUBLIC SEACASExodus) diff --git a/src/egads_lite_adapt_test.cpp b/src/egads_lite_adapt_test.cpp index 5ba41369f..db69be1fc 100644 --- a/src/egads_lite_adapt_test.cpp +++ b/src/egads_lite_adapt_test.cpp @@ -9,6 +9,8 @@ #include #endif +#include + #include static void compute_implied_metric(Omega_h::Mesh* mesh) { @@ -26,8 +28,34 @@ static void compute_target_metric(Omega_h::Mesh* mesh) { Omega_h::symm_ncomps(mesh->dim()), metrics); } +void checkCudaError(int line) { +#ifdef __NVCC__ + cudaError_t code = cudaDeviceSynchronize(); + const char * errorMessage = cudaGetErrorString(code); + if( code != cudaSuccess ) { + fprintf(stderr, "CUDA error on line %d Error code: %d (%s)\n", line, code, errorMessage); + } + assert(code == cudaSuccess); +#endif +} + +void setCudaStackSz() { + size_t stackLimit; + cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); + checkCudaError(__LINE__); + printf("original stack limit %d\n", stackLimit); + stackLimit=8*1024; + cuCtxSetLimit(CU_LIMIT_STACK_SIZE,stackLimit); + checkCudaError(__LINE__); + cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); + checkCudaError(__LINE__); + printf("new stack limit %d\n", stackLimit); + printf("stack limit %d\n", stackLimit); +} + int main(int argc, char** argv) { auto lib = Omega_h::Library(&argc, &argv); + setCudaStackSz(); Omega_h::CmdLine cmdline; cmdline.add_arg("mesh_in.meshb"); cmdline.add_arg("mesh_out.meshb"); @@ -58,6 +86,7 @@ int main(int argc, char** argv) { opts.egads_lite_model = eg; } #endif + fprintf(stderr, "numverts %d\n", mesh.nents(0)); auto has_viz = cmdline.parsed("--viz"); Omega_h::vtk::Writer writer; if (has_viz) { From fd730b7e9e365a34394b73df7c70a60e238a76ff Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 21 Feb 2022 20:53:14 -0500 Subject: [PATCH 27/40] egadslite: remove debug prints and vtx isolation --- src/Omega_h_egads_lite.cpp | 15 +++------------ 1 file changed, 3 insertions(+), 12 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 0bf767a44..5082bd613 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -435,18 +435,9 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto index2 = EG_indexBodyTopo(egBody_d, g); assert(index2 > 0); OMEGA_H_CHECK(index2 == index + 1); - int debug = 0; - if(i == 22 && index2 == 5 && class_dim == 2) { - debug=1; - int isEdge = (g->oclass == EGADS_EDGE); - int isFace = (g->oclass == EGADS_FACE); - printf("vtx %d class_id %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", - i, index2, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); - auto b = get_closest_point(g, a); - printf("clPt %.3f %.3f %.3f\n", b[0], b[1], b[2]); - clPt = b; - d = b - a; - } + auto b = get_closest_point(g, a); + clPt = b; + d = b - a; } set_vector(warp, i, d); set_vector(closePts, i, clPt); From fa55e2c717f384e9a0d4cbaa3975ae79d1626e2b Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 21 Feb 2022 22:21:04 -0500 Subject: [PATCH 28/40] egadslite: finish the previous commit... --- src/Omega_h_egads_lite.cpp | 18 +++++++++++++----- 1 file changed, 13 insertions(+), 5 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 5082bd613..85b9abffd 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -419,7 +419,6 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto egCounts_d = eg->counts_d; auto egBody_d = eg->body; auto calc_warp = OMEGA_H_LAMBDA(LO i) { - i=22; auto a = get_vector<3>(coords, i); Int class_dim = class_dims[i]; OMEGA_H_CHECK(class_dim >= 0); @@ -435,14 +434,23 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto index2 = EG_indexBodyTopo(egBody_d, g); assert(index2 > 0); OMEGA_H_CHECK(index2 == index + 1); - auto b = get_closest_point(g, a); - clPt = b; - d = b - a; + int debug = 0; + if(i == 22 && index2 == 5 && class_dim == 2) { + debug=1; + int isEdge = (g->oclass == EGADS_EDGE); + int isFace = (g->oclass == EGADS_FACE); + printf("vtx %d class_id %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", + i, index2, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); + auto b = get_closest_point(g, a); + printf("clPt %.3f %.3f %.3f\n", b[0], b[1], b[2]); + clPt = b; + d = b - a; + } } set_vector(warp, i, d); set_vector(closePts, i, clPt); }; - parallel_for(1, std::move(calc_warp), "calc_warp"); + parallel_for(mesh->nverts(), std::move(calc_warp), "calc_warp"); assert(cudaSuccess == cudaDeviceSynchronize()); auto t1 = now(); mesh->add_tag(0, "warpVec", 3, read(warp)); From 8dfdac6d8227c4563b5dde179e2dea6cf129f55c Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Mon, 21 Feb 2022 22:22:10 -0500 Subject: [PATCH 29/40] debug: write pt data to csv --- src/Omega_h_egads_lite.cpp | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 85b9abffd..6a4b1cc66 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -418,6 +418,7 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { egoPtrToGo(eg->entities_d[2])}; auto egCounts_d = eg->counts_d; auto egBody_d = eg->body; + printf("vtx,class_id,class_dim,isEdge,isFace,ptX,ptY,ptZ,clPtX,clPtY,clPtZ,warpX,warpY,warpZ\n"); auto calc_warp = OMEGA_H_LAMBDA(LO i) { auto a = get_vector<3>(coords, i); Int class_dim = class_dims[i]; @@ -434,17 +435,16 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto index2 = EG_indexBodyTopo(egBody_d, g); assert(index2 > 0); OMEGA_H_CHECK(index2 == index + 1); - int debug = 0; - if(i == 22 && index2 == 5 && class_dim == 2) { - debug=1; - int isEdge = (g->oclass == EGADS_EDGE); - int isFace = (g->oclass == EGADS_FACE); - printf("vtx %d class_id %d class_dim %d oclass %d isEdge %d isFace %d pt %.3f %.3f %.3f\n", - i, index2, class_dim, g->oclass, isEdge, isFace, a[0], a[1], a[2]); - auto b = get_closest_point(g, a); - printf("clPt %.3f %.3f %.3f\n", b[0], b[1], b[2]); - clPt = b; - d = b - a; + auto b = get_closest_point(g, a); + clPt = b; + d = b - a; + { + int isEdge = (g->oclass == EGADS_EDGE); + int isFace = (g->oclass == EGADS_FACE); + printf("%d,%d,%d,%d,%d,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f,%.3f\n", + i, index2, class_dim, isEdge, isFace, a[0], a[1], a[2], + clPt[0], clPt[1], clPt[2], + d[0],d[1],d[2]); } } set_vector(warp, i, d); From 3090ede488ca0d7c104fdacb8c919e132c6f8406 Mon Sep 17 00:00:00 2001 From: Cameron Smith Date: Wed, 23 Feb 2022 14:33:58 -0500 Subject: [PATCH 30/40] egadslite: hack classification for wacky topology --- src/egads_lite_adapt_test.cpp | 34 ++++++++++++++++++++++++++++++++++ 1 file changed, 34 insertions(+) diff --git a/src/egads_lite_adapt_test.cpp b/src/egads_lite_adapt_test.cpp index db69be1fc..1e54ffa13 100644 --- a/src/egads_lite_adapt_test.cpp +++ b/src/egads_lite_adapt_test.cpp @@ -4,6 +4,7 @@ #include #include #include +#include #ifdef OMEGA_H_USE_EGADSLITE #include @@ -39,6 +40,37 @@ void checkCudaError(int line) { #endif } +void hackClassification(Omega_h::Mesh* mesh) { + fprintf(stderr, "hacking classification\n"); + OMEGA_H_CHECK(mesh->dim() == 3); + auto vtx_class_dims = mesh->get_array(Omega_h::VERT, "class_dim"); + auto vtx_class_ids_r = mesh->get_array(Omega_h::VERT, "class_id"); + auto vtx_class_ids_w = Omega_h::deep_copy(vtx_class_ids_r, "vtxClassIds_w"); + auto setVtxClass = OMEGA_H_LAMBDA(int i) { + if(vtx_class_dims[i] == 1 && vtx_class_ids_w[i] == 1) { + printf("vtx %i reclassified\n",i); + vtx_class_ids_w[i] = 7; + } + }; + Omega_h::parallel_for(mesh->nents(0), setVtxClass, "setVtxClass"); + fprintf(stderr, "done hacking vtx classification\n"); + mesh->set_tag(0, "class_id", Omega_h::read(vtx_class_ids_w)); + + auto edge_class_dims = mesh->get_array(Omega_h::EDGE, "class_dim"); + auto edge_class_ids_r = mesh->get_array(Omega_h::EDGE, "class_id"); + auto edge_class_ids_w = Omega_h::deep_copy(edge_class_ids_r, "edgeClassIds_w"); + auto setEdgeClass = OMEGA_H_LAMBDA(int i) { + if(edge_class_dims[i] == 1 && edge_class_ids_w[i] == 1) { + printf("edge %i reclassified\n",i); + edge_class_ids_w[i] = 7; + } + }; + Omega_h::parallel_for(mesh->nents(1), setEdgeClass, "setEdgeClass"); + fprintf(stderr, "done hacking edge classification\n"); + mesh->set_tag(1, "class_id", Omega_h::read(edge_class_ids_w)); + +} + void setCudaStackSz() { size_t stackLimit; cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); @@ -84,6 +116,8 @@ int main(int argc, char** argv) { auto eg = Omega_h::egads_lite_load(model_path); Omega_h::egads_lite_reclassify(&mesh, eg); opts.egads_lite_model = eg; + hackClassification(&mesh); //there are problems... + Omega_h::vtk::write_parallel("postHack", &mesh, mesh.dim()); } #endif fprintf(stderr, "numverts %d\n", mesh.nents(0)); From 12cacb9e4053ad0f2cc48f94b50e8fc3bdde4412 Mon Sep 17 00:00:00 2001 From: joshia5 Date: Thu, 7 Apr 2022 12:18:17 -0400 Subject: [PATCH 31/40] fix compilation bug assumes the write_parallel was for debugging and not required for functionality --- src/Omega_h_egads_lite.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index 6a4b1cc66..bfe101775 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -404,7 +404,7 @@ OMEGA_H_INLINE Vector<3> get_closest_point(ego g, Vector<3> in) { Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { fprintf(stderr, "numverts %d\n", mesh->nverts()); - Omega_h::vtk::write_parallel("preWarp", mesh, mesh->dim()); + //Omega_h::vtk::write_parallel("preWarp", mesh, mesh->dim()); OMEGA_H_CHECK(mesh->dim() == 3); if (verbose) std::cout << "Querying closest points for surface vertices...\n"; auto t0 = now(); @@ -455,7 +455,7 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { auto t1 = now(); mesh->add_tag(0, "warpVec", 3, read(warp)); mesh->add_tag(0, "closePts", 3, read(closePts)); - Omega_h::vtk::write_parallel("warpVec", mesh, mesh->dim()); + //Omega_h::vtk::write_parallel("warpVec", mesh, mesh->dim()); if (verbose) { std::cout << "Querying closest points for surface vertices took " << (t1 - t0) << " seconds\n"; From 8930b2326e2a669daec4226d4b798c8f377871b8 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Fri, 31 Jul 2026 13:46:41 -0400 Subject: [PATCH 32/40] compile changes --- src/Omega_h_reduce.hpp | 1 + src/Omega_h_sort.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/src/Omega_h_reduce.hpp b/src/Omega_h_reduce.hpp index 166c19890..18fbeb5bf 100644 --- a/src/Omega_h_reduce.hpp +++ b/src/Omega_h_reduce.hpp @@ -15,6 +15,7 @@ #endif #include #include +#include #include #ifdef __GNUC__ #pragma GCC diagnostic pop diff --git a/src/Omega_h_sort.cpp b/src/Omega_h_sort.cpp index a0046849d..a117d79c3 100644 --- a/src/Omega_h_sort.cpp +++ b/src/Omega_h_sort.cpp @@ -116,7 +116,7 @@ LO number_same_values( auto const last = IntIterator(a.size()); auto const result = tmp_perm.begin() + 1; auto const op = plus(); - auto transform = OMEGA_H_LAMBDA(LO i)->LO { + auto transform = [=] __host__ __device__ (LO i)->LO { return a[i] == value ? LO(1) : LO(0); }; transform_inclusive_scan(first, last, result, op, std::move(transform)); From 421a7ba444e0315bc8434d1c9891e44f6fbef160 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Wed, 5 Aug 2026 18:12:51 -0400 Subject: [PATCH 33/40] fixed compile errors --- CMakeLists.txt | 5 +++++ src/CMakeLists.txt | 20 ++++---------------- 2 files changed, 9 insertions(+), 16 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 3c10f360f..7861180a5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -64,6 +64,11 @@ if (Omega_h_USE_Kokkos) set_target_properties(Kokkos::kokkos PROPERTIES INTERFACE_COMPILE_OPTIONS "") endif() +if (Omega_h_USE_EGADSlite) + enable_language(CUDA) + install(PROGRAMS ${CMAKE_CURRENT_SOURCE_DIR}/bin/nvcc_wrapper DESTINATION bin) +endif() + set(libMeshb_REQUIRED_VERSION 7.24) set(Omega_h_USE_libMeshb_DEFAULT OFF) bob_add_dependency(PUBLIC NAME libMeshb TARGETS libMeshb::Meshb.7) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 883063031..860addab7 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -187,20 +187,6 @@ if(Omega_h_USE_EGADSlite) install(TARGETS ohEgadslite EXPORT omega_h-target ARCHIVE) #HACK endif() -if(Omega_h_USE_EGADSlite) - find_package(egadslite REQUIRED) - add_library(ohEgadslite Omega_h_egads_lite.cpp) - set_source_files_properties(Omega_h_egads_lite.cpp PROPERTIES LANGUAGE CUDA) - set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) - set_target_properties(ohEgadslite PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) - target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite) - target_include_directories(ohEgadslite PRIVATE - ${CMAKE_CURRENT_SOURCE_DIR} - ${CMAKE_CURRENT_BINARY_DIR}) - target_compile_options(ohEgadslite PRIVATE "--expt-extended-lambda") - install(TARGETS ohEgadslite EXPORT omega_h-target ARCHIVE) #HACK -endif() - if(Omega_h_USE_SEACASExodus) list(APPEND Omega_h_SOURCES Omega_h_exodus.cpp) endif() @@ -236,7 +222,6 @@ if (Omega_h_USE_OpenMP) target_compile_options(omega_h PUBLIC -fopenmp) endif() - bob_link_dependency(omega_h PUBLIC Kokkos) bob_link_dependency(omega_h PUBLIC libMeshb) @@ -307,11 +292,14 @@ if(Omega_h_USE_CTAGS) endif() function(osh_add_exe EXE_NAME) + if (Omega_h_USE_EGADSlite) + set_source_files_properties(${EXE_NAME}.cpp PROPERTIES LANGUAGE CUDA) + endif() add_executable(${EXE_NAME} ${EXE_NAME}.cpp) set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD "17") set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD_REQUIRED ON) set_property(TARGET ${EXE_NAME} PROPERTY CXX_EXTENSIONS OFF) - set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) + # set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) target_link_libraries(${EXE_NAME} PRIVATE omega_h) if (OMP_EXTRA_LIBS) file(GLOB OMP_OBJS ${OMP_EXTRA_LIBS}) From 5b845c40bb9ab0b04c72e0c41fb6249bffc76477 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Thu, 6 Aug 2026 17:40:38 -0400 Subject: [PATCH 34/40] fixed runtime tests --- src/CMakeLists.txt | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 860addab7..b266f91a0 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -173,13 +173,15 @@ if(Omega_h_USE_ADIOS2) find_package(ADIOS2 REQUIRED) endif() +cmake_path(REPLACE_FILENAME MPI_CXX_COMPILER "../include" OUTPUT_VARIABLE MPI_INCLUDE) + if(Omega_h_USE_EGADSlite) find_package(egadslite REQUIRED) add_library(ohEgadslite Omega_h_egads_lite.cpp) set_source_files_properties(Omega_h_egads_lite.cpp PROPERTIES LANGUAGE CUDA) set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) - set_target_properties(ohEgadslite PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) - target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite) + target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite Kokkos::kokkos) + target_include_directories(ohEgadslite PRIVATE ${MPI_INCLUDE}) target_include_directories(ohEgadslite PRIVATE ${CMAKE_CURRENT_SOURCE_DIR} ${CMAKE_CURRENT_BINARY_DIR}) @@ -259,6 +261,7 @@ bob_link_dependency(omega_h PUBLIC ZLIB) if (Omega_h_USE_MPI) target_link_libraries(omega_h PUBLIC MPI::MPI_CXX) + target_include_directories(omega_h PRIVATE ${MPI_INCLUDE}) endif() if (Omega_h_USE_dwarf) @@ -299,7 +302,9 @@ function(osh_add_exe EXE_NAME) set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD "17") set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD_REQUIRED ON) set_property(TARGET ${EXE_NAME} PROPERTY CXX_EXTENSIONS OFF) - # set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) + set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) + set_property(TARGET ${EXE_NAME} PROPERTY CUDA_SEPARABLE_COMPILATION ON) + target_include_directories(${EXE_NAME} PRIVATE ${MPI_INCLUDE}) target_link_libraries(${EXE_NAME} PRIVATE omega_h) if (OMP_EXTRA_LIBS) file(GLOB OMP_OBJS ${OMP_EXTRA_LIBS}) From 4a76c1bcc2cbe4aebf058aacc6315084ea2629ce Mon Sep 17 00:00:00 2001 From: Angelyr Date: Fri, 14 Aug 2026 15:33:37 -0400 Subject: [PATCH 35/40] export egadslite --- CMakeLists.txt | 5 +++-- src/CMakeLists.txt | 17 ++++++++++------- src/Omega_h_adapt.cpp | 4 ++-- 3 files changed, 15 insertions(+), 11 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 7861180a5..2bd95d396 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -20,7 +20,7 @@ bob_option(Omega_h_CHECK_BOUNDS "Check array bounds when running on host (makes bob_option(Omega_h_THROW "Errors throw exceptions instead of abort" ${USE_XSDK_DEFAULTS}) bob_input(Omega_h_DATA "" PATH "Path to omega_h-data test files") bob_option(Omega_h_USE_EGADS "Use EGADS from ESP for geometry" OFF) -bob_option(Omega_h_USE_EGADSlite "Use EGADSlite from ESP for geometry" OFF) +bob_option(Omega_h_USE_egadslite "Use EGADSlite from ESP for geometry" OFF) bob_input(EGADS_PREFIX "" PATH "EGADS (or ESP) installation directory") bob_input(EGADSlite_PREFIX "" PATH "EGADSlite (or ESP) installation directory") bob_option(Omega_h_USE_Kokkos "Use Kokkos as a backend" OFF) @@ -64,8 +64,9 @@ if (Omega_h_USE_Kokkos) set_target_properties(Kokkos::kokkos PROPERTIES INTERFACE_COMPILE_OPTIONS "") endif() -if (Omega_h_USE_EGADSlite) +if (Omega_h_USE_egadslite) enable_language(CUDA) + bob_add_dependency(PUBLIC NAME egadslite TARGETS egadslite::egadslite) install(PROGRAMS ${CMAKE_CURRENT_SOURCE_DIR}/bin/nvcc_wrapper DESTINATION bin) endif() diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index b266f91a0..d2b880367 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -175,11 +175,12 @@ endif() cmake_path(REPLACE_FILENAME MPI_CXX_COMPILER "../include" OUTPUT_VARIABLE MPI_INCLUDE) -if(Omega_h_USE_EGADSlite) +if(Omega_h_USE_egadslite) find_package(egadslite REQUIRED) add_library(ohEgadslite Omega_h_egads_lite.cpp) set_source_files_properties(Omega_h_egads_lite.cpp PROPERTIES LANGUAGE CUDA) set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) + # set_target_properties(ohEgadslite PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite Kokkos::kokkos) target_include_directories(ohEgadslite PRIVATE ${MPI_INCLUDE}) target_include_directories(ohEgadslite PRIVATE @@ -230,12 +231,14 @@ bob_link_dependency(omega_h PUBLIC libMeshb) bob_link_dependency(omega_h PUBLIC Gmsh) +bob_link_dependency(omega_h PUBLIC egadslite) + if(Omega_h_USE_EGADS) target_include_directories(omega_h PUBLIC "${EGADS_INCLUDE_DIR}") target_link_libraries(omega_h PUBLIC "${EGADS_LIBRARY}") endif() -if(Omega_h_USE_EGADSlite) +if(Omega_h_USE_egadslite) target_link_libraries(omega_h PUBLIC ohEgadslite) target_link_libraries(omega_h PUBLIC cuda) endif() @@ -295,17 +298,17 @@ if(Omega_h_USE_CTAGS) endif() function(osh_add_exe EXE_NAME) - if (Omega_h_USE_EGADSlite) - set_source_files_properties(${EXE_NAME}.cpp PROPERTIES LANGUAGE CUDA) - endif() add_executable(${EXE_NAME} ${EXE_NAME}.cpp) set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD "17") set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD_REQUIRED ON) set_property(TARGET ${EXE_NAME} PROPERTY CXX_EXTENSIONS OFF) set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) - set_property(TARGET ${EXE_NAME} PROPERTY CUDA_SEPARABLE_COMPILATION ON) target_include_directories(${EXE_NAME} PRIVATE ${MPI_INCLUDE}) target_link_libraries(${EXE_NAME} PRIVATE omega_h) + if (Omega_h_USE_egadslite) + set_source_files_properties(${EXE_NAME}.cpp PROPERTIES LANGUAGE CUDA) + set_property(TARGET ${EXE_NAME} PROPERTY CUDA_SEPARABLE_COMPILATION ON) + endif() if (OMP_EXTRA_LIBS) file(GLOB OMP_OBJS ${OMP_EXTRA_LIBS}) message(STATUS "adding OMP_EXTRA_LIBS ${OMP_OBJS} to omega_h") @@ -671,7 +674,7 @@ if(BUILD_TESTING) if(Omega_h_USE_EGADS) osh_add_exe(egads_test) endif() - if(Omega_h_USE_EGADSlite) + if(Omega_h_USE_egadslite) osh_add_exe(egads_lite_adapt_test) endif() osh_add_exe(advect2d_test) diff --git a/src/Omega_h_adapt.cpp b/src/Omega_h_adapt.cpp index 9aa9d92a7..ed0eeb453 100644 --- a/src/Omega_h_adapt.cpp +++ b/src/Omega_h_adapt.cpp @@ -82,7 +82,7 @@ AdaptOpts::AdaptOpts(Int dim) { #ifdef OMEGA_H_USE_EGADS egads_model = nullptr; #endif -#ifdef OMEGA_H_USE_EGADSlite +#ifdef OMEGA_H_USE_EGADSLITE egads_lite_model = nullptr; #endif should_refine = true; @@ -299,7 +299,7 @@ static void post_adapt( #ifdef OMEGA_H_USE_EGADS if (opts.egads_model) std::cout << "snapping while "; #endif -#ifdef OMEGA_H_USE_EGADSlite +#ifdef OMEGA_H_USE_EGADSLITE if (opts.egads_lite_model) std::cout << "snapping while "; #endif std::cout << "addressing element qualities took " << (t3 - t2); From d8bc4060db19e39055d5572ca88939d984e5cad2 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Thu, 27 Aug 2026 13:48:44 -0400 Subject: [PATCH 36/40] export egadslite --- CMakeLists.txt | 2 +- src/CMakeLists.txt | 1 + 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 2bd95d396..aabdff69a 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -159,7 +159,7 @@ set(Omega_h_KEY_BOOLS Omega_h_USE_ZLIB Omega_h_USE_libMeshb Omega_h_USE_EGADS - Omega_h_USE_EGADSlite + Omega_h_USE_egadslite Omega_h_USE_SEACASExodus Omega_h_USE_SimModSuite Omega_h_USE_SimDiscrete diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index d2b880367..e7649e167 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -803,6 +803,7 @@ set(Omega_h_HEADERS Omega_h_dist.hpp Omega_h_eigen.hpp Omega_h_egads.hpp + Omega_h_egads_lite.hpp Omega_h_element.hpp Omega_h_expr.hpp Omega_h_fail.hpp From 0cec0feeb899911a83b35e7db0ba27539c222215 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Fri, 28 Aug 2026 19:00:46 -0400 Subject: [PATCH 37/40] removed unnecessary --- CMakeLists.txt | 1 - src/CMakeLists.txt | 8 +------- 2 files changed, 1 insertion(+), 8 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index aabdff69a..a17e3cd81 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -65,7 +65,6 @@ if (Omega_h_USE_Kokkos) endif() if (Omega_h_USE_egadslite) - enable_language(CUDA) bob_add_dependency(PUBLIC NAME egadslite TARGETS egadslite::egadslite) install(PROGRAMS ${CMAKE_CURRENT_SOURCE_DIR}/bin/nvcc_wrapper DESTINATION bin) endif() diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index e7649e167..3efedf15d 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -178,9 +178,7 @@ cmake_path(REPLACE_FILENAME MPI_CXX_COMPILER "../include" OUTPUT_VARIABLE MPI_IN if(Omega_h_USE_egadslite) find_package(egadslite REQUIRED) add_library(ohEgadslite Omega_h_egads_lite.cpp) - set_source_files_properties(Omega_h_egads_lite.cpp PROPERTIES LANGUAGE CUDA) set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) - # set_target_properties(ohEgadslite PROPERTIES CUDA_RESOLVE_DEVICE_SYMBOLS ON) target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite Kokkos::kokkos) target_include_directories(ohEgadslite PRIVATE ${MPI_INCLUDE}) target_include_directories(ohEgadslite PRIVATE @@ -240,7 +238,6 @@ endif() if(Omega_h_USE_egadslite) target_link_libraries(omega_h PUBLIC ohEgadslite) - target_link_libraries(omega_h PUBLIC cuda) endif() if(Omega_h_USE_SimModSuite) @@ -305,10 +302,6 @@ function(osh_add_exe EXE_NAME) set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) target_include_directories(${EXE_NAME} PRIVATE ${MPI_INCLUDE}) target_link_libraries(${EXE_NAME} PRIVATE omega_h) - if (Omega_h_USE_egadslite) - set_source_files_properties(${EXE_NAME}.cpp PROPERTIES LANGUAGE CUDA) - set_property(TARGET ${EXE_NAME} PROPERTY CUDA_SEPARABLE_COMPILATION ON) - endif() if (OMP_EXTRA_LIBS) file(GLOB OMP_OBJS ${OMP_EXTRA_LIBS}) message(STATUS "adding OMP_EXTRA_LIBS ${OMP_OBJS} to omega_h") @@ -676,6 +669,7 @@ if(BUILD_TESTING) endif() if(Omega_h_USE_egadslite) osh_add_exe(egads_lite_adapt_test) + set_property(TARGET egads_lite_adapt_test PROPERTY CUDA_SEPARABLE_COMPILATION ON) endif() osh_add_exe(advect2d_test) if(Omega_h_DATA) From 758aa4b3903b7d9cff9b4298c60224ff78443359 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Fri, 28 Aug 2026 20:28:09 -0400 Subject: [PATCH 38/40] simplify build --- CMakeLists.txt | 7 ++----- cmake/bob.cmake | 3 +-- src/CMakeLists.txt | 20 +------------------- 3 files changed, 4 insertions(+), 26 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index a17e3cd81..f750d2243 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -22,7 +22,6 @@ bob_input(Omega_h_DATA "" PATH "Path to omega_h-data test files") bob_option(Omega_h_USE_EGADS "Use EGADS from ESP for geometry" OFF) bob_option(Omega_h_USE_egadslite "Use EGADSlite from ESP for geometry" OFF) bob_input(EGADS_PREFIX "" PATH "EGADS (or ESP) installation directory") -bob_input(EGADSlite_PREFIX "" PATH "EGADSlite (or ESP) installation directory") bob_option(Omega_h_USE_Kokkos "Use Kokkos as a backend" OFF) bob_input(Kokkos_PREFIX "" PATH "Path to Kokkos install") bob_option(OMEGA_H_USE_GPU_AWARE_MPI "Assume MPI is GPU-aware, make use of that" OFF) @@ -64,10 +63,8 @@ if (Omega_h_USE_Kokkos) set_target_properties(Kokkos::kokkos PROPERTIES INTERFACE_COMPILE_OPTIONS "") endif() -if (Omega_h_USE_egadslite) - bob_add_dependency(PUBLIC NAME egadslite TARGETS egadslite::egadslite) - install(PROGRAMS ${CMAKE_CURRENT_SOURCE_DIR}/bin/nvcc_wrapper DESTINATION bin) -endif() +set(Omega_h_USE_egadslite_DEFAULT OFF) +bob_add_dependency(PUBLIC NAME egadslite TARGETS egadslite::egadslite) set(libMeshb_REQUIRED_VERSION 7.24) set(Omega_h_USE_libMeshb_DEFAULT OFF) diff --git a/cmake/bob.cmake b/cmake/bob.cmake index f6dbdef86..6c57b5a65 100644 --- a/cmake/bob.cmake +++ b/cmake/bob.cmake @@ -506,8 +506,7 @@ function(bob_get_link_libs tgt var) endif() set(link_libs) foreach(lib IN LISTS sublibs) - if (${lib} STREQUAL "ohEgadslite") #HACK - elseif (TARGET ${lib}) + if (TARGET ${lib}) get_target_property(subtgt_type "${lib}" TYPE) if (subtgt_type MATCHES "STATIC_LIBRARY|SHARED_LIBRARY") get_target_property(sublibtgt_loc "${lib}" LOCATION) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 3efedf15d..32b90b816 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -173,19 +173,8 @@ if(Omega_h_USE_ADIOS2) find_package(ADIOS2 REQUIRED) endif() -cmake_path(REPLACE_FILENAME MPI_CXX_COMPILER "../include" OUTPUT_VARIABLE MPI_INCLUDE) - if(Omega_h_USE_egadslite) - find_package(egadslite REQUIRED) - add_library(ohEgadslite Omega_h_egads_lite.cpp) - set_target_properties(ohEgadslite PROPERTIES CUDA_SEPARABLE_COMPILATION ON) - target_link_libraries(ohEgadslite PUBLIC egadslite::egadslite Kokkos::kokkos) - target_include_directories(ohEgadslite PRIVATE ${MPI_INCLUDE}) - target_include_directories(ohEgadslite PRIVATE - ${CMAKE_CURRENT_SOURCE_DIR} - ${CMAKE_CURRENT_BINARY_DIR}) - target_compile_options(ohEgadslite PRIVATE "--expt-extended-lambda") - install(TARGETS ohEgadslite EXPORT omega_h-target ARCHIVE) #HACK + list(APPEND Omega_h_SOURCES Omega_h_egads_lite.cpp) endif() if(Omega_h_USE_SEACASExodus) @@ -236,10 +225,6 @@ if(Omega_h_USE_EGADS) target_link_libraries(omega_h PUBLIC "${EGADS_LIBRARY}") endif() -if(Omega_h_USE_egadslite) - target_link_libraries(omega_h PUBLIC ohEgadslite) -endif() - if(Omega_h_USE_SimModSuite) target_include_directories(omega_h PUBLIC "${SIMMODSUITE_INCLUDE_DIR}") target_link_libraries(omega_h PUBLIC "${SIMMODSUITE_LIBS}") @@ -261,7 +246,6 @@ bob_link_dependency(omega_h PUBLIC ZLIB) if (Omega_h_USE_MPI) target_link_libraries(omega_h PUBLIC MPI::MPI_CXX) - target_include_directories(omega_h PRIVATE ${MPI_INCLUDE}) endif() if (Omega_h_USE_dwarf) @@ -300,7 +284,6 @@ function(osh_add_exe EXE_NAME) set_property(TARGET ${EXE_NAME} PROPERTY CXX_STANDARD_REQUIRED ON) set_property(TARGET ${EXE_NAME} PROPERTY CXX_EXTENSIONS OFF) set_property(TARGET ${EXE_NAME} PROPERTY CUDA_ARCHITECTURES ${Omega_h_CUDA_ARCH}) - target_include_directories(${EXE_NAME} PRIVATE ${MPI_INCLUDE}) target_link_libraries(${EXE_NAME} PRIVATE omega_h) if (OMP_EXTRA_LIBS) file(GLOB OMP_OBJS ${OMP_EXTRA_LIBS}) @@ -669,7 +652,6 @@ if(BUILD_TESTING) endif() if(Omega_h_USE_egadslite) osh_add_exe(egads_lite_adapt_test) - set_property(TARGET egads_lite_adapt_test PROPERTY CUDA_SEPARABLE_COMPILATION ON) endif() osh_add_exe(advect2d_test) if(Omega_h_DATA) From ec55331a029a24dc52d9a99d7202e039661b9bf9 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Fri, 28 Aug 2026 20:52:12 -0400 Subject: [PATCH 39/40] removed duplicate code --- src/Omega_h_adapt.cpp | 46 ++++++----------------------------- src/Omega_h_adapt.hpp | 5 +--- src/egads_lite_adapt_test.cpp | 4 +-- 3 files changed, 11 insertions(+), 44 deletions(-) diff --git a/src/Omega_h_adapt.cpp b/src/Omega_h_adapt.cpp index ed0eeb453..6b722c13c 100644 --- a/src/Omega_h_adapt.cpp +++ b/src/Omega_h_adapt.cpp @@ -79,11 +79,8 @@ AdaptOpts::AdaptOpts(Int dim) { should_smooth_snap = true; snap_smooth_tolerance = 1e-2; allow_snap_failure = false; -#ifdef OMEGA_H_USE_EGADS +#if defined(OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) egads_model = nullptr; -#endif -#ifdef OMEGA_H_USE_EGADSLITE - egads_lite_model = nullptr; #endif should_refine = true; should_coarsen = true; @@ -218,7 +215,7 @@ static bool satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { } static void snap_and_satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { -#ifdef OMEGA_H_USE_EGADS +#if defined(OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) if (opts.egads_model) { ScopedTimer snap_timer("snap"); @@ -228,38 +225,14 @@ static void snap_and_satisfy_quality(Mesh* mesh, AdaptOpts const& opts) { //mesh->change_all_rcFieldsToMesh(); + #ifdef OMEGA_H_USE_EGADSLITE + auto warp = egads_lite_get_snap_warp( + mesh, opts.egads_model, opts.verbosity >= EACH_REBUILD); + #else auto warp = egads_get_snap_warp( mesh, opts.egads_model, opts.verbosity >= EACH_REBUILD); - if (opts.should_smooth_snap) { - if (opts.verbosity >= EACH_REBUILD) { - std::cout << "Solving Laplacian of warp field...\n"; - } - auto t0 = now(); - warp = - solve_laplacian(mesh, warp, mesh->dim(), opts.snap_smooth_tolerance); - auto t1 = now(); - if (opts.verbosity >= EACH_REBUILD) { - std::cout << "Solving Laplacian of warp field took " << (t1 - t0) - << " seconds\n"; - } - } - mesh->add_tag(VERT, "warp", mesh->dim(), warp); - while (warp_to_limit(mesh, opts, opts.allow_snap_failure)) { - if (!satisfy_quality(mesh, opts)) { - mesh->remove_tag(VERT, "warp"); - break; - } - } - } else -#endif -#ifdef OMEGA_H_USE_EGADSLITE - if (opts.egads_lite_model) { - ScopedTimer snap_timer("snap"); - - mesh->set_parting(OMEGA_H_GHOSTED); + #endif - auto warp = egads_lite_get_snap_warp( - mesh, opts.egads_lite_model, opts.verbosity >= EACH_REBUILD); if (opts.should_smooth_snap) { if (opts.verbosity >= EACH_REBUILD) { std::cout << "Solving Laplacian of warp field...\n"; @@ -296,11 +269,8 @@ static void post_adapt( std::cout << "addressing edge lengths took " << (t2 - t1) << " seconds\n"; } if (opts.verbosity > SILENT && !mesh->comm()->rank()) { -#ifdef OMEGA_H_USE_EGADS +#if defined(OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) if (opts.egads_model) std::cout << "snapping while "; -#endif -#ifdef OMEGA_H_USE_EGADSLITE - if (opts.egads_lite_model) std::cout << "snapping while "; #endif std::cout << "addressing element qualities took " << (t3 - t2); std::cout << " seconds\n"; diff --git a/src/Omega_h_adapt.hpp b/src/Omega_h_adapt.hpp index 273dc0f21..4eb9c348a 100644 --- a/src/Omega_h_adapt.hpp +++ b/src/Omega_h_adapt.hpp @@ -171,11 +171,8 @@ struct AdaptOpts { bool should_smooth_snap; Real snap_smooth_tolerance; bool allow_snap_failure; -#ifdef OMEGA_H_USE_EGADS +#if defined(OMEGA_H_USE_EGADS) || defined(OMEGA_H_USE_EGADSLITE) Egads* egads_model; -#endif -#ifdef OMEGA_H_USE_EGADSLITE - Egads* egads_lite_model; #endif bool should_refine; bool should_coarsen; diff --git a/src/egads_lite_adapt_test.cpp b/src/egads_lite_adapt_test.cpp index 1e54ffa13..dd885ce87 100644 --- a/src/egads_lite_adapt_test.cpp +++ b/src/egads_lite_adapt_test.cpp @@ -115,7 +115,7 @@ int main(int argc, char** argv) { std::cout << "reading in " << model_path << '\n'; auto eg = Omega_h::egads_lite_load(model_path); Omega_h::egads_lite_reclassify(&mesh, eg); - opts.egads_lite_model = eg; + opts.egads_model = eg; hackClassification(&mesh); //there are problems... Omega_h::vtk::write_parallel("postHack", &mesh, mesh.dim()); } @@ -139,7 +139,7 @@ int main(int argc, char** argv) { Omega_h::meshb::write(&mesh, path_out); #ifdef OMEGA_H_USE_EGADSLITE if (has_model) { - Omega_h::egads_lite_free(opts.egads_lite_model); + Omega_h::egads_lite_free(opts.egads_model); } #endif } From 195888ec450d2d7e02aa4f88d7d533d8559933c0 Mon Sep 17 00:00:00 2001 From: Angelyr Date: Fri, 28 Aug 2026 21:24:02 -0400 Subject: [PATCH 40/40] reuse classification and cudastacksize --- src/Omega_h_egads_lite.cpp | 56 ++++++++++++++++++++++++++++++++ src/Omega_h_egads_lite.hpp | 2 ++ src/egads_lite_adapt_test.cpp | 60 ++--------------------------------- 3 files changed, 60 insertions(+), 58 deletions(-) diff --git a/src/Omega_h_egads_lite.cpp b/src/Omega_h_egads_lite.cpp index bfe101775..902d06a7a 100644 --- a/src/Omega_h_egads_lite.cpp +++ b/src/Omega_h_egads_lite.cpp @@ -463,4 +463,60 @@ Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose) { return warp; } +void checkCudaError(int line) { +#ifdef __NVCC__ + cudaError_t code = cudaDeviceSynchronize(); + const char * errorMessage = cudaGetErrorString(code); + if( code != cudaSuccess ) { + fprintf(stderr, "CUDA error on line %d Error code: %d (%s)\n", line, code, errorMessage); + } + assert(code == cudaSuccess); +#endif +} + +void hackClassification(Omega_h::Mesh* mesh) { + fprintf(stderr, "hacking classification\n"); + OMEGA_H_CHECK(mesh->dim() == 3); + auto vtx_class_dims = mesh->get_array(Omega_h::VERT, "class_dim"); + auto vtx_class_ids_r = mesh->get_array(Omega_h::VERT, "class_id"); + auto vtx_class_ids_w = Omega_h::deep_copy(vtx_class_ids_r, "vtxClassIds_w"); + auto setVtxClass = OMEGA_H_LAMBDA(int i) { + if(vtx_class_dims[i] == 1 && vtx_class_ids_w[i] == 1) { + printf("vtx %i reclassified\n",i); + vtx_class_ids_w[i] = 7; + } + }; + Omega_h::parallel_for(mesh->nents(0), setVtxClass, "setVtxClass"); + fprintf(stderr, "done hacking vtx classification\n"); + mesh->set_tag(0, "class_id", Omega_h::read(vtx_class_ids_w)); + + auto edge_class_dims = mesh->get_array(Omega_h::EDGE, "class_dim"); + auto edge_class_ids_r = mesh->get_array(Omega_h::EDGE, "class_id"); + auto edge_class_ids_w = Omega_h::deep_copy(edge_class_ids_r, "edgeClassIds_w"); + auto setEdgeClass = OMEGA_H_LAMBDA(int i) { + if(edge_class_dims[i] == 1 && edge_class_ids_w[i] == 1) { + printf("edge %i reclassified\n",i); + edge_class_ids_w[i] = 7; + } + }; + Omega_h::parallel_for(mesh->nents(1), setEdgeClass, "setEdgeClass"); + fprintf(stderr, "done hacking edge classification\n"); + mesh->set_tag(1, "class_id", Omega_h::read(edge_class_ids_w)); + +} + +void setCudaStackSz() { + size_t stackLimit; + cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); + checkCudaError(__LINE__); + printf("original stack limit %d\n", stackLimit); + stackLimit=8*1024; + cuCtxSetLimit(CU_LIMIT_STACK_SIZE,stackLimit); + checkCudaError(__LINE__); + cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); + checkCudaError(__LINE__); + printf("new stack limit %d\n", stackLimit); + printf("stack limit %d\n", stackLimit); +} + } // namespace Omega_h diff --git a/src/Omega_h_egads_lite.hpp b/src/Omega_h_egads_lite.hpp index 64a61c576..948821657 100644 --- a/src/Omega_h_egads_lite.hpp +++ b/src/Omega_h_egads_lite.hpp @@ -15,6 +15,8 @@ void egads_lite_classify(Egads* eg, int nadj_faces, int const adj_face_ids[], void egads_lite_free(Egads* eg); void egads_lite_reclassify(Mesh* mesh, Egads* eg); Reals egads_lite_get_snap_warp(Mesh* mesh, Egads* eg, bool verbose); +void hackClassification(Omega_h::Mesh* mesh); +void setCudaStackSz(); } // namespace Omega_h diff --git a/src/egads_lite_adapt_test.cpp b/src/egads_lite_adapt_test.cpp index dd885ce87..da180129f 100644 --- a/src/egads_lite_adapt_test.cpp +++ b/src/egads_lite_adapt_test.cpp @@ -29,65 +29,9 @@ static void compute_target_metric(Omega_h::Mesh* mesh) { Omega_h::symm_ncomps(mesh->dim()), metrics); } -void checkCudaError(int line) { -#ifdef __NVCC__ - cudaError_t code = cudaDeviceSynchronize(); - const char * errorMessage = cudaGetErrorString(code); - if( code != cudaSuccess ) { - fprintf(stderr, "CUDA error on line %d Error code: %d (%s)\n", line, code, errorMessage); - } - assert(code == cudaSuccess); -#endif -} - -void hackClassification(Omega_h::Mesh* mesh) { - fprintf(stderr, "hacking classification\n"); - OMEGA_H_CHECK(mesh->dim() == 3); - auto vtx_class_dims = mesh->get_array(Omega_h::VERT, "class_dim"); - auto vtx_class_ids_r = mesh->get_array(Omega_h::VERT, "class_id"); - auto vtx_class_ids_w = Omega_h::deep_copy(vtx_class_ids_r, "vtxClassIds_w"); - auto setVtxClass = OMEGA_H_LAMBDA(int i) { - if(vtx_class_dims[i] == 1 && vtx_class_ids_w[i] == 1) { - printf("vtx %i reclassified\n",i); - vtx_class_ids_w[i] = 7; - } - }; - Omega_h::parallel_for(mesh->nents(0), setVtxClass, "setVtxClass"); - fprintf(stderr, "done hacking vtx classification\n"); - mesh->set_tag(0, "class_id", Omega_h::read(vtx_class_ids_w)); - - auto edge_class_dims = mesh->get_array(Omega_h::EDGE, "class_dim"); - auto edge_class_ids_r = mesh->get_array(Omega_h::EDGE, "class_id"); - auto edge_class_ids_w = Omega_h::deep_copy(edge_class_ids_r, "edgeClassIds_w"); - auto setEdgeClass = OMEGA_H_LAMBDA(int i) { - if(edge_class_dims[i] == 1 && edge_class_ids_w[i] == 1) { - printf("edge %i reclassified\n",i); - edge_class_ids_w[i] = 7; - } - }; - Omega_h::parallel_for(mesh->nents(1), setEdgeClass, "setEdgeClass"); - fprintf(stderr, "done hacking edge classification\n"); - mesh->set_tag(1, "class_id", Omega_h::read(edge_class_ids_w)); - -} - -void setCudaStackSz() { - size_t stackLimit; - cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); - checkCudaError(__LINE__); - printf("original stack limit %d\n", stackLimit); - stackLimit=8*1024; - cuCtxSetLimit(CU_LIMIT_STACK_SIZE,stackLimit); - checkCudaError(__LINE__); - cuCtxGetLimit(&stackLimit, CU_LIMIT_STACK_SIZE); - checkCudaError(__LINE__); - printf("new stack limit %d\n", stackLimit); - printf("stack limit %d\n", stackLimit); -} - int main(int argc, char** argv) { auto lib = Omega_h::Library(&argc, &argv); - setCudaStackSz(); + Omega_h::setCudaStackSz(); Omega_h::CmdLine cmdline; cmdline.add_arg("mesh_in.meshb"); cmdline.add_arg("mesh_out.meshb"); @@ -116,7 +60,7 @@ int main(int argc, char** argv) { auto eg = Omega_h::egads_lite_load(model_path); Omega_h::egads_lite_reclassify(&mesh, eg); opts.egads_model = eg; - hackClassification(&mesh); //there are problems... + Omega_h::hackClassification(&mesh); //there are problems... Omega_h::vtk::write_parallel("postHack", &mesh, mesh.dim()); } #endif