From b28f3c2db34740df3884da575e975c4aa2200623 Mon Sep 17 00:00:00 2001 From: ahmed Date: Wed, 9 Oct 2024 15:59:57 -0400 Subject: [PATCH] simplify reading attributes to glm --- apps/ARAP/arap.cu | 8 +++--- apps/Delaunay/delaunay_rxmesh.cuh | 8 +++--- apps/Delaunay/mcf_rxmesh_kernel.cuh | 14 +++++------ .../gaussian_curvature_kernel.cuh | 6 ++--- apps/Geodesic/geodesic_kernel.cuh | 7 +++--- apps/MCF/mcf_kernels.cuh | 14 +++++------ apps/Remesh/collapse.cuh | 25 +++++++++---------- apps/Remesh/smoothing.cuh | 7 +++--- apps/Remesh/split.cuh | 21 ++++++++-------- apps/Remesh/util.cuh | 4 +-- apps/SCP/scp.cu | 8 +++--- apps/SECHistogram/sec_kernels.cuh | 12 ++++----- apps/VertexNormal/vertex_normal.cu | 2 +- apps/VertexNormal/vertex_normal_kernel.cuh | 6 ++--- apps/XPBD/xpbd.cu | 8 +++--- include/rxmesh/attribute.h | 17 +++++++++++++ include/rxmesh/types.h | 23 ++++++++++------- 17 files changed, 106 insertions(+), 84 deletions(-) diff --git a/apps/ARAP/arap.cu b/apps/ARAP/arap.cu index c8fff3d9..8cff74e4 100644 --- a/apps/ARAP/arap.cu +++ b/apps/ARAP/arap.cu @@ -30,10 +30,10 @@ __global__ static void calc_edge_weights_mat( return; } - const vec3 p(coords(p_id, 0), coords(p_id, 1), coords(p_id, 2)); - const vec3 r(coords(r_id, 0), coords(r_id, 1), coords(r_id, 2)); - const vec3 q(coords(q_id, 0), coords(q_id, 1), coords(q_id, 2)); - const vec3 s(coords(s_id, 0), coords(s_id, 1), coords(s_id, 2)); + const vec3 p = coords.to_glm<3>(p_id); + const vec3 r = coords.to_glm<3>(r_id); + const vec3 q = coords.to_glm<3>(q_id); + const vec3 s = coords.to_glm<3>(s_id); // cotans[(v1, v2)] =np.dot(e1, e2) / np.linalg.norm(np.cross(e1, e2)) diff --git a/apps/Delaunay/delaunay_rxmesh.cuh b/apps/Delaunay/delaunay_rxmesh.cuh index 113691d7..eb9066b8 100644 --- a/apps/Delaunay/delaunay_rxmesh.cuh +++ b/apps/Delaunay/delaunay_rxmesh.cuh @@ -67,10 +67,10 @@ __global__ static void __launch_bounds__(blockThreads) constexpr T PII = 3.14159265358979323f; - const vec3 V0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 V1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); - const vec3 V2(coords(v2, 0), coords(v2, 1), coords(v2, 2)); - const vec3 V3(coords(v3, 0), coords(v3, 1), coords(v3, 2)); + const vec3 V0 = coords.to_glm<3>(v0); + const vec3 V1 = coords.to_glm<3>(v1); + const vec3 V2 = coords.to_glm<3>(v2); + const vec3 V3 = coords.to_glm<3>(v3); // find the angle between S, M, Q vertices (i.e., angle at M) auto angle_between_three_vertices = [](const vec3& S, diff --git a/apps/Delaunay/mcf_rxmesh_kernel.cuh b/apps/Delaunay/mcf_rxmesh_kernel.cuh index 990f0277..4aa2326c 100644 --- a/apps/Delaunay/mcf_rxmesh_kernel.cuh +++ b/apps/Delaunay/mcf_rxmesh_kernel.cuh @@ -20,10 +20,10 @@ edge_cotan_weight(const rxmesh::VertexHandle& p_id, // q and s composes the diamond around p-r using namespace rxmesh; - const vec3 p(X(p_id, 0), X(p_id, 1), X(p_id, 2)); - const vec3 r(X(r_id, 0), X(r_id, 1), X(r_id, 2)); - const vec3 q(X(q_id, 0), X(q_id, 1), X(q_id, 2)); - const vec3 s(X(s_id, 0), X(s_id, 1), X(s_id, 2)); + const vec3 p = X.to_glm<3>(p_id); + const vec3 r = X.to_glm<3>(r_id); + const vec3 q = X.to_glm<3>(q_id); + const vec3 s = X.to_glm<3>(s_id); return edge_cotan_weight(p, r, q, s); } @@ -42,9 +42,9 @@ partial_voronoi_area(const rxmesh::VertexHandle& p_id, // center // the triangle p->q->r (oriented ccw) using namespace rxmesh; - const vec3 p(X(p_id, 0), X(p_id, 1), X(p_id, 2)); - const vec3 q(X(q_id, 0), X(q_id, 1), X(q_id, 2)); - const vec3 r(X(r_id, 0), X(r_id, 1), X(r_id, 2)); + const vec3 p = X.to_glm<3>(p_id); + const vec3 q = X.to_glm<3>(q_id); + const vec3 r = X.to_glm<3>(r_id); return partial_voronoi_area(p, q, r); } diff --git a/apps/GaussianCurvature/gaussian_curvature_kernel.cuh b/apps/GaussianCurvature/gaussian_curvature_kernel.cuh index c1a265ae..22135d74 100644 --- a/apps/GaussianCurvature/gaussian_curvature_kernel.cuh +++ b/apps/GaussianCurvature/gaussian_curvature_kernel.cuh @@ -17,9 +17,9 @@ __global__ static void compute_gaussian_curvature( auto gc_lambda = [&](FaceHandle face_id, VertexIterator& fv) { // get the face's three vertices coordinates - vec3 c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2)); - vec3 c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2)); - vec3 c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2)); + const vec3 c0 = coords.to_glm<3>(fv[0]); + const vec3 c1 = coords.to_glm<3>(fv[1]); + const vec3 c2 = coords.to_glm<3>(fv[2]); // the three edges length vec3 l(glm::distance2(c0, c1), diff --git a/apps/Geodesic/geodesic_kernel.cuh b/apps/Geodesic/geodesic_kernel.cuh index 3de8d5ca..5b34fe59 100644 --- a/apps/Geodesic/geodesic_kernel.cuh +++ b/apps/Geodesic/geodesic_kernel.cuh @@ -18,9 +18,10 @@ __device__ __inline__ T update_step( const T infinity_val) { using namespace rxmesh; - const vec3 v0(coords(v0_id, 0), coords(v0_id, 1), coords(v0_id, 2)); - const vec3 v1(coords(v1_id, 0), coords(v1_id, 1), coords(v1_id, 2)); - const vec3 v2(coords(v2_id, 0), coords(v2_id, 1), coords(v2_id, 2)); + const vec3 v0 = coords.to_glm<3>(v0_id); + const vec3 v1 = coords.to_glm<3>(v1_id); + const vec3 v2 = coords.to_glm<3>(v2_id); + const vec3 x0 = v1 - v0; const vec3 x1 = v2 - v0; diff --git a/apps/MCF/mcf_kernels.cuh b/apps/MCF/mcf_kernels.cuh index 990f0277..4aa2326c 100644 --- a/apps/MCF/mcf_kernels.cuh +++ b/apps/MCF/mcf_kernels.cuh @@ -20,10 +20,10 @@ edge_cotan_weight(const rxmesh::VertexHandle& p_id, // q and s composes the diamond around p-r using namespace rxmesh; - const vec3 p(X(p_id, 0), X(p_id, 1), X(p_id, 2)); - const vec3 r(X(r_id, 0), X(r_id, 1), X(r_id, 2)); - const vec3 q(X(q_id, 0), X(q_id, 1), X(q_id, 2)); - const vec3 s(X(s_id, 0), X(s_id, 1), X(s_id, 2)); + const vec3 p = X.to_glm<3>(p_id); + const vec3 r = X.to_glm<3>(r_id); + const vec3 q = X.to_glm<3>(q_id); + const vec3 s = X.to_glm<3>(s_id); return edge_cotan_weight(p, r, q, s); } @@ -42,9 +42,9 @@ partial_voronoi_area(const rxmesh::VertexHandle& p_id, // center // the triangle p->q->r (oriented ccw) using namespace rxmesh; - const vec3 p(X(p_id, 0), X(p_id, 1), X(p_id, 2)); - const vec3 q(X(q_id, 0), X(q_id, 1), X(q_id, 2)); - const vec3 r(X(r_id, 0), X(r_id, 1), X(r_id, 2)); + const vec3 p = X.to_glm<3>(p_id); + const vec3 q = X.to_glm<3>(q_id); + const vec3 r = X.to_glm<3>(r_id); return partial_voronoi_area(p, q, r); } diff --git a/apps/Remesh/collapse.cuh b/apps/Remesh/collapse.cuh index 8340b1f1..c34006bf 100644 --- a/apps/Remesh/collapse.cuh +++ b/apps/Remesh/collapse.cuh @@ -79,8 +79,8 @@ __global__ static void __launch_bounds__(blockThreads) v2 == v3) { return; } - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); const T edge_len_sq = glm::distance2(p0, p1); if (edge_len_sq < low_edge_len_sq) { @@ -124,8 +124,8 @@ __global__ static void __launch_bounds__(blockThreads) cavity.get_vertices(src, v0, v1); - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); const vec3 new_p((p0[0] + p1[0]) * T(0.5), (p0[1] + p1[1]) * T(0.5), @@ -139,8 +139,7 @@ __global__ static void __launch_bounds__(blockThreads) const VertexHandle vvv = cavity.get_cavity_vertex(c, i); - const vec3 vp( - coords(vvv, 0), coords(vvv, 1), coords(vvv, 2)); + const vec3 vp = coords.to_glm<3>(vvv); const T edge_len_sq = glm::distance2(vp, new_p); @@ -300,9 +299,10 @@ __global__ static void __launch_bounds__(blockThreads) return; } - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); - const T edge_len_sq = glm::distance2(p0, p1); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); + + const T edge_len_sq = glm::distance2(p0, p1); if (edge_len_sq < low_edge_len_sq) { @@ -391,8 +391,8 @@ __global__ static void __launch_bounds__(blockThreads) cavity.get_vertices(src, v0, v1); - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); const vec3 new_p((p0[0] + p1[0]) * T(0.5), (p0[1] + p1[1]) * T(0.5), @@ -404,8 +404,7 @@ __global__ static void __launch_bounds__(blockThreads) for (uint16_t i = 0; i < size; ++i) { const VertexHandle vvv = cavity.get_cavity_vertex(c, i); - const vec3 vp( - coords(vvv, 0), coords(vvv, 1), coords(vvv, 2)); + const vec3 vp = coords.to_glm<3>(vvv); const T edge_len_sq = glm::distance2(vp, new_p); if (edge_len_sq > high_edge_len_sq) { diff --git a/apps/Remesh/smoothing.cuh b/apps/Remesh/smoothing.cuh index 726a8f7a..851d0585 100644 --- a/apps/Remesh/smoothing.cuh +++ b/apps/Remesh/smoothing.cuh @@ -22,7 +22,7 @@ __global__ static void __launch_bounds__(blockThreads) return; } - const vec3 v(coords(v_id, 0), coords(v_id, 1), coords(v_id, 2)); + const vec3 v = coords.to_glm<3>(v_id); // compute both vertex normal and the new position // the new position is the average of the one-ring @@ -33,7 +33,8 @@ __global__ static void __launch_bounds__(blockThreads) // this is the last vertex in the one-ring (before r_id) VertexHandle q_id = iter.back(); - vec3 q(coords(q_id, 0), coords(q_id, 1), coords(q_id, 2)); + + vec3 q = coords.to_glm<3>(q_id); vec3 new_v(0.0, 0.0, 0.0); vec3 v_normal(0.0, 0.0, 0.0); @@ -44,7 +45,7 @@ __global__ static void __launch_bounds__(blockThreads) // the current one ring vertex const VertexHandle r_id = iter[i]; - const vec3 r(coords(r_id, 0), coords(r_id, 1), coords(r_id, 2)); + const vec3 r = coords.to_glm<3>(r_id); vec3 c = glm::cross(q - v, r - v); diff --git a/apps/Remesh/split.cuh b/apps/Remesh/split.cuh index 0df19a47..906c9b75 100644 --- a/apps/Remesh/split.cuh +++ b/apps/Remesh/split.cuh @@ -59,8 +59,8 @@ __global__ static void edge_split(rxmesh::Context context, edge_status(eh) = SKIP; return; } - const vec3 pa(coords(va, 0), coords(va, 1), coords(va, 2)); - const vec3 pb(coords(vb, 0), coords(vb, 1), coords(vb, 2)); + const vec3 pa = coords.to_glm<3>(va); + const vec3 pb = coords.to_glm<3>(vb); const T edge_len = glm::distance2(pa, pb); @@ -68,8 +68,8 @@ __global__ static void edge_split(rxmesh::Context context, vec3 p_new = (pa + pb) * T(0.5); - vec3 pc(coords(vc, 0), coords(vc, 1), coords(vc, 2)); - vec3 pd(coords(vd, 0), coords(vd, 1), coords(vd, 2)); + vec3 pc = coords.to_glm<3>(vc); + vec3 pd = coords.to_glm<3>(vd); T min_new_edge_len = std::numeric_limits::max(); @@ -124,14 +124,13 @@ __global__ static void edge_split(rxmesh::Context context, #ifndef NDEBUG // sanity check: we don't introduce small edges - const vec3 p_new( - coords(new_v, 0), coords(new_v, 1), coords(new_v, 2)); + const vec3 p_new = coords.to_glm<3>(new_v); for (int i = 0; i < 4; ++i) { const VertexHandle v = cavity.get_cavity_vertex(c, i); - const vec3 p(coords(v, 0), coords(v, 1), coords(v, 2)); + const vec3 p = coords.to_glm<3>(v); assert(glm::distance2(p_new, p) >= low_edge_len_sq); } @@ -241,19 +240,19 @@ inline void split_long_edges(rxmesh::RXMeshDynamic& rx, high_edge_len_sq, low_edge_len_sq, d_buffer); - + timers.stop("Split"); timers.start("SplitCleanup"); - rx.cleanup(); + rx.cleanup(); timers.stop("SplitCleanup"); timers.start("SplitSlice"); - rx.slice_patches(*coords, *edge_status); + rx.slice_patches(*coords, *edge_status); timers.stop("SplitSlice"); timers.start("SplitCleanup"); - rx.cleanup(); + rx.cleanup(); timers.stop("SplitCleanup"); bool show = false; diff --git a/apps/Remesh/util.cuh b/apps/Remesh/util.cuh index c19cbc11..d1c4d68e 100644 --- a/apps/Remesh/util.cuh +++ b/apps/Remesh/util.cuh @@ -29,8 +29,8 @@ __global__ static void stats_kernel(const rxmesh::Context context, ShmemAllocator shrd_alloc; auto compute_edge_len = [&](const EdgeHandle eh, const VertexIterator& ev) { - const vec3 v0(coords(ev[0], 0), coords(ev[0], 1), coords(ev[0], 2)); - const vec3 v1(coords(ev[1], 0), coords(ev[1], 1), coords(ev[1], 2)); + const vec3 v0 = coords.to_glm<3>(ev[0]); + const vec3 v1 = coords.to_glm<3>(ev[1]); T len = glm::distance(v0, v1); diff --git a/apps/SCP/scp.cu b/apps/SCP/scp.cu index 93b84454..ea161963 100644 --- a/apps/SCP/scp.cu +++ b/apps/SCP/scp.cu @@ -61,17 +61,17 @@ __global__ static void conformal_energy(const Context context, assert(p.is_valid() && q.is_valid()); assert(o0.is_valid() || o1.is_valid()); - const vec3 P(coord(p, 0), coord(p, 1), coord(p, 2)); - const vec3 Q(coord(q, 0), coord(q, 1), coord(q, 2)); + const vec3 P = coord.to_glm<3>(p); + const vec3 Q = coord.to_glm<3>(q); float coef = 0; if (o0.is_valid()) { - const vec3 O0(coord(o0, 0), coord(o0, 1), coord(o0, 2)); + const vec3 O0 = coord.to_glm<3>(o0); coef += weight(P, Q, O0); } if (o1.is_valid()) { - const vec3 O1(coord(o1, 0), coord(o1, 1), coord(o1, 2)); + const vec3 O1 = coord.to_glm<3>(o1); coef += weight(P, Q, O1); } diff --git a/apps/SECHistogram/sec_kernels.cuh b/apps/SECHistogram/sec_kernels.cuh index d6f15ad9..fc527416 100644 --- a/apps/SECHistogram/sec_kernels.cuh +++ b/apps/SECHistogram/sec_kernels.cuh @@ -48,8 +48,8 @@ __global__ static void sec(rxmesh::Context context, const VertexHandle v0 = iter[0]; const VertexHandle v1 = iter[1]; - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); T len2 = logf(glm::distance2(p0, p1)); @@ -170,8 +170,8 @@ __global__ static void compute_min_max_cost( const VertexHandle v0 = iter[0]; const VertexHandle v1 = iter[1]; - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); T len2 = logf(glm::distance2(p0, p1)); @@ -198,8 +198,8 @@ __global__ static void populate_histogram( const VertexHandle v0 = iter[0]; const VertexHandle v1 = iter[1]; - const vec3 p0(coords(v0, 0), coords(v0, 1), coords(v0, 2)); - const vec3 p1(coords(v1, 0), coords(v1, 1), coords(v1, 2)); + const vec3 p0 = coords.to_glm<3>(v0); + const vec3 p1 = coords.to_glm<3>(v1); T len2 = logf(glm::distance2(p0, p1)); diff --git a/apps/VertexNormal/vertex_normal.cu b/apps/VertexNormal/vertex_normal.cu index 358574af..d3354879 100644 --- a/apps/VertexNormal/vertex_normal.cu +++ b/apps/VertexNormal/vertex_normal.cu @@ -88,7 +88,7 @@ void vertex_normal_rxmesh(rxmesh::RXMeshStatic& rx, // Verify v_normals->move(rxmesh::DEVICE, rxmesh::HOST); - rx.for_each_vertex(HOST, [&](const VertexHandle& vh) { + rx.for_each_vertex(HOST, [&](const VertexHandle& vh) { uint32_t v_id = rx.map_to_global(vh); for (uint32_t i = 0; i < 3; ++i) { diff --git a/apps/VertexNormal/vertex_normal_kernel.cuh b/apps/VertexNormal/vertex_normal_kernel.cuh index 01bf77c1..b48a5b1d 100644 --- a/apps/VertexNormal/vertex_normal_kernel.cuh +++ b/apps/VertexNormal/vertex_normal_kernel.cuh @@ -15,9 +15,9 @@ __global__ static void compute_vertex_normal(const rxmesh::Context context, using namespace rxmesh; auto vn_lambda = [&](FaceHandle face_id, VertexIterator& fv) { // get the face's three vertices coordinates - vec3 c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2)); - vec3 c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2)); - vec3 c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2)); + vec3 c0 = coords.to_glm<3>(fv[0]); + vec3 c1 = coords.to_glm<3>(fv[1]); + vec3 c2 = coords.to_glm<3>(fv[2]); // compute the face normal vec3 n = cross(c1 - c0, c2 - c0); diff --git a/apps/XPBD/xpbd.cu b/apps/XPBD/xpbd.cu index 534ad499..e4db77c9 100644 --- a/apps/XPBD/xpbd.cu +++ b/apps/XPBD/xpbd.cu @@ -19,8 +19,8 @@ void __global__ init_edges(const Context context, auto v0 = iter[0]; auto v1 = iter[1]; - const glm::fvec3 x0(x(v0, 0), x(v0, 1), x(v0, 2)); - const glm::fvec3 x1(x(v1, 0), x(v1, 1), x(v1, 2)); + const glm::fvec3 x0 = x.to_glm<3>(v0); + const glm::fvec3 x1 = x.to_glm<3>(v1); rest_len(eh, 0) = glm::length(x0 - x1); }; @@ -49,8 +49,8 @@ void __global__ solve_stretch(const Context context, auto v0 = iter[0]; auto v1 = iter[1]; - const glm::fvec3 x0(new_x(v0, 0), new_x(v0, 1), new_x(v0, 2)); - const glm::fvec3 x1(new_x(v1, 0), new_x(v1, 1), new_x(v1, 2)); + const glm::fvec3 x0 = new_x.to_glm<3>(v0); + const glm::fvec3 x1 = new_x.to_glm<3>(v1); const float w1(invM(v0, 0)), w2(invM(v1, 0)); diff --git a/include/rxmesh/attribute.h b/include/rxmesh/attribute.h index 3ff0b8b0..843e61f2 100644 --- a/include/rxmesh/attribute.h +++ b/include/rxmesh/attribute.h @@ -657,6 +657,23 @@ class Attribute : public AttributeBase } + /** + * @brief Accessing the attribute a glm vector. This is used for read only + * since the return result is a copy + */ + template + __host__ __device__ __inline__ vec to_glm(const HandleT& handle) const + { + assert(N == get_num_attributes()); + + vec ret; + + for (int i = 0; i < N; ++i) { + ret[i] = this->operator()(handle, i); + } + return ret; + } + /** * @brief Accessing an attribute using a handle to the mesh element * @param handle input handle diff --git a/include/rxmesh/types.h b/include/rxmesh/types.h index e529ca53..d5e72d36 100644 --- a/include/rxmesh/types.h +++ b/include/rxmesh/types.h @@ -7,30 +7,35 @@ namespace rxmesh { -template -using vec1 = glm::vec<1, T, glm::defaultp>; +template +using vec = glm::vec; template -using vec2 = glm::vec<2, T, glm::defaultp>; +using vec1 = vec; template -using vec3 = glm::vec<3, T, glm::defaultp>; +using vec2 = vec; template -using mat2x2 = glm::mat<2, 2, T, glm::defaultp>; +using vec3 = vec; + +template +using mat = glm::mat; template -using mat3x2 = glm::mat<3, 2, T, glm::defaultp>; +using mat2x2 = mat; template -using mat2x3 = glm::mat<2, 3, T, glm::defaultp>; +using mat3x2 = mat; template -using mat3x3 = glm::mat<3, 3, T, glm::defaultp>; +using mat2x3 = mat; template -using mat4x4 = glm::mat<4, 4, T, glm::defaultp>; +using mat3x3 = mat; +template +using mat4x4 = mat; /** * @brief Flags for where data resides. Used with Attributes