Skip to content

Commit 4c49be4

Browse files
authored
Merge branch 'main' into dyn
2 parents f82c9b9 + 89011c8 commit 4c49be4

File tree

99 files changed

+15224
-7981
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

99 files changed

+15224
-7981
lines changed

.gitignore

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@ input/*
1010
!input/cloth.obj
1111
!input/plane_5.obj
1212
!input/sphere1.obj
13+
!input/bunnyhead.obj
1314
build/
1415
build_debug/
1516
include/rxmesh/util/git_sha1.cpp

CMakeLists.txt

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -187,6 +187,12 @@ find_package(CUDAToolkit REQUIRED)
187187
target_link_libraries(RXMesh INTERFACE CUDA::cusparse)
188188
target_link_libraries(RXMesh INTERFACE CUDA::cusolver)
189189

190+
#Eigen
191+
include("cmake/eigen.cmake")
192+
target_link_libraries(RXMesh INTERFACE Eigen3::Eigen)
193+
# https://eigen.tuxfamily.org/dox/TopicCUDA.html
194+
target_compile_definitions(RXMesh INTERFACE "EIGEN_DEFAULT_DENSE_INDEX_TYPE=int")
195+
190196
include(GoogleTest)
191197
add_subdirectory(apps)
192198
add_subdirectory(tests)

README.md

Lines changed: 84 additions & 44 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@
1313
* [**Structures**](#structures)
1414
* [**Computation**](#computation)
1515
* [**Viewer**](#viewer)
16+
* [**Matrices and Vectors**](#matrices-and-vectors)
1617
- [**Replicability**](#replicability)
1718
- [**Bibtex**](#bibtex)
1819

@@ -25,16 +26,23 @@ RXMesh is a surface triangle mesh data structure and programming model for proce
2526

2627
- *[RXMesh: A High-performance Mesh Data Structure and Programming Model on the GPU [S41051]](https://www.nvidia.com/gtc/session-catalog/?tab.scheduledorondemand=1583520458947001NJiE&search=rxmesh#/session/1633891051385001Q9SE)—NVIDIA GTC 2022*
2728

28-
This repository provides 1) source code to reproduce the results presented in the paper (git tag [`v0.1.0`](https://github.com/owensgroup/RXMesh/tree/v0.1.0)) and 2) ongoing development of RXMesh. For 1), all input models used in the paper can be found [here](https://ucdavis365-my.sharepoint.com/:f:/g/personal/ahmahmoud_ucdavis_edu/En-vEpIdSGBHqvCIa-MVXRQBg5g7GfM3P3RwZBHL4Hby3w?e=2EVnJd). Models were collected from [Thingi10K](https://ten-thousand-models.appspot.com/) and [Smithsonian 3D](https://3d.si.edu/explore) repository.
29+
The library also features a sparse and dense matrix infrastructure that is tightly coupled with the mesh data structure. We expose various [cuSolver](https://docs.nvidia.com/cuda/cusolver/index.html), [cuSparse](https://docs.nvidia.com/cuda/cusparse/), and [cuBlas](https://docs.nvidia.com/cuda/cublas/) operations through the sparse and dense matrices, tailored for geometry processing applications.
30+
31+
This repository provides 1) source code to reproduce the results presented in the paper (git tag [`v0.1.0`](https://github.com/owensgroup/RXMesh/tree/v0.1.0)) and 2) ongoing development of RXMesh.
2932

3033
## **Compilation**
31-
The code can be compiled on Ubuntu (GCC 9) and Windows (Visual Studio 2019) providing that CUDA (>=11.1.0) is installed. To run the executable(s), an NVIDIA GPU should be installed on the machine.
34+
The code can be compiled on Ubuntu, Windows, and WSL providing that CUDA (>=11.1.0) is installed. To run the executable(s), an NVIDIA GPU should be installed on the machine.
3235

3336
### **Dependencies**
3437
- [OpenMesh](https://www.graphics.rwth-aachen.de:9000/OpenMesh/OpenMesh) to verify the applications against reference CPU implementation
3538
- [RapidJson](https://github.com/Tencent/rapidjson) to report the results in JSON file(s)
3639
- [GoogleTest](https://github.com/google/googletest) for unit tests
3740
- [spdlog](https://github.com/gabime/spdlog) for logging
41+
- [glm](https://github.com/g-truc/glm.git) for small vectors and matrices operations
42+
- [Eigen](https://gitlab.com/libeigen/eigen) for small vectors and matrices operations
43+
- [Polyscope ](https://github.com/nmwsharp/polyscope) for visualization
44+
- [cereal](https://github.com/USCiLab/cereal.git) for serialization
45+
3846

3947
All the dependencies are installed automatically! To compile the code:
4048

@@ -109,7 +117,7 @@ The goal of defining a programming model is to make it easy to write applicatio
109117
vertex_color(vh, 2) = 0.6;
110118
```
111119

112-
- **Iterators** are used during query operations to iterate over the output of the query operation. The type of iterator defines the type of mesh element iterated on e.g., `VertexIterator` iterates over vertices which is the output of `VV`, `EV`, or `FV` query operations. Since query operations are only supported on the device, iterators can be only used inside the kernel. Iterators are usually populated internally.
120+
- **Iterators** are used during query operations to iterate over the output of the query operation. The type of iterator defines the type of mesh element iterated on e.g., `VertexIterator` iterates over vertices which is the output of `VV`, `EV`, or `FV` query operations. Since query operations are only supported on the device, iterators can be only used inside the GPU kernel. Iterators are usually populated internally.
113121

114122
- Example: Iterating over faces
115123
```c++
@@ -140,7 +148,7 @@ The goal of defining a programming model is to make it easy to write applicatio
140148
vertex_color(vh, 2) = 0.9;
141149
});
142150
```
143-
Alternatively, `for_each` operations could be written the same way as Queries operations (see below) using `for_each_dispatcher()`. This might be useful if the user would like to combine a `for_each` with queries operations in the same kernel. For more examples, checkout [`ForEach`](/tests/RXMesh_test/test_for_each.cuh) unit test.
151+
Alternatively, `for_each` operations could be written the same way as Queries operations (see below). This might be useful if the user would like to combine a `for_each` with queries operations in the same kernel. For more examples, checkout [`ForEach`](/tests/RXMesh_test/test_for_each.cuh) unit test.
144152

145153
- **Queries** operations supported by RXMesh with description are listed below
146154

@@ -161,30 +169,39 @@ The goal of defining a programming model is to make it easy to write applicatio
161169
```cpp
162170
template<uint32_t blockSize>
163171
__global__ void vertex_normal (Context context){
164-
auto compute_vn = [&](FaceHandle face_id, VertexIterator& fv) {
172+
auto compute_vn = [&](const FaceHandle face_id, const VertexIterator& fv) {
165173
//This thread is assigned to face_id
166174

167175
// get the face's three vertices coordinates
168-
Vector<3, T> c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2));
169-
Vector<3, T> c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2));
170-
Vector<3, T> c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2));
176+
vec3<T> c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2));
177+
vec3<T> c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2));
178+
vec3<T> c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2));
171179

172180
//compute face normal
173-
Vector<3, T> n = cross(c1 - c0, c2 - c0);
181+
vec3<T> n = cross(c1 - c0, c2 - c0);
174182

175183
// add the face's normal to its vertices
176184
for (uint32_t v = 0; v < 3; ++v) // for every vertex in this face
177185
for (uint32_t i = 0; i < 3; ++i) // for the vertex 3 coordinates
178186
atomicAdd(&normals(fv[v], i), n[i]);
179187
};
180188

181-
//Query dispatcher must be called by all threads in the block.
182-
//Dispatcher will first perform the query, store the results in shared memory, then
183-
//run the user-defined computation i.e., compute_vn
184-
query_block_dispatcher<Op::FV, blockSize>(context, compute_vn);
189+
//Query must be called by all threads in the block. Thus, we create this cooperative_group
190+
//that uses all threads in the block and pass to the Query
191+
auto block = cooperative_groups::this_thread_block();
192+
193+
Query<blockThreads> query(context);
194+
195+
//Qeury will first perform the query, store the results in shared memory. ShmemAllocator is
196+
//passed to the function to make sure we don't over-allocate or overwrite user-allocated shared
197+
//memory
198+
ShmemAllocator shrd_alloc;
199+
200+
//Finally, we run the user-defined computation i.e., compute_vn
201+
query.dispatch<Op::FV>(block, shrd_alloc, compute_vn);
185202
}
186203
```
187-
To save computation, `query_block_dispatcher` could be run on a subset of the input mesh element i.e., _active set_. The user can define the active set using a lambda function that returns true if the input mesh element is in the active set.
204+
To save computation, `query.dispatch` could be run on a subset of the input mesh element i.e., _active set_. The user can define the active set using a lambda function that returns true if the input mesh element is in the active set.
188205

189206
- Example: defining active set
190207
```cpp
@@ -194,11 +211,11 @@ The goal of defining a programming model is to make it easy to write applicatio
194211
// ....
195212
};
196213

197-
auto computation = [&](FaceHandle face_id, VertexIterator& fv) {
214+
auto computation = [&](const FaceHandle face_id, const VertexIterator& fv) {
198215
// ....
199216
};
200217

201-
query_block_dispatcher<Op::FV, blockSize>(context, computation, active_set);
218+
query.dispatch<Op::FV, blockSize>(context, computation, active_set);
202219
}
203220
```
204221

@@ -242,13 +259,10 @@ Starting v0.2.1, RXMesh integrates [Polyscope](https://polyscope.run) as a mesh
242259
> cd build
243260
> cmake -DUSE_POLYSCOPE=True ../
244261
```
245-
By default, the parameter is set to True on Windows and False on Linux machines. RXMesh implements the necessary functionalities to pass attributes to Polyscope—thanks to its [data adaptors](https://polyscope.run/data_adaptors/). However, this needs attributes to be moved to the host first before passing it to Polyscope. For more information about Polyscope's different visualization options, please checkout Polyscope's [Surface Mesh documentation](https://polyscope.run/structures/surface_mesh/basics/).
262+
By default, the parameter is set to True. RXMesh implements the necessary functionalities to pass attributes to Polyscope—thanks to its [data adaptors](https://polyscope.run/data_adaptors/). However, this needs attributes to be moved to the host first before passing it to Polyscope. For more information about Polyscope's different visualization options, please checkout Polyscope's [Surface Mesh documentation](https://polyscope.run/structures/surface_mesh/basics/).
246263

247264
- Example: [render vertex color](./tests/Polyscope_test/test_polyscope.cu)
248265
```cpp
249-
//initialize polyscope
250-
polyscope::init();
251-
252266
RXMeshStatic rx("dragon.obj");
253267

254268
//vertex color attribute
@@ -273,36 +287,62 @@ By default, the parameter is set to True on Windows and False on Linux machines.
273287
<img src="./assets/polyscope_dragon.PNG" width="80%"><br>
274288
</p>
275289

290+
### **Matrices and Vectors**
291+
- **Large Matrices:** RXMesh has built-in support for large sparse and dense matrices built on top of [cuSparse](https://docs.nvidia.com/cuda/cusparse/) and [cuBlas](https://docs.nvidia.com/cuda/cublas/), respectively. For example, attributes can be converted to dense matrices as follows
276292

277-
## **Replicability**
278-
This repo was awarded the [replicability stamp](http://www.replicabilitystamp.org#https-github-com-owensgroup-rxmesh) by the Graphics Replicability Stamp Initiative (GRSI) :tada:
293+
```cpp
294+
295+
RXMeshStatic rx("input.obj");
296+
297+
//Input mesh coordinates as VertexAttribute
298+
std::shared_ptr<VertexAttribute<float>> x = rx.get_input_vertex_coordinates();
299+
300+
//Convert the attributes to a (#vertices x 3) dense matrix
301+
std::shared_ptr<DenseMatrix<float>> x_mat = x->to_matrix();
302+
303+
//do something with x_mat
304+
//....
279305

280-
The scripts used to generate the data shown in the paper can be found under
281-
* [Figure 6](https://github.com/owensgroup/RXMesh/blob/main/tests/RXMesh_test/benchmark.sh)
282-
* [Figure 8 (a)](https://github.com/owensgroup/RXMesh/blob/main/apps/MCF/benchmark.sh)
283-
* [Figure 8 (b)](https://github.com/owensgroup/RXMesh/blob/main/apps/Geodesic/benchmark.sh)
284-
* [Figure 8 (c)](https://github.com/owensgroup/RXMesh/blob/main/apps/Filtering/benchmark.sh)
285-
* [Figure 8 (d)](https://github.com/owensgroup/RXMesh/blob/main/apps/VertexNormal/benchmark.sh)
306+
//Populate the VertexAttribute coordinates back with the content of the dense matrix
307+
x->from_matrix(x_mat.get());
286308

287-
Each script should be run from the script's containing directory after compiling the code in the `build/` directory. The only input parameter needed is the path to the input OBJ files. The resulting JSON files will be written to the `output/` directory.
309+
```
310+
Dense matrices can be accessed using the usual row and column indices or via the mesh element handle (Vertex/Edge/FaceHandle) as a row index. This allows for easy access to the correct row associated with a specific vertex, edge, or face. Dense matrices support various operations such as absolute sum, AXPY, dot products, norm2, scaling, and swapping.
311+
312+
RXMesh supports sparse matrices, where the sparsity pattern matches the query operations. For example, it is often necessary to build a sparse matrix of size #V x #V with non-zero values at (i, j) only if the vertex corresponding to row i is connected by an edge to the vertex corresponding to column j. Currently, we only support the VV sparsity pattern, but we are working on expanding to all other types of queries.
313+
314+
The sparse matrix can be used to solve a linear system via Cholesky, LU, or QR factorization (relying on [cuSolver](https://docs.nvidia.com/cuda/cusolver/index.html))). The solver offers two APIs. The high-level API reorders the input sparse matrix (to reduce non-zero fill-in after matrix factorization) and allocates the additional memory needed to solve the system. Repeated calls to this API will reorder the matrix and allocate/deallocate the temporary memory with each call. For scenarios where the matrix remains unchanged but multiple right-hand sides need to be solved, users can utilize the low-level API, which splits the solve method into pre_solve() and solve(). The former reorders the matrix and allocates temporary memory only once. The low-level API is currently only supported for Cholesky-based factorization. Check out the MCF application for an example of how to set up and use the solver.
315+
316+
Similar to dense matrices, sparse matrices also support accessing the matrix using the VertexHandle and multiplication by dense matrices.
317+
318+
- **Small Matrices:**
319+
It is often necessary to perform operations on small matrices as part of geometry processing applications, such as computing the SVD of a 3x3 matrix or normalizing a 1x3 vector. For this purpose, RXMesh attributes can be converted into glm or Eigen matrices, as demonstrated in the vertex_normal example above. Both glm and Eigen support small matrix operations inside the GPU kernel.
320+
321+
322+
323+
## **Replicability**
324+
This repo was awarded the [replicability stamp](http://www.replicabilitystamp.org#https-github-com-owensgroup-rxmesh) by the Graphics Replicability Stamp Initiative (GRSI) :tada:. Visit git tag [`v0.1.0`](https://github.com/owensgroup/RXMesh/tree/v0.1.0) for more information about replicability scripts.
288325
289326
## **Bibtex**
290327
```
291328
@article{Mahmoud:2021:RAG,
292-
author = {Mahmoud, Ahmed H. and Porumbescu, Serban D. and Owens, John D.},
293-
title = {{RXM}esh: A {GPU} Mesh Data Structure},
294-
journal = {ACM Transactions on Graphics},
295-
year = 2021,
296-
volume = 40,
297-
number = 4,
298-
month = aug,
299-
issue_date = {August 2021},
300-
articleno = 104,
301-
numpages = 16,
302-
pages = {104:1--104:16},
303-
url = {https://escholarship.org/uc/item/8r5848vp},
304-
full_talk = {https://youtu.be/Se_cNAol4hY},
305-
short_talk = {https://youtu.be/V_SHMXnCVws},
306-
doi = {10.1145/3450626.3459748}
329+
author = {Ahmed H. Mahmoud and Serban D. Porumbescu and John D. Owens},
330+
title = {{RXM}esh: A {GPU} Mesh Data Structure},
331+
journal = {ACM Transactions on Graphics},
332+
year = 2021,
333+
volume = 40,
334+
number = 4,
335+
month = aug,
336+
issue_date = {August 2021},
337+
articleno = 104,
338+
numpages = 16,
339+
pages = {104:1--104:16},
340+
url = {https://escholarship.org/uc/item/8r5848vp},
341+
full_talk = {https://youtu.be/Se_cNAol4hY},
342+
short_talk = {https://youtu.be/V_SHMXnCVws},
343+
doi = {10.1145/3450626.3459748},
344+
acmauthorize = {https://dl.acm.org/doi/10.1145/3450626.3459748?cid=81100458295},
345+
acceptance = {149/444 (33.6\%)},
346+
ucdcite = {a140}
307347
}
308348
```

apps/ARAP/CMakeLists.txt

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
add_executable(ARAP)
2+
3+
set(SOURCE_LIST
4+
arap.cu
5+
)
6+
7+
target_sources(ARAP
8+
PRIVATE
9+
${SOURCE_LIST}
10+
)
11+
12+
set_target_properties(ARAP PROPERTIES FOLDER "apps")
13+
14+
set_property(TARGET ARAP PROPERTY CUDA_SEPARABLE_COMPILATION ON)
15+
16+
source_group(TREE ${CMAKE_CURRENT_LIST_DIR} PREFIX "ARAP" FILES ${SOURCE_LIST})
17+
18+
target_link_libraries(ARAP
19+
PRIVATE RXMesh
20+
)
21+
22+
#gtest_discover_tests( ARAP )

0 commit comments

Comments
 (0)