{"id":469,"date":"2024-08-02T15:28:38","date_gmt":"2024-08-02T15:28:38","guid":{"rendered":"https:\/\/summergeometry.org\/sgi2024\/?p=469"},"modified":"2024-08-02T15:28:40","modified_gmt":"2024-08-02T15:28:40","slug":"an-introduction-to-rxmesh","status":"publish","type":"post","link":"https:\/\/summergeometry.org\/sgi2024\/an-introduction-to-rxmesh\/","title":{"rendered":"An Introduction to RXMesh"},"content":{"rendered":"\n

This post serves as a tutorial for using RXMesh<\/a>.<\/p>\n\n\n\n

RXMesh is a framework for processing triangle mesh on the GPU. It allows developers to easily use the GPU’s massive parallelism to perform common geometry processing tasks.<\/p>\n\n\n\n

This tutorial will cover the following topics:<\/p>\n\n\n\n

    \n
  1. Setting up RXMesh<\/li>\n\n\n\n
  2. RXMesh basic workflow<\/li>\n\n\n\n
  3. Writing kernels for queries<\/li>\n\n\n\n
  4. Example: Visualizing face normals<\/li>\n<\/ol>\n\n\n\n

    Setting up RXMesh<\/h2>\n\n\n\n

    For the sake of this tutorial, you can set up a fork of RXMesh and then create your own small section to implement your own tutorial examples in the fork. <\/p>\n\n\n\n

    Go to https:\/\/github.com\/owensgroup\/RXMesh<\/a> and fork the repository. You can create your own directory with a simple CUDA script to start using RXMesh. To do so, go to https:\/\/github.com\/owensgroup\/RXMesh\/blob\/main\/apps\/CMakeLists.txt<\/a> in your fork and add a line:<\/p>\n\n\n\n

    add_subdirectory(introduction)<\/code><\/code><\/pre>\n\n\n\n
    In the apps<\/a> directory on your local fork, create a folder called “introduction”. There should be two files in this folder.<\/p>\n\n\n\n
      \n
    1. A .cu<\/em> file to run the code. You can call this introduction.cu<\/em>. To start, it can have the following contents:<\/li>\n<\/ol>\n\n\n\n
      #include \"rxmesh\/query.cuh\"\n#include \"rxmesh\/rxmesh_static.h\"\n\n#include \"rxmesh\/matrix\/sparse_matrix.cuh\"\n\nusing namespace rxmesh;\n\nint main(int argc, char** argv)\n{\n    Log::init();\n\n    const uint32_t device_id = 0;\n    cuda_query(device_id);\n\n    RXMeshStatic rx(STRINGIFY(INPUT_DIR) \"sphere3.obj\");\n\n#if USE_POLYSCOPE\n    polyscope::show();\n#endif\n}<\/code><\/pre>\n\n\n\n
      2. A CMake file called CMakeLists.txt<\/em>. It must have the following contents:<\/p>\n\n\n\n
      add_executable(introduction)\n\nset(SOURCE_LIST\n    introduction.cu\t\n)\n\ntarget_sources(introduction\n    PRIVATE\n    ${SOURCE_LIST}\n)\n\nset_target_properties(introduction PROPERTIES FOLDER \"apps\")\n\nset_property(TARGET ARAP PROPERTY CUDA_SEPARABLE_COMPILATION ON)\n\nsource_group(TREE ${CMAKE_CURRENT_LIST_DIR} PREFIX \"introduction\" FILES ${SOURCE_LIST})\n\ntarget_link_libraries(introduction\n    PRIVATE RXMesh\n)\n\n#gtest_discover_tests( introduction )<\/code><\/pre>\n\n\n\n
      Once you have done the above steps, building the project with CMake is simple. Refer to the README in https:\/\/github.com\/owensgroup\/RXMesh<\/a> to know the requirements and steps to build the project for your device.<\/p>\n\n\n\n
      Once the project has been set up, we can begin to test and run our introduction.cu<\/em> file. If you were to run it, you should get an output of a 3D sphere.<\/p>\n\n\n\n
      \"\"<\/figure>\n\n\n\n
      This means the project set up has been successful, we can now begin learning RXMesh.<\/p>\n\n\n\n
      RXMesh basic workflow<\/h2>\n\n\n\n
      RXMesh provides custom data structures and functions to handle various low-level computational tasks that would otherwise require low-level implementation. Let’s explore the fundamentals of how this workflow operates.<\/p>\n\n\n\n
      To start, look at this piece of code.<\/p>\n\n\n\n
      \n
      \n
      int main(int argc, char** argv)\n{\n    Log::init();\n    const uint32_t device_id = 0;\n    rxmesh::cuda_query(device_id);\n\n    rxmesh::RXMeshStatic rx(STRINGIFY(INPUT_DIR) \"dragon.obj\");\n    auto polyscope_mesh = rx.get_polyscope_mesh();\n\n    auto vertex_pos   = *rx.get_input_vertex_coordinates();\n    auto vertex_color = *rx.add_vertex_attribute<float>(\"vColor\", 3);\n\n    rx.for_each_vertex(rxmesh::DEVICE,\n    [vertex_color, vertex_pos] __device__(const rxmesh::VertexHandle vh)\n    {\n            vertex_color(vh, 0) = 0.9;\n            vertex_color(vh, 1) = vertex_pos(vh, 1);\n            vertex_color(vh, 2) = 0.9;\n    });\n\n    vertex_color.move(rxmesh::DEVICE, rxmesh::HOST);\n\n    polyscope_mesh->addVertexColorQuantity(\"vColor\", vertex_color);\n    polyscope::show();\n\n    return 0;\n}<\/code><\/pre>\n<\/div>\n<\/div>\n\n\n\n
      The above is a slightly simplified version of what a for_each<\/code> computation program would look like using RXMesh. for_each<\/code> involves accessing every mesh element<\/em> of a specific type (vertex, edge, or face) and performing some process on\/with it.<\/p>\n\n\n\n
      In the code example above, the computation adjusts the green component of the vertex’s color based on the vertex’s Y coordinate in 3D space.<\/p>\n\n\n\n
      But how do we comprehend this code? We will look at it line by line:<\/p>\n\n\n\n
          rxmesh::RXMeshStatic rx(STRINGIFY(INPUT_DIR) \"dragon.obj\");<\/code><\/pre>\n\n\n\n
      The above line declares the RXMeshStatic<\/code> object. Static here stands for a static mesh (one whose connectivity remains constant for the duration of the program’s runtime). Using the RXMesh’s Input directory, we can use some meshes that come with it. In this case, we pass in the dragon.obj file, which holds the mesh data for a dragon.<\/p>\n\n\n\n
          auto polyscope_mesh = rx.get_polyscope_mesh();<\/code><\/pre>\n\n\n\n
      This line returns the data needed for Polyscope <\/a>to display the mesh along with any attribute content we add to it for visualization.<\/p>\n\n\n\n
          auto vertex_pos = *rx.get_input_vertex_coordinates();<\/code><\/pre>\n\n\n\n
      This line is a special function that directly gives us the coordinates of the input mesh.<\/p>\n\n\n\n
      auto vertex_color = *rx.add_vertex_attribute<float>(\"vColor\", 3);<\/code><\/pre>\n\n\n\n
      This line gives vertex attributes called “vColor”. Attributes are simply data that lives on top of vertices, edges, or faces. To learn more about how to handle attributes in RXMesh, check out this<\/a>. In this case, we associate three float-point numbers to each vertex in the mesh.<\/p>\n\n\n\n
      rx.for_each_vertex(rxmesh::DEVICE,\n    [vertex_color, vertex_pos] __device__(const rxmesh::VertexHandle vh)\n    {\n            vertex_color(vh, 0) = 0.9;\n            vertex_color(vh, 1) = vertex_pos(vh, 1);\n            vertex_color(vh, 2) = 0.9;\n    });<\/code><\/pre>\n\n\n\n
      The above lines represent a lambda function that is utilized by the for_each<\/code> computation. In this case, for_each_vertex<\/code> accesses each of the vertices of the mesh associated with our dragon. We pass in the arguments:<\/p>\n\n\n\n
        \n
      1. rxmesh::DEVICE<\/code> – to let RXMesh know this is happening on the device i.e., the GPU.<\/li>\n\n\n\n
      2. [vertex_color, vertex_pos]<\/code> – represents the data we are passing to the function.<\/li>\n\n\n\n
      3. (const rxmesh::VertexHandle vh)<\/code> – is the handle that is used in any RXMesh computation. A handle allows us to access data associated to individual geometric elements. In this case, we have a handle that allows us to access each vertex in the mesh.<\/li>\n<\/ol>\n\n\n\n
        What do these lines mean? <\/p>\n\n\n\n
                    vertex_color(vh, 0) = 0.9;\n            vertex_color(vh, 1) = vertex_pos(vh, 1);\n            vertex_color(vh, 2) = 0.9;<\/code><\/pre>\n\n\n\n
        These 3 lines bring together a lot of the declarations made earlier. Through the kernel, we are accessing the attribute data we defined before. Since we need to know which vertex we are accessing its color, we pass in vh<\/code> as the first argument. Since a vertex has three components (standing for RGB), we also need to pass which index in the attribute’s vector (you can think of it as a 1D array too) we are accessing. Hence vertex_color(vh, 0) = 0.9;<\/code> which stands for “in the vertex color associated with the handle vh (which for the kernel represents a specific vertex on the mesh), the value of the first component is 0.9”. Note that this “first component” represents red for us.<\/p>\n\n\n\n
        \"\"<\/figure>\n\n\n\n
        What about vertex_color(vh, 1) = vertex_pos(vh, 1)<\/code>? This line, similar to the previous one, is accessing the second component associated with the color, in the vertex the handle is associated with.<\/p>\n\n\n\n
        But what is on the right-hand side? We are accessing vertex_pos (our coordinates of each vertex in the mesh) and we are accessing it the same way we access our color. In this case, the line is telling us that we are accessing the 2nd positional coordinate (y coordinate) associated with our vertex (that our handle gives to the kernel).<\/p>\n\n\n\n
            vertex_color.move(rxmesh::DEVICE, rxmesh::HOST);<\/code><\/pre>\n\n\n\n
        This line moves the attribute data from the GPU to the CPU.  <\/p>\n\n\n\n
        polyscope_mesh->addVertexColorQuantity(\"vColor\", vertex_color);<\/code><\/pre>\n\n\n\n
        This line uses Polyscope to add a “quantity” to Polyscope’s visualization of the mesh. In this case, when we add the VertexColorQuantity and pass vertex_color, Polyscope will now visualize the per-vertex color information we calculated in the lambda function.<\/p>\n\n\n\n
            polyscope::show();<\/code><\/pre>\n\n\n\n
        We finally render the entire mesh with our added quantities using the show()<\/code> function from Polyscope.<\/p>\n\n\n\n
        Throughout this basic workflow, you may have noticed it works similarly to any other program where we receive some input data, perform some processing, and then render\/output it in some form.<\/p>\n\n\n\n
        More specifically, we use RXMesh to read that data, set up our attributes, and then pass that into a kernel to perform some processing. Once we move that data from the device to the host, we can either perform further processing or render it using Polyscope. <\/p>\n\n\n\n
        It is important to look at the kernel as computing per geometric element.<\/em> This means we only <\/em>need to think of computation on a single mesh element of a certain type (i.e., vertex, edge, or face) since RXMesh then takes this computation and runs it in parallel on all mesh elements of that type.<\/p>\n\n\n\n
        Writing kernels for queries<\/h2>\n\n\n\n
        While our basic workflow covers how to perform for_each<\/code> operation using the GPU, we may often require geometric connectivity information for different geometry processing tasks. To do that, RXMesh implements various query operations. To understand the different types of queries, check out this<\/a> part of the README.<\/p>\n\n\n\n
        We can say there are two parts to run a query.<\/p>\n\n\n\n
        The first part consists of creating the launchbox <\/code>for the query, which defines the threads and (shared) memory allocated for the process along with calling the function from the host to run on the device.<\/p>\n\n\n\n
        The second part consists of actually defining what the kernel <\/em>looks like.<\/p>\n\n\n\n
        We will look at these one by one.<\/p>\n\n\n\n
        Here’s an example of how a launchbox <\/code>is created and used for a process where we want to find all the vertex normals:<\/p>\n\n\n\n
        \/\/ Vertex Normal\nauto vertex_normals = rx.add_vertex_attribute(\"vNormals\", 3);\nconstexpr uint32_t CUDABlockSize = 256;\n\nrxmesh::LaunchBox<CUDABlockSize> launch_box;\n\nrx.prepare_launch_box(\n  {rxmesh::Op::FV},\n  launch_box,\n  (void*)compute_vertex_normal<float,CUDABlockSize>);\n\ncompute_vertex_normal<float, CUDABlockSize>\n<<<launch_box.blocks, \n   launch_box.num_threads, \n   launch_box.smem_bytes_dyn>>>\n(rx.get_context(), vertex_pos, *vertex_normals);\n\nvertex_normals->move(rxmesh::DEVICE, rxmesh::HOST);<\/code><\/pre>\n\n\n\n
        Notice how, in replacing the for_each part of our basic workflow, we instead declare the launchbox and call our function compute_vertex_normal <\/code>(which we will look at next) from our main function.<\/p>\n\n\n\n
        We must also define the kernel which will run on the device.<\/p>\n\n\n\n
        template <typename T, uint32_t blockThreads>\n__global__ static void compute_vertex_normal(const rxmesh::Context      context,\nrxmesh::VertexAttribute<T> coords,\nrxmesh::VertexAttribute<T> normals)\n{\n\n\n    auto vn_lambda = [&](FaceHandle face_id, VertexIterator& fv){\n        \/\/ get the face's three vertices coordinates\n        glm::fvec3 c0(coords(fv[0], 0), coords(fv[0], 1), coords(fv[0], 2));\n        glm::fvec3 c1(coords(fv[1], 0), coords(fv[1], 1), coords(fv[1], 2));\n        glm::fvec3 c2(coords(fv[2], 0), coords(fv[2], 1), coords(fv[2], 2));\n\n        \/\/ compute the face normal\n        glm::fvec3 n = cross(c1 - c0, c2 - c0);\n\n        \/\/ the three edges length\n        glm::fvec3 l(glm::distance2(c0, c1),\n                     glm::distance2(c1, c2),\n                     glm::distance2(c2, c0));\n\n        \/\/ add the face's normal to its vertices\n        for (uint32_t v = 0; v < 3; ++v) {      \n\/\/ for every vertex in this face\n            for (uint32_t i = 0; i < 3; ++i) {  \n\/\/ for the vertex 3 coordinates\n                atomicAdd(&normals(fv[v], i), n[i] \/ (l[v] + l[(v + 2) % 3]));\n            }\n        }\n    };\n\n    auto block = cooperative_groups::this_thread_block();\n\n    Query<blockThreads> query(context);\n    ShmemAllocator      shrd_alloc;\n    query.dispatch<Op::FV>(block, shrd_alloc, vn_lambda);\n}<\/code><\/pre>\n\n\n\n
        A few things to note about our kernel<\/p>\n\n\n\n
          \n
        1. We pass in a few arguments to our kernel. \n