• About
• Compilation
  • Dependencies
• Organization
• Programming Model
  • Structures
  • Computation
  • Viewer
• Replicability
• Bibtex

RXMesh is a surface triangle mesh data structure and programming model for processing static meshes on the GPU. RXMesh aims at provides a high-performance, generic, and compact data structure that can handle meshes regardless of their quality (e.g., non-manifold). The programming model helps to hide the complexity of the data structure and provides an intuitive access model for different use cases. For more details, please check out our paper and GTC talk:

• RXMesh: A GPU Mesh Data Structure
  Ahmed H. Mahmoud, Serban D. Porumbescu, and John D. Owens
  ACM Transaction on Graphics (Proceedings of SIGGRAPH 2021)

• RXMesh: A High-performance Mesh Data Structure and Programming Model on the GPU [S41051]—NVIDIA GTC 2022

This repository provides 1) source code to reproduce the results presented in the paper (git tag v0.1.0) and 2) ongoing development of RXMesh. For 1), all input models used in the paper can be found here. Models were collected from Thingi10K and Smithsonian 3D repository.

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.

• OpenMesh to verify the applications against reference CPU implementation
• RapidJson to report the results in JSON file(s)
• GoogleTest for unit tests
• spdlog for logging

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

> git clone https://github.com/owensgroup/RXMesh.git
> cd RXMesh
> mkdir build 
> cd build 
> cmake ../

Depending on the system, this will generate either a .sln project on Windows or a make file for a Linux system.

RXMesh is a CUDA/C++ header-only library. All unit tests are under the tests/ folder. This includes the unit test for some basic functionalities along with the unit test for the query operations. All applications are under the apps/ folder.

The goal of defining a programming model is to make it easy to write applications using RXMesh without getting into the nuances of the data structure. Applications written using RXMesh are composed of one or more of the high-level building blocks defined under Computation. To use these building blocks, the user would have to interact with data structures specific to RXMesh discussed under Structures. Finally, RXMesh integrates Polyscope as a mesh Viewer which the user can use to render their final results or for debugging purposes.

• Attributes are the metadata (geometry information) attached to vertices, edges, or faces. Allocation of the attributes is per-patch basis and managed internally by RXMesh. The allocation could be done on the host, device, or both. Allocating attributes on the host is only beneficial for I/O operations or initializing attributes and then eventually moving them to the device.

  • Example: allocation

    RXMeshStatic rx("input.obj");
    auto vertex_color = 
      rx.add_vertex_attribute<float>("vColor", //Unique name 
                                     3,        //Number of attribute per vertex 
                                     DEVICE,   //Allocation place 
                                     SoA);     //Memory layout (SoA vs. AoS)                                 
    

  • Example: reading from std::vector

    RXMeshStatic rx("input.obj");
    std::vector<std::vector<float>> face_color_vector;
    //....
    
    auto face_color = 
      rx.add_face_attribute<int>(face_color_vector,//Input attribute where number of attributes per face is inferred 
                                 "fColor",         //Unique name                                
                                 SoA);             //Memory layout (SoA vs. AoS)                                  

  • Example: move, reset, and copy

    //By default, attributes are allocated on both host and device     
    auto edge_attr = rx.add_edge_attribute<float>("eAttr", 1);  
    //Initialize edge_attr on the host 
    // ..... 
    
    //Move attributes from host to device 
    edge_attr.move(HOST, DEVICE);
    
    //Reset all entries to zero
    edge_attr.reset(0, DEVICE);
    
    auto edge_attr_1 = rx.add_edge_attribute<float>("eAttr1", 1);  
    
    //Copy from another attribute. 
    //Here, what is on the host sde of edge_attr will be copied into the device side of edge_attr_1
    edge_attr_1.copy_from(edge_attr, HOST, DEVICE);

• Handles are the unique identifiers for vertices, edges, and faces. They are usually internally populated by RXMesh (by concatenating the patch ID and mesh element index within the patch). Handles can be used to access attributes, for_each operations, and query operations.

  • Example: Setting vertex attribute using vertex handle

    auto vertex_color = ...    
    VertexHandle vh; 
    //...
    
    vertex_color(vh, 0) = 0.9;
    vertex_color(vh, 1) = 0.5;
    vertex_color(vh, 2) = 0.6;

• 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.

  • Example: Iterating over faces

    FaceIterator f_iter; 
    //...
    
    for (uint32_t f = 0; f < f_iter.size(); ++f) {	
      FaceHandle fh = f_iter[f];
      //do something with fh ....
    }

• for_each runs a computation over all vertices, edges, or faces without requiring information from neighbor mesh elements. The computation run on each mesh element is defined as a lambda function that takes a handle as an input. The lambda function could run either on the host, device, or both. On the host, we parallelize the computation using OpenMP. Care must be taken for lambda function on the device since it needs to be annotated using __device__ and it can only capture by value. More about lambda function in CUDA can be found here

  • Example: using for_each to initialize attributes

    RXMeshStatic rx("input.obj");
    auto vertex_pos   = rx.get_input_vertex_coordinates();                   //vertex position 
    auto vertex_color = rx.add_vertex_attribute<float>("vColor", 3, DEVICE); //vertex color 
    
    //This function will be executed on the device 
    rx.for_each_vertex(
        DEVICE,
        [vertex_color, vertex_pos] __device__(const VertexHandle vh) {
            vertex_color(vh, 0) = 0.9;
            vertex_color(vh, 1) = vertex_pos(vh, 1);
            vertex_color(vh, 2) = 0.9;
        });

• Queries operations supported by RXMesh with description are listed below

  Query	Description
  VV	For vertex V, return its adjacent vertices
  VE	For vertex V, return its incident edges
  VF	For vertex V, return its incident faces
  EV	For edge E, return its incident vertices
  EF	For edge E, return its incident faces
  FV	For face F, return its incident vertices
  FE	For face F, return its incident edges
  FF	For face F, return its adjacent faces

  Queries are only supported on the device. RXMesh API for queries takes a lambda function along with the type of query. The lambda function defines the computation that will be run on the query output.

  • Example: vertex normal computation

    template<uint32_t blockSize>
    __global__ void vertex_normal (Context context){      
        auto compute_vn = [&](FaceHandle face_id, VertexIterator& fv) {
        	//This thread is assigned to face_id
    
        	// get the face's three vertices coordinates
        	Vector<3, T> c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2));
        	Vector<3, T> c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2));
            Vector<3, T> c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2));
    
          //compute face normal
          Vector<3, T> n = cross(c1 - c0, c2 - c0);
    
            // add the face's normal to its vertices
        		for (uint32_t v = 0; v < 3; ++v)     // for every vertex in this face
                for (uint32_t i = 0; i < 3; ++i)   // for the vertex 3 coordinates
        		        atomicAdd(&normals(fv[v], i), n[i]);          
        };
    
        //Query dispatcher must be called by all threads in the block. 
        //Dispatcher will first perform the query, store the results in shared memory, then 
        //run the user-defined computation i.e., compute_vn
        query_block_dispatcher<Op::FV, blockSize>(context, compute_vn);
    } 

  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.

  • Example: defining active set

    template<uint32_t blockSize>
    __global__ void active_set_query (Context context){
      auto active_set = [&](FaceHandle face_id) -> bool{ 
        // ....         
        };
    
        auto computation = [&](FaceHandle face_id, VertexIterator& fv) {          
        // ....         
        };
    
        query_block_dispatcher<Op::FV, blockSize>(context, computation, active_set);
    } 

• Reduction operations apply a binary associative operation on the input attributes. RXMesh provides dot products between two attributes (of the same type), L2 norm of an input attribute, and user-defined reduction operation on an input attribute. For user-defined reduction operation, the user needs to pass a binary reduction functor with member __device__ T operator()(const T &a, const T &b) or use on of CUB's thread operators e.g., cub::Max(). Reduction operations require allocation of temporary buffers which we abstract away using ReduceHandle.

  • Example: dot product, L2 norm, user-defined reduction

    RXMeshStatic rx("input.obj");
    auto vertex_attr1 = rx.add_vertex_attribute<float>("v_attr1", 3, DEVICE);
    auto vertex_attr2 = rx.add_vertex_attribute<float>("v_attr2", 3, DEVICE);
    
    // Populate vertex_attr1 and vertex_attr2 
    //....
    
    //Reduction handle 
    ReduceHandle reduce(v1_attr);
    
    //Dot product between two attributes. Results are returned on the host 
    float dot_product = reduce.dot(v1_attr, v2_attr);
    
    cudaStream_t stream; 
    //init stream 
    //...
    
    //Reduction operation could be performed on specific attribute and using specific stream 
    float l2_norm = reduce.norm2(v1_attr, //input attribute 
                                 1,       //attribute ID. If not specified, reduction is run on all attributes 
                                 stream); //stream used for reduction. 
    
    
    //User-defined reduction operation 
    float l2_norm = reduce.reduce(v1_attr,                               //input attribute 
                                  cub::Max(),                            //binary reduction functor 
                                  std::numeric_limits<float>::lowest()); //initial value 

Starting v0.2.1, RXMesh integrates Polyscope as a mesh viewer. To use it, make sure to turn on the CMake parameter USE_POLYSCOPE i.e.,

> cd build 
> cmake -DUSE_POLYSCOPE=True ../

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. 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.

• Example: render vertex color

  //initialize polyscope 
  polyscope::init();
  
  RXMeshStatic rx("dragon.obj");
  
  //vertex color attribute 
  auto vertex_color = rx.add_vertex_attribute<float>("vColor", 3);
  
  //Populate vertex color on the device
  //....
  
  //Move vertex color to the host 
  vertex_color.move(DEVICE, HOST);
  
  //polyscope instance associated with rx 
  auto polyscope_mesh = rx.get_polyscope_mesh();
  
  //pass vertex color to polyscope 
  polyscope_mesh->addVertexColorQuantity("vColor", vertex_color);
  
  //render 
  polyscope::show();

  
  

This repo was awarded the replicability stamp by the Graphics Replicability Stamp Initiative (GRSI) 🎉

The scripts used to generate the data shown in the paper can be found under

• Figure 6
• Figure 8 (a)
• Figure 8 (b)
• Figure 8 (c)
• Figure 8 (d)

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.

@article{Mahmoud:2021:RAG,
  author = {Mahmoud, Ahmed H. and Porumbescu, Serban D. and Owens, John D.},
  title = {{RXM}esh: A {GPU} Mesh Data Structure},
  journal = {ACM Transactions on Graphics},
  year = 2021,
  volume = 40,
  number = 4,
  month = aug,
  issue_date = {August 2021},
  articleno = 104,
  numpages = 16,
  pages = {104:1--104:16},
  url = {https://escholarship.org/uc/item/8r5848vp},
  full_talk = {https://youtu.be/Se_cNAol4hY},
  short_talk = {https://youtu.be/V_SHMXnCVws},
  doi = {10.1145/3450626.3459748}
}