Format code and change app.py.

.gitignore

@@ -0,0 +1 @@
+__pycache__/

.pre-commit-config.yaml

@@ -0,0 +1,11 @@
+repos:
+  - repo: https://github.com/pre-commit/pre-commit-hooks
+    rev: v2.3.0
+    hooks:
+      - id: check-yaml
+      - id: end-of-file-fixer
+      - id: trailing-whitespace
+  - repo: https://github.com/psf/black
+    rev: 22.10.0
+    hooks:
+      - id: black

README.md

@@ -13,4 +13,4 @@ short_description: Scalable and Versatile 3D Generation from images
 
 Check out the configuration reference at https://huggingface.co/docs/hub/spaces-config-reference
 
-Paper: https://huggingface.co/papers/2412.01506
+Paper: https://huggingface.co/papers/2412.01506

app.py

@@ -4,7 +4,8 @@ from gradio_litmodel3d import LitModel3D
 
 import os
 import shutil
-os.environ['SPCONV_ALGO'] = 'native'
+
+os.environ["SPCONV_ALGO"] = "native"
 from typing import *
 import torch
 import numpy as np
@@ -17,15 +18,24 @@ from trellis.utils import render_utils, postprocessing_utils
 
 
 MAX_SEED = np.iinfo(np.int32).max
-TMP_DIR = os.path.join(os.path.dirname(os.path.abspath(__file__)), 'tmp')
+TMP_DIR = os.path.join(os.path.dirname(os.path.abspath(__file__)), "tmp")
 os.makedirs(TMP_DIR, exist_ok=True)
 
+pipeline = TrellisImageTo3DPipeline.from_pretrained("JeffreyXiang/TRELLIS-image-large")
+pipeline.cuda()
+try:
+    pipeline.preprocess_image(
+        Image.fromarray(np.zeros((512, 512, 3), dtype=np.uint8))
+    )  # Preload rembg
+except:
+    pass
+
 
 def start_session(req: gr.Request):
     user_dir = os.path.join(TMP_DIR, str(req.session_hash))
     os.makedirs(user_dir, exist_ok=True)
-    
-    
+
+
 def end_session(req: gr.Request):
     user_dir = os.path.join(TMP_DIR, str(req.session_hash))
     shutil.rmtree(user_dir)
@@ -48,10 +58,10 @@ def preprocess_image(image: Image.Image) -> Image.Image:
 def preprocess_images(images: List[Tuple[Image.Image, str]]) -> List[Image.Image]:
     """
     Preprocess a list of input images.
-    
+
     Args:
         images (List[Tuple[Image.Image, str]]): The input images.
-        
+
     Returns:
         List[Image.Image]: The preprocessed images.
     """
@@ -62,41 +72,41 @@ def preprocess_images(images: List[Tuple[Image.Image, str]]) -> List[Image.Image
 
 def pack_state(gs: Gaussian, mesh: MeshExtractResult) -> dict:
     return {
-        'gaussian': {
+        "gaussian": {
             **gs.init_params,
-            '_xyz': gs._xyz.cpu().numpy(),
-            '_features_dc': gs._features_dc.cpu().numpy(),
-            '_scaling': gs._scaling.cpu().numpy(),
-            '_rotation': gs._rotation.cpu().numpy(),
-            '_opacity': gs._opacity.cpu().numpy(),
+            "_xyz": gs._xyz.cpu().numpy(),
+            "_features_dc": gs._features_dc.cpu().numpy(),
+            "_scaling": gs._scaling.cpu().numpy(),
+            "_rotation": gs._rotation.cpu().numpy(),
+            "_opacity": gs._opacity.cpu().numpy(),
         },
-        'mesh': {
-            'vertices': mesh.vertices.cpu().numpy(),
-            'faces': mesh.faces.cpu().numpy(),
+        "mesh": {
+            "vertices": mesh.vertices.cpu().numpy(),
+            "faces": mesh.faces.cpu().numpy(),
         },
     }
-    
-    
+
+
 def unpack_state(state: dict) -> Tuple[Gaussian, edict, str]:
     gs = Gaussian(
-        aabb=state['gaussian']['aabb'],
-        sh_degree=state['gaussian']['sh_degree'],
-        mininum_kernel_size=state['gaussian']['mininum_kernel_size'],
-        scaling_bias=state['gaussian']['scaling_bias'],
-        opacity_bias=state['gaussian']['opacity_bias'],
-        scaling_activation=state['gaussian']['scaling_activation'],
+        aabb=state["gaussian"]["aabb"],
+        sh_degree=state["gaussian"]["sh_degree"],
+        mininum_kernel_size=state["gaussian"]["mininum_kernel_size"],
+        scaling_bias=state["gaussian"]["scaling_bias"],
+        opacity_bias=state["gaussian"]["opacity_bias"],
+        scaling_activation=state["gaussian"]["scaling_activation"],
     )
-    gs._xyz = torch.tensor(state['gaussian']['_xyz'], device='cuda')
-    gs._features_dc = torch.tensor(state['gaussian']['_features_dc'], device='cuda')
-    gs._scaling = torch.tensor(state['gaussian']['_scaling'], device='cuda')
-    gs._rotation = torch.tensor(state['gaussian']['_rotation'], device='cuda')
-    gs._opacity = torch.tensor(state['gaussian']['_opacity'], device='cuda')
-    
+    gs._xyz = torch.tensor(state["gaussian"]["_xyz"], device="cuda")
+    gs._features_dc = torch.tensor(state["gaussian"]["_features_dc"], device="cuda")
+    gs._scaling = torch.tensor(state["gaussian"]["_scaling"], device="cuda")
+    gs._rotation = torch.tensor(state["gaussian"]["_rotation"], device="cuda")
+    gs._opacity = torch.tensor(state["gaussian"]["_opacity"], device="cuda")
+
     mesh = edict(
-        vertices=torch.tensor(state['mesh']['vertices'], device='cuda'),
-        faces=torch.tensor(state['mesh']['faces'], device='cuda'),
+        vertices=torch.tensor(state["mesh"]["vertices"], device="cuda"),
+        faces=torch.tensor(state["mesh"]["faces"], device="cuda"),
     )
-    
+
     return gs, mesh
 
 
@@ -170,12 +180,14 @@ def image_to_3d(
             },
             mode=multiimage_algo,
         )
-    video = render_utils.render_video(outputs['gaussian'][0], num_frames=120)['color']
-    video_geo = render_utils.render_video(outputs['mesh'][0], num_frames=120)['normal']
-    video = [np.concatenate([video[i], video_geo[i]], axis=1) for i in range(len(video))]
-    video_path = os.path.join(user_dir, 'sample.mp4')
+    video = render_utils.render_video(outputs["gaussian"][0], num_frames=120)["color"]
+    video_geo = render_utils.render_video(outputs["mesh"][0], num_frames=120)["normal"]
+    video = [
+        np.concatenate([video[i], video_geo[i]], axis=1) for i in range(len(video))
+    ]
+    video_path = os.path.join(user_dir, "sample.mp4")
     imageio.mimsave(video_path, video, fps=15)
-    state = pack_state(outputs['gaussian'][0], outputs['mesh'][0])
+    state = pack_state(outputs["gaussian"][0], outputs["mesh"][0])
     torch.cuda.empty_cache()
     return state, video_path
 
@@ -200,8 +212,10 @@ def extract_glb(
     """
     user_dir = os.path.join(TMP_DIR, str(req.session_hash))
     gs, mesh = unpack_state(state)
-    glb = postprocessing_utils.to_glb(gs, mesh, simplify=mesh_simplify, texture_size=texture_size, verbose=False)
-    glb_path = os.path.join(user_dir, 'sample.glb')
+    glb = postprocessing_utils.to_glb(
+        gs, mesh, simplify=mesh_simplify, texture_size=texture_size, verbose=False
+    )
+    glb_path = os.path.join(user_dir, "sample.glb")
     glb.export(glb_path)
     torch.cuda.empty_cache()
     return glb_path, glb_path
@@ -220,19 +234,21 @@ def extract_gaussian(state: dict, req: gr.Request) -> Tuple[str, str]:
     """
     user_dir = os.path.join(TMP_DIR, str(req.session_hash))
     gs, _ = unpack_state(state)
-    gaussian_path = os.path.join(user_dir, 'sample.ply')
+    gaussian_path = os.path.join(user_dir, "sample.ply")
     gs.save_ply(gaussian_path)
     torch.cuda.empty_cache()
     return gaussian_path, gaussian_path
 
 
 def prepare_multi_example() -> List[Image.Image]:
-    multi_case = list(set([i.split('_')[0] for i in os.listdir("assets/example_multi_image")]))
+    multi_case = list(
+        set([i.split("_")[0] for i in os.listdir("assets/example_multi_image")])
+    )
     images = []
     for case in multi_case:
         _images = []
         for i in range(1, 4):
-            img = Image.open(f'assets/example_multi_image/{case}_{i}.png')
+            img = Image.open(f"assets/example_multi_image/{case}_{i}.png")
             W, H = img.size
             img = img.resize((int(W / H * 512), 512))
             _images.append(np.array(img))
@@ -246,71 +262,113 @@ def split_image(image: Image.Image) -> List[Image.Image]:
     """
     image = np.array(image)
     alpha = image[..., 3]
-    alpha = np.any(alpha>0, axis=0)
+    alpha = np.any(alpha > 0, axis=0)
     start_pos = np.where(~alpha[:-1] & alpha[1:])[0].tolist()
     end_pos = np.where(alpha[:-1] & ~alpha[1:])[0].tolist()
     images = []
     for s, e in zip(start_pos, end_pos):
-        images.append(Image.fromarray(image[:, s:e+1]))
+        images.append(Image.fromarray(image[:, s : e + 1]))
     return [preprocess_image(image) for image in images]
 
 
 with gr.Blocks(delete_cache=(600, 600)) as demo:
-    gr.Markdown("""
+    gr.Markdown(
+        """
     ## Image to 3D Asset with [TRELLIS](https://trellis3d.github.io/)
     * Upload an image and click "Generate" to create a 3D asset. If the image has alpha channel, it be used as the mask. Otherwise, we use `rembg` to remove the background.
     * If you find the generated 3D asset satisfactory, click "Extract GLB" to extract the GLB file and download it.
-    
+
     ✨New: 1) Experimental multi-image support. 2) Gaussian file extraction.
-    """)
-    
+    """
+    )
+
     with gr.Row():
         with gr.Column():
             with gr.Tabs() as input_tabs:
                 with gr.Tab(label="Single Image", id=0) as single_image_input_tab:
-                    image_prompt = gr.Image(label="Image Prompt", format="png", image_mode="RGBA", type="pil", height=300)
+                    image_prompt = gr.Image(
+                        label="Image Prompt",
+                        format="png",
+                        image_mode="RGBA",
+                        type="pil",
+                        height=300,
+                    )
                 with gr.Tab(label="Multiple Images", id=1) as multiimage_input_tab:
-                    multiimage_prompt = gr.Gallery(label="Image Prompt", format="png", type="pil", height=300, columns=3)
-                    gr.Markdown("""
-                        Input different views of the object in separate images. 
-                        
+                    multiimage_prompt = gr.Gallery(
+                        label="Image Prompt",
+                        format="png",
+                        type="pil",
+                        height=300,
+                        columns=3,
+                    )
+                    gr.Markdown(
+                        """
+                        Input different views of the object in separate images.
+
                         *NOTE: this is an experimental algorithm without training a specialized model. It may not produce the best results for all images, especially those having different poses or inconsistent details.*
-                    """)
-            
+                    """
+                    )
+
             with gr.Accordion(label="Generation Settings", open=False):
                 seed = gr.Slider(0, MAX_SEED, label="Seed", value=0, step=1)
                 randomize_seed = gr.Checkbox(label="Randomize Seed", value=True)
                 gr.Markdown("Stage 1: Sparse Structure Generation")
                 with gr.Row():
-                    ss_guidance_strength = gr.Slider(0.0, 10.0, label="Guidance Strength", value=7.5, step=0.1)
-                    ss_sampling_steps = gr.Slider(1, 50, label="Sampling Steps", value=12, step=1)
+                    ss_guidance_strength = gr.Slider(
+                        0.0, 10.0, label="Guidance Strength", value=7.5, step=0.1
+                    )
+                    ss_sampling_steps = gr.Slider(
+                        1, 50, label="Sampling Steps", value=12, step=1
+                    )
                 gr.Markdown("Stage 2: Structured Latent Generation")
                 with gr.Row():
-                    slat_guidance_strength = gr.Slider(0.0, 10.0, label="Guidance Strength", value=3.0, step=0.1)
-                    slat_sampling_steps = gr.Slider(1, 50, label="Sampling Steps", value=12, step=1)
-                multiimage_algo = gr.Radio(["stochastic", "multidiffusion"], label="Multi-image Algorithm", value="stochastic")
+                    slat_guidance_strength = gr.Slider(
+                        0.0, 10.0, label="Guidance Strength", value=3.0, step=0.1
+                    )
+                    slat_sampling_steps = gr.Slider(
+                        1, 50, label="Sampling Steps", value=12, step=1
+                    )
+                multiimage_algo = gr.Radio(
+                    ["stochastic", "multidiffusion"],
+                    label="Multi-image Algorithm",
+                    value="stochastic",
+                )
 
             generate_btn = gr.Button("Generate")
-            
+
             with gr.Accordion(label="GLB Extraction Settings", open=False):
-                mesh_simplify = gr.Slider(0.0, 0.98, label="Simplify", value=0.0, step=0.01)
-                texture_size = gr.Slider(512, 2048, label="Texture Size", value=2048, step=512)
-            
+                mesh_simplify = gr.Slider(
+                    0.0, 0.98, label="Simplify", value=0.0, step=0.01
+                )
+                texture_size = gr.Slider(
+                    512, 2048, label="Texture Size", value=2048, step=512
+                )
+
             with gr.Row():
                 extract_glb_btn = gr.Button("Extract GLB", interactive=False)
                 extract_gs_btn = gr.Button("Extract Gaussian", interactive=False)
-            gr.Markdown("""
+            gr.Markdown(
+                """
                         *NOTE: Gaussian file can be very large (~50MB), it will take a while to display and download.*
-                        """)
+                        """
+            )
 
         with gr.Column():
-            video_output = gr.Video(label="Generated 3D Asset", autoplay=True, loop=True, height=300)
-            model_output = LitModel3D(label="Extracted GLB/Gaussian", exposure=10.0, height=300)
-            
+            video_output = gr.Video(
+                label="Generated 3D Asset", autoplay=True, loop=True, height=300
+            )
+            model_output = LitModel3D(
+                label="Extracted GLB/Gaussian", exposure=10.0, height=300
+            )
+
             with gr.Row():
-                download_glb = gr.DownloadButton(label="Download GLB", interactive=False)
-                download_gs = gr.DownloadButton(label="Download Gaussian", interactive=False)  
-    
+                download_glb = gr.DownloadButton(
+                    label="Download GLB", interactive=False
+                )
+                download_gs = gr.DownloadButton(
+                    label="Download Gaussian", interactive=False
+                )
+
     is_multiimage = gr.State(False)
     output_buf = gr.State()
 
@@ -318,7 +376,7 @@ with gr.Blocks(delete_cache=(600, 600)) as demo:
     with gr.Row() as single_image_example:
         examples = gr.Examples(
             examples=[
-                f'assets/example_image/{image}'
+                f"assets/example_image/{image}"
                 for image in os.listdir("assets/example_image")
             ],
             inputs=[image_prompt],
@@ -340,16 +398,20 @@ with gr.Blocks(delete_cache=(600, 600)) as demo:
     # Handlers
     demo.load(start_session)
     demo.unload(end_session)
-    
+
     single_image_input_tab.select(
-        lambda: tuple([False, gr.Row.update(visible=True), gr.Row.update(visible=False)]),
-        outputs=[is_multiimage, single_image_example, multiimage_example]
+        lambda: tuple(
+            [False, gr.Row.update(visible=True), gr.Row.update(visible=False)]
+        ),
+        outputs=[is_multiimage, single_image_example, multiimage_example],
     )
     multiimage_input_tab.select(
-        lambda: tuple([True, gr.Row.update(visible=False), gr.Row.update(visible=True)]),
-        outputs=[is_multiimage, single_image_example, multiimage_example]
+        lambda: tuple(
+            [True, gr.Row.update(visible=False), gr.Row.update(visible=True)]
+        ),
+        outputs=[is_multiimage, single_image_example, multiimage_example],
     )
-    
+
     image_prompt.upload(
         preprocess_image,
         inputs=[image_prompt],
@@ -361,13 +423,19 @@ with gr.Blocks(delete_cache=(600, 600)) as demo:
         outputs=[multiimage_prompt],
     )
 
-    generate_btn.click(
-        get_seed,
-        inputs=[randomize_seed, seed],
-        outputs=[seed],
-    ).then(
+    generate_btn.click(get_seed, inputs=[randomize_seed, seed], outputs=[seed],).then(
         image_to_3d,
-        inputs=[image_prompt, multiimage_prompt, is_multiimage, seed, ss_guidance_strength, ss_sampling_steps, slat_guidance_strength, slat_sampling_steps, multiimage_algo],
+        inputs=[
+            image_prompt,
+            multiimage_prompt,
+            is_multiimage,
+            seed,
+            ss_guidance_strength,
+            ss_sampling_steps,
+            slat_guidance_strength,
+            slat_sampling_steps,
+            multiimage_algo,
+        ],
         outputs=[output_buf, video_output],
     ).then(
         lambda: tuple([gr.Button(interactive=True), gr.Button(interactive=True)]),
@@ -387,7 +455,7 @@ with gr.Blocks(delete_cache=(600, 600)) as demo:
         lambda: gr.Button(interactive=True),
         outputs=[download_glb],
     )
-    
+
     extract_gs_btn.click(
         extract_gaussian,
         inputs=[output_buf],
@@ -401,14 +469,8 @@ with gr.Blocks(delete_cache=(600, 600)) as demo:
         lambda: gr.Button(interactive=False),
         outputs=[download_glb],
     )
-    
+
 
 # Launch the Gradio app
 if __name__ == "__main__":
-    pipeline = TrellisImageTo3DPipeline.from_pretrained("JeffreyXiang/TRELLIS-image-large")
-    pipeline.cuda()
-    try:
-        pipeline.preprocess_image(Image.fromarray(np.zeros((512, 512, 3), dtype=np.uint8)))    # Preload rembg
-    except:
-        pass
     demo.launch()

extensions/nvdiffrast/README.md

@@ -21,7 +21,7 @@ We do not currently accept outside code contributions in the form of pull reques
 
 Environment map stored as part of `samples/data/envphong.npz` is derived from a Wave Engine
 [sample material](https://github.com/WaveEngine/Samples-2.5/tree/master/Materials/EnvironmentMap/Content/Assets/CubeMap.cubemap)
-originally shared under 
+originally shared under
 [MIT License](https://github.com/WaveEngine/Samples-2.5/blob/master/LICENSE.md).
 Mesh and texture stored as part of `samples/data/earth.npz` are derived from
 [3D Earth Photorealistic 2K](https://www.turbosquid.com/3d-models/3d-realistic-earth-photorealistic-2k-1279125)

extensions/nvdiffrast/nvdiffrast/__init__.py

@@ -6,4 +6,4 @@
 # distribution of this software and related documentation without an express
 # license agreement from NVIDIA CORPORATION is strictly prohibited.
 
-__version__ = '0.3.3'
+__version__ = "0.3.3"

extensions/nvdiffrast/nvdiffrast/common/antialias.cu

@@ -112,7 +112,7 @@ static __device__ __forceinline__ void evhash_insert_vertex(const AntialiasKerne
 {
     if (va == vb)
         return;
-    
+
     uint64_t v0 = (uint32_t)min(va, vb) + 1; // canonical vertex order
     uint64_t v1 = (uint32_t)max(va, vb) + 1;
     uint64_t vk = v0 | (v1 << 32); // hash key

extensions/nvdiffrast/nvdiffrast/common/cudaraster/CudaRaster.hpp

@@ -60,4 +60,3 @@ private:
 
 //------------------------------------------------------------------------
 } // namespace CR
-

extensions/nvdiffrast/nvdiffrast/common/cudaraster/impl/RasterImpl.hpp

@@ -99,4 +99,3 @@ private:
 
 //------------------------------------------------------------------------
 } // namespace CR
-

extensions/nvdiffrast/nvdiffrast/common/interpolate.cu

@@ -94,9 +94,9 @@ static __forceinline__ __device__ void InterpolateFwdKernelTemplate(const Interp
     float dvdx = db.z;
     float dvdy = db.w;
 
-    // Calculate the pixel differentials of chosen attributes.    
+    // Calculate the pixel differentials of chosen attributes.
     for (int i=0; i < p.numDiffAttr; i++)
-    {   
+    {
         // Input attribute index.
         int j = p.diff_attrs_all ? i : p.diffAttrs[i];
         if (j < 0)

extensions/nvdiffrast/nvdiffrast/common/texture.cpp

@@ -47,7 +47,7 @@ void raiseMipSizeError(NVDR_CTX_ARGS, const TextureKernelParams& p)
 
         // Append level size to error message.
         snprintf(buf, bufsz, "mip %-2d ", level);
-        msg += buf; 
+        msg += buf;
         if (ew) snprintf(buf, bufsz, "  err  ");
         else    snprintf(buf, bufsz, "%5d  ", w);
         msg += buf;

extensions/nvdiffrast/nvdiffrast/tensorflow/ops.py

@@ -11,22 +11,26 @@ import numpy as np
 import os
 from . import plugin_loader
 
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 # Helpers.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 
 # OpenGL-related linker options depending on platform.
 def _get_gl_opts():
     libs = {
-        'posix': ['GL', 'EGL'],
-        'nt':    ['gdi32', 'opengl32', 'user32', 'setgpu'],
+        "posix": ["GL", "EGL"],
+        "nt": ["gdi32", "opengl32", "user32", "setgpu"],
     }
-    return ['-l' + x for x in libs[os.name]]
+    return ["-l" + x for x in libs[os.name]]
+
 
 # Load the cpp plugin.
 def _get_plugin():
-    fn = os.path.join(os.path.dirname(__file__), 'tf_all.cu')
-    return plugin_loader.get_plugin(fn, extra_nvcc_options=_get_gl_opts() + ['-DNVDR_TENSORFLOW'])
+    fn = os.path.join(os.path.dirname(__file__), "tf_all.cu")
+    return plugin_loader.get_plugin(
+        fn, extra_nvcc_options=_get_gl_opts() + ["-DNVDR_TENSORFLOW"]
+    )
+
 
 # Convert parameter to a numpy array if possible.
 def _get_constant(x, dtype):
@@ -35,19 +39,24 @@ def _get_constant(x, dtype):
     except (TypeError, ValueError):
         return None
 
+
 # Tests for a construction-time constantness instead of tf.constant node because
 # the latter can be overridden in Session.run() feed_dict at evaluation time.
 def _is_constant(x, dtype):
     if isinstance(x, np.ndarray):
-        return np.can_cast(x.dtype, dtype, 'unsafe')
+        return np.can_cast(x.dtype, dtype, "unsafe")
     else:
         return _get_constant(x, dtype) is not None
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Rasterize.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 
-def rasterize(pos, tri, resolution, ranges=None, tri_const=False, output_db=True, grad_db=True):
+
+def rasterize(
+    pos, tri, resolution, ranges=None, tri_const=False, output_db=True, grad_db=True
+):
     assert tri_const is True or tri_const is False
     assert output_db is True or output_db is False
 
@@ -63,15 +72,19 @@ def rasterize(pos, tri, resolution, ranges=None, tri_const=False, output_db=True
     pos = tf.convert_to_tensor(pos, dtype=tf.float32)
     resolution = tf.convert_to_tensor(resolution, dtype=tf.int32)
     if ranges is None:
-        ranges = tf.convert_to_tensor(np.zeros(shape=[0, 2], dtype=np.int32)) # Empty tensor.
+        ranges = tf.convert_to_tensor(
+            np.zeros(shape=[0, 2], dtype=np.int32)
+        )  # Empty tensor.
     else:
-        ranges = tf.convert_to_tensor(ranges, dtype=tf.int32) # Convert input to tensor.
+        ranges = tf.convert_to_tensor(
+            ranges, dtype=tf.int32
+        )  # Convert input to tensor.
 
     # Infer as much about the output shape as possible.
     out_shape = [None, None, None, 4]
-    if pos.shape.rank == 3: # Instanced mode.
+    if pos.shape.rank == 3:  # Instanced mode.
         out_shape[0] = pos.shape[0].value
-    elif pos.shape.rank == 2: # Range mode.
+    elif pos.shape.rank == 2:  # Range mode.
         if ranges.shape.rank not in [None, 0]:
             out_shape[0] = ranges.shape[0].value
     if resolution_c is not None:
@@ -81,24 +94,32 @@ def rasterize(pos, tri, resolution, ranges=None, tri_const=False, output_db=True
     # Output pixel differentials.
     @tf.custom_gradient
     def func_db(pos):
-        out, out_db = _get_plugin().rasterize_fwd(pos, tri, resolution, ranges, 1, tri_const)
+        out, out_db = _get_plugin().rasterize_fwd(
+            pos, tri, resolution, ranges, 1, tri_const
+        )
         out.set_shape(out_shape)
         out_db.set_shape(out_shape)
+
         def grad(dy, ddb):
             if grad_db:
                 return _get_plugin().rasterize_grad_db(pos, tri, out, dy, ddb)
             else:
                 return _get_plugin().rasterize_grad(pos, tri, out, dy)
+
         return (out, out_db), grad
 
     # Do not output pixel differentials.
     @tf.custom_gradient
     def func(pos):
-        out, out_db = _get_plugin().rasterize_fwd(pos, tri, resolution, ranges, 0, tri_const)
+        out, out_db = _get_plugin().rasterize_fwd(
+            pos, tri, resolution, ranges, 0, tri_const
+        )
         out.set_shape(out_shape)
-        out_db.set_shape(out_shape[:-1] + [0]) # Zero channels in out_db.
+        out_db.set_shape(out_shape[:-1] + [0])  # Zero channels in out_db.
+
         def grad(dy, _):
             return _get_plugin().rasterize_grad(pos, tri, out, dy)
+
         return (out, out_db), grad
 
     # Choose stub.
@@ -107,15 +128,17 @@ def rasterize(pos, tri, resolution, ranges=None, tri_const=False, output_db=True
     else:
         return func(pos)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Interpolate.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     # Sanitize the list of pixel differential attributes.
     if diff_attrs is None:
         diff_attrs = []
-    elif diff_attrs != 'all':
+    elif diff_attrs != "all":
         diff_attrs = _get_constant(diff_attrs, np.int32)
         assert (diff_attrs is not None) and len(diff_attrs.shape) == 1
         diff_attrs = diff_attrs.tolist()
@@ -130,16 +153,23 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     # Infer output shape.
     out_shape = [None, None, None, None]
     if rast.shape.rank is not None:
-        out_shape = [rast.shape[0].value, rast.shape[1].value, rast.shape[2].value, None]
+        out_shape = [
+            rast.shape[0].value,
+            rast.shape[1].value,
+            rast.shape[2].value,
+            None,
+        ]
     if attr.shape.rank in [2, 3]:
         out_shape[3] = attr.shape[-1].value
 
     # Output pixel differentials for at least some attributes.
     @tf.custom_gradient
     def func_da(attr, rast, rast_db):
-        diff_attrs_all = int(diff_attrs == 'all')
+        diff_attrs_all = int(diff_attrs == "all")
         diff_attrs_list = [] if diff_attrs_all else diff_attrs
-        out, out_da = _get_plugin().interpolate_fwd_da(attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list)
+        out, out_da = _get_plugin().interpolate_fwd_da(
+            attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list
+        )
 
         # Infer number of channels in out_da.
         if not diff_attrs_all:
@@ -154,7 +184,10 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
         out_da.set_shape([out_shape[0], out_shape[1], out_shape[2], da_channels])
 
         def grad(dy, dda):
-            return _get_plugin().interpolate_grad_da(attr, rast, tri, dy, rast_db, dda, diff_attrs_all, diff_attrs_list)
+            return _get_plugin().interpolate_grad_da(
+                attr, rast, tri, dy, rast_db, dda, diff_attrs_all, diff_attrs_list
+            )
+
         return (out, out_da), grad
 
     # No pixel differentials for any attribute.
@@ -162,9 +195,11 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     def func(attr, rast):
         out, out_da = _get_plugin().interpolate_fwd(attr, rast, tri)
         out.set_shape(out_shape)
-        out_da.set_shape(out_shape[:-1] + [0]) # Zero channels in out_da.
+        out_da.set_shape(out_shape[:-1] + [0])  # Zero channels in out_da.
+
         def grad(dy, _):
             return _get_plugin().interpolate_grad(attr, rast, tri, dy)
+
         return (out, out_da), grad
 
     # Choose stub.
@@ -173,16 +208,26 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     else:
         return func(attr, rast)
 
-#----------------------------------------------------------------------------
-# Texture.
-#----------------------------------------------------------------------------
 
-def texture(tex, uv, uv_da=None, filter_mode='auto', boundary_mode='wrap', tex_const=False, max_mip_level=None):
+# ----------------------------------------------------------------------------
+# Texture.
+# ----------------------------------------------------------------------------
+
+
+def texture(
+    tex,
+    uv,
+    uv_da=None,
+    filter_mode="auto",
+    boundary_mode="wrap",
+    tex_const=False,
+    max_mip_level=None,
+):
     assert tex_const is True or tex_const is False
 
     # Default filter mode.
-    if filter_mode == 'auto':
-        filter_mode = 'linear-mipmap-linear' if (uv_da is not None) else 'linear'
+    if filter_mode == "auto":
+        filter_mode = "linear-mipmap-linear" if (uv_da is not None) else "linear"
 
     # Known constant texture?
     tex_const = tex_const or _is_constant(tex, np.float32)
@@ -198,7 +243,7 @@ def texture(tex, uv, uv_da=None, filter_mode='auto', boundary_mode='wrap', tex_c
     # Convert inputs to tensors.
     tex = tf.convert_to_tensor(tex, dtype=tf.float32)
     uv = tf.convert_to_tensor(uv, dtype=tf.float32)
-    if 'mipmap' in filter_mode:
+    if "mipmap" in filter_mode:
         uv_da = tf.convert_to_tensor(uv_da, dtype=tf.float32)
 
     # Infer output shape.
@@ -207,37 +252,83 @@ def texture(tex, uv, uv_da=None, filter_mode='auto', boundary_mode='wrap', tex_c
         assert uv.shape.rank == 4
         out_shape = [uv.shape[0].value, uv.shape[1].value, uv.shape[2].value, None]
     if tex.shape.rank is not None:
-        assert tex.shape.rank == (5 if boundary_mode == 'cube' else 4)
+        assert tex.shape.rank == (5 if boundary_mode == "cube" else 4)
         out_shape[-1] = tex.shape[-1].value
 
     # If mipping disabled via max level=0, we may as well use simpler filtering internally.
-    if max_mip_level == 0 and filter_mode in ['linear-mipmap-nearest', 'linear-mipmap-linear']:
-        filter_mode = 'linear'
+    if max_mip_level == 0 and filter_mode in [
+        "linear-mipmap-nearest",
+        "linear-mipmap-linear",
+    ]:
+        filter_mode = "linear"
 
     # Convert filter mode to internal enumeration.
-    filter_mode_dict = {'nearest': 0, 'linear': 1, 'linear-mipmap-nearest': 2, 'linear-mipmap-linear': 3}
+    filter_mode_dict = {
+        "nearest": 0,
+        "linear": 1,
+        "linear-mipmap-nearest": 2,
+        "linear-mipmap-linear": 3,
+    }
     filter_mode_enum = filter_mode_dict[filter_mode]
 
     # Convert boundary mode to internal enumeration.
-    boundary_mode_dict = {'cube': 0, 'wrap': 1, 'clamp': 2, 'zero': 3}
+    boundary_mode_dict = {"cube": 0, "wrap": 1, "clamp": 2, "zero": 3}
     boundary_mode_enum = boundary_mode_dict[boundary_mode]
 
     # Linear-mipmap-linear: Mipmaps enabled, all gradients active.
     @tf.custom_gradient
     def func_linear_mipmap_linear(tex, uv, uv_da):
-        out, mip = _get_plugin().texture_fwd_mip(tex, uv, uv_da, filter_mode_enum, boundary_mode_enum, tex_const, max_mip_level)
+        out, mip = _get_plugin().texture_fwd_mip(
+            tex,
+            uv,
+            uv_da,
+            filter_mode_enum,
+            boundary_mode_enum,
+            tex_const,
+            max_mip_level,
+        )
         out.set_shape(out_shape)
+
         def grad(dy):
-            return _get_plugin().texture_grad_linear_mipmap_linear(tex, uv, dy, uv_da, mip, filter_mode_enum, boundary_mode_enum, max_mip_level)
+            return _get_plugin().texture_grad_linear_mipmap_linear(
+                tex,
+                uv,
+                dy,
+                uv_da,
+                mip,
+                filter_mode_enum,
+                boundary_mode_enum,
+                max_mip_level,
+            )
+
         return out, grad
 
     # Linear-mipmap-nearest: Mipmaps enabled, no gradients to uv_da.
     @tf.custom_gradient
     def func_linear_mipmap_nearest(tex, uv):
-        out, mip = _get_plugin().texture_fwd_mip(tex, uv, uv_da, filter_mode_enum, boundary_mode_enum, tex_const, max_mip_level)
+        out, mip = _get_plugin().texture_fwd_mip(
+            tex,
+            uv,
+            uv_da,
+            filter_mode_enum,
+            boundary_mode_enum,
+            tex_const,
+            max_mip_level,
+        )
         out.set_shape(out_shape)
+
         def grad(dy):
-            return _get_plugin().texture_grad_linear_mipmap_nearest(tex, uv, dy, uv_da, mip, filter_mode_enum, boundary_mode_enum, max_mip_level)
+            return _get_plugin().texture_grad_linear_mipmap_nearest(
+                tex,
+                uv,
+                dy,
+                uv_da,
+                mip,
+                filter_mode_enum,
+                boundary_mode_enum,
+                max_mip_level,
+            )
+
         return out, grad
 
     # Linear: Mipmaps disabled, no uv_da, no gradients to uv_da.
@@ -245,8 +336,12 @@ def texture(tex, uv, uv_da=None, filter_mode='auto', boundary_mode='wrap', tex_c
     def func_linear(tex, uv):
         out = _get_plugin().texture_fwd(tex, uv, filter_mode_enum, boundary_mode_enum)
         out.set_shape(out_shape)
+
         def grad(dy):
-            return _get_plugin().texture_grad_linear(tex, uv, dy, filter_mode_enum, boundary_mode_enum)
+            return _get_plugin().texture_grad_linear(
+                tex, uv, dy, filter_mode_enum, boundary_mode_enum
+            )
+
         return out, grad
 
     # Nearest: Mipmaps disabled, no uv_da, no gradients to uv_da or uv.
@@ -254,23 +349,29 @@ def texture(tex, uv, uv_da=None, filter_mode='auto', boundary_mode='wrap', tex_c
     def func_nearest(tex):
         out = _get_plugin().texture_fwd(tex, uv, filter_mode_enum, boundary_mode_enum)
         out.set_shape(out_shape)
+
         def grad(dy):
-            return _get_plugin().texture_grad_nearest(tex, uv, dy, filter_mode_enum, boundary_mode_enum)
+            return _get_plugin().texture_grad_nearest(
+                tex, uv, dy, filter_mode_enum, boundary_mode_enum
+            )
+
         return out, grad
 
     # Choose stub.
-    if filter_mode == 'linear-mipmap-linear':
+    if filter_mode == "linear-mipmap-linear":
         return func_linear_mipmap_linear(tex, uv, uv_da)
-    elif filter_mode == 'linear-mipmap-nearest':
+    elif filter_mode == "linear-mipmap-nearest":
         return func_linear_mipmap_nearest(tex, uv)
-    elif filter_mode == 'linear':
+    elif filter_mode == "linear":
         return func_linear(tex, uv)
-    elif filter_mode == 'nearest':
+    elif filter_mode == "nearest":
         return func_nearest(tex)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Antialias.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 def antialias(color, rast, pos, tri, tri_const=False, pos_gradient_boost=1.0):
     assert tri_const is True or tri_const is False
@@ -289,15 +390,22 @@ def antialias(color, rast, pos, tri, tri_const=False, pos_gradient_boost=1.0):
 
     @tf.custom_gradient
     def func(color, pos):
-        color_out, work_buffer = _get_plugin().antialias_fwd(color, rast, pos, tri, tri_const)
+        color_out, work_buffer = _get_plugin().antialias_fwd(
+            color, rast, pos, tri, tri_const
+        )
         color_out.set_shape(color.shape)
+
         def grad(dy):
-            grad_color, grad_pos = _get_plugin().antialias_grad(color, rast, pos, tri, dy, work_buffer)
+            grad_color, grad_pos = _get_plugin().antialias_grad(
+                color, rast, pos, tri, dy, work_buffer
+            )
             if pos_gradient_boost != 1.0:
                 grad_pos = grad_pos * pos_gradient_boost
             return grad_color, grad_pos
+
         return color_out, grad
 
     return func(color, pos)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------

extensions/nvdiffrast/nvdiffrast/tensorflow/plugin_loader.py

@@ -14,15 +14,16 @@ import hashlib
 import tempfile
 import shutil
 import tensorflow as tf
-from tensorflow.python.client import device_lib # pylint: disable=no-name-in-module
+from tensorflow.python.client import device_lib  # pylint: disable=no-name-in-module
 
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 # Global options.
 
 _nvdiffrast_cache_dir = None
 
+
 def set_cache_dir(path: str) -> None:
-    '''Set CUDA kernel compilation temp dir.
+    """Set CUDA kernel compilation temp dir.
 
     If `set_cache_dir` is not called, the cache directory will default to
     one of the below:
@@ -33,103 +34,164 @@ def set_cache_dir(path: str) -> None:
 
     Args:
       path: Where to save CUDA kernel build temporaries
-    '''
+    """
     global _nvdiffrast_cache_dir
     _nvdiffrast_cache_dir = path
 
+
 def make_cache_dir_path(*paths: str) -> str:
     if _nvdiffrast_cache_dir is not None:
         return os.path.join(_nvdiffrast_cache_dir, *paths)
-    if 'NVDIFFRAST_CACHE_DIR' in os.environ:
-        return os.path.join(os.environ['NVDIFFRAST_CACHE_DIR'], *paths)
-    if 'HOME' in os.environ:
-        return os.path.join(os.environ['HOME'], '.cache', 'nvdiffrast', *paths)
-    if 'USERPROFILE' in os.environ:
-        return os.path.join(os.environ['USERPROFILE'], '.cache', 'nvdiffrast', *paths)
-    return os.path.join(tempfile.gettempdir(), '.cache', 'nvdiffrast', *paths)
-
-cuda_cache_version_tag = 'v1'
-do_not_hash_included_headers = False # Speed up compilation by assuming that headers included by the CUDA code never change. Unsafe!
-verbose = True # Print status messages to stdout.
-
-#----------------------------------------------------------------------------
+    if "NVDIFFRAST_CACHE_DIR" in os.environ:
+        return os.path.join(os.environ["NVDIFFRAST_CACHE_DIR"], *paths)
+    if "HOME" in os.environ:
+        return os.path.join(os.environ["HOME"], ".cache", "nvdiffrast", *paths)
+    if "USERPROFILE" in os.environ:
+        return os.path.join(os.environ["USERPROFILE"], ".cache", "nvdiffrast", *paths)
+    return os.path.join(tempfile.gettempdir(), ".cache", "nvdiffrast", *paths)
+
+
+cuda_cache_version_tag = "v1"
+do_not_hash_included_headers = False  # Speed up compilation by assuming that headers included by the CUDA code never change. Unsafe!
+verbose = True  # Print status messages to stdout.
+
+# ----------------------------------------------------------------------------
 # Internal helper funcs.
 
+
 def _find_compiler_bindir():
-    hostx64_paths = sorted(glob.glob('C:/Program Files/Microsoft Visual Studio/*/Enterprise/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files/Microsoft Visual Studio/*/Enterprise/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files (x86)/Microsoft Visual Studio/*/Enterprise/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files (x86)/Microsoft Visual Studio/*/Enterprise/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files/Microsoft Visual Studio/*/Professional/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files/Microsoft Visual Studio/*/Professional/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files (x86)/Microsoft Visual Studio/*/Professional/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files (x86)/Microsoft Visual Studio/*/Professional/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files/Microsoft Visual Studio/*/BuildTools/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files/Microsoft Visual Studio/*/BuildTools/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files (x86)/Microsoft Visual Studio/*/BuildTools/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files (x86)/Microsoft Visual Studio/*/BuildTools/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files/Microsoft Visual Studio/*/Community/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files/Microsoft Visual Studio/*/Community/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    hostx64_paths = sorted(glob.glob('C:/Program Files (x86)/Microsoft Visual Studio/*/Community/VC/Tools/MSVC/*/bin/Hostx64/x64'), reverse=True)
+    hostx64_paths = sorted(
+        glob.glob(
+            "C:/Program Files (x86)/Microsoft Visual Studio/*/Community/VC/Tools/MSVC/*/bin/Hostx64/x64"
+        ),
+        reverse=True,
+    )
     if hostx64_paths != []:
         return hostx64_paths[0]
-    vc_bin_dir = 'C:/Program Files (x86)/Microsoft Visual Studio 14.0/vc/bin'
+    vc_bin_dir = "C:/Program Files (x86)/Microsoft Visual Studio 14.0/vc/bin"
     if os.path.isdir(vc_bin_dir):
         return vc_bin_dir
     return None
 
+
 def _get_compute_cap(device):
     caps_str = device.physical_device_desc
-    m = re.search('compute capability: (\\d+).(\\d+)', caps_str)
+    m = re.search("compute capability: (\\d+).(\\d+)", caps_str)
     major = m.group(1)
     minor = m.group(2)
     return (major, minor)
 
+
 def _get_cuda_gpu_arch_string():
-    gpus = [x for x in device_lib.list_local_devices() if x.device_type == 'GPU']
+    gpus = [x for x in device_lib.list_local_devices() if x.device_type == "GPU"]
     if len(gpus) == 0:
-        raise RuntimeError('No GPU devices found')
+        raise RuntimeError("No GPU devices found")
     (major, minor) = _get_compute_cap(gpus[0])
-    return 'sm_%s%s' % (major, minor)
+    return "sm_%s%s" % (major, minor)
+
 
 def _run_cmd(cmd):
     with os.popen(cmd) as pipe:
         output = pipe.read()
         status = pipe.close()
     if status is not None:
-        raise RuntimeError('NVCC returned an error. See below for full command line and output log:\n\n%s\n\n%s' % (cmd, output))
+        raise RuntimeError(
+            "NVCC returned an error. See below for full command line and output log:\n\n%s\n\n%s"
+            % (cmd, output)
+        )
+
 
 def _prepare_nvcc_cli(opts):
-    cmd = 'nvcc ' + opts.strip()
-    cmd += ' --disable-warnings'
+    cmd = "nvcc " + opts.strip()
+    cmd += " --disable-warnings"
     cmd += ' --include-path "%s"' % tf.sysconfig.get_include()
-    cmd += ' --include-path "%s"' % os.path.join(tf.sysconfig.get_include(), 'external', 'protobuf_archive', 'src')
-    cmd += ' --include-path "%s"' % os.path.join(tf.sysconfig.get_include(), 'external', 'com_google_absl')
-    cmd += ' --include-path "%s"' % os.path.join(tf.sysconfig.get_include(), 'external', 'eigen_archive')
+    cmd += ' --include-path "%s"' % os.path.join(
+        tf.sysconfig.get_include(), "external", "protobuf_archive", "src"
+    )
+    cmd += ' --include-path "%s"' % os.path.join(
+        tf.sysconfig.get_include(), "external", "com_google_absl"
+    )
+    cmd += ' --include-path "%s"' % os.path.join(
+        tf.sysconfig.get_include(), "external", "eigen_archive"
+    )
 
     compiler_bindir = _find_compiler_bindir()
     if compiler_bindir is None:
         # Require that _find_compiler_bindir succeeds on Windows.  Allow
         # nvcc to use whatever is the default on Linux.
-        if os.name == 'nt':
-            raise RuntimeError('Could not find MSVC/GCC/CLANG installation on this computer. Check compiler_bindir_search_path list in "%s".' % __file__)
+        if os.name == "nt":
+            raise RuntimeError(
+                'Could not find MSVC/GCC/CLANG installation on this computer. Check compiler_bindir_search_path list in "%s".'
+                % __file__
+            )
     else:
         cmd += ' --compiler-bindir "%s"' % compiler_bindir
-    cmd += ' 2>&1'
+    cmd += " 2>&1"
     return cmd
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Main entry point.
 
 _plugin_cache = dict()
 
+
 def get_plugin(cuda_file, extra_nvcc_options=[]):
     cuda_file_base = os.path.basename(cuda_file)
     cuda_file_name, cuda_file_ext = os.path.splitext(cuda_file_base)
@@ -140,80 +202,112 @@ def get_plugin(cuda_file, extra_nvcc_options=[]):
 
     # Setup plugin.
     if verbose:
-        print('Setting up TensorFlow plugin "%s": ' % cuda_file_base, end='', flush=True)
+        print(
+            'Setting up TensorFlow plugin "%s": ' % cuda_file_base, end="", flush=True
+        )
     try:
         # Hash CUDA source.
         md5 = hashlib.md5()
-        with open(cuda_file, 'rb') as f:
+        with open(cuda_file, "rb") as f:
             md5.update(f.read())
-        md5.update(b'\n')
+        md5.update(b"\n")
 
         # Hash headers included by the CUDA code by running it through the preprocessor.
         if not do_not_hash_included_headers:
             if verbose:
-                print('Preprocessing... ', end='', flush=True)
+                print("Preprocessing... ", end="", flush=True)
             with tempfile.TemporaryDirectory() as tmp_dir:
-                tmp_file = os.path.join(tmp_dir, cuda_file_name + '_tmp' + cuda_file_ext)
-                _run_cmd(_prepare_nvcc_cli('"%s" --preprocess -o "%s" --keep --keep-dir "%s"' % (cuda_file, tmp_file, tmp_dir)))
-                with open(tmp_file, 'rb') as f:
-                    bad_file_str = ('"' + cuda_file.replace('\\', '/') + '"').encode('utf-8') # __FILE__ in error check macros
-                    good_file_str = ('"' + cuda_file_base + '"').encode('utf-8')
+                tmp_file = os.path.join(
+                    tmp_dir, cuda_file_name + "_tmp" + cuda_file_ext
+                )
+                _run_cmd(
+                    _prepare_nvcc_cli(
+                        '"%s" --preprocess -o "%s" --keep --keep-dir "%s"'
+                        % (cuda_file, tmp_file, tmp_dir)
+                    )
+                )
+                with open(tmp_file, "rb") as f:
+                    bad_file_str = ('"' + cuda_file.replace("\\", "/") + '"').encode(
+                        "utf-8"
+                    )  # __FILE__ in error check macros
+                    good_file_str = ('"' + cuda_file_base + '"').encode("utf-8")
                     for ln in f:
-                        if not ln.startswith(b'# ') and not ln.startswith(b'#line '): # ignore line number pragmas
+                        if not ln.startswith(b"# ") and not ln.startswith(
+                            b"#line "
+                        ):  # ignore line number pragmas
                             ln = ln.replace(bad_file_str, good_file_str)
                             md5.update(ln)
-                    md5.update(b'\n')
+                    md5.update(b"\n")
 
         # Select compiler options.
-        compile_opts = ''
-        if os.name == 'nt':
-            compile_opts += '"%s"' % os.path.join(tf.sysconfig.get_lib(), 'python', '_pywrap_tensorflow_internal.lib')
-            compile_opts += ' --library-path="%s"' % (os.path.dirname(__file__) + r"\..\lib") # Find libraries during compilation.
-        elif os.name == 'posix':
-            compile_opts += '"%s"' % os.path.join(tf.sysconfig.get_lib(), 'python', '_pywrap_tensorflow_internal.so')
-            compile_opts += ' --compiler-options \'-fPIC -D_GLIBCXX_USE_CXX11_ABI=0\''
+        compile_opts = ""
+        if os.name == "nt":
+            compile_opts += '"%s"' % os.path.join(
+                tf.sysconfig.get_lib(), "python", "_pywrap_tensorflow_internal.lib"
+            )
+            compile_opts += ' --library-path="%s"' % (
+                os.path.dirname(__file__) + r"\..\lib"
+            )  # Find libraries during compilation.
+        elif os.name == "posix":
+            compile_opts += '"%s"' % os.path.join(
+                tf.sysconfig.get_lib(), "python", "_pywrap_tensorflow_internal.so"
+            )
+            compile_opts += " --compiler-options '-fPIC -D_GLIBCXX_USE_CXX11_ABI=0'"
         else:
-            assert False # not Windows or Linux, w00t?
-        compile_opts += ' --gpu-architecture=%s' % _get_cuda_gpu_arch_string()
-        compile_opts += ' --use_fast_math'
+            assert False  # not Windows or Linux, w00t?
+        compile_opts += " --gpu-architecture=%s" % _get_cuda_gpu_arch_string()
+        compile_opts += " --use_fast_math"
         for opt in extra_nvcc_options:
-            compile_opts += ' ' + opt
+            compile_opts += " " + opt
         nvcc_cmd = _prepare_nvcc_cli(compile_opts)
 
         # Hash build configuration.
-        md5.update(('nvcc_cmd: ' + nvcc_cmd).encode('utf-8') + b'\n')
-        md5.update(('tf.VERSION: ' + tf.VERSION).encode('utf-8') + b'\n')
-        md5.update(('cuda_cache_version_tag: ' + cuda_cache_version_tag).encode('utf-8') + b'\n')
+        md5.update(("nvcc_cmd: " + nvcc_cmd).encode("utf-8") + b"\n")
+        md5.update(("tf.VERSION: " + tf.VERSION).encode("utf-8") + b"\n")
+        md5.update(
+            ("cuda_cache_version_tag: " + cuda_cache_version_tag).encode("utf-8")
+            + b"\n"
+        )
 
         # Compile if not already compiled.
-        bin_file_ext = '.dll' if os.name == 'nt' else '.so'
+        bin_file_ext = ".dll" if os.name == "nt" else ".so"
         cuda_cache_path = make_cache_dir_path()
-        bin_file = os.path.join(make_cache_dir_path(), cuda_file_name + '_' + md5.hexdigest() + bin_file_ext)
+        bin_file = os.path.join(
+            make_cache_dir_path(), cuda_file_name + "_" + md5.hexdigest() + bin_file_ext
+        )
         if not os.path.isfile(bin_file):
             if verbose:
-                print('Compiling... ', end='', flush=True)
+                print("Compiling... ", end="", flush=True)
             with tempfile.TemporaryDirectory() as tmp_dir:
-                tmp_file = os.path.join(tmp_dir, cuda_file_name + '_tmp' + bin_file_ext)
-                _run_cmd(nvcc_cmd + ' "%s" --shared -o "%s" --keep --keep-dir "%s"' % (cuda_file, tmp_file, tmp_dir))
+                tmp_file = os.path.join(tmp_dir, cuda_file_name + "_tmp" + bin_file_ext)
+                _run_cmd(
+                    nvcc_cmd
+                    + ' "%s" --shared -o "%s" --keep --keep-dir "%s"'
+                    % (cuda_file, tmp_file, tmp_dir)
+                )
                 os.makedirs(cuda_cache_path, exist_ok=True)
-                intermediate_file = os.path.join(cuda_cache_path, cuda_file_name + '_' + uuid.uuid4().hex + '_tmp' + bin_file_ext)
+                intermediate_file = os.path.join(
+                    cuda_cache_path,
+                    cuda_file_name + "_" + uuid.uuid4().hex + "_tmp" + bin_file_ext,
+                )
                 shutil.copyfile(tmp_file, intermediate_file)
-                os.rename(intermediate_file, bin_file) # atomic
+                os.rename(intermediate_file, bin_file)  # atomic
 
         # Load.
         if verbose:
-            print('Loading... ', end='', flush=True)
+            print("Loading... ", end="", flush=True)
         plugin = tf.load_op_library(bin_file)
 
         # Add to cache.
         _plugin_cache[cuda_file] = plugin
         if verbose:
-            print('Done.', flush=True)
+            print("Done.", flush=True)
         return plugin
 
     except:
         if verbose:
-            print('Failed!', flush=True)
+            print("Failed!", flush=True)
         raise
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------

extensions/nvdiffrast/nvdiffrast/tensorflow/tf_antialias.cu

@@ -100,13 +100,13 @@ struct AntialiasFwdOp : public OpKernel
 
         // (Re-)calculate opposite vertex hash.
         if (!p.evHash || !p.tri_const)
-        {            
+        {
             if (p.allocTriangles < p.numTriangles)
             {
                 p.allocTriangles = max(p.allocTriangles, 64);
                 while (p.allocTriangles < p.numTriangles)
                     p.allocTriangles <<= 1; // Must be power of two.
-               
+
                 // (Re-)allocate memory for the hash.
                 OP_CHECK_CUDA_ERROR(ctx, cudaFree(p.evHash));
                 OP_CHECK_CUDA_ERROR(ctx, cudaMalloc(&p.evHash, p.allocTriangles * AA_HASH_ELEMENTS_PER_TRIANGLE(p.allocTriangles) * sizeof(uint4)));

extensions/nvdiffrast/nvdiffrast/tensorflow/tf_interpolate.cu

@@ -112,7 +112,7 @@ struct InterpolateFwdOp : public OpKernel
 
         // Verify that buffers are aligned to allow float2/float4 operations.
         OP_REQUIRES(ctx, !((uintptr_t)p.rast   & 15), errors::Internal("rast input tensor not aligned to float4"));
-        OP_REQUIRES(ctx, !((uintptr_t)p.rastDB & 15), errors::Internal("rast_db input tensor not aligned to float4"));        
+        OP_REQUIRES(ctx, !((uintptr_t)p.rastDB & 15), errors::Internal("rast_db input tensor not aligned to float4"));
         if (ENABLE_DA)
             OP_REQUIRES(ctx, !((uintptr_t)p.outDA & 7), errors::Internal("out_da output tensor not aligned to float2"));
 
@@ -158,7 +158,7 @@ struct InterpolateGradOp : public OpKernel
     InterpolateGradOp(OpKernelConstruction* ctx): OpKernel(ctx)
     {
         memset(&m_attribs, 0, sizeof(m_attribs));
-        interpolateParseOpAttributes(ctx, m_attribs, ENABLE_DA);      
+        interpolateParseOpAttributes(ctx, m_attribs, ENABLE_DA);
     }
 
     void Compute(OpKernelContext* ctx)
@@ -247,7 +247,7 @@ struct InterpolateGradOp : public OpKernel
             OP_REQUIRES_OK(ctx, ctx->allocate_output(2, grad_rast_shape, &grad_rast_db_tensor));
             p.gradRasterDB = grad_rast_db_tensor->flat<float>().data();
         }
-        
+
         // Clear attribute gradients.
         cudaMemsetAsync(p.gradAttr, 0, attr_depth * p.numVertices * p.numAttr * sizeof(float), stream);
 
@@ -257,10 +257,10 @@ struct InterpolateGradOp : public OpKernel
         if (ENABLE_DA)
         {
             OP_REQUIRES(ctx, !((uintptr_t)p.dda & 7), errors::Internal("dda input tensor not aligned to float2"));
-            OP_REQUIRES(ctx, !((uintptr_t)p.rastDB & 15), errors::Internal("rast_db input tensor not aligned to float4"));        
+            OP_REQUIRES(ctx, !((uintptr_t)p.rastDB & 15), errors::Internal("rast_db input tensor not aligned to float4"));
             OP_REQUIRES(ctx, !((uintptr_t)p.gradRasterDB & 15), errors::Internal("grad_rast_db output tensor not aligned to float4"));
         }
-    
+
         // Choose launch parameters.
         dim3 blockSize = getLaunchBlockSize(IP_GRAD_MAX_KERNEL_BLOCK_WIDTH, IP_GRAD_MAX_KERNEL_BLOCK_HEIGHT, p.width, p.height);
         dim3 gridSize  = getLaunchGridSize(blockSize, p.width, p.height, p.depth);

extensions/nvdiffrast/nvdiffrast/tensorflow/tf_texture.cu

@@ -503,7 +503,7 @@ REGISTER_OP("TextureGradLinearMipmapNearest")
     .Attr       ("filter_mode: int")
     .Attr       ("boundary_mode: int")
     .Attr       ("max_mip_level: int");
-    
+
 REGISTER_OP("TextureGradLinearMipmapLinear")
     .Input      ("tex: float")
     .Input      ("uv: float")
@@ -516,10 +516,10 @@ REGISTER_OP("TextureGradLinearMipmapLinear")
     .Attr       ("filter_mode: int")
     .Attr       ("boundary_mode: int")
     .Attr       ("max_mip_level: int");
-    
+
 REGISTER_KERNEL_BUILDER(Name("TextureGradNearest")            .Device(DEVICE_GPU), TextureGradOp);
 REGISTER_KERNEL_BUILDER(Name("TextureGradLinear")             .Device(DEVICE_GPU), TextureGradOp);
 REGISTER_KERNEL_BUILDER(Name("TextureGradLinearMipmapNearest").Device(DEVICE_GPU), TextureGradOp);
 REGISTER_KERNEL_BUILDER(Name("TextureGradLinearMipmapLinear") .Device(DEVICE_GPU), TextureGradOp);
-        
+
 //------------------------------------------------------------------------

extensions/nvdiffrast/nvdiffrast/torch/__init__.py

@@ -6,5 +6,30 @@
 # distribution of this software and related documentation without an express
 # license agreement from NVIDIA CORPORATION is strictly prohibited.
 
-from .ops import RasterizeCudaContext, RasterizeGLContext, get_log_level, set_log_level, rasterize, DepthPeeler, interpolate, texture, texture_construct_mip, antialias, antialias_construct_topology_hash
-__all__ = ["RasterizeCudaContext", "RasterizeGLContext", "get_log_level", "set_log_level", "rasterize", "DepthPeeler", "interpolate", "texture", "texture_construct_mip", "antialias", "antialias_construct_topology_hash"]
+from .ops import (
+    RasterizeCudaContext,
+    RasterizeGLContext,
+    get_log_level,
+    set_log_level,
+    rasterize,
+    DepthPeeler,
+    interpolate,
+    texture,
+    texture_construct_mip,
+    antialias,
+    antialias_construct_topology_hash,
+)
+
+__all__ = [
+    "RasterizeCudaContext",
+    "RasterizeGLContext",
+    "get_log_level",
+    "set_log_level",
+    "rasterize",
+    "DepthPeeler",
+    "interpolate",
+    "texture",
+    "texture_construct_mip",
+    "antialias",
+    "antialias_construct_topology_hash",
+]

extensions/nvdiffrast/nvdiffrast/torch/ops.py

@@ -14,13 +14,15 @@ import torch
 import torch.utils.cpp_extension
 from . import _C
 
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 # C++/Cuda plugin compiler/loader.
 
 _cached_plugin = {}
+
+
 def _get_plugin(gl=False):
     assert isinstance(gl, bool)
-    
+
     # Modified with precompiled torch CUDA extension
     if not gl:
         return _C
@@ -30,16 +32,27 @@ def _get_plugin(gl=False):
         return _cached_plugin[gl]
 
     # Make sure we can find the necessary compiler and libary binaries.
-    if os.name == 'nt':
+    if os.name == "nt":
         lib_dir = os.path.dirname(__file__) + r"\..\lib"
+
         def find_cl_path():
             import glob
+
             def get_sort_key(x):
                 # Primary criterion is VS version, secondary is edition, third is internal MSVC version.
-                x = x.split('\\')[3:]
-                x[1] = {'BuildTools': '~0', 'Community': '~1', 'Pro': '~2', 'Professional': '~3', 'Enterprise': '~4'}.get(x[1], x[1])
+                x = x.split("\\")[3:]
+                x[1] = {
+                    "BuildTools": "~0",
+                    "Community": "~1",
+                    "Pro": "~2",
+                    "Professional": "~3",
+                    "Enterprise": "~4",
+                }.get(x[1], x[1])
                 return x
-            vs_relative_path = r"\Microsoft Visual Studio\*\*\VC\Tools\MSVC\*\bin\Hostx64\x64"
+
+            vs_relative_path = (
+                r"\Microsoft Visual Studio\*\*\VC\Tools\MSVC\*\bin\Hostx64\x64"
+            )
             paths = glob.glob(r"C:\Program Files" + vs_relative_path)
             paths += glob.glob(r"C:\Program Files (x86)" + vs_relative_path)
             if paths:
@@ -49,104 +62,126 @@ def _get_plugin(gl=False):
         if os.system("where cl.exe >nul 2>nul") != 0:
             cl_path = find_cl_path()
             if cl_path is None:
-                raise RuntimeError("Could not locate a supported Microsoft Visual C++ installation")
-            os.environ['PATH'] += ';' + cl_path
+                raise RuntimeError(
+                    "Could not locate a supported Microsoft Visual C++ installation"
+                )
+            os.environ["PATH"] += ";" + cl_path
 
     # Compiler options.
-    common_opts = ['-DNVDR_TORCH']
+    common_opts = ["-DNVDR_TORCH"]
     cc_opts = []
-    if os.name == 'nt':
-        cc_opts += ['/wd4067', '/wd4624'] # Disable warnings in torch headers.
+    if os.name == "nt":
+        cc_opts += ["/wd4067", "/wd4624"]  # Disable warnings in torch headers.
 
     # Linker options for the GL-interfacing plugin.
     ldflags = []
     if gl:
-        if os.name == 'posix':
-            ldflags = ['-lGL', '-lEGL']
-        elif os.name == 'nt':
-            libs = ['gdi32', 'opengl32', 'user32', 'setgpu']
-            ldflags = ['/LIBPATH:' + lib_dir] + ['/DEFAULTLIB:' + x for x in libs]
+        if os.name == "posix":
+            ldflags = ["-lGL", "-lEGL"]
+        elif os.name == "nt":
+            libs = ["gdi32", "opengl32", "user32", "setgpu"]
+            ldflags = ["/LIBPATH:" + lib_dir] + ["/DEFAULTLIB:" + x for x in libs]
 
     # List of source files.
     if gl:
         source_files = [
-            '../common/common.cpp',
-            '../common/glutil.cpp',
-            '../common/rasterize_gl.cpp',
-            'torch_bindings_gl.cpp',
-            'torch_rasterize_gl.cpp',
+            "../common/common.cpp",
+            "../common/glutil.cpp",
+            "../common/rasterize_gl.cpp",
+            "torch_bindings_gl.cpp",
+            "torch_rasterize_gl.cpp",
         ]
     else:
         source_files = [
-            '../common/cudaraster/impl/Buffer.cpp',
-            '../common/cudaraster/impl/CudaRaster.cpp',
-            '../common/cudaraster/impl/RasterImpl.cu',
-            '../common/cudaraster/impl/RasterImpl.cpp',
-            '../common/common.cpp',
-            '../common/rasterize.cu',
-            '../common/interpolate.cu',
-            '../common/texture.cu',
-            '../common/texture.cpp',
-            '../common/antialias.cu',
-            'torch_bindings.cpp',
-            'torch_rasterize.cpp',
-            'torch_interpolate.cpp',
-            'torch_texture.cpp',
-            'torch_antialias.cpp',
+            "../common/cudaraster/impl/Buffer.cpp",
+            "../common/cudaraster/impl/CudaRaster.cpp",
+            "../common/cudaraster/impl/RasterImpl.cu",
+            "../common/cudaraster/impl/RasterImpl.cpp",
+            "../common/common.cpp",
+            "../common/rasterize.cu",
+            "../common/interpolate.cu",
+            "../common/texture.cu",
+            "../common/texture.cpp",
+            "../common/antialias.cu",
+            "torch_bindings.cpp",
+            "torch_rasterize.cpp",
+            "torch_interpolate.cpp",
+            "torch_texture.cpp",
+            "torch_antialias.cpp",
         ]
 
     # Some containers set this to contain old architectures that won't compile. We only need the one installed in the machine.
-    os.environ['TORCH_CUDA_ARCH_LIST'] = ''
+    os.environ["TORCH_CUDA_ARCH_LIST"] = ""
 
     # On Linux, show a warning if GLEW is being forcibly loaded when compiling the GL plugin.
-    if gl and (os.name == 'posix') and ('libGLEW' in os.environ.get('LD_PRELOAD', '')):
-        logging.getLogger('nvdiffrast').warning("Warning: libGLEW is being loaded via LD_PRELOAD, and will probably conflict with the OpenGL plugin")
+    if gl and (os.name == "posix") and ("libGLEW" in os.environ.get("LD_PRELOAD", "")):
+        logging.getLogger("nvdiffrast").warning(
+            "Warning: libGLEW is being loaded via LD_PRELOAD, and will probably conflict with the OpenGL plugin"
+        )
 
     # Try to detect if a stray lock file is left in cache directory and show a warning. This sometimes happens on Windows if the build is interrupted at just the right moment.
-    plugin_name = 'nvdiffrast_plugin' + ('_gl' if gl else '')
+    plugin_name = "nvdiffrast_plugin" + ("_gl" if gl else "")
     try:
-        lock_fn = os.path.join(torch.utils.cpp_extension._get_build_directory(plugin_name, False), 'lock')
+        lock_fn = os.path.join(
+            torch.utils.cpp_extension._get_build_directory(plugin_name, False), "lock"
+        )
         if os.path.exists(lock_fn):
-            logging.getLogger('nvdiffrast').warning("Lock file exists in build directory: '%s'" % lock_fn)
+            logging.getLogger("nvdiffrast").warning(
+                "Lock file exists in build directory: '%s'" % lock_fn
+            )
     except:
         pass
 
     # Speed up compilation on Windows.
-    if os.name == 'nt':
+    if os.name == "nt":
         # Skip telemetry sending step in vcvarsall.bat
-        os.environ['VSCMD_SKIP_SENDTELEMETRY'] = '1'
+        os.environ["VSCMD_SKIP_SENDTELEMETRY"] = "1"
 
         # Opportunistically patch distutils to cache MSVC environments.
         try:
             import distutils._msvccompiler
             import functools
-            if not hasattr(distutils._msvccompiler._get_vc_env, '__wrapped__'):
-                distutils._msvccompiler._get_vc_env = functools.lru_cache()(distutils._msvccompiler._get_vc_env)
+
+            if not hasattr(distutils._msvccompiler._get_vc_env, "__wrapped__"):
+                distutils._msvccompiler._get_vc_env = functools.lru_cache()(
+                    distutils._msvccompiler._get_vc_env
+                )
         except:
             pass
 
     # Compile and load.
     source_paths = [os.path.join(os.path.dirname(__file__), fn) for fn in source_files]
-    torch.utils.cpp_extension.load(name=plugin_name, sources=source_paths, extra_cflags=common_opts+cc_opts, extra_cuda_cflags=common_opts+['-lineinfo'], extra_ldflags=ldflags, with_cuda=True, verbose=False)
+    torch.utils.cpp_extension.load(
+        name=plugin_name,
+        sources=source_paths,
+        extra_cflags=common_opts + cc_opts,
+        extra_cuda_cflags=common_opts + ["-lineinfo"],
+        extra_ldflags=ldflags,
+        with_cuda=True,
+        verbose=False,
+    )
 
     # Import, cache, and return the compiled module.
     _cached_plugin[gl] = importlib.import_module(plugin_name)
     return _cached_plugin[gl]
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Log level.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 def get_log_level():
-    '''Get current log level.
+    """Get current log level.
 
     Returns:
       Current log level in nvdiffrast. See `set_log_level()` for possible values.
-    '''
+    """
     return _get_plugin().get_log_level()
 
+
 def set_log_level(level):
-    '''Set log level.
+    """Set log level.
 
     Log levels follow the convention on the C++ side of Torch:
       0 = Info,
@@ -156,19 +191,21 @@ def set_log_level(level):
     The default log level is 1.
 
     Args:
-      level: New log level as integer. Internal nvdiffrast messages of this 
+      level: New log level as integer. Internal nvdiffrast messages of this
              severity or higher will be printed, while messages of lower
              severity will be silent.
-    '''
+    """
     _get_plugin().set_log_level(level)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # CudaRaster state wrapper.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 class RasterizeCudaContext:
     def __init__(self, device=None):
-        '''Create a new Cuda rasterizer context.
+        """Create a new Cuda rasterizer context.
 
         The context is deleted and internal storage is released when the object is
         destroyed.
@@ -180,7 +217,7 @@ class RasterizeCudaContext:
                              device.
         Returns:
           The newly created Cuda rasterizer context.
-        '''
+        """
         if device is None:
             cuda_device_idx = torch.cuda.current_device()
         else:
@@ -190,13 +227,15 @@ class RasterizeCudaContext:
         self.output_db = True
         self.active_depth_peeler = None
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # GL state wrapper.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 class RasterizeGLContext:
-    def __init__(self, output_db=True, mode='automatic', device=None):
-        '''Create a new OpenGL rasterizer context.
+    def __init__(self, output_db=True, mode="automatic", device=None):
+        """Create a new OpenGL rasterizer context.
 
         Creating an OpenGL context is a slow operation so you should usually reuse the same
         context in all calls to `rasterize()` on the same CPU thread. The OpenGL context
@@ -220,9 +259,9 @@ class RasterizeGLContext:
                              device.
         Returns:
           The newly created OpenGL rasterizer context.
-        '''
+        """
         assert output_db is True or output_db is False
-        assert mode in ['automatic', 'manual']
+        assert mode in ["automatic", "manual"]
         self.output_db = output_db
         self.mode = mode
         if device is None:
@@ -230,34 +269,42 @@ class RasterizeGLContext:
         else:
             with torch.cuda.device(device):
                 cuda_device_idx = torch.cuda.current_device()
-        self.cpp_wrapper = _get_plugin(gl=True).RasterizeGLStateWrapper(output_db, mode == 'automatic', cuda_device_idx)
-        self.active_depth_peeler = None # For error checking only.
+        self.cpp_wrapper = _get_plugin(gl=True).RasterizeGLStateWrapper(
+            output_db, mode == "automatic", cuda_device_idx
+        )
+        self.active_depth_peeler = None  # For error checking only.
 
     def set_context(self):
-        '''Set (activate) OpenGL context in the current CPU thread.
-           Only available if context was created in manual mode.
-        '''
-        assert self.mode == 'manual'
+        """Set (activate) OpenGL context in the current CPU thread.
+        Only available if context was created in manual mode.
+        """
+        assert self.mode == "manual"
         self.cpp_wrapper.set_context()
 
     def release_context(self):
-        '''Release (deactivate) currently active OpenGL context.
-           Only available if context was created in manual mode.
-        '''
-        assert self.mode == 'manual'
+        """Release (deactivate) currently active OpenGL context.
+        Only available if context was created in manual mode.
+        """
+        assert self.mode == "manual"
         self.cpp_wrapper.release_context()
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Rasterize.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 class _rasterize_func(torch.autograd.Function):
     @staticmethod
     def forward(ctx, raster_ctx, pos, tri, resolution, ranges, grad_db, peeling_idx):
         if isinstance(raster_ctx, RasterizeGLContext):
-            out, out_db = _get_plugin(gl=True).rasterize_fwd_gl(raster_ctx.cpp_wrapper, pos, tri, resolution, ranges, peeling_idx)
+            out, out_db = _get_plugin(gl=True).rasterize_fwd_gl(
+                raster_ctx.cpp_wrapper, pos, tri, resolution, ranges, peeling_idx
+            )
         else:
-            out, out_db = _get_plugin().rasterize_fwd_cuda(raster_ctx.cpp_wrapper, pos, tri, resolution, ranges, peeling_idx)
+            out, out_db = _get_plugin().rasterize_fwd_cuda(
+                raster_ctx.cpp_wrapper, pos, tri, resolution, ranges, peeling_idx
+            )
         ctx.save_for_backward(pos, tri, out)
         ctx.saved_grad_db = grad_db
         return out, out_db
@@ -271,9 +318,10 @@ class _rasterize_func(torch.autograd.Function):
             g_pos = _get_plugin().rasterize_grad(pos, tri, out, dy)
         return None, g_pos, None, None, None, None, None
 
+
 # Op wrapper.
 def rasterize(glctx, pos, tri, resolution, ranges=None, grad_db=True):
-    '''Rasterize triangles.
+    """Rasterize triangles.
 
     All input tensors must be contiguous and reside in GPU memory except for
     the `ranges` tensor that, if specified, has to reside in CPU memory. The
@@ -301,7 +349,7 @@ def rasterize(glctx, pos, tri, resolution, ranges=None, grad_db=True):
         [minibatch_size, height, width, 4] and contain said derivatives in order
         (du/dX, du/dY, dv/dX, dv/dY). Otherwise it will be an empty tensor with shape
         [minibatch_size, height, width, 0].
-    '''
+    """
     assert isinstance(glctx, (RasterizeGLContext, RasterizeCudaContext))
     assert grad_db is True or grad_db is False
     grad_db = grad_db and glctx.output_db
@@ -310,30 +358,34 @@ def rasterize(glctx, pos, tri, resolution, ranges=None, grad_db=True):
     assert isinstance(pos, torch.Tensor) and isinstance(tri, torch.Tensor)
     resolution = tuple(resolution)
     if ranges is None:
-        ranges = torch.empty(size=(0, 2), dtype=torch.int32, device='cpu')
+        ranges = torch.empty(size=(0, 2), dtype=torch.int32, device="cpu")
     else:
         assert isinstance(ranges, torch.Tensor)
 
     # Check that context is not currently reserved for depth peeling.
     if glctx.active_depth_peeler is not None:
-        return RuntimeError("Cannot call rasterize() during depth peeling operation, use rasterize_next_layer() instead")
+        return RuntimeError(
+            "Cannot call rasterize() during depth peeling operation, use rasterize_next_layer() instead"
+        )
 
     # Instantiate the function.
     return _rasterize_func.apply(glctx, pos, tri, resolution, ranges, grad_db, -1)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Depth peeler context manager for rasterizing multiple depth layers.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 class DepthPeeler:
     def __init__(self, glctx, pos, tri, resolution, ranges=None, grad_db=True):
-        '''Create a depth peeler object for rasterizing multiple depth layers.
+        """Create a depth peeler object for rasterizing multiple depth layers.
 
         Arguments are the same as in `rasterize()`.
 
         Returns:
           The newly created depth peeler.
-        '''
+        """
         assert isinstance(glctx, (RasterizeGLContext, RasterizeCudaContext))
         assert grad_db is True or grad_db is False
         grad_db = grad_db and glctx.output_db
@@ -342,7 +394,7 @@ class DepthPeeler:
         assert isinstance(pos, torch.Tensor) and isinstance(tri, torch.Tensor)
         resolution = tuple(resolution)
         if ranges is None:
-            ranges = torch.empty(size=(0, 2), dtype=torch.int32, device='cpu')
+            ranges = torch.empty(size=(0, 2), dtype=torch.int32, device="cpu")
         else:
             assert isinstance(ranges, torch.Tensor)
 
@@ -359,7 +411,9 @@ class DepthPeeler:
         if self.raster_ctx is None:
             raise RuntimeError("Cannot re-enter a terminated depth peeling operation")
         if self.raster_ctx.active_depth_peeler is not None:
-            raise RuntimeError("Cannot have multiple depth peelers active simultaneously in a rasterization context")
+            raise RuntimeError(
+                "Cannot have multiple depth peelers active simultaneously in a rasterization context"
+            )
         self.raster_ctx.active_depth_peeler = self
         self.peeling_idx = 0
         return self
@@ -367,7 +421,9 @@ class DepthPeeler:
     def __exit__(self, *args):
         assert self.raster_ctx.active_depth_peeler is self
         self.raster_ctx.active_depth_peeler = None
-        self.raster_ctx = None # Remove all references to input tensor so they're not left dangling.
+        self.raster_ctx = (
+            None  # Remove all references to input tensor so they're not left dangling.
+        )
         self.pos = None
         self.tri = None
         self.resolution = None
@@ -377,29 +433,40 @@ class DepthPeeler:
         return None
 
     def rasterize_next_layer(self):
-        '''Rasterize next depth layer.
+        """Rasterize next depth layer.
 
         Operation is equivalent to `rasterize()` except that previously reported
         surface points are culled away.
 
         Returns:
           A tuple of two tensors as in `rasterize()`.
-        '''
+        """
         assert self.raster_ctx.active_depth_peeler is self
         assert self.peeling_idx >= 0
-        result = _rasterize_func.apply(self.raster_ctx, self.pos, self.tri, self.resolution, self.ranges, self.grad_db, self.peeling_idx)
+        result = _rasterize_func.apply(
+            self.raster_ctx,
+            self.pos,
+            self.tri,
+            self.resolution,
+            self.ranges,
+            self.grad_db,
+            self.peeling_idx,
+        )
         self.peeling_idx += 1
         return result
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Interpolate.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 
 # Output pixel differentials for at least some attributes.
 class _interpolate_func_da(torch.autograd.Function):
     @staticmethod
     def forward(ctx, attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list):
-        out, out_da = _get_plugin().interpolate_fwd_da(attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list)
+        out, out_da = _get_plugin().interpolate_fwd_da(
+            attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list
+        )
         ctx.save_for_backward(attr, rast, tri, rast_db)
         ctx.saved_misc = diff_attrs_all, diff_attrs_list
         return out, out_da
@@ -408,9 +475,12 @@ class _interpolate_func_da(torch.autograd.Function):
     def backward(ctx, dy, dda):
         attr, rast, tri, rast_db = ctx.saved_tensors
         diff_attrs_all, diff_attrs_list = ctx.saved_misc
-        g_attr, g_rast, g_rast_db = _get_plugin().interpolate_grad_da(attr, rast, tri, dy, rast_db, dda, diff_attrs_all, diff_attrs_list)
+        g_attr, g_rast, g_rast_db = _get_plugin().interpolate_grad_da(
+            attr, rast, tri, dy, rast_db, dda, diff_attrs_all, diff_attrs_list
+        )
         return g_attr, g_rast, None, g_rast_db, None, None
 
+
 # No pixel differential for any attribute.
 class _interpolate_func(torch.autograd.Function):
     @staticmethod
@@ -425,6 +495,7 @@ class _interpolate_func(torch.autograd.Function):
         g_attr, g_rast = _get_plugin().interpolate_grad(attr, rast, tri, dy)
         return g_attr, g_rast, None
 
+
 # Op wrapper.
 def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     """Interpolate vertex attributes.
@@ -433,13 +504,13 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     will be contiguous and reside in GPU memory.
 
     Args:
-        attr: Attribute tensor with dtype `torch.float32`. 
-              Shape is [num_vertices, num_attributes] in range mode, or 
+        attr: Attribute tensor with dtype `torch.float32`.
+              Shape is [num_vertices, num_attributes] in range mode, or
               [minibatch_size, num_vertices, num_attributes] in instanced mode.
               Broadcasting is supported along the minibatch axis.
         rast: Main output tensor from `rasterize()`.
         tri: Triangle tensor with shape [num_triangles, 3] and dtype `torch.int32`.
-        rast_db: (Optional) Tensor containing image-space derivatives of barycentrics, 
+        rast_db: (Optional) Tensor containing image-space derivatives of barycentrics,
                  i.e., the second output tensor from `rasterize()`. Enables computing
                  image-space derivatives of attributes.
         diff_attrs: (Optional) List of attribute indices for which image-space
@@ -459,12 +530,12 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
     # Sanitize the list of pixel differential attributes.
     if diff_attrs is None:
         diff_attrs = []
-    elif diff_attrs != 'all':
+    elif diff_attrs != "all":
         diff_attrs = np.asarray(diff_attrs, np.int32)
         assert len(diff_attrs.shape) == 1
         diff_attrs = diff_attrs.tolist()
 
-    diff_attrs_all = int(diff_attrs == 'all')
+    diff_attrs_all = int(diff_attrs == "all")
     diff_attrs_list = [] if diff_attrs_all else diff_attrs
 
     # Check inputs.
@@ -474,18 +545,32 @@ def interpolate(attr, rast, tri, rast_db=None, diff_attrs=None):
 
     # Choose stub.
     if diff_attrs:
-        return _interpolate_func_da.apply(attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list)
+        return _interpolate_func_da.apply(
+            attr, rast, tri, rast_db, diff_attrs_all, diff_attrs_list
+        )
     else:
         return _interpolate_func.apply(attr, rast, tri)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Texture
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
 
 # Linear-mipmap-linear and linear-mipmap-nearest: Mipmaps enabled.
 class _texture_func_mip(torch.autograd.Function):
     @staticmethod
-    def forward(ctx, filter_mode, tex, uv, uv_da, mip_level_bias, mip_wrapper, filter_mode_enum, boundary_mode_enum, *mip_stack):
+    def forward(
+        ctx,
+        filter_mode,
+        tex,
+        uv,
+        uv_da,
+        mip_level_bias,
+        mip_wrapper,
+        filter_mode_enum,
+        boundary_mode_enum,
+        *mip_stack
+    ):
         empty = torch.tensor([])
         if uv_da is None:
             uv_da = empty
@@ -493,7 +578,16 @@ class _texture_func_mip(torch.autograd.Function):
             mip_level_bias = empty
         if mip_wrapper is None:
             mip_wrapper = _get_plugin().TextureMipWrapper()
-        out = _get_plugin().texture_fwd_mip(tex, uv, uv_da, mip_level_bias, mip_wrapper, mip_stack, filter_mode_enum, boundary_mode_enum)
+        out = _get_plugin().texture_fwd_mip(
+            tex,
+            uv,
+            uv_da,
+            mip_level_bias,
+            mip_wrapper,
+            mip_stack,
+            filter_mode_enum,
+            boundary_mode_enum,
+        )
         ctx.save_for_backward(tex, uv, uv_da, mip_level_bias, *mip_stack)
         ctx.saved_misc = filter_mode, mip_wrapper, filter_mode_enum, boundary_mode_enum
         return out
@@ -502,12 +596,50 @@ class _texture_func_mip(torch.autograd.Function):
     def backward(ctx, dy):
         tex, uv, uv_da, mip_level_bias, *mip_stack = ctx.saved_tensors
         filter_mode, mip_wrapper, filter_mode_enum, boundary_mode_enum = ctx.saved_misc
-        if filter_mode == 'linear-mipmap-linear':
-            g_tex, g_uv, g_uv_da, g_mip_level_bias, g_mip_stack = _get_plugin().texture_grad_linear_mipmap_linear(tex, uv, dy, uv_da, mip_level_bias, mip_wrapper, mip_stack, filter_mode_enum, boundary_mode_enum)
-            return (None, g_tex, g_uv, g_uv_da, g_mip_level_bias, None, None, None) + tuple(g_mip_stack)
-        else: # linear-mipmap-nearest
-            g_tex, g_uv, g_mip_stack = _get_plugin().texture_grad_linear_mipmap_nearest(tex, uv, dy, uv_da, mip_level_bias, mip_wrapper, mip_stack, filter_mode_enum, boundary_mode_enum)
-            return (None, g_tex, g_uv, None, None, None, None, None) + tuple(g_mip_stack)
+        if filter_mode == "linear-mipmap-linear":
+            (
+                g_tex,
+                g_uv,
+                g_uv_da,
+                g_mip_level_bias,
+                g_mip_stack,
+            ) = _get_plugin().texture_grad_linear_mipmap_linear(
+                tex,
+                uv,
+                dy,
+                uv_da,
+                mip_level_bias,
+                mip_wrapper,
+                mip_stack,
+                filter_mode_enum,
+                boundary_mode_enum,
+            )
+            return (
+                None,
+                g_tex,
+                g_uv,
+                g_uv_da,
+                g_mip_level_bias,
+                None,
+                None,
+                None,
+            ) + tuple(g_mip_stack)
+        else:  # linear-mipmap-nearest
+            g_tex, g_uv, g_mip_stack = _get_plugin().texture_grad_linear_mipmap_nearest(
+                tex,
+                uv,
+                dy,
+                uv_da,
+                mip_level_bias,
+                mip_wrapper,
+                mip_stack,
+                filter_mode_enum,
+                boundary_mode_enum,
+            )
+            return (None, g_tex, g_uv, None, None, None, None, None) + tuple(
+                g_mip_stack
+            )
+
 
 # Linear and nearest: Mipmaps disabled.
 class _texture_func(torch.autograd.Function):
@@ -522,15 +654,29 @@ class _texture_func(torch.autograd.Function):
     def backward(ctx, dy):
         tex, uv = ctx.saved_tensors
         filter_mode, filter_mode_enum, boundary_mode_enum = ctx.saved_misc
-        if filter_mode == 'linear':
-            g_tex, g_uv = _get_plugin().texture_grad_linear(tex, uv, dy, filter_mode_enum, boundary_mode_enum)
+        if filter_mode == "linear":
+            g_tex, g_uv = _get_plugin().texture_grad_linear(
+                tex, uv, dy, filter_mode_enum, boundary_mode_enum
+            )
             return None, g_tex, g_uv, None, None
-        else: # nearest
-            g_tex = _get_plugin().texture_grad_nearest(tex, uv, dy, filter_mode_enum, boundary_mode_enum)
+        else:  # nearest
+            g_tex = _get_plugin().texture_grad_nearest(
+                tex, uv, dy, filter_mode_enum, boundary_mode_enum
+            )
             return None, g_tex, None, None, None
 
+
 # Op wrapper.
-def texture(tex, uv, uv_da=None, mip_level_bias=None, mip=None, filter_mode='auto', boundary_mode='wrap', max_mip_level=None):
+def texture(
+    tex,
+    uv,
+    uv_da=None,
+    mip_level_bias=None,
+    mip=None,
+    filter_mode="auto",
+    boundary_mode="wrap",
+    max_mip_level=None,
+):
     """Perform texture sampling.
 
     All input tensors must be contiguous and reside in GPU memory. The output tensor
@@ -580,8 +726,12 @@ def texture(tex, uv, uv_da=None, mip_level_bias=None, mip=None, filter_mode='aut
     """
 
     # Default filter mode.
-    if filter_mode == 'auto':
-        filter_mode = 'linear-mipmap-linear' if (uv_da is not None or mip_level_bias is not None) else 'linear'
+    if filter_mode == "auto":
+        filter_mode = (
+            "linear-mipmap-linear"
+            if (uv_da is not None or mip_level_bias is not None)
+            else "linear"
+        )
 
     # Sanitize inputs.
     if max_mip_level is None:
@@ -592,23 +742,33 @@ def texture(tex, uv, uv_da=None, mip_level_bias=None, mip=None, filter_mode='aut
 
     # Check inputs.
     assert isinstance(tex, torch.Tensor) and isinstance(uv, torch.Tensor)
-    if 'mipmap' in filter_mode:
-        assert isinstance(uv_da, torch.Tensor) or isinstance(mip_level_bias, torch.Tensor)
+    if "mipmap" in filter_mode:
+        assert isinstance(uv_da, torch.Tensor) or isinstance(
+            mip_level_bias, torch.Tensor
+        )
 
     # If mipping disabled via max level=0, we may as well use simpler filtering internally.
-    if max_mip_level == 0 and filter_mode in ['linear-mipmap-nearest', 'linear-mipmap-linear']:
-        filter_mode = 'linear'
+    if max_mip_level == 0 and filter_mode in [
+        "linear-mipmap-nearest",
+        "linear-mipmap-linear",
+    ]:
+        filter_mode = "linear"
 
     # Convert filter mode to internal enumeration.
-    filter_mode_dict = {'nearest': 0, 'linear': 1, 'linear-mipmap-nearest': 2, 'linear-mipmap-linear': 3}
+    filter_mode_dict = {
+        "nearest": 0,
+        "linear": 1,
+        "linear-mipmap-nearest": 2,
+        "linear-mipmap-linear": 3,
+    }
     filter_mode_enum = filter_mode_dict[filter_mode]
 
     # Convert boundary mode to internal enumeration.
-    boundary_mode_dict = {'cube': 0, 'wrap': 1, 'clamp': 2, 'zero': 3}
+    boundary_mode_dict = {"cube": 0, "wrap": 1, "clamp": 2, "zero": 3}
     boundary_mode_enum = boundary_mode_dict[boundary_mode]
 
     # Construct a mipmap if necessary.
-    if 'mipmap' in filter_mode:
+    if "mipmap" in filter_mode:
         mip_wrapper, mip_stack = None, []
         if mip is not None:
             assert isinstance(mip, (_get_plugin().TextureMipWrapper, list))
@@ -618,13 +778,28 @@ def texture(tex, uv, uv_da=None, mip_level_bias=None, mip=None, filter_mode='aut
             else:
                 mip_wrapper = mip
         else:
-            mip_wrapper = _get_plugin().texture_construct_mip(tex, max_mip_level, boundary_mode == 'cube')
+            mip_wrapper = _get_plugin().texture_construct_mip(
+                tex, max_mip_level, boundary_mode == "cube"
+            )
 
     # Choose stub.
-    if filter_mode == 'linear-mipmap-linear' or filter_mode == 'linear-mipmap-nearest':
-        return _texture_func_mip.apply(filter_mode, tex, uv, uv_da, mip_level_bias, mip_wrapper, filter_mode_enum, boundary_mode_enum, *mip_stack)
+    if filter_mode == "linear-mipmap-linear" or filter_mode == "linear-mipmap-nearest":
+        return _texture_func_mip.apply(
+            filter_mode,
+            tex,
+            uv,
+            uv_da,
+            mip_level_bias,
+            mip_wrapper,
+            filter_mode_enum,
+            boundary_mode_enum,
+            *mip_stack
+        )
     else:
-        return _texture_func.apply(filter_mode, tex, uv, filter_mode_enum, boundary_mode_enum)
+        return _texture_func.apply(
+            filter_mode, tex, uv, filter_mode_enum, boundary_mode_enum
+        )
+
 
 # Mipmap precalculation for cases where the texture stays constant.
 def texture_construct_mip(tex, max_mip_level=None, cube_mode=False):
@@ -639,7 +814,7 @@ def texture_construct_mip(tex, max_mip_level=None, cube_mode=False):
         cube_mode: Must be set to True if `tex` specifies a cube map texture.
 
     Returns:
-        An opaque object containing the mipmap stack. This can be supplied in a call to `texture()` 
+        An opaque object containing the mipmap stack. This can be supplied in a call to `texture()`
         in the `mip` argument.
     """
 
@@ -652,14 +827,18 @@ def texture_construct_mip(tex, max_mip_level=None, cube_mode=False):
         assert max_mip_level >= 0
     return _get_plugin().texture_construct_mip(tex, max_mip_level, cube_mode)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------
 # Antialias.
-#----------------------------------------------------------------------------
+# ----------------------------------------------------------------------------
+
 
 class _antialias_func(torch.autograd.Function):
     @staticmethod
     def forward(ctx, color, rast, pos, tri, topology_hash, pos_gradient_boost):
-        out, work_buffer = _get_plugin().antialias_fwd(color, rast, pos, tri, topology_hash)
+        out, work_buffer = _get_plugin().antialias_fwd(
+            color, rast, pos, tri, topology_hash
+        )
         ctx.save_for_backward(color, rast, pos, tri)
         ctx.saved_misc = pos_gradient_boost, work_buffer
         return out
@@ -668,11 +847,14 @@ class _antialias_func(torch.autograd.Function):
     def backward(ctx, dy):
         color, rast, pos, tri = ctx.saved_tensors
         pos_gradient_boost, work_buffer = ctx.saved_misc
-        g_color, g_pos = _get_plugin().antialias_grad(color, rast, pos, tri, dy, work_buffer)
+        g_color, g_pos = _get_plugin().antialias_grad(
+            color, rast, pos, tri, dy, work_buffer
+        )
         if pos_gradient_boost != 1.0:
             g_pos = g_pos * pos_gradient_boost
         return g_color, None, g_pos, None, None, None
 
+
 # Op wrapper.
 def antialias(color, rast, pos, tri, topology_hash=None, pos_gradient_boost=1.0):
     """Perform antialiasing.
@@ -711,13 +893,16 @@ def antialias(color, rast, pos, tri, topology_hash=None, pos_gradient_boost=1.0)
         topology_hash = _get_plugin().antialias_construct_topology_hash(tri)
 
     # Instantiate the function.
-    return _antialias_func.apply(color, rast, pos, tri, topology_hash, pos_gradient_boost)
+    return _antialias_func.apply(
+        color, rast, pos, tri, topology_hash, pos_gradient_boost
+    )
+
 
 # Topology hash precalculation for cases where the triangle array stays constant.
 def antialias_construct_topology_hash(tri):
     """Construct a topology hash for a triangle tensor.
 
-    This function can be used for constructing a topology hash for a triangle tensor that is 
+    This function can be used for constructing a topology hash for a triangle tensor that is
     known to remain constant. This avoids reconstructing it every time `antialias()` is called.
 
     Args:
@@ -725,10 +910,11 @@ def antialias_construct_topology_hash(tri):
              GPU memory.
 
     Returns:
-        An opaque object containing the topology hash. This can be supplied in a call to 
+        An opaque object containing the topology hash. This can be supplied in a call to
         `antialias()` in the `topology_hash` argument.
     """
     assert isinstance(tri, torch.Tensor)
     return _get_plugin().antialias_construct_topology_hash(tri)
 
-#----------------------------------------------------------------------------
+
+# ----------------------------------------------------------------------------

extensions/nvdiffrast/setup copy.py

@@ -24,28 +24,31 @@ setuptools.setup(
     url="https://github.com/NVlabs/nvdiffrast",
     packages=setuptools.find_packages(),
     package_data={
-        'nvdiffrast': [
-            'common/*.h',
-            'common/*.inl',
-            'common/*.cu',
-            'common/*.cpp',
-            'common/cudaraster/*.hpp',
-            'common/cudaraster/impl/*.cpp',
-            'common/cudaraster/impl/*.hpp',
-            'common/cudaraster/impl/*.inl',
-            'common/cudaraster/impl/*.cu',
-            'lib/*.h',
-            'torch/*.h',
-            'torch/*.inl',
-            'torch/*.cpp',
-            'tensorflow/*.cu',
-        ] + (['lib/*.lib'] if os.name == 'nt' else [])
+        "nvdiffrast": [
+            "common/*.h",
+            "common/*.inl",
+            "common/*.cu",
+            "common/*.cpp",
+            "common/cudaraster/*.hpp",
+            "common/cudaraster/impl/*.cpp",
+            "common/cudaraster/impl/*.hpp",
+            "common/cudaraster/impl/*.inl",
+            "common/cudaraster/impl/*.cu",
+            "lib/*.h",
+            "torch/*.h",
+            "torch/*.inl",
+            "torch/*.cpp",
+            "tensorflow/*.cu",
+        ]
+        + (["lib/*.lib"] if os.name == "nt" else [])
     },
     include_package_data=True,
-    install_requires=['numpy'],  # note: can't require torch here as it will install torch even for a TensorFlow container
+    install_requires=[
+        "numpy"
+    ],  # note: can't require torch here as it will install torch even for a TensorFlow container
     classifiers=[
         "Programming Language :: Python :: 3",
         "Operating System :: OS Independent",
     ],
-    python_requires='>=3.6',
+    python_requires=">=3.6",
 )

extensions/nvdiffrast/setup.py

@@ -48,35 +48,35 @@ setuptools.setup(
         CUDAExtension(
             name="nvdiffrast.torch._C",
             sources=[
-                'nvdiffrast/common/cudaraster/impl/Buffer.cpp',
-                'nvdiffrast/common/cudaraster/impl/CudaRaster.cpp',
-                'nvdiffrast/common/cudaraster/impl/RasterImpl_.cu',
-                'nvdiffrast/common/cudaraster/impl/RasterImpl.cpp',
-                'nvdiffrast/common/common.cpp',
-                'nvdiffrast/common/rasterize.cu',
-                'nvdiffrast/common/interpolate.cu',
-                'nvdiffrast/common/texture_.cu',
-                'nvdiffrast/common/texture.cpp',
-                'nvdiffrast/common/antialias.cu',
-                'nvdiffrast/torch/torch_bindings.cpp',
-                'nvdiffrast/torch/torch_rasterize.cpp',
-                'nvdiffrast/torch/torch_interpolate.cpp',
-                'nvdiffrast/torch/torch_texture.cpp',
-                'nvdiffrast/torch/torch_antialias.cpp',
+                "nvdiffrast/common/cudaraster/impl/Buffer.cpp",
+                "nvdiffrast/common/cudaraster/impl/CudaRaster.cpp",
+                "nvdiffrast/common/cudaraster/impl/RasterImpl_.cu",
+                "nvdiffrast/common/cudaraster/impl/RasterImpl.cpp",
+                "nvdiffrast/common/common.cpp",
+                "nvdiffrast/common/rasterize.cu",
+                "nvdiffrast/common/interpolate.cu",
+                "nvdiffrast/common/texture_.cu",
+                "nvdiffrast/common/texture.cpp",
+                "nvdiffrast/common/antialias.cu",
+                "nvdiffrast/torch/torch_bindings.cpp",
+                "nvdiffrast/torch/torch_rasterize.cpp",
+                "nvdiffrast/torch/torch_interpolate.cpp",
+                "nvdiffrast/torch/torch_texture.cpp",
+                "nvdiffrast/torch/torch_antialias.cpp",
             ],
             extra_compile_args={
-                'cxx': ['-DNVDR_TORCH'],
-                'nvcc': ['-DNVDR_TORCH', '-lineinfo'],
+                "cxx": ["-DNVDR_TORCH"],
+                "nvcc": ["-DNVDR_TORCH", "-lineinfo"],
             },
         )
     ],
-    cmdclass={
-        'build_ext': BuildExtension
-    },
-    install_requires=['numpy'],  # note: can't require torch here as it will install torch even for a TensorFlow container
+    cmdclass={"build_ext": BuildExtension},
+    install_requires=[
+        "numpy"
+    ],  # note: can't require torch here as it will install torch even for a TensorFlow container
     classifiers=[
         "Programming Language :: Python :: 3",
         "Operating System :: OS Independent",
     ],
-    python_requires='>=3.6',
+    python_requires=">=3.6",
 )

requirements.txt

@@ -26,4 +26,4 @@ https://huggingface.co/spaces/JeffreyXiang/TRELLIS/resolve/main/wheels/diff_gaus
 https://huggingface.co/spaces/JeffreyXiang/TRELLIS/resolve/main/wheels/nvdiffrast-0.3.3-cp310-cp310-linux_x86_64.whl?download=true
 spaces
 plyfile==1.1
-utils3d
+utils3d

trellis/models/__init__.py

@@ -1,20 +1,21 @@
 import importlib
 
 __attributes = {
-    'SparseStructureEncoder': 'sparse_structure_vae',
-    'SparseStructureDecoder': 'sparse_structure_vae',
-    'SparseStructureFlowModel': 'sparse_structure_flow',
-    'SLatEncoder': 'structured_latent_vae',
-    'SLatGaussianDecoder': 'structured_latent_vae',
-    'SLatRadianceFieldDecoder': 'structured_latent_vae',
-    'SLatMeshDecoder': 'structured_latent_vae',
-    'SLatFlowModel': 'structured_latent_flow',
+    "SparseStructureEncoder": "sparse_structure_vae",
+    "SparseStructureDecoder": "sparse_structure_vae",
+    "SparseStructureFlowModel": "sparse_structure_flow",
+    "SLatEncoder": "structured_latent_vae",
+    "SLatGaussianDecoder": "structured_latent_vae",
+    "SLatRadianceFieldDecoder": "structured_latent_vae",
+    "SLatMeshDecoder": "structured_latent_vae",
+    "SLatFlowModel": "structured_latent_flow",
 }
 
 __submodules = []
 
 __all__ = list(__attributes.keys()) + __submodules
 
+
 def __getattr__(name):
     if name not in globals():
         if name in __attributes:
@@ -41,6 +42,7 @@ def from_pretrained(path: str, **kwargs):
     import os
     import json
     from safetensors.torch import load_file
+
     is_local = os.path.exists(f"{path}.json") and os.path.exists(f"{path}.safetensors")
 
     if is_local:
@@ -48,23 +50,29 @@ def from_pretrained(path: str, **kwargs):
         model_file = f"{path}.safetensors"
     else:
         from huggingface_hub import hf_hub_download
-        path_parts = path.split('/')
-        repo_id = f'{path_parts[0]}/{path_parts[1]}'
-        model_name = '/'.join(path_parts[2:])
+
+        path_parts = path.split("/")
+        repo_id = f"{path_parts[0]}/{path_parts[1]}"
+        model_name = "/".join(path_parts[2:])
         config_file = hf_hub_download(repo_id, f"{model_name}.json")
         model_file = hf_hub_download(repo_id, f"{model_name}.safetensors")
 
-    with open(config_file, 'r') as f:
+    with open(config_file, "r") as f:
         config = json.load(f)
-    model = __getattr__(config['name'])(**config['args'], **kwargs)
+    model = __getattr__(config["name"])(**config["args"], **kwargs)
     model.load_state_dict(load_file(model_file))
 
     return model
 
 
 # For Pylance
-if __name__ == '__main__':
+if __name__ == "__main__":
     from .sparse_structure_vae import SparseStructureEncoder, SparseStructureDecoder
     from .sparse_structure_flow import SparseStructureFlowModel
-    from .structured_latent_vae import SLatEncoder, SLatGaussianDecoder, SLatRadianceFieldDecoder, SLatMeshDecoder
+    from .structured_latent_vae import (
+        SLatEncoder,
+        SLatGaussianDecoder,
+        SLatRadianceFieldDecoder,
+        SLatMeshDecoder,
+    )
     from .structured_latent_flow import SLatFlowModel

trellis/models/sparse_structure_flow.py

@@ -4,7 +4,10 @@ import torch.nn as nn
 import torch.nn.functional as F
 import numpy as np
 from ..modules.utils import convert_module_to_f16, convert_module_to_f32
-from ..modules.transformer import AbsolutePositionEmbedder, ModulatedTransformerCrossBlock
+from ..modules.transformer import (
+    AbsolutePositionEmbedder,
+    ModulatedTransformerCrossBlock,
+)
 from ..modules.spatial import patchify, unpatchify
 
 
@@ -12,6 +15,7 @@ class TimestepEmbedder(nn.Module):
     """
     Embeds scalar timesteps into vector representations.
     """
+
     def __init__(self, hidden_size, frequency_embedding_size=256):
         super().__init__()
         self.mlp = nn.Sequential(
@@ -38,12 +42,16 @@ class TimestepEmbedder(nn.Module):
         # https://github.com/openai/glide-text2im/blob/main/glide_text2im/nn.py
         half = dim // 2
         freqs = torch.exp(
-            -np.log(max_period) * torch.arange(start=0, end=half, dtype=torch.float32) / half
+            -np.log(max_period)
+            * torch.arange(start=0, end=half, dtype=torch.float32)
+            / half
         ).to(device=t.device)
         args = t[:, None].float() * freqs[None]
         embedding = torch.cat([torch.cos(args), torch.sin(args)], dim=-1)
         if dim % 2:
-            embedding = torch.cat([embedding, torch.zeros_like(embedding[:, :1])], dim=-1)
+            embedding = torch.cat(
+                [embedding, torch.zeros_like(embedding[:, :1])], dim=-1
+            )
         return embedding
 
     def forward(self, t):
@@ -93,34 +101,41 @@ class SparseStructureFlowModel(nn.Module):
         self.t_embedder = TimestepEmbedder(model_channels)
         if share_mod:
             self.adaLN_modulation = nn.Sequential(
-                nn.SiLU(),
-                nn.Linear(model_channels, 6 * model_channels, bias=True)
+                nn.SiLU(), nn.Linear(model_channels, 6 * model_channels, bias=True)
             )
 
         if pe_mode == "ape":
             pos_embedder = AbsolutePositionEmbedder(model_channels, 3)
-            coords = torch.meshgrid(*[torch.arange(res, device=self.device) for res in [resolution // patch_size] * 3], indexing='ij')
+            coords = torch.meshgrid(
+                *[
+                    torch.arange(res, device=self.device)
+                    for res in [resolution // patch_size] * 3
+                ],
+                indexing="ij",
+            )
             coords = torch.stack(coords, dim=-1).reshape(-1, 3)
             pos_emb = pos_embedder(coords)
             self.register_buffer("pos_emb", pos_emb)
 
         self.input_layer = nn.Linear(in_channels * patch_size**3, model_channels)
-            
-        self.blocks = nn.ModuleList([
-            ModulatedTransformerCrossBlock(
-                model_channels,
-                cond_channels,
-                num_heads=self.num_heads,
-                mlp_ratio=self.mlp_ratio,
-                attn_mode='full',
-                use_checkpoint=self.use_checkpoint,
-                use_rope=(pe_mode == "rope"),
-                share_mod=share_mod,
-                qk_rms_norm=self.qk_rms_norm,
-                qk_rms_norm_cross=self.qk_rms_norm_cross,
-            )
-            for _ in range(num_blocks)
-        ])
+
+        self.blocks = nn.ModuleList(
+            [
+                ModulatedTransformerCrossBlock(
+                    model_channels,
+                    cond_channels,
+                    num_heads=self.num_heads,
+                    mlp_ratio=self.mlp_ratio,
+                    attn_mode="full",
+                    use_checkpoint=self.use_checkpoint,
+                    use_rope=(pe_mode == "rope"),
+                    share_mod=share_mod,
+                    qk_rms_norm=self.qk_rms_norm,
+                    qk_rms_norm_cross=self.qk_rms_norm_cross,
+                )
+                for _ in range(num_blocks)
+            ]
+        )
 
         self.out_layer = nn.Linear(model_channels, out_channels * patch_size**3)
 
@@ -154,6 +169,7 @@ class SparseStructureFlowModel(nn.Module):
                 torch.nn.init.xavier_uniform_(module.weight)
                 if module.bias is not None:
                     nn.init.constant_(module.bias, 0)
+
         self.apply(_basic_init)
 
         # Initialize timestep embedding MLP:
@@ -173,9 +189,14 @@ class SparseStructureFlowModel(nn.Module):
         nn.init.constant_(self.out_layer.weight, 0)
         nn.init.constant_(self.out_layer.bias, 0)
 
-    def forward(self, x: torch.Tensor, t: torch.Tensor, cond: torch.Tensor) -> torch.Tensor:
-        assert [*x.shape] == [x.shape[0], self.in_channels, *[self.resolution] * 3], \
-                f"Input shape mismatch, got {x.shape}, expected {[x.shape[0], self.in_channels, *[self.resolution] * 3]}"
+    def forward(
+        self, x: torch.Tensor, t: torch.Tensor, cond: torch.Tensor
+    ) -> torch.Tensor:
+        assert [*x.shape] == [
+            x.shape[0],
+            self.in_channels,
+            *[self.resolution] * 3,
+        ], f"Input shape mismatch, got {x.shape}, expected {[x.shape[0], self.in_channels, *[self.resolution] * 3]}"
 
         h = patchify(x, self.patch_size)
         h = h.view(*h.shape[:2], -1).permute(0, 2, 1).contiguous()
@@ -194,7 +215,9 @@ class SparseStructureFlowModel(nn.Module):
         h = F.layer_norm(h, h.shape[-1:])
         h = self.out_layer(h)
 
-        h = h.permute(0, 2, 1).view(h.shape[0], h.shape[2], *[self.resolution // self.patch_size] * 3)
+        h = h.permute(0, 2, 1).view(
+            h.shape[0], h.shape[2], *[self.resolution // self.patch_size] * 3
+        )
         h = unpatchify(h, self.patch_size).contiguous()
 
         return h

trellis/models/sparse_structure_vae.py

@@ -33,9 +33,15 @@ class ResBlock3d(nn.Module):
         self.norm1 = norm_layer(norm_type, channels)
         self.norm2 = norm_layer(norm_type, self.out_channels)
         self.conv1 = nn.Conv3d(channels, self.out_channels, 3, padding=1)
-        self.conv2 = zero_module(nn.Conv3d(self.out_channels, self.out_channels, 3, padding=1))
-        self.skip_connection = nn.Conv3d(channels, self.out_channels, 1) if channels != self.out_channels else nn.Identity()
-    
+        self.conv2 = zero_module(
+            nn.Conv3d(self.out_channels, self.out_channels, 3, padding=1)
+        )
+        self.skip_connection = (
+            nn.Conv3d(channels, self.out_channels, 1)
+            if channels != self.out_channels
+            else nn.Identity()
+        )
+
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         h = self.norm1(x)
         h = F.silu(h)
@@ -63,7 +69,9 @@ class DownsampleBlock3d(nn.Module):
         if mode == "conv":
             self.conv = nn.Conv3d(in_channels, out_channels, 2, stride=2)
         elif mode == "avgpool":
-            assert in_channels == out_channels, "Pooling mode requires in_channels to be equal to out_channels"
+            assert (
+                in_channels == out_channels
+            ), "Pooling mode requires in_channels to be equal to out_channels"
 
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         if hasattr(self, "conv"):
@@ -86,9 +94,11 @@ class UpsampleBlock3d(nn.Module):
         self.out_channels = out_channels
 
         if mode == "conv":
-            self.conv = nn.Conv3d(in_channels, out_channels*8, 3, padding=1)
+            self.conv = nn.Conv3d(in_channels, out_channels * 8, 3, padding=1)
         elif mode == "nearest":
-            assert in_channels == out_channels, "Nearest mode requires in_channels to be equal to out_channels"
+            assert (
+                in_channels == out_channels
+            ), "Nearest mode requires in_channels to be equal to out_channels"
 
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         if hasattr(self, "conv"):
@@ -96,12 +106,12 @@ class UpsampleBlock3d(nn.Module):
             return pixel_shuffle_3d(x, 2)
         else:
             return F.interpolate(x, scale_factor=2, mode="nearest")
-        
+
 
 class SparseStructureEncoder(nn.Module):
     """
     Encoder for Sparse Structure (\mathcal{E}_S in the paper Sec. 3.3).
-    
+
     Args:
         in_channels (int): Channels of the input.
         latent_channels (int): Channels of the latent representation.
@@ -111,6 +121,7 @@ class SparseStructureEncoder(nn.Module):
         norm_type (Literal["group", "layer"]): Type of normalization layer.
         use_fp16 (bool): Whether to use FP16.
     """
+
     def __init__(
         self,
         in_channels: int,
@@ -135,24 +146,21 @@ class SparseStructureEncoder(nn.Module):
 
         self.blocks = nn.ModuleList([])
         for i, ch in enumerate(channels):
-            self.blocks.extend([
-                ResBlock3d(ch, ch)
-                for _ in range(num_res_blocks)
-            ])
+            self.blocks.extend([ResBlock3d(ch, ch) for _ in range(num_res_blocks)])
             if i < len(channels) - 1:
-                self.blocks.append(
-                    DownsampleBlock3d(ch, channels[i+1])
-                )
-        
-        self.middle_block = nn.Sequential(*[
-            ResBlock3d(channels[-1], channels[-1])
-            for _ in range(num_res_blocks_middle)
-        ])
+                self.blocks.append(DownsampleBlock3d(ch, channels[i + 1]))
+
+        self.middle_block = nn.Sequential(
+            *[
+                ResBlock3d(channels[-1], channels[-1])
+                for _ in range(num_res_blocks_middle)
+            ]
+        )
 
         self.out_layer = nn.Sequential(
             norm_layer(norm_type, channels[-1]),
             nn.SiLU(),
-            nn.Conv3d(channels[-1], latent_channels*2, 3, padding=1)
+            nn.Conv3d(channels[-1], latent_channels * 2, 3, padding=1),
         )
 
         if use_fp16:
@@ -183,7 +191,9 @@ class SparseStructureEncoder(nn.Module):
         self.blocks.apply(convert_module_to_f32)
         self.middle_block.apply(convert_module_to_f32)
 
-    def forward(self, x: torch.Tensor, sample_posterior: bool = False, return_raw: bool = False) -> torch.Tensor:
+    def forward(
+        self, x: torch.Tensor, sample_posterior: bool = False, return_raw: bool = False
+    ) -> torch.Tensor:
         h = self.input_layer(x)
         h = h.type(self.dtype)
 
@@ -201,16 +211,16 @@ class SparseStructureEncoder(nn.Module):
             z = mean + std * torch.randn_like(std)
         else:
             z = mean
-            
+
         if return_raw:
             return z, mean, logvar
         return z
-        
+
 
 class SparseStructureDecoder(nn.Module):
     """
     Decoder for Sparse Structure (\mathcal{D}_S in the paper Sec. 3.3).
-    
+
     Args:
         out_channels (int): Channels of the output.
         latent_channels (int): Channels of the latent representation.
@@ -219,7 +229,8 @@ class SparseStructureDecoder(nn.Module):
         num_res_blocks_middle (int): Number of residual blocks in the middle.
         norm_type (Literal["group", "layer"]): Type of normalization layer.
         use_fp16 (bool): Whether to use FP16.
-    """ 
+    """
+
     def __init__(
         self,
         out_channels: int,
@@ -242,26 +253,23 @@ class SparseStructureDecoder(nn.Module):
 
         self.input_layer = nn.Conv3d(latent_channels, channels[0], 3, padding=1)
 
-        self.middle_block = nn.Sequential(*[
-            ResBlock3d(channels[0], channels[0])
-            for _ in range(num_res_blocks_middle)
-        ])
+        self.middle_block = nn.Sequential(
+            *[
+                ResBlock3d(channels[0], channels[0])
+                for _ in range(num_res_blocks_middle)
+            ]
+        )
 
         self.blocks = nn.ModuleList([])
         for i, ch in enumerate(channels):
-            self.blocks.extend([
-                ResBlock3d(ch, ch)
-                for _ in range(num_res_blocks)
-            ])
+            self.blocks.extend([ResBlock3d(ch, ch) for _ in range(num_res_blocks)])
             if i < len(channels) - 1:
-                self.blocks.append(
-                    UpsampleBlock3d(ch, channels[i+1])
-                )
+                self.blocks.append(UpsampleBlock3d(ch, channels[i + 1]))
 
         self.out_layer = nn.Sequential(
             norm_layer(norm_type, channels[-1]),
             nn.SiLU(),
-            nn.Conv3d(channels[-1], out_channels, 3, padding=1)
+            nn.Conv3d(channels[-1], out_channels, 3, padding=1),
         )
 
         if use_fp16:
@@ -273,7 +281,7 @@ class SparseStructureDecoder(nn.Module):
         Return the device of the model.
         """
         return next(self.parameters()).device
-    
+
     def convert_to_fp16(self) -> None:
         """
         Convert the torso of the model to float16.
@@ -291,12 +299,12 @@ class SparseStructureDecoder(nn.Module):
         self.dtype = torch.float32
         self.blocks.apply(convert_module_to_f32)
         self.middle_block.apply(convert_module_to_f32)
-    
+
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         h = self.input_layer(x)
-        
+
         h = h.type(self.dtype)
-                
+
         h = self.middle_block(h)
         for block in self.blocks:
             h = block(h)

trellis/models/structured_latent_flow.py

@@ -26,18 +26,26 @@ class SparseResBlock3d(nn.Module):
         self.out_channels = out_channels or channels
         self.downsample = downsample
         self.upsample = upsample
-        
-        assert not (downsample and upsample), "Cannot downsample and upsample at the same time"
+
+        assert not (
+            downsample and upsample
+        ), "Cannot downsample and upsample at the same time"
 
         self.norm1 = LayerNorm32(channels, elementwise_affine=True, eps=1e-6)
         self.norm2 = LayerNorm32(self.out_channels, elementwise_affine=False, eps=1e-6)
         self.conv1 = sp.SparseConv3d(channels, self.out_channels, 3)
-        self.conv2 = zero_module(sp.SparseConv3d(self.out_channels, self.out_channels, 3))
+        self.conv2 = zero_module(
+            sp.SparseConv3d(self.out_channels, self.out_channels, 3)
+        )
         self.emb_layers = nn.Sequential(
             nn.SiLU(),
             nn.Linear(emb_channels, 2 * self.out_channels, bias=True),
         )
-        self.skip_connection = sp.SparseLinear(channels, self.out_channels) if channels != self.out_channels else nn.Identity()
+        self.skip_connection = (
+            sp.SparseLinear(channels, self.out_channels)
+            if channels != self.out_channels
+            else nn.Identity()
+        )
         self.updown = None
         if self.downsample:
             self.updown = sp.SparseDownsample(2)
@@ -63,7 +71,7 @@ class SparseResBlock3d(nn.Module):
         h = h + self.skip_connection(x)
 
         return h
-    
+
 
 class SLatFlowModel(nn.Module):
     def __init__(
@@ -109,14 +117,17 @@ class SLatFlowModel(nn.Module):
         self.qk_rms_norm_cross = qk_rms_norm_cross
         self.dtype = torch.float16 if use_fp16 else torch.float32
 
-        assert int(np.log2(patch_size)) == np.log2(patch_size), "Patch size must be a power of 2"
-        assert np.log2(patch_size) == len(io_block_channels), "Number of IO ResBlocks must match the number of stages"
+        assert int(np.log2(patch_size)) == np.log2(
+            patch_size
+        ), "Patch size must be a power of 2"
+        assert np.log2(patch_size) == len(
+            io_block_channels
+        ), "Number of IO ResBlocks must match the number of stages"
 
         self.t_embedder = TimestepEmbedder(model_channels)
         if share_mod:
             self.adaLN_modulation = nn.Sequential(
-                nn.SiLU(),
-                nn.Linear(model_channels, 6 * model_channels, bias=True)
+                nn.SiLU(), nn.Linear(model_channels, 6 * model_channels, bias=True)
             )
 
         if pe_mode == "ape":
@@ -124,15 +135,19 @@ class SLatFlowModel(nn.Module):
 
         self.input_layer = sp.SparseLinear(in_channels, io_block_channels[0])
         self.input_blocks = nn.ModuleList([])
-        for chs, next_chs in zip(io_block_channels, io_block_channels[1:] + [model_channels]):
-            self.input_blocks.extend([
-                SparseResBlock3d(
-                    chs,
-                    model_channels,
-                    out_channels=chs,
-                )
-                for _ in range(num_io_res_blocks-1)
-            ])
+        for chs, next_chs in zip(
+            io_block_channels, io_block_channels[1:] + [model_channels]
+        ):
+            self.input_blocks.extend(
+                [
+                    SparseResBlock3d(
+                        chs,
+                        model_channels,
+                        out_channels=chs,
+                    )
+                    for _ in range(num_io_res_blocks - 1)
+                ]
+            )
             self.input_blocks.append(
                 SparseResBlock3d(
                     chs,
@@ -141,25 +156,30 @@ class SLatFlowModel(nn.Module):
                     downsample=True,
                 )
             )
-            
-        self.blocks = nn.ModuleList([
-            ModulatedSparseTransformerCrossBlock(
-                model_channels,
-                cond_channels,
-                num_heads=self.num_heads,
-                mlp_ratio=self.mlp_ratio,
-                attn_mode='full',
-                use_checkpoint=self.use_checkpoint,
-                use_rope=(pe_mode == "rope"),
-                share_mod=self.share_mod,
-                qk_rms_norm=self.qk_rms_norm,
-                qk_rms_norm_cross=self.qk_rms_norm_cross,
-            )
-            for _ in range(num_blocks)
-        ])
+
+        self.blocks = nn.ModuleList(
+            [
+                ModulatedSparseTransformerCrossBlock(
+                    model_channels,
+                    cond_channels,
+                    num_heads=self.num_heads,
+                    mlp_ratio=self.mlp_ratio,
+                    attn_mode="full",
+                    use_checkpoint=self.use_checkpoint,
+                    use_rope=(pe_mode == "rope"),
+                    share_mod=self.share_mod,
+                    qk_rms_norm=self.qk_rms_norm,
+                    qk_rms_norm_cross=self.qk_rms_norm_cross,
+                )
+                for _ in range(num_blocks)
+            ]
+        )
 
         self.out_blocks = nn.ModuleList([])
-        for chs, prev_chs in zip(reversed(io_block_channels), [model_channels] + list(reversed(io_block_channels[1:]))):
+        for chs, prev_chs in zip(
+            reversed(io_block_channels),
+            [model_channels] + list(reversed(io_block_channels[1:])),
+        ):
             self.out_blocks.append(
                 SparseResBlock3d(
                     prev_chs * 2 if self.use_skip_connection else prev_chs,
@@ -168,14 +188,16 @@ class SLatFlowModel(nn.Module):
                     upsample=True,
                 )
             )
-            self.out_blocks.extend([
-                SparseResBlock3d(
-                    chs * 2 if self.use_skip_connection else chs,
-                    model_channels,
-                    out_channels=chs,
-                )
-                for _ in range(num_io_res_blocks-1)
-            ])
+            self.out_blocks.extend(
+                [
+                    SparseResBlock3d(
+                        chs * 2 if self.use_skip_connection else chs,
+                        model_channels,
+                        out_channels=chs,
+                    )
+                    for _ in range(num_io_res_blocks - 1)
+                ]
+            )
         self.out_layer = sp.SparseLinear(io_block_channels[0], out_channels)
 
         self.initialize_weights()
@@ -212,6 +234,7 @@ class SLatFlowModel(nn.Module):
                 torch.nn.init.xavier_uniform_(module.weight)
                 if module.bias is not None:
                     nn.init.constant_(module.bias, 0)
+
         self.apply(_basic_init)
 
         # Initialize timestep embedding MLP:
@@ -231,7 +254,9 @@ class SLatFlowModel(nn.Module):
         nn.init.constant_(self.out_layer.weight, 0)
         nn.init.constant_(self.out_layer.bias, 0)
 
-    def forward(self, x: sp.SparseTensor, t: torch.Tensor, cond: torch.Tensor) -> sp.SparseTensor:
+    def forward(
+        self, x: sp.SparseTensor, t: torch.Tensor, cond: torch.Tensor
+    ) -> sp.SparseTensor:
         h = self.input_layer(x).type(self.dtype)
         t_emb = self.t_embedder(t)
         if self.share_mod:
@@ -244,7 +269,7 @@ class SLatFlowModel(nn.Module):
         for block in self.input_blocks:
             h = block(h, t_emb)
             skips.append(h.feats)
-        
+
         if self.pe_mode == "ape":
             h = h + self.pos_embedder(h.coords[:, 1:]).type(self.dtype)
         for block in self.blocks:

trellis/models/structured_latent_vae/base.py

@@ -13,15 +13,23 @@ def block_attn_config(self):
     """
     for i in range(self.num_blocks):
         if self.attn_mode == "shift_window":
-            yield "serialized", self.window_size, 0, (16 * (i % 2),) * 3, sp.SerializeMode.Z_ORDER
+            yield "serialized", self.window_size, 0, (
+                16 * (i % 2),
+            ) * 3, sp.SerializeMode.Z_ORDER
         elif self.attn_mode == "shift_sequence":
-            yield "serialized", self.window_size, self.window_size // 2 * (i % 2), (0, 0, 0), sp.SerializeMode.Z_ORDER
+            yield "serialized", self.window_size, self.window_size // 2 * (i % 2), (
+                0,
+                0,
+                0,
+            ), sp.SerializeMode.Z_ORDER
         elif self.attn_mode == "shift_order":
             yield "serialized", self.window_size, 0, (0, 0, 0), sp.SerializeModes[i % 4]
         elif self.attn_mode == "full":
             yield "full", None, None, None, None
         elif self.attn_mode == "swin":
-            yield "windowed", self.window_size, None, self.window_size // 2 * (i % 2), None
+            yield "windowed", self.window_size, None, self.window_size // 2 * (
+                i % 2
+            ), None
 
 
 class SparseTransformerBase(nn.Module):
@@ -29,6 +37,7 @@ class SparseTransformerBase(nn.Module):
     Sparse Transformer without output layers.
     Serve as the base class for encoder and decoder.
     """
+
     def __init__(
         self,
         in_channels: int,
@@ -37,7 +46,9 @@ class SparseTransformerBase(nn.Module):
         num_heads: Optional[int] = None,
         num_head_channels: Optional[int] = 64,
         mlp_ratio: float = 4.0,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "full",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "full",
         window_size: Optional[int] = None,
         pe_mode: Literal["ape", "rope"] = "ape",
         use_fp16: bool = False,
@@ -62,22 +73,26 @@ class SparseTransformerBase(nn.Module):
             self.pos_embedder = AbsolutePositionEmbedder(model_channels)
 
         self.input_layer = sp.SparseLinear(in_channels, model_channels)
-        self.blocks = nn.ModuleList([
-            SparseTransformerBlock(
-                model_channels,
-                num_heads=self.num_heads,
-                mlp_ratio=self.mlp_ratio,
-                attn_mode=attn_mode,
-                window_size=window_size,
-                shift_sequence=shift_sequence,
-                shift_window=shift_window,
-                serialize_mode=serialize_mode,
-                use_checkpoint=self.use_checkpoint,
-                use_rope=(pe_mode == "rope"),
-                qk_rms_norm=self.qk_rms_norm,
-            )
-            for attn_mode, window_size, shift_sequence, shift_window, serialize_mode in block_attn_config(self)
-        ])
+        self.blocks = nn.ModuleList(
+            [
+                SparseTransformerBlock(
+                    model_channels,
+                    num_heads=self.num_heads,
+                    mlp_ratio=self.mlp_ratio,
+                    attn_mode=attn_mode,
+                    window_size=window_size,
+                    shift_sequence=shift_sequence,
+                    shift_window=shift_window,
+                    serialize_mode=serialize_mode,
+                    use_checkpoint=self.use_checkpoint,
+                    use_rope=(pe_mode == "rope"),
+                    qk_rms_norm=self.qk_rms_norm,
+                )
+                for attn_mode, window_size, shift_sequence, shift_window, serialize_mode in block_attn_config(
+                    self
+                )
+            ]
+        )
 
     @property
     def device(self) -> torch.device:
@@ -105,6 +120,7 @@ class SparseTransformerBase(nn.Module):
                 torch.nn.init.xavier_uniform_(module.weight)
                 if module.bias is not None:
                     nn.init.constant_(module.bias, 0)
+
         self.apply(_basic_init)
 
     def forward(self, x: sp.SparseTensor) -> sp.SparseTensor:

trellis/models/structured_latent_vae/decoder_gs.py

@@ -18,7 +18,9 @@ class SLatGaussianDecoder(SparseTransformerBase):
         num_heads: Optional[int] = None,
         num_head_channels: Optional[int] = 64,
         mlp_ratio: float = 4,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "swin",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "swin",
         window_size: int = 8,
         pe_mode: Literal["ape", "rope"] = "ape",
         use_fp16: bool = False,
@@ -57,26 +59,44 @@ class SLatGaussianDecoder(SparseTransformerBase):
         nn.init.constant_(self.out_layer.bias, 0)
 
     def _build_perturbation(self) -> None:
-        perturbation = [hammersley_sequence(3, i, self.rep_config['num_gaussians']) for i in range(self.rep_config['num_gaussians'])]
+        perturbation = [
+            hammersley_sequence(3, i, self.rep_config["num_gaussians"])
+            for i in range(self.rep_config["num_gaussians"])
+        ]
         perturbation = torch.tensor(perturbation).float() * 2 - 1
-        perturbation = perturbation / self.rep_config['voxel_size']
+        perturbation = perturbation / self.rep_config["voxel_size"]
         perturbation = torch.atanh(perturbation).to(self.device)
-        self.register_buffer('offset_perturbation', perturbation)
+        self.register_buffer("offset_perturbation", perturbation)
 
     def _calc_layout(self) -> None:
         self.layout = {
-            '_xyz' : {'shape': (self.rep_config['num_gaussians'], 3), 'size': self.rep_config['num_gaussians'] * 3},
-            '_features_dc' : {'shape': (self.rep_config['num_gaussians'], 1, 3), 'size': self.rep_config['num_gaussians'] * 3},
-            '_scaling' : {'shape': (self.rep_config['num_gaussians'], 3), 'size': self.rep_config['num_gaussians'] * 3},
-            '_rotation' : {'shape': (self.rep_config['num_gaussians'], 4), 'size': self.rep_config['num_gaussians'] * 4},
-            '_opacity' : {'shape': (self.rep_config['num_gaussians'], 1), 'size': self.rep_config['num_gaussians']},
+            "_xyz": {
+                "shape": (self.rep_config["num_gaussians"], 3),
+                "size": self.rep_config["num_gaussians"] * 3,
+            },
+            "_features_dc": {
+                "shape": (self.rep_config["num_gaussians"], 1, 3),
+                "size": self.rep_config["num_gaussians"] * 3,
+            },
+            "_scaling": {
+                "shape": (self.rep_config["num_gaussians"], 3),
+                "size": self.rep_config["num_gaussians"] * 3,
+            },
+            "_rotation": {
+                "shape": (self.rep_config["num_gaussians"], 4),
+                "size": self.rep_config["num_gaussians"] * 4,
+            },
+            "_opacity": {
+                "shape": (self.rep_config["num_gaussians"], 1),
+                "size": self.rep_config["num_gaussians"],
+            },
         }
         start = 0
         for k, v in self.layout.items():
-            v['range'] = (start, start + v['size'])
-            start += v['size']
+            v["range"] = (start, start + v["size"])
+            start += v["size"]
         self.out_channels = start
-    
+
     def to_representation(self, x: sp.SparseTensor) -> List[Gaussian]:
         """
         Convert a batch of network outputs to 3D representations.
@@ -92,24 +112,35 @@ class SLatGaussianDecoder(SparseTransformerBase):
             representation = Gaussian(
                 sh_degree=0,
                 aabb=[-0.5, -0.5, -0.5, 1.0, 1.0, 1.0],
-                mininum_kernel_size = self.rep_config['3d_filter_kernel_size'],
-                scaling_bias = self.rep_config['scaling_bias'],
-                opacity_bias = self.rep_config['opacity_bias'],
-                scaling_activation = self.rep_config['scaling_activation']
+                mininum_kernel_size=self.rep_config["3d_filter_kernel_size"],
+                scaling_bias=self.rep_config["scaling_bias"],
+                opacity_bias=self.rep_config["opacity_bias"],
+                scaling_activation=self.rep_config["scaling_activation"],
             )
             xyz = (x.coords[x.layout[i]][:, 1:].float() + 0.5) / self.resolution
             for k, v in self.layout.items():
-                if k == '_xyz':
-                    offset = x.feats[x.layout[i]][:, v['range'][0]:v['range'][1]].reshape(-1, *v['shape'])
-                    offset = offset * self.rep_config['lr'][k]
-                    if self.rep_config['perturb_offset']:
+                if k == "_xyz":
+                    offset = x.feats[x.layout[i]][
+                        :, v["range"][0] : v["range"][1]
+                    ].reshape(-1, *v["shape"])
+                    offset = offset * self.rep_config["lr"][k]
+                    if self.rep_config["perturb_offset"]:
                         offset = offset + self.offset_perturbation
-                    offset = torch.tanh(offset) / self.resolution * 0.5 * self.rep_config['voxel_size']
+                    offset = (
+                        torch.tanh(offset)
+                        / self.resolution
+                        * 0.5
+                        * self.rep_config["voxel_size"]
+                    )
                     _xyz = xyz.unsqueeze(1) + offset
                     setattr(representation, k, _xyz.flatten(0, 1))
                 else:
-                    feats = x.feats[x.layout[i]][:, v['range'][0]:v['range'][1]].reshape(-1, *v['shape']).flatten(0, 1)
-                    feats = feats * self.rep_config['lr'][k]
+                    feats = (
+                        x.feats[x.layout[i]][:, v["range"][0] : v["range"][1]]
+                        .reshape(-1, *v["shape"])
+                        .flatten(0, 1)
+                    )
+                    feats = feats * self.rep_config["lr"][k]
                     setattr(representation, k, feats)
             ret.append(representation)
         return ret

trellis/models/structured_latent_vae/decoder_mesh.py

@@ -19,12 +19,13 @@ class SparseSubdivideBlock3d(nn.Module):
         out_channels: if specified, the number of output channels.
         num_groups: the number of groups for the group norm.
     """
+
     def __init__(
         self,
         channels: int,
         resolution: int,
         out_channels: Optional[int] = None,
-        num_groups: int = 32
+        num_groups: int = 32,
     ):
         super().__init__()
         self.channels = channels
@@ -33,24 +34,34 @@ class SparseSubdivideBlock3d(nn.Module):
         self.out_channels = out_channels or channels
 
         self.act_layers = nn.Sequential(
-            sp.SparseGroupNorm32(num_groups, channels),
-            sp.SparseSiLU()
+            sp.SparseGroupNorm32(num_groups, channels), sp.SparseSiLU()
         )
-        
+
         self.sub = sp.SparseSubdivide()
-        
+
         self.out_layers = nn.Sequential(
-            sp.SparseConv3d(channels, self.out_channels, 3, indice_key=f"res_{self.out_resolution}"),
+            sp.SparseConv3d(
+                channels, self.out_channels, 3, indice_key=f"res_{self.out_resolution}"
+            ),
             sp.SparseGroupNorm32(num_groups, self.out_channels),
             sp.SparseSiLU(),
-            zero_module(sp.SparseConv3d(self.out_channels, self.out_channels, 3, indice_key=f"res_{self.out_resolution}")),
+            zero_module(
+                sp.SparseConv3d(
+                    self.out_channels,
+                    self.out_channels,
+                    3,
+                    indice_key=f"res_{self.out_resolution}",
+                )
+            ),
         )
-        
+
         if self.out_channels == channels:
             self.skip_connection = nn.Identity()
         else:
-            self.skip_connection = sp.SparseConv3d(channels, self.out_channels, 1, indice_key=f"res_{self.out_resolution}")
-        
+            self.skip_connection = sp.SparseConv3d(
+                channels, self.out_channels, 1, indice_key=f"res_{self.out_resolution}"
+            )
+
     def forward(self, x: sp.SparseTensor) -> sp.SparseTensor:
         """
         Apply the block to a Tensor, conditioned on a timestep embedding.
@@ -78,7 +89,9 @@ class SLatMeshDecoder(SparseTransformerBase):
         num_heads: Optional[int] = None,
         num_head_channels: Optional[int] = 64,
         mlp_ratio: float = 4,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "swin",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "swin",
         window_size: int = 8,
         pe_mode: Literal["ape", "rope"] = "ape",
         use_fp16: bool = False,
@@ -102,20 +115,24 @@ class SLatMeshDecoder(SparseTransformerBase):
         )
         self.resolution = resolution
         self.rep_config = representation_config
-        self.mesh_extractor = SparseFeatures2Mesh(res=self.resolution*4, use_color=self.rep_config.get('use_color', False))
+        self.mesh_extractor = SparseFeatures2Mesh(
+            res=self.resolution * 4, use_color=self.rep_config.get("use_color", False)
+        )
         self.out_channels = self.mesh_extractor.feats_channels
-        self.upsample = nn.ModuleList([
-            SparseSubdivideBlock3d(
-                channels=model_channels,
-                resolution=resolution,
-                out_channels=model_channels // 4
-            ),
-            SparseSubdivideBlock3d(
-                channels=model_channels // 4,
-                resolution=resolution * 2,
-                out_channels=model_channels // 8
-            )
-        ])
+        self.upsample = nn.ModuleList(
+            [
+                SparseSubdivideBlock3d(
+                    channels=model_channels,
+                    resolution=resolution,
+                    out_channels=model_channels // 4,
+                ),
+                SparseSubdivideBlock3d(
+                    channels=model_channels // 4,
+                    resolution=resolution * 2,
+                    out_channels=model_channels // 8,
+                ),
+            ]
+        )
         self.out_layer = sp.SparseLinear(model_channels // 8, self.out_channels)
 
         self.initialize_weights()
@@ -140,8 +157,8 @@ class SLatMeshDecoder(SparseTransformerBase):
         Convert the torso of the model to float32.
         """
         super().convert_to_fp32()
-        self.upsample.apply(convert_module_to_f32)  
-    
+        self.upsample.apply(convert_module_to_f32)
+
     def to_representation(self, x: sp.SparseTensor) -> List[MeshExtractResult]:
         """
         Convert a batch of network outputs to 3D representations.

trellis/models/structured_latent_vae/decoder_rf.py

@@ -18,7 +18,9 @@ class SLatRadianceFieldDecoder(SparseTransformerBase):
         num_heads: Optional[int] = None,
         num_head_channels: Optional[int] = 64,
         mlp_ratio: float = 4,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "swin",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "swin",
         window_size: int = 8,
         pe_mode: Literal["ape", "rope"] = "ape",
         use_fp16: bool = False,
@@ -57,16 +59,25 @@ class SLatRadianceFieldDecoder(SparseTransformerBase):
 
     def _calc_layout(self) -> None:
         self.layout = {
-            'trivec': {'shape': (self.rep_config['rank'], 3, self.rep_config['dim']), 'size': self.rep_config['rank'] * 3 * self.rep_config['dim']},
-            'density': {'shape': (self.rep_config['rank'],), 'size': self.rep_config['rank']},
-            'features_dc': {'shape': (self.rep_config['rank'], 1, 3), 'size': self.rep_config['rank'] * 3},
+            "trivec": {
+                "shape": (self.rep_config["rank"], 3, self.rep_config["dim"]),
+                "size": self.rep_config["rank"] * 3 * self.rep_config["dim"],
+            },
+            "density": {
+                "shape": (self.rep_config["rank"],),
+                "size": self.rep_config["rank"],
+            },
+            "features_dc": {
+                "shape": (self.rep_config["rank"], 1, 3),
+                "size": self.rep_config["rank"] * 3,
+            },
         }
         start = 0
         for k, v in self.layout.items():
-            v['range'] = (start, start + v['size'])
-            start += v['size']
-        self.out_channels = start    
-    
+            v["range"] = (start, start + v["size"])
+            start += v["size"]
+        self.out_channels = start
+
     def to_representation(self, x: sp.SparseTensor) -> List[Strivec]:
         """
         Convert a batch of network outputs to 3D representations.
@@ -83,15 +94,28 @@ class SLatRadianceFieldDecoder(SparseTransformerBase):
                 sh_degree=0,
                 resolution=self.resolution,
                 aabb=[-0.5, -0.5, -0.5, 1, 1, 1],
-                rank=self.rep_config['rank'],
-                dim=self.rep_config['dim'],
-                device='cuda',
+                rank=self.rep_config["rank"],
+                dim=self.rep_config["dim"],
+                device="cuda",
             )
             representation.density_shift = 0.0
-            representation.position = (x.coords[x.layout[i]][:, 1:].float() + 0.5) / self.resolution
-            representation.depth = torch.full((representation.position.shape[0], 1), int(np.log2(self.resolution)), dtype=torch.uint8, device='cuda')
+            representation.position = (
+                x.coords[x.layout[i]][:, 1:].float() + 0.5
+            ) / self.resolution
+            representation.depth = torch.full(
+                (representation.position.shape[0], 1),
+                int(np.log2(self.resolution)),
+                dtype=torch.uint8,
+                device="cuda",
+            )
             for k, v in self.layout.items():
-                setattr(representation, k, x.feats[x.layout[i]][:, v['range'][0]:v['range'][1]].reshape(-1, *v['shape']))
+                setattr(
+                    representation,
+                    k,
+                    x.feats[x.layout[i]][:, v["range"][0] : v["range"][1]].reshape(
+                        -1, *v["shape"]
+                    ),
+                )
             representation.trivec = representation.trivec + 1
             ret.append(representation)
         return ret

trellis/models/structured_latent_vae/encoder.py

@@ -17,7 +17,9 @@ class SLatEncoder(SparseTransformerBase):
         num_heads: Optional[int] = None,
         num_head_channels: Optional[int] = 64,
         mlp_ratio: float = 4,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "swin",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "swin",
         window_size: int = 8,
         pe_mode: Literal["ape", "rope"] = "ape",
         use_fp16: bool = False,
@@ -56,7 +58,7 @@ class SLatEncoder(SparseTransformerBase):
         h = h.type(x.dtype)
         h = h.replace(F.layer_norm(h.feats, h.feats.shape[-1:]))
         h = self.out_layer(h)
-        
+
         # Sample from the posterior distribution
         mean, logvar = h.feats.chunk(2, dim=-1)
         if sample_posterior:
@@ -65,7 +67,7 @@ class SLatEncoder(SparseTransformerBase):
         else:
             z = mean
         z = h.replace(z)
-            
+
         if return_raw:
             return z, mean, logvar
         else:

trellis/modules/attention/__init__.py

@@ -1,32 +1,39 @@
 from typing import *
 
-BACKEND = 'flash_attn' 
+BACKEND = "flash_attn"
 DEBUG = False
 
+
 def __from_env():
     import os
-    
+
     global BACKEND
     global DEBUG
-    
-    env_attn_backend = os.environ.get('ATTN_BACKEND')
-    env_sttn_debug = os.environ.get('ATTN_DEBUG')
-    
-    if env_attn_backend is not None and env_attn_backend in ['xformers', 'flash_attn', 'sdpa', 'naive']:
+
+    env_attn_backend = os.environ.get("ATTN_BACKEND")
+    env_sttn_debug = os.environ.get("ATTN_DEBUG")
+
+    if env_attn_backend is not None and env_attn_backend in [
+        "xformers",
+        "flash_attn",
+        "sdpa",
+        "naive",
+    ]:
         BACKEND = env_attn_backend
     if env_sttn_debug is not None:
-        DEBUG = env_sttn_debug == '1'
+        DEBUG = env_sttn_debug == "1"
 
     print(f"[ATTENTION] Using backend: {BACKEND}")
-        
+
 
 __from_env()
-    
 
-def set_backend(backend: Literal['xformers', 'flash_attn']):
+
+def set_backend(backend: Literal["xformers", "flash_attn"]):
     global BACKEND
     BACKEND = backend
 
+
 def set_debug(debug: bool):
     global DEBUG
     DEBUG = debug

trellis/modules/attention/full_attn.py

@@ -3,20 +3,20 @@ import torch
 import math
 from . import DEBUG, BACKEND
 
-if BACKEND == 'xformers':
+if BACKEND == "xformers":
     import xformers.ops as xops
-elif BACKEND == 'flash_attn':
+elif BACKEND == "flash_attn":
     import flash_attn
-elif BACKEND == 'sdpa':
+elif BACKEND == "sdpa":
     from torch.nn.functional import scaled_dot_product_attention as sdpa
-elif BACKEND == 'naive':
+elif BACKEND == "naive":
     pass
 else:
     raise ValueError(f"Unknown attention backend: {BACKEND}")
 
 
 __all__ = [
-    'scaled_dot_product_attention',
+    "scaled_dot_product_attention",
 ]
 
 
@@ -24,14 +24,14 @@ def _naive_sdpa(q, k, v):
     """
     Naive implementation of scaled dot product attention.
     """
-    q = q.permute(0, 2, 1, 3)   # [N, H, L, C]
-    k = k.permute(0, 2, 1, 3)   # [N, H, L, C]
-    v = v.permute(0, 2, 1, 3)   # [N, H, L, C]
+    q = q.permute(0, 2, 1, 3)  # [N, H, L, C]
+    k = k.permute(0, 2, 1, 3)  # [N, H, L, C]
+    v = v.permute(0, 2, 1, 3)  # [N, H, L, C]
     scale_factor = 1 / math.sqrt(q.size(-1))
     attn_weight = q @ k.transpose(-2, -1) * scale_factor
     attn_weight = torch.softmax(attn_weight, dim=-1)
     out = attn_weight @ v
-    out = out.permute(0, 2, 1, 3)   # [N, L, H, C]
+    out = out.permute(0, 2, 1, 3)  # [N, L, H, C]
     return out
 
 
@@ -45,6 +45,7 @@ def scaled_dot_product_attention(qkv: torch.Tensor) -> torch.Tensor:
     """
     ...
 
+
 @overload
 def scaled_dot_product_attention(q: torch.Tensor, kv: torch.Tensor) -> torch.Tensor:
     """
@@ -56,8 +57,11 @@ def scaled_dot_product_attention(q: torch.Tensor, kv: torch.Tensor) -> torch.Ten
     """
     ...
 
+
 @overload
-def scaled_dot_product_attention(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor) -> torch.Tensor:
+def scaled_dot_product_attention(
+    q: torch.Tensor, k: torch.Tensor, v: torch.Tensor
+) -> torch.Tensor:
     """
     Apply scaled dot product attention.
 
@@ -71,64 +75,79 @@ def scaled_dot_product_attention(q: torch.Tensor, k: torch.Tensor, v: torch.Tens
     """
     ...
 
+
 def scaled_dot_product_attention(*args, **kwargs):
-    arg_names_dict = {
-        1: ['qkv'],
-        2: ['q', 'kv'],
-        3: ['q', 'k', 'v']
-    }
+    arg_names_dict = {1: ["qkv"], 2: ["q", "kv"], 3: ["q", "k", "v"]}
     num_all_args = len(args) + len(kwargs)
-    assert num_all_args in arg_names_dict, f"Invalid number of arguments, got {num_all_args}, expected 1, 2, or 3"
-    for key in arg_names_dict[num_all_args][len(args):]:
+    assert (
+        num_all_args in arg_names_dict
+    ), f"Invalid number of arguments, got {num_all_args}, expected 1, 2, or 3"
+    for key in arg_names_dict[num_all_args][len(args) :]:
         assert key in kwargs, f"Missing argument {key}"
 
     if num_all_args == 1:
-        qkv = args[0] if len(args) > 0 else kwargs['qkv']
-        assert len(qkv.shape) == 5 and qkv.shape[2] == 3, f"Invalid shape for qkv, got {qkv.shape}, expected [N, L, 3, H, C]"
+        qkv = args[0] if len(args) > 0 else kwargs["qkv"]
+        assert (
+            len(qkv.shape) == 5 and qkv.shape[2] == 3
+        ), f"Invalid shape for qkv, got {qkv.shape}, expected [N, L, 3, H, C]"
         device = qkv.device
 
     elif num_all_args == 2:
-        q = args[0] if len(args) > 0 else kwargs['q']
-        kv = args[1] if len(args) > 1 else kwargs['kv']
-        assert q.shape[0] == kv.shape[0], f"Batch size mismatch, got {q.shape[0]} and {kv.shape[0]}"
-        assert len(q.shape) == 4, f"Invalid shape for q, got {q.shape}, expected [N, L, H, C]"
-        assert len(kv.shape) == 5, f"Invalid shape for kv, got {kv.shape}, expected [N, L, 2, H, C]"
+        q = args[0] if len(args) > 0 else kwargs["q"]
+        kv = args[1] if len(args) > 1 else kwargs["kv"]
+        assert (
+            q.shape[0] == kv.shape[0]
+        ), f"Batch size mismatch, got {q.shape[0]} and {kv.shape[0]}"
+        assert (
+            len(q.shape) == 4
+        ), f"Invalid shape for q, got {q.shape}, expected [N, L, H, C]"
+        assert (
+            len(kv.shape) == 5
+        ), f"Invalid shape for kv, got {kv.shape}, expected [N, L, 2, H, C]"
         device = q.device
 
     elif num_all_args == 3:
-        q = args[0] if len(args) > 0 else kwargs['q']
-        k = args[1] if len(args) > 1 else kwargs['k']
-        v = args[2] if len(args) > 2 else kwargs['v']
-        assert q.shape[0] == k.shape[0] == v.shape[0], f"Batch size mismatch, got {q.shape[0]}, {k.shape[0]}, and {v.shape[0]}"
-        assert len(q.shape) == 4, f"Invalid shape for q, got {q.shape}, expected [N, L, H, Ci]"
-        assert len(k.shape) == 4, f"Invalid shape for k, got {k.shape}, expected [N, L, H, Ci]"
-        assert len(v.shape) == 4, f"Invalid shape for v, got {v.shape}, expected [N, L, H, Co]"
-        device = q.device    
-
-    if BACKEND == 'xformers':
+        q = args[0] if len(args) > 0 else kwargs["q"]
+        k = args[1] if len(args) > 1 else kwargs["k"]
+        v = args[2] if len(args) > 2 else kwargs["v"]
+        assert (
+            q.shape[0] == k.shape[0] == v.shape[0]
+        ), f"Batch size mismatch, got {q.shape[0]}, {k.shape[0]}, and {v.shape[0]}"
+        assert (
+            len(q.shape) == 4
+        ), f"Invalid shape for q, got {q.shape}, expected [N, L, H, Ci]"
+        assert (
+            len(k.shape) == 4
+        ), f"Invalid shape for k, got {k.shape}, expected [N, L, H, Ci]"
+        assert (
+            len(v.shape) == 4
+        ), f"Invalid shape for v, got {v.shape}, expected [N, L, H, Co]"
+        device = q.device
+
+    if BACKEND == "xformers":
         if num_all_args == 1:
             q, k, v = qkv.unbind(dim=2)
         elif num_all_args == 2:
             k, v = kv.unbind(dim=2)
         out = xops.memory_efficient_attention(q, k, v)
-    elif BACKEND == 'flash_attn':
+    elif BACKEND == "flash_attn":
         if num_all_args == 1:
             out = flash_attn.flash_attn_qkvpacked_func(qkv)
         elif num_all_args == 2:
             out = flash_attn.flash_attn_kvpacked_func(q, kv)
         elif num_all_args == 3:
             out = flash_attn.flash_attn_func(q, k, v)
-    elif BACKEND == 'sdpa':
+    elif BACKEND == "sdpa":
         if num_all_args == 1:
             q, k, v = qkv.unbind(dim=2)
         elif num_all_args == 2:
             k, v = kv.unbind(dim=2)
-        q = q.permute(0, 2, 1, 3)   # [N, H, L, C]
-        k = k.permute(0, 2, 1, 3)   # [N, H, L, C]
-        v = v.permute(0, 2, 1, 3)   # [N, H, L, C]
-        out = sdpa(q, k, v)         # [N, H, L, C]
-        out = out.permute(0, 2, 1, 3)   # [N, L, H, C]
-    elif BACKEND == 'naive':
+        q = q.permute(0, 2, 1, 3)  # [N, H, L, C]
+        k = k.permute(0, 2, 1, 3)  # [N, H, L, C]
+        v = v.permute(0, 2, 1, 3)  # [N, H, L, C]
+        out = sdpa(q, k, v)  # [N, H, L, C]
+        out = out.permute(0, 2, 1, 3)  # [N, L, H, C]
+    elif BACKEND == "naive":
         if num_all_args == 1:
             q, k, v = qkv.unbind(dim=2)
         elif num_all_args == 2:
@@ -136,5 +155,5 @@ def scaled_dot_product_attention(*args, **kwargs):
         out = _naive_sdpa(q, k, v)
     else:
         raise ValueError(f"Unknown attention module: {BACKEND}")
-    
+
     return out

trellis/modules/attention/modules.py

@@ -8,11 +8,11 @@ from .full_attn import scaled_dot_product_attention
 class MultiHeadRMSNorm(nn.Module):
     def __init__(self, dim: int, heads: int):
         super().__init__()
-        self.scale = dim ** 0.5
+        self.scale = dim**0.5
         self.gamma = nn.Parameter(torch.ones(heads, dim))
 
     def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return (F.normalize(x.float(), dim = -1) * self.gamma * self.scale).to(x.dtype)
+        return (F.normalize(x.float(), dim=-1) * self.gamma * self.scale).to(x.dtype)
 
 
 class RotaryPositionEmbedder(nn.Module):
@@ -23,21 +23,25 @@ class RotaryPositionEmbedder(nn.Module):
         self.in_channels = in_channels
         self.freq_dim = hidden_size // in_channels // 2
         self.freqs = torch.arange(self.freq_dim, dtype=torch.float32) / self.freq_dim
-        self.freqs = 1.0 / (10000 ** self.freqs)
-        
+        self.freqs = 1.0 / (10000**self.freqs)
+
     def _get_phases(self, indices: torch.Tensor) -> torch.Tensor:
         self.freqs = self.freqs.to(indices.device)
         phases = torch.outer(indices, self.freqs)
         phases = torch.polar(torch.ones_like(phases), phases)
         return phases
-        
+
     def _rotary_embedding(self, x: torch.Tensor, phases: torch.Tensor) -> torch.Tensor:
         x_complex = torch.view_as_complex(x.float().reshape(*x.shape[:-1], -1, 2))
         x_rotated = x_complex * phases
-        x_embed = torch.view_as_real(x_rotated).reshape(*x_rotated.shape[:-1], -1).to(x.dtype)
+        x_embed = (
+            torch.view_as_real(x_rotated).reshape(*x_rotated.shape[:-1], -1).to(x.dtype)
+        )
         return x_embed
-        
-    def forward(self, q: torch.Tensor, k: torch.Tensor, indices: Optional[torch.Tensor] = None) -> Tuple[torch.Tensor, torch.Tensor]:
+
+    def forward(
+        self, q: torch.Tensor, k: torch.Tensor, indices: Optional[torch.Tensor] = None
+    ) -> Tuple[torch.Tensor, torch.Tensor]:
         """
         Args:
             q (sp.SparseTensor): [..., N, D] tensor of queries
@@ -48,24 +52,38 @@ class RotaryPositionEmbedder(nn.Module):
             indices = torch.arange(q.shape[-2], device=q.device)
             if len(q.shape) > 2:
                 indices = indices.unsqueeze(0).expand(q.shape[:-2] + (-1,))
-        
+
         phases = self._get_phases(indices.reshape(-1)).reshape(*indices.shape[:-1], -1)
         if phases.shape[1] < self.hidden_size // 2:
-            phases = torch.cat([phases, torch.polar(
-                torch.ones(*phases.shape[:-1], self.hidden_size // 2 - phases.shape[1], device=phases.device),
-                torch.zeros(*phases.shape[:-1], self.hidden_size // 2 - phases.shape[1], device=phases.device)
-            )], dim=-1)
+            phases = torch.cat(
+                [
+                    phases,
+                    torch.polar(
+                        torch.ones(
+                            *phases.shape[:-1],
+                            self.hidden_size // 2 - phases.shape[1],
+                            device=phases.device,
+                        ),
+                        torch.zeros(
+                            *phases.shape[:-1],
+                            self.hidden_size // 2 - phases.shape[1],
+                            device=phases.device,
+                        ),
+                    ),
+                ],
+                dim=-1,
+            )
         q_embed = self._rotary_embedding(q, phases)
         k_embed = self._rotary_embedding(k, phases)
         return q_embed, k_embed
-    
+
 
 class MultiHeadAttention(nn.Module):
     def __init__(
         self,
         channels: int,
         num_heads: int,
-        ctx_channels: Optional[int]=None,
+        ctx_channels: Optional[int] = None,
         type: Literal["self", "cross"] = "self",
         attn_mode: Literal["full", "windowed"] = "full",
         window_size: Optional[int] = None,
@@ -78,11 +96,13 @@ class MultiHeadAttention(nn.Module):
         assert channels % num_heads == 0
         assert type in ["self", "cross"], f"Invalid attention type: {type}"
         assert attn_mode in ["full", "windowed"], f"Invalid attention mode: {attn_mode}"
-        assert type == "self" or attn_mode == "full", "Cross-attention only supports full attention"
-        
+        assert (
+            type == "self" or attn_mode == "full"
+        ), "Cross-attention only supports full attention"
+
         if attn_mode == "windowed":
             raise NotImplementedError("Windowed attention is not yet implemented")
-        
+
         self.channels = channels
         self.head_dim = channels // num_heads
         self.ctx_channels = ctx_channels if ctx_channels is not None else channels
@@ -99,17 +119,22 @@ class MultiHeadAttention(nn.Module):
         else:
             self.to_q = nn.Linear(channels, channels, bias=qkv_bias)
             self.to_kv = nn.Linear(self.ctx_channels, channels * 2, bias=qkv_bias)
-            
+
         if self.qk_rms_norm:
             self.q_rms_norm = MultiHeadRMSNorm(self.head_dim, num_heads)
             self.k_rms_norm = MultiHeadRMSNorm(self.head_dim, num_heads)
-            
+
         self.to_out = nn.Linear(channels, channels)
 
         if use_rope:
             self.rope = RotaryPositionEmbedder(channels)
-    
-    def forward(self, x: torch.Tensor, context: Optional[torch.Tensor] = None, indices: Optional[torch.Tensor] = None) -> torch.Tensor:
+
+    def forward(
+        self,
+        x: torch.Tensor,
+        context: Optional[torch.Tensor] = None,
+        indices: Optional[torch.Tensor] = None,
+    ) -> torch.Tensor:
         B, L, C = x.shape
         if self._type == "self":
             qkv = self.to_qkv(x)

trellis/modules/norm.py

@@ -5,21 +5,21 @@ import torch.nn as nn
 class LayerNorm32(nn.LayerNorm):
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         return super().forward(x.float()).type(x.dtype)
-    
+
 
 class GroupNorm32(nn.GroupNorm):
     """
     A GroupNorm layer that converts to float32 before the forward pass.
     """
+
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         return super().forward(x.float()).type(x.dtype)
-    
-    
+
+
 class ChannelLayerNorm32(LayerNorm32):
     def forward(self, x: torch.Tensor) -> torch.Tensor:
         DIM = x.dim()
         x = x.permute(0, *range(2, DIM), 1).contiguous()
         x = super().forward(x)
-        x = x.permute(0, DIM-1, *range(1, DIM-1)).contiguous()
+        x = x.permute(0, DIM - 1, *range(1, DIM - 1)).contiguous()
         return x
-    

trellis/modules/sparse/__init__.py

@@ -1,81 +1,88 @@
 from typing import *
 
-BACKEND = 'spconv' 
+BACKEND = "spconv"
 DEBUG = False
-ATTN = 'flash_attn'
+ATTN = "flash_attn"
+
 
 def __from_env():
     import os
-    
+
     global BACKEND
     global DEBUG
     global ATTN
-    
-    env_sparse_backend = os.environ.get('SPARSE_BACKEND')
-    env_sparse_debug = os.environ.get('SPARSE_DEBUG')
-    env_sparse_attn = os.environ.get('SPARSE_ATTN_BACKEND')
+
+    env_sparse_backend = os.environ.get("SPARSE_BACKEND")
+    env_sparse_debug = os.environ.get("SPARSE_DEBUG")
+    env_sparse_attn = os.environ.get("SPARSE_ATTN_BACKEND")
     if env_sparse_attn is None:
-        env_sparse_attn = os.environ.get('ATTN_BACKEND')
+        env_sparse_attn = os.environ.get("ATTN_BACKEND")
 
-    if env_sparse_backend is not None and env_sparse_backend in ['spconv', 'torchsparse']:
+    if env_sparse_backend is not None and env_sparse_backend in [
+        "spconv",
+        "torchsparse",
+    ]:
         BACKEND = env_sparse_backend
     if env_sparse_debug is not None:
-        DEBUG = env_sparse_debug == '1'
-    if env_sparse_attn is not None and env_sparse_attn in ['xformers', 'flash_attn']:
+        DEBUG = env_sparse_debug == "1"
+    if env_sparse_attn is not None and env_sparse_attn in ["xformers", "flash_attn"]:
         ATTN = env_sparse_attn
-        
+
     print(f"[SPARSE] Backend: {BACKEND}, Attention: {ATTN}")
-        
+
 
 __from_env()
-    
 
-def set_backend(backend: Literal['spconv', 'torchsparse']):
+
+def set_backend(backend: Literal["spconv", "torchsparse"]):
     global BACKEND
     BACKEND = backend
 
+
 def set_debug(debug: bool):
     global DEBUG
     DEBUG = debug
 
-def set_attn(attn: Literal['xformers', 'flash_attn']):
+
+def set_attn(attn: Literal["xformers", "flash_attn"]):
     global ATTN
     ATTN = attn
-    
-    
+
+
 import importlib
 
 __attributes = {
-    'SparseTensor': 'basic',
-    'sparse_batch_broadcast': 'basic',
-    'sparse_batch_op': 'basic',
-    'sparse_cat': 'basic',
-    'sparse_unbind': 'basic',
-    'SparseGroupNorm': 'norm',
-    'SparseLayerNorm': 'norm',
-    'SparseGroupNorm32': 'norm',
-    'SparseLayerNorm32': 'norm',
-    'SparseReLU': 'nonlinearity',
-    'SparseSiLU': 'nonlinearity',
-    'SparseGELU': 'nonlinearity',
-    'SparseActivation': 'nonlinearity',
-    'SparseLinear': 'linear',
-    'sparse_scaled_dot_product_attention': 'attention',
-    'SerializeMode': 'attention',
-    'sparse_serialized_scaled_dot_product_self_attention': 'attention',
-    'sparse_windowed_scaled_dot_product_self_attention': 'attention',
-    'SparseMultiHeadAttention': 'attention',
-    'SparseConv3d': 'conv',
-    'SparseInverseConv3d': 'conv',
-    'SparseDownsample': 'spatial',
-    'SparseUpsample': 'spatial',
-    'SparseSubdivide' : 'spatial'
+    "SparseTensor": "basic",
+    "sparse_batch_broadcast": "basic",
+    "sparse_batch_op": "basic",
+    "sparse_cat": "basic",
+    "sparse_unbind": "basic",
+    "SparseGroupNorm": "norm",
+    "SparseLayerNorm": "norm",
+    "SparseGroupNorm32": "norm",
+    "SparseLayerNorm32": "norm",
+    "SparseReLU": "nonlinearity",
+    "SparseSiLU": "nonlinearity",
+    "SparseGELU": "nonlinearity",
+    "SparseActivation": "nonlinearity",
+    "SparseLinear": "linear",
+    "sparse_scaled_dot_product_attention": "attention",
+    "SerializeMode": "attention",
+    "sparse_serialized_scaled_dot_product_self_attention": "attention",
+    "sparse_windowed_scaled_dot_product_self_attention": "attention",
+    "SparseMultiHeadAttention": "attention",
+    "SparseConv3d": "conv",
+    "SparseInverseConv3d": "conv",
+    "SparseDownsample": "spatial",
+    "SparseUpsample": "spatial",
+    "SparseSubdivide": "spatial",
 }
 
-__submodules = ['transformer']
+__submodules = ["transformer"]
 
 __all__ = list(__attributes.keys()) + __submodules
 
+
 def __getattr__(name):
     if name not in globals():
         if name in __attributes:
@@ -91,7 +98,7 @@ def __getattr__(name):
 
 
 # For Pylance
-if __name__ == '__main__':
+if __name__ == "__main__":
     from .basic import *
     from .norm import *
     from .nonlinearity import *

trellis/modules/sparse/attention/full_attn.py

@@ -3,16 +3,16 @@ import torch
 from .. import SparseTensor
 from .. import DEBUG, ATTN
 
-if ATTN == 'xformers':
+if ATTN == "xformers":
     import xformers.ops as xops
-elif ATTN == 'flash_attn':
+elif ATTN == "flash_attn":
     import flash_attn
 else:
     raise ValueError(f"Unknown attention module: {ATTN}")
 
 
 __all__ = [
-    'sparse_scaled_dot_product_attention',
+    "sparse_scaled_dot_product_attention",
 ]
 
 
@@ -26,8 +26,11 @@ def sparse_scaled_dot_product_attention(qkv: SparseTensor) -> SparseTensor:
     """
     ...
 
+
 @overload
-def sparse_scaled_dot_product_attention(q: SparseTensor, kv: Union[SparseTensor, torch.Tensor]) -> SparseTensor:
+def sparse_scaled_dot_product_attention(
+    q: SparseTensor, kv: Union[SparseTensor, torch.Tensor]
+) -> SparseTensor:
     """
     Apply scaled dot product attention to a sparse tensor.
 
@@ -37,8 +40,11 @@ def sparse_scaled_dot_product_attention(q: SparseTensor, kv: Union[SparseTensor,
     """
     ...
 
+
 @overload
-def sparse_scaled_dot_product_attention(q: torch.Tensor, kv: SparseTensor) -> torch.Tensor:
+def sparse_scaled_dot_product_attention(
+    q: torch.Tensor, kv: SparseTensor
+) -> torch.Tensor:
     """
     Apply scaled dot product attention to a sparse tensor.
 
@@ -48,8 +54,11 @@ def sparse_scaled_dot_product_attention(q: torch.Tensor, kv: SparseTensor) -> to
     """
     ...
 
+
 @overload
-def sparse_scaled_dot_product_attention(q: SparseTensor, k: SparseTensor, v: SparseTensor) -> SparseTensor:
+def sparse_scaled_dot_product_attention(
+    q: SparseTensor, k: SparseTensor, v: SparseTensor
+) -> SparseTensor:
     """
     Apply scaled dot product attention to a sparse tensor.
 
@@ -63,8 +72,11 @@ def sparse_scaled_dot_product_attention(q: SparseTensor, k: SparseTensor, v: Spa
     """
     ...
 
+
 @overload
-def sparse_scaled_dot_product_attention(q: SparseTensor, k: torch.Tensor, v: torch.Tensor) -> SparseTensor:
+def sparse_scaled_dot_product_attention(
+    q: SparseTensor, k: torch.Tensor, v: torch.Tensor
+) -> SparseTensor:
     """
     Apply scaled dot product attention to a sparse tensor.
 
@@ -75,8 +87,11 @@ def sparse_scaled_dot_product_attention(q: SparseTensor, k: torch.Tensor, v: tor
     """
     ...
 
+
 @overload
-def sparse_scaled_dot_product_attention(q: torch.Tensor, k: SparseTensor, v: SparseTensor) -> torch.Tensor:
+def sparse_scaled_dot_product_attention(
+    q: torch.Tensor, k: SparseTensor, v: SparseTensor
+) -> torch.Tensor:
     """
     Apply scaled dot product attention to a sparse tensor.
 
@@ -87,106 +102,158 @@ def sparse_scaled_dot_product_attention(q: torch.Tensor, k: SparseTensor, v: Spa
     """
     ...
 
+
 def sparse_scaled_dot_product_attention(*args, **kwargs):
-    arg_names_dict = {
-        1: ['qkv'],
-        2: ['q', 'kv'],
-        3: ['q', 'k', 'v']
-    }
+    arg_names_dict = {1: ["qkv"], 2: ["q", "kv"], 3: ["q", "k", "v"]}
     num_all_args = len(args) + len(kwargs)
-    assert num_all_args in arg_names_dict, f"Invalid number of arguments, got {num_all_args}, expected 1, 2, or 3"
-    for key in arg_names_dict[num_all_args][len(args):]:
+    assert (
+        num_all_args in arg_names_dict
+    ), f"Invalid number of arguments, got {num_all_args}, expected 1, 2, or 3"
+    for key in arg_names_dict[num_all_args][len(args) :]:
         assert key in kwargs, f"Missing argument {key}"
 
     if num_all_args == 1:
-        qkv = args[0] if len(args) > 0 else kwargs['qkv']
-        assert isinstance(qkv, SparseTensor), f"qkv must be a SparseTensor, got {type(qkv)}"
-        assert len(qkv.shape) == 4 and qkv.shape[1] == 3, f"Invalid shape for qkv, got {qkv.shape}, expected [N, *, 3, H, C]"
+        qkv = args[0] if len(args) > 0 else kwargs["qkv"]
+        assert isinstance(
+            qkv, SparseTensor
+        ), f"qkv must be a SparseTensor, got {type(qkv)}"
+        assert (
+            len(qkv.shape) == 4 and qkv.shape[1] == 3
+        ), f"Invalid shape for qkv, got {qkv.shape}, expected [N, *, 3, H, C]"
         device = qkv.device
 
         s = qkv
-        q_seqlen = [qkv.layout[i].stop - qkv.layout[i].start for i in range(qkv.shape[0])]
+        q_seqlen = [
+            qkv.layout[i].stop - qkv.layout[i].start for i in range(qkv.shape[0])
+        ]
         kv_seqlen = q_seqlen
-        qkv = qkv.feats     # [T, 3, H, C]
+        qkv = qkv.feats  # [T, 3, H, C]
 
     elif num_all_args == 2:
-        q = args[0] if len(args) > 0 else kwargs['q']
-        kv = args[1] if len(args) > 1 else kwargs['kv']
-        assert isinstance(q, SparseTensor) and isinstance(kv, (SparseTensor, torch.Tensor)) or \
-               isinstance(q, torch.Tensor) and isinstance(kv, SparseTensor), \
-               f"Invalid types, got {type(q)} and {type(kv)}"
-        assert q.shape[0] == kv.shape[0], f"Batch size mismatch, got {q.shape[0]} and {kv.shape[0]}"
+        q = args[0] if len(args) > 0 else kwargs["q"]
+        kv = args[1] if len(args) > 1 else kwargs["kv"]
+        assert (
+            isinstance(q, SparseTensor)
+            and isinstance(kv, (SparseTensor, torch.Tensor))
+            or isinstance(q, torch.Tensor)
+            and isinstance(kv, SparseTensor)
+        ), f"Invalid types, got {type(q)} and {type(kv)}"
+        assert (
+            q.shape[0] == kv.shape[0]
+        ), f"Batch size mismatch, got {q.shape[0]} and {kv.shape[0]}"
         device = q.device
 
         if isinstance(q, SparseTensor):
-            assert len(q.shape) == 3, f"Invalid shape for q, got {q.shape}, expected [N, *, H, C]"
+            assert (
+                len(q.shape) == 3
+            ), f"Invalid shape for q, got {q.shape}, expected [N, *, H, C]"
             s = q
             q_seqlen = [q.layout[i].stop - q.layout[i].start for i in range(q.shape[0])]
-            q = q.feats     # [T_Q, H, C]
+            q = q.feats  # [T_Q, H, C]
         else:
-            assert len(q.shape) == 4, f"Invalid shape for q, got {q.shape}, expected [N, L, H, C]"
+            assert (
+                len(q.shape) == 4
+            ), f"Invalid shape for q, got {q.shape}, expected [N, L, H, C]"
             s = None
             N, L, H, C = q.shape
             q_seqlen = [L] * N
-            q = q.reshape(N * L, H, C)   # [T_Q, H, C]
+            q = q.reshape(N * L, H, C)  # [T_Q, H, C]
 
         if isinstance(kv, SparseTensor):
-            assert len(kv.shape) == 4 and kv.shape[1] == 2, f"Invalid shape for kv, got {kv.shape}, expected [N, *, 2, H, C]"
-            kv_seqlen = [kv.layout[i].stop - kv.layout[i].start for i in range(kv.shape[0])]
-            kv = kv.feats     # [T_KV, 2, H, C]
+            assert (
+                len(kv.shape) == 4 and kv.shape[1] == 2
+            ), f"Invalid shape for kv, got {kv.shape}, expected [N, *, 2, H, C]"
+            kv_seqlen = [
+                kv.layout[i].stop - kv.layout[i].start for i in range(kv.shape[0])
+            ]
+            kv = kv.feats  # [T_KV, 2, H, C]
         else:
-            assert len(kv.shape) == 5, f"Invalid shape for kv, got {kv.shape}, expected [N, L, 2, H, C]"
+            assert (
+                len(kv.shape) == 5
+            ), f"Invalid shape for kv, got {kv.shape}, expected [N, L, 2, H, C]"
             N, L, _, H, C = kv.shape
             kv_seqlen = [L] * N
-            kv = kv.reshape(N * L, 2, H, C)   # [T_KV, 2, H, C]
+            kv = kv.reshape(N * L, 2, H, C)  # [T_KV, 2, H, C]
 
     elif num_all_args == 3:
-        q = args[0] if len(args) > 0 else kwargs['q']
-        k = args[1] if len(args) > 1 else kwargs['k']
-        v = args[2] if len(args) > 2 else kwargs['v']
-        assert isinstance(q, SparseTensor) and isinstance(k, (SparseTensor, torch.Tensor)) and type(k) == type(v) or \
-               isinstance(q, torch.Tensor) and isinstance(k, SparseTensor) and isinstance(v, SparseTensor), \
-               f"Invalid types, got {type(q)}, {type(k)}, and {type(v)}"
-        assert q.shape[0] == k.shape[0] == v.shape[0], f"Batch size mismatch, got {q.shape[0]}, {k.shape[0]}, and {v.shape[0]}"
+        q = args[0] if len(args) > 0 else kwargs["q"]
+        k = args[1] if len(args) > 1 else kwargs["k"]
+        v = args[2] if len(args) > 2 else kwargs["v"]
+        assert (
+            isinstance(q, SparseTensor)
+            and isinstance(k, (SparseTensor, torch.Tensor))
+            and type(k) == type(v)
+            or isinstance(q, torch.Tensor)
+            and isinstance(k, SparseTensor)
+            and isinstance(v, SparseTensor)
+        ), f"Invalid types, got {type(q)}, {type(k)}, and {type(v)}"
+        assert (
+            q.shape[0] == k.shape[0] == v.shape[0]
+        ), f"Batch size mismatch, got {q.shape[0]}, {k.shape[0]}, and {v.shape[0]}"
         device = q.device
 
         if isinstance(q, SparseTensor):
-            assert len(q.shape) == 3, f"Invalid shape for q, got {q.shape}, expected [N, *, H, Ci]"
+            assert (
+                len(q.shape) == 3
+            ), f"Invalid shape for q, got {q.shape}, expected [N, *, H, Ci]"
             s = q
             q_seqlen = [q.layout[i].stop - q.layout[i].start for i in range(q.shape[0])]
-            q = q.feats     # [T_Q, H, Ci]
+            q = q.feats  # [T_Q, H, Ci]
         else:
-            assert len(q.shape) == 4, f"Invalid shape for q, got {q.shape}, expected [N, L, H, Ci]"
+            assert (
+                len(q.shape) == 4
+            ), f"Invalid shape for q, got {q.shape}, expected [N, L, H, Ci]"
             s = None
             N, L, H, CI = q.shape
             q_seqlen = [L] * N
             q = q.reshape(N * L, H, CI)  # [T_Q, H, Ci]
 
         if isinstance(k, SparseTensor):
-            assert len(k.shape) == 3, f"Invalid shape for k, got {k.shape}, expected [N, *, H, Ci]"
-            assert len(v.shape) == 3, f"Invalid shape for v, got {v.shape}, expected [N, *, H, Co]"
-            kv_seqlen = [k.layout[i].stop - k.layout[i].start for i in range(k.shape[0])]
-            k = k.feats     # [T_KV, H, Ci]
-            v = v.feats     # [T_KV, H, Co]
+            assert (
+                len(k.shape) == 3
+            ), f"Invalid shape for k, got {k.shape}, expected [N, *, H, Ci]"
+            assert (
+                len(v.shape) == 3
+            ), f"Invalid shape for v, got {v.shape}, expected [N, *, H, Co]"
+            kv_seqlen = [
+                k.layout[i].stop - k.layout[i].start for i in range(k.shape[0])
+            ]
+            k = k.feats  # [T_KV, H, Ci]
+            v = v.feats  # [T_KV, H, Co]
         else:
-            assert len(k.shape) == 4, f"Invalid shape for k, got {k.shape}, expected [N, L, H, Ci]"
-            assert len(v.shape) == 4, f"Invalid shape for v, got {v.shape}, expected [N, L, H, Co]"
+            assert (
+                len(k.shape) == 4
+            ), f"Invalid shape for k, got {k.shape}, expected [N, L, H, Ci]"
+            assert (
+                len(v.shape) == 4
+            ), f"Invalid shape for v, got {v.shape}, expected [N, L, H, Co]"
             N, L, H, CI, CO = *k.shape, v.shape[-1]
             kv_seqlen = [L] * N
-            k = k.reshape(N * L, H, CI)     # [T_KV, H, Ci]
-            v = v.reshape(N * L, H, CO)     # [T_KV, H, Co]
+            k = k.reshape(N * L, H, CI)  # [T_KV, H, Ci]
+            v = v.reshape(N * L, H, CO)  # [T_KV, H, Co]
 
     if DEBUG:
         if s is not None:
             for i in range(s.shape[0]):
-                assert (s.coords[s.layout[i]] == i).all(), f"SparseScaledDotProductSelfAttention: batch index mismatch"
+                assert (
+                    s.coords[s.layout[i]] == i
+                ).all(), f"SparseScaledDotProductSelfAttention: batch index mismatch"
         if num_all_args in [2, 3]:
-            assert q.shape[:2] == [1, sum(q_seqlen)], f"SparseScaledDotProductSelfAttention: q shape mismatch"
+            assert q.shape[:2] == [
+                1,
+                sum(q_seqlen),
+            ], f"SparseScaledDotProductSelfAttention: q shape mismatch"
         if num_all_args == 3:
-            assert k.shape[:2] == [1, sum(kv_seqlen)], f"SparseScaledDotProductSelfAttention: k shape mismatch"
-            assert v.shape[:2] == [1, sum(kv_seqlen)], f"SparseScaledDotProductSelfAttention: v shape mismatch"
+            assert k.shape[:2] == [
+                1,
+                sum(kv_seqlen),
+            ], f"SparseScaledDotProductSelfAttention: k shape mismatch"
+            assert v.shape[:2] == [
+                1,
+                sum(kv_seqlen),
+            ], f"SparseScaledDotProductSelfAttention: v shape mismatch"
 
-    if ATTN == 'xformers':
+    if ATTN == "xformers":
         if num_all_args == 1:
             q, k, v = qkv.unbind(dim=1)
         elif num_all_args == 2:
@@ -196,19 +263,35 @@ def sparse_scaled_dot_product_attention(*args, **kwargs):
         v = v.unsqueeze(0)
         mask = xops.fmha.BlockDiagonalMask.from_seqlens(q_seqlen, kv_seqlen)
         out = xops.memory_efficient_attention(q, k, v, mask)[0]
-    elif ATTN == 'flash_attn':
-        cu_seqlens_q = torch.cat([torch.tensor([0]), torch.cumsum(torch.tensor(q_seqlen), dim=0)]).int().to(device)
+    elif ATTN == "flash_attn":
+        cu_seqlens_q = (
+            torch.cat([torch.tensor([0]), torch.cumsum(torch.tensor(q_seqlen), dim=0)])
+            .int()
+            .to(device)
+        )
         if num_all_args in [2, 3]:
-            cu_seqlens_kv = torch.cat([torch.tensor([0]), torch.cumsum(torch.tensor(kv_seqlen), dim=0)]).int().to(device)
+            cu_seqlens_kv = (
+                torch.cat(
+                    [torch.tensor([0]), torch.cumsum(torch.tensor(kv_seqlen), dim=0)]
+                )
+                .int()
+                .to(device)
+            )
         if num_all_args == 1:
-            out = flash_attn.flash_attn_varlen_qkvpacked_func(qkv, cu_seqlens_q, max(q_seqlen))
+            out = flash_attn.flash_attn_varlen_qkvpacked_func(
+                qkv, cu_seqlens_q, max(q_seqlen)
+            )
         elif num_all_args == 2:
-            out = flash_attn.flash_attn_varlen_kvpacked_func(q, kv, cu_seqlens_q, cu_seqlens_kv, max(q_seqlen), max(kv_seqlen))
+            out = flash_attn.flash_attn_varlen_kvpacked_func(
+                q, kv, cu_seqlens_q, cu_seqlens_kv, max(q_seqlen), max(kv_seqlen)
+            )
         elif num_all_args == 3:
-            out = flash_attn.flash_attn_varlen_func(q, k, v, cu_seqlens_q, cu_seqlens_kv, max(q_seqlen), max(kv_seqlen))
+            out = flash_attn.flash_attn_varlen_func(
+                q, k, v, cu_seqlens_q, cu_seqlens_kv, max(q_seqlen), max(kv_seqlen)
+            )
     else:
         raise ValueError(f"Unknown attention module: {ATTN}")
-    
+
     if s is not None:
         return s.replace(out)
     else:

trellis/modules/sparse/attention/modules.py

@@ -4,7 +4,10 @@ import torch.nn as nn
 import torch.nn.functional as F
 from .. import SparseTensor
 from .full_attn import sparse_scaled_dot_product_attention
-from .serialized_attn import SerializeMode, sparse_serialized_scaled_dot_product_self_attention
+from .serialized_attn import (
+    SerializeMode,
+    sparse_serialized_scaled_dot_product_self_attention,
+)
 from .windowed_attn import sparse_windowed_scaled_dot_product_self_attention
 from ...attention import RotaryPositionEmbedder
 
@@ -12,16 +15,18 @@ from ...attention import RotaryPositionEmbedder
 class SparseMultiHeadRMSNorm(nn.Module):
     def __init__(self, dim: int, heads: int):
         super().__init__()
-        self.scale = dim ** 0.5
+        self.scale = dim**0.5
         self.gamma = nn.Parameter(torch.ones(heads, dim))
 
-    def forward(self, x: Union[SparseTensor, torch.Tensor]) -> Union[SparseTensor, torch.Tensor]:
+    def forward(
+        self, x: Union[SparseTensor, torch.Tensor]
+    ) -> Union[SparseTensor, torch.Tensor]:
         x_type = x.dtype
         x = x.float()
         if isinstance(x, SparseTensor):
             x = x.replace(F.normalize(x.feats, dim=-1))
         else:
-            x = F.normalize(x, dim=-1)            
+            x = F.normalize(x, dim=-1)
         return (x * self.gamma * self.scale).to(x_type)
 
 
@@ -44,9 +49,17 @@ class SparseMultiHeadAttention(nn.Module):
         super().__init__()
         assert channels % num_heads == 0
         assert type in ["self", "cross"], f"Invalid attention type: {type}"
-        assert attn_mode in ["full", "serialized", "windowed"], f"Invalid attention mode: {attn_mode}"
-        assert type == "self" or attn_mode == "full", "Cross-attention only supports full attention"
-        assert type == "self" or use_rope is False, "Rotary position embeddings only supported for self-attention"
+        assert attn_mode in [
+            "full",
+            "serialized",
+            "windowed",
+        ], f"Invalid attention mode: {attn_mode}"
+        assert (
+            type == "self" or attn_mode == "full"
+        ), "Cross-attention only supports full attention"
+        assert (
+            type == "self" or use_rope is False
+        ), "Rotary position embeddings only supported for self-attention"
         self.channels = channels
         self.ctx_channels = ctx_channels if ctx_channels is not None else channels
         self.num_heads = num_heads
@@ -64,31 +77,37 @@ class SparseMultiHeadAttention(nn.Module):
         else:
             self.to_q = nn.Linear(channels, channels, bias=qkv_bias)
             self.to_kv = nn.Linear(self.ctx_channels, channels * 2, bias=qkv_bias)
-        
+
         if self.qk_rms_norm:
             self.q_rms_norm = SparseMultiHeadRMSNorm(channels // num_heads, num_heads)
             self.k_rms_norm = SparseMultiHeadRMSNorm(channels // num_heads, num_heads)
-            
+
         self.to_out = nn.Linear(channels, channels)
 
         if use_rope:
             self.rope = RotaryPositionEmbedder(channels)
 
     @staticmethod
-    def _linear(module: nn.Linear, x: Union[SparseTensor, torch.Tensor]) -> Union[SparseTensor, torch.Tensor]:
+    def _linear(
+        module: nn.Linear, x: Union[SparseTensor, torch.Tensor]
+    ) -> Union[SparseTensor, torch.Tensor]:
         if isinstance(x, SparseTensor):
             return x.replace(module(x.feats))
         else:
             return module(x)
 
     @staticmethod
-    def _reshape_chs(x: Union[SparseTensor, torch.Tensor], shape: Tuple[int, ...]) -> Union[SparseTensor, torch.Tensor]:
+    def _reshape_chs(
+        x: Union[SparseTensor, torch.Tensor], shape: Tuple[int, ...]
+    ) -> Union[SparseTensor, torch.Tensor]:
         if isinstance(x, SparseTensor):
             return x.reshape(*shape)
         else:
             return x.reshape(*x.shape[:2], *shape)
 
-    def _fused_pre(self, x: Union[SparseTensor, torch.Tensor], num_fused: int) -> Union[SparseTensor, torch.Tensor]:
+    def _fused_pre(
+        self, x: Union[SparseTensor, torch.Tensor], num_fused: int
+    ) -> Union[SparseTensor, torch.Tensor]:
         if isinstance(x, SparseTensor):
             x_feats = x.feats.unsqueeze(0)
         else:
@@ -97,12 +116,16 @@ class SparseMultiHeadAttention(nn.Module):
         return x.replace(x_feats.squeeze(0)) if isinstance(x, SparseTensor) else x_feats
 
     def _rope(self, qkv: SparseTensor) -> SparseTensor:
-        q, k, v = qkv.feats.unbind(dim=1)   # [T, H, C]
+        q, k, v = qkv.feats.unbind(dim=1)  # [T, H, C]
         q, k = self.rope(q, k, qkv.coords[:, 1:])
-        qkv = qkv.replace(torch.stack([q, k, v], dim=1)) 
+        qkv = qkv.replace(torch.stack([q, k, v], dim=1))
         return qkv
-    
-    def forward(self, x: Union[SparseTensor, torch.Tensor], context: Optional[Union[SparseTensor, torch.Tensor]] = None) -> Union[SparseTensor, torch.Tensor]:
+
+    def forward(
+        self,
+        x: Union[SparseTensor, torch.Tensor],
+        context: Optional[Union[SparseTensor, torch.Tensor]] = None,
+    ) -> Union[SparseTensor, torch.Tensor]:
         if self._type == "self":
             qkv = self._linear(self.to_qkv, x)
             qkv = self._fused_pre(qkv, num_fused=3)
@@ -117,7 +140,11 @@ class SparseMultiHeadAttention(nn.Module):
                 h = sparse_scaled_dot_product_attention(qkv)
             elif self.attn_mode == "serialized":
                 h = sparse_serialized_scaled_dot_product_self_attention(
-                    qkv, self.window_size, serialize_mode=self.serialize_mode, shift_sequence=self.shift_sequence, shift_window=self.shift_window
+                    qkv,
+                    self.window_size,
+                    serialize_mode=self.serialize_mode,
+                    shift_sequence=self.shift_sequence,
+                    shift_window=self.shift_window,
                 )
             elif self.attn_mode == "windowed":
                 h = sparse_windowed_scaled_dot_product_self_attention(

trellis/modules/sparse/attention/serialized_attn.py

@@ -5,16 +5,16 @@ import math
 from .. import SparseTensor
 from .. import DEBUG, ATTN
 
-if ATTN == 'xformers':
+if ATTN == "xformers":
     import xformers.ops as xops
-elif ATTN == 'flash_attn':
+elif ATTN == "flash_attn":
     import flash_attn
 else:
     raise ValueError(f"Unknown attention module: {ATTN}")
 
 
 __all__ = [
-    'sparse_serialized_scaled_dot_product_self_attention',
+    "sparse_serialized_scaled_dot_product_self_attention",
 ]
 
 
@@ -29,7 +29,7 @@ SerializeModes = [
     SerializeMode.Z_ORDER,
     SerializeMode.Z_ORDER_TRANSPOSED,
     SerializeMode.HILBERT,
-    SerializeMode.HILBERT_TRANSPOSED
+    SerializeMode.HILBERT_TRANSPOSED,
 ]
 
 
@@ -38,7 +38,7 @@ def calc_serialization(
     window_size: int,
     serialize_mode: SerializeMode = SerializeMode.Z_ORDER,
     shift_sequence: int = 0,
-    shift_window: Tuple[int, int, int] = (0, 0, 0)
+    shift_window: Tuple[int, int, int] = (0, 0, 0),
 ) -> Tuple[torch.Tensor, torch.Tensor, List[int]]:
     """
     Calculate serialization and partitioning for a set of coordinates.
@@ -58,32 +58,38 @@ def calc_serialization(
     seq_lens = []
     seq_batch_indices = []
     offsets = [0]
-    
-    if 'vox2seq' not in globals():
+
+    if "vox2seq" not in globals():
         import vox2seq
 
     # Serialize the input
     serialize_coords = tensor.coords[:, 1:].clone()
-    serialize_coords += torch.tensor(shift_window, dtype=torch.int32, device=tensor.device).reshape(1, 3)
+    serialize_coords += torch.tensor(
+        shift_window, dtype=torch.int32, device=tensor.device
+    ).reshape(1, 3)
     if serialize_mode == SerializeMode.Z_ORDER:
-        code = vox2seq.encode(serialize_coords, mode='z_order', permute=[0, 1, 2])
+        code = vox2seq.encode(serialize_coords, mode="z_order", permute=[0, 1, 2])
     elif serialize_mode == SerializeMode.Z_ORDER_TRANSPOSED:
-        code = vox2seq.encode(serialize_coords, mode='z_order', permute=[1, 0, 2])
+        code = vox2seq.encode(serialize_coords, mode="z_order", permute=[1, 0, 2])
     elif serialize_mode == SerializeMode.HILBERT:
-        code = vox2seq.encode(serialize_coords, mode='hilbert', permute=[0, 1, 2])
+        code = vox2seq.encode(serialize_coords, mode="hilbert", permute=[0, 1, 2])
     elif serialize_mode == SerializeMode.HILBERT_TRANSPOSED:
-        code = vox2seq.encode(serialize_coords, mode='hilbert', permute=[1, 0, 2])
+        code = vox2seq.encode(serialize_coords, mode="hilbert", permute=[1, 0, 2])
     else:
         raise ValueError(f"Unknown serialize mode: {serialize_mode}")
-    
+
     for bi, s in enumerate(tensor.layout):
         num_points = s.stop - s.start
         num_windows = (num_points + window_size - 1) // window_size
         valid_window_size = num_points / num_windows
-        to_ordered = torch.argsort(code[s.start:s.stop])
+        to_ordered = torch.argsort(code[s.start : s.stop])
         if num_windows == 1:
             fwd_indices.append(to_ordered)
-            bwd_indices.append(torch.zeros_like(to_ordered).scatter_(0, to_ordered, torch.arange(num_points, device=tensor.device)))
+            bwd_indices.append(
+                torch.zeros_like(to_ordered).scatter_(
+                    0, to_ordered, torch.arange(num_points, device=tensor.device)
+                )
+            )
             fwd_indices[-1] += s.start
             bwd_indices[-1] += offsets[-1]
             seq_lens.append(num_points)
@@ -92,18 +98,39 @@ def calc_serialization(
         else:
             # Partition the input
             offset = 0
-            mids = [(i + 0.5) * valid_window_size + shift_sequence for i in range(num_windows)]
-            split = [math.floor(i * valid_window_size + shift_sequence) for i in range(num_windows + 1)]
-            bwd_index = torch.zeros((num_points,), dtype=torch.int64, device=tensor.device)
+            mids = [
+                (i + 0.5) * valid_window_size + shift_sequence
+                for i in range(num_windows)
+            ]
+            split = [
+                math.floor(i * valid_window_size + shift_sequence)
+                for i in range(num_windows + 1)
+            ]
+            bwd_index = torch.zeros(
+                (num_points,), dtype=torch.int64, device=tensor.device
+            )
             for i in range(num_windows):
                 mid = mids[i]
                 valid_start = split[i]
                 valid_end = split[i + 1]
                 padded_start = math.floor(mid - 0.5 * window_size)
                 padded_end = padded_start + window_size
-                fwd_indices.append(to_ordered[torch.arange(padded_start, padded_end, device=tensor.device) % num_points])
+                fwd_indices.append(
+                    to_ordered[
+                        torch.arange(padded_start, padded_end, device=tensor.device)
+                        % num_points
+                    ]
+                )
                 offset += valid_start - padded_start
-                bwd_index.scatter_(0, fwd_indices[-1][valid_start-padded_start:valid_end-padded_start], torch.arange(offset, offset + valid_end - valid_start, device=tensor.device))
+                bwd_index.scatter_(
+                    0,
+                    fwd_indices[-1][
+                        valid_start - padded_start : valid_end - padded_start
+                    ],
+                    torch.arange(
+                        offset, offset + valid_end - valid_start, device=tensor.device
+                    ),
+                )
                 offset += padded_end - valid_start
                 fwd_indices[-1] += s.start
             seq_lens.extend([window_size] * num_windows)
@@ -115,14 +142,14 @@ def calc_serialization(
     bwd_indices = torch.cat(bwd_indices)
 
     return fwd_indices, bwd_indices, seq_lens, seq_batch_indices
-    
+
 
 def sparse_serialized_scaled_dot_product_self_attention(
     qkv: SparseTensor,
     window_size: int,
     serialize_mode: SerializeMode = SerializeMode.Z_ORDER,
     shift_sequence: int = 0,
-    shift_window: Tuple[int, int, int] = (0, 0, 0)
+    shift_window: Tuple[int, int, int] = (0, 0, 0),
 ) -> SparseTensor:
     """
     Apply serialized scaled dot product self attention to a sparse tensor.
@@ -135,59 +162,89 @@ def sparse_serialized_scaled_dot_product_self_attention(
         shift_window (Tuple[int, int, int]): The shift of serialized coordinates.
         shift (int): The shift to use.
     """
-    assert len(qkv.shape) == 4 and qkv.shape[1] == 3, f"Invalid shape for qkv, got {qkv.shape}, expected [N, *, 3, H, C]"
-
-    serialization_spatial_cache_name = f'serialization_{serialize_mode}_{window_size}_{shift_sequence}_{shift_window}'
-    serialization_spatial_cache = qkv.get_spatial_cache(serialization_spatial_cache_name)
+    assert (
+        len(qkv.shape) == 4 and qkv.shape[1] == 3
+    ), f"Invalid shape for qkv, got {qkv.shape}, expected [N, *, 3, H, C]"
+
+    serialization_spatial_cache_name = (
+        f"serialization_{serialize_mode}_{window_size}_{shift_sequence}_{shift_window}"
+    )
+    serialization_spatial_cache = qkv.get_spatial_cache(
+        serialization_spatial_cache_name
+    )
     if serialization_spatial_cache is None:
-        fwd_indices, bwd_indices, seq_lens, seq_batch_indices = calc_serialization(qkv, window_size, serialize_mode, shift_sequence, shift_window)
-        qkv.register_spatial_cache(serialization_spatial_cache_name, (fwd_indices, bwd_indices, seq_lens, seq_batch_indices))
+        fwd_indices, bwd_indices, seq_lens, seq_batch_indices = calc_serialization(
+            qkv, window_size, serialize_mode, shift_sequence, shift_window
+        )
+        qkv.register_spatial_cache(
+            serialization_spatial_cache_name,
+            (fwd_indices, bwd_indices, seq_lens, seq_batch_indices),
+        )
     else:
-        fwd_indices, bwd_indices, seq_lens, seq_batch_indices = serialization_spatial_cache
+        (
+            fwd_indices,
+            bwd_indices,
+            seq_lens,
+            seq_batch_indices,
+        ) = serialization_spatial_cache
 
     M = fwd_indices.shape[0]
     T = qkv.feats.shape[0]
     H = qkv.feats.shape[2]
     C = qkv.feats.shape[3]
-    
-    qkv_feats = qkv.feats[fwd_indices]      # [M, 3, H, C]
+
+    qkv_feats = qkv.feats[fwd_indices]  # [M, 3, H, C]
 
     if DEBUG:
         start = 0
         qkv_coords = qkv.coords[fwd_indices]
         for i in range(len(seq_lens)):
-            assert (qkv_coords[start:start+seq_lens[i], 0] == seq_batch_indices[i]).all(), f"SparseWindowedScaledDotProductSelfAttention: batch index mismatch"
+            assert (
+                qkv_coords[start : start + seq_lens[i], 0] == seq_batch_indices[i]
+            ).all(), (
+                f"SparseWindowedScaledDotProductSelfAttention: batch index mismatch"
+            )
             start += seq_lens[i]
 
     if all([seq_len == window_size for seq_len in seq_lens]):
         B = len(seq_lens)
         N = window_size
         qkv_feats = qkv_feats.reshape(B, N, 3, H, C)
-        if ATTN == 'xformers':
-            q, k, v = qkv_feats.unbind(dim=2)                       # [B, N, H, C]
-            out = xops.memory_efficient_attention(q, k, v)          # [B, N, H, C]
-        elif ATTN == 'flash_attn':
-            out = flash_attn.flash_attn_qkvpacked_func(qkv_feats)   # [B, N, H, C]
+        if ATTN == "xformers":
+            q, k, v = qkv_feats.unbind(dim=2)  # [B, N, H, C]
+            out = xops.memory_efficient_attention(q, k, v)  # [B, N, H, C]
+        elif ATTN == "flash_attn":
+            out = flash_attn.flash_attn_qkvpacked_func(qkv_feats)  # [B, N, H, C]
         else:
             raise ValueError(f"Unknown attention module: {ATTN}")
-        out = out.reshape(B * N, H, C)                              # [M, H, C]
+        out = out.reshape(B * N, H, C)  # [M, H, C]
     else:
-        if ATTN == 'xformers':
-            q, k, v = qkv_feats.unbind(dim=1)                       # [M, H, C]
-            q = q.unsqueeze(0)                                      # [1, M, H, C]
-            k = k.unsqueeze(0)                                      # [1, M, H, C]
-            v = v.unsqueeze(0)                                      # [1, M, H, C]
+        if ATTN == "xformers":
+            q, k, v = qkv_feats.unbind(dim=1)  # [M, H, C]
+            q = q.unsqueeze(0)  # [1, M, H, C]
+            k = k.unsqueeze(0)  # [1, M, H, C]
+            v = v.unsqueeze(0)  # [1, M, H, C]
             mask = xops.fmha.BlockDiagonalMask.from_seqlens(seq_lens)
-            out = xops.memory_efficient_attention(q, k, v, mask)[0] # [M, H, C]
-        elif ATTN == 'flash_attn':
-            cu_seqlens = torch.cat([torch.tensor([0]), torch.cumsum(torch.tensor(seq_lens), dim=0)], dim=0) \
-                        .to(qkv.device).int()
-            out = flash_attn.flash_attn_varlen_qkvpacked_func(qkv_feats, cu_seqlens, max(seq_lens)) # [M, H, C]
-
-    out = out[bwd_indices]      # [T, H, C]
+            out = xops.memory_efficient_attention(q, k, v, mask)[0]  # [M, H, C]
+        elif ATTN == "flash_attn":
+            cu_seqlens = (
+                torch.cat(
+                    [torch.tensor([0]), torch.cumsum(torch.tensor(seq_lens), dim=0)],
+                    dim=0,
+                )
+                .to(qkv.device)
+                .int()
+            )
+            out = flash_attn.flash_attn_varlen_qkvpacked_func(
+                qkv_feats, cu_seqlens, max(seq_lens)
+            )  # [M, H, C]
+
+    out = out[bwd_indices]  # [T, H, C]
 
     if DEBUG:
         qkv_coords = qkv_coords[bwd_indices]
-        assert torch.equal(qkv_coords, qkv.coords), "SparseWindowedScaledDotProductSelfAttention: coordinate mismatch"
+        assert torch.equal(
+            qkv_coords, qkv.coords
+        ), "SparseWindowedScaledDotProductSelfAttention: coordinate mismatch"
 
     return qkv.replace(out)

trellis/modules/sparse/attention/windowed_attn.py

@@ -4,23 +4,23 @@ import math
 from .. import SparseTensor
 from .. import DEBUG, ATTN
 
-if ATTN == 'xformers':
+if ATTN == "xformers":
     import xformers.ops as xops
-elif ATTN == 'flash_attn':
+elif ATTN == "flash_attn":
     import flash_attn
 else:
     raise ValueError(f"Unknown attention module: {ATTN}")
 
 
 __all__ = [
-    'sparse_windowed_scaled_dot_product_self_attention',
+    "sparse_windowed_scaled_dot_product_self_attention",
 ]
 
 
 def calc_window_partition(
     tensor: SparseTensor,
     window_size: Union[int, Tuple[int, ...]],
-    shift_window: Union[int, Tuple[int, ...]] = 0
+    shift_window: Union[int, Tuple[int, ...]] = 0,
 ) -> Tuple[torch.Tensor, torch.Tensor, List[int], List[int]]:
     """
     Calculate serialization and partitioning for a set of coordinates.
@@ -37,33 +37,43 @@ def calc_window_partition(
         (List[int]): Sequence batch indices.
     """
     DIM = tensor.coords.shape[1] - 1
-    shift_window = (shift_window,) * DIM if isinstance(shift_window, int) else shift_window
+    shift_window = (
+        (shift_window,) * DIM if isinstance(shift_window, int) else shift_window
+    )
     window_size = (window_size,) * DIM if isinstance(window_size, int) else window_size
     shifted_coords = tensor.coords.clone().detach()
-    shifted_coords[:, 1:] += torch.tensor(shift_window, device=tensor.device, dtype=torch.int32).unsqueeze(0)
+    shifted_coords[:, 1:] += torch.tensor(
+        shift_window, device=tensor.device, dtype=torch.int32
+    ).unsqueeze(0)
 
     MAX_COORDS = shifted_coords[:, 1:].max(dim=0).values.tolist()
     NUM_WINDOWS = [math.ceil((mc + 1) / ws) for mc, ws in zip(MAX_COORDS, window_size)]
     OFFSET = torch.cumprod(torch.tensor([1] + NUM_WINDOWS[::-1]), dim=0).tolist()[::-1]
 
-    shifted_coords[:, 1:] //= torch.tensor(window_size, device=tensor.device, dtype=torch.int32).unsqueeze(0)
-    shifted_indices = (shifted_coords * torch.tensor(OFFSET, device=tensor.device, dtype=torch.int32).unsqueeze(0)).sum(dim=1)
+    shifted_coords[:, 1:] //= torch.tensor(
+        window_size, device=tensor.device, dtype=torch.int32
+    ).unsqueeze(0)
+    shifted_indices = (
+        shifted_coords
+        * torch.tensor(OFFSET, device=tensor.device, dtype=torch.int32).unsqueeze(0)
+    ).sum(dim=1)
     fwd_indices = torch.argsort(shifted_indices)
     bwd_indices = torch.empty_like(fwd_indices)
     bwd_indices[fwd_indices] = torch.arange(fwd_indices.shape[0], device=tensor.device)
     seq_lens = torch.bincount(shifted_indices)
-    seq_batch_indices = torch.arange(seq_lens.shape[0], device=tensor.device, dtype=torch.int32) // OFFSET[0]
+    seq_batch_indices = (
+        torch.arange(seq_lens.shape[0], device=tensor.device, dtype=torch.int32)
+        // OFFSET[0]
+    )
     mask = seq_lens != 0
     seq_lens = seq_lens[mask].tolist()
     seq_batch_indices = seq_batch_indices[mask].tolist()
 
     return fwd_indices, bwd_indices, seq_lens, seq_batch_indices
-    
+
 
 def sparse_windowed_scaled_dot_product_self_attention(
-    qkv: SparseTensor,
-    window_size: int,
-    shift_window: Tuple[int, int, int] = (0, 0, 0)
+    qkv: SparseTensor, window_size: int, shift_window: Tuple[int, int, int] = (0, 0, 0)
 ) -> SparseTensor:
     """
     Apply windowed scaled dot product self attention to a sparse tensor.
@@ -74,62 +84,95 @@ def sparse_windowed_scaled_dot_product_self_attention(
         shift_window (Tuple[int, int, int]): The shift of serialized coordinates.
         shift (int): The shift to use.
     """
-    assert len(qkv.shape) == 4 and qkv.shape[1] == 3, f"Invalid shape for qkv, got {qkv.shape}, expected [N, *, 3, H, C]"
-
-    serialization_spatial_cache_name = f'window_partition_{window_size}_{shift_window}'
-    serialization_spatial_cache = qkv.get_spatial_cache(serialization_spatial_cache_name)
+    assert (
+        len(qkv.shape) == 4 and qkv.shape[1] == 3
+    ), f"Invalid shape for qkv, got {qkv.shape}, expected [N, *, 3, H, C]"
+
+    serialization_spatial_cache_name = f"window_partition_{window_size}_{shift_window}"
+    serialization_spatial_cache = qkv.get_spatial_cache(
+        serialization_spatial_cache_name
+    )
     if serialization_spatial_cache is None:
-        fwd_indices, bwd_indices, seq_lens, seq_batch_indices = calc_window_partition(qkv, window_size, shift_window)
-        qkv.register_spatial_cache(serialization_spatial_cache_name, (fwd_indices, bwd_indices, seq_lens, seq_batch_indices))
+        fwd_indices, bwd_indices, seq_lens, seq_batch_indices = calc_window_partition(
+            qkv, window_size, shift_window
+        )
+        qkv.register_spatial_cache(
+            serialization_spatial_cache_name,
+            (fwd_indices, bwd_indices, seq_lens, seq_batch_indices),
+        )
     else:
-        fwd_indices, bwd_indices, seq_lens, seq_batch_indices = serialization_spatial_cache
+        (
+            fwd_indices,
+            bwd_indices,
+            seq_lens,
+            seq_batch_indices,
+        ) = serialization_spatial_cache
 
     M = fwd_indices.shape[0]
     T = qkv.feats.shape[0]
     H = qkv.feats.shape[2]
     C = qkv.feats.shape[3]
-    
-    qkv_feats = qkv.feats[fwd_indices]      # [M, 3, H, C]
+
+    qkv_feats = qkv.feats[fwd_indices]  # [M, 3, H, C]
 
     if DEBUG:
         start = 0
         qkv_coords = qkv.coords[fwd_indices]
         for i in range(len(seq_lens)):
-            seq_coords = qkv_coords[start:start+seq_lens[i]]
-            assert (seq_coords[:, 0] == seq_batch_indices[i]).all(), f"SparseWindowedScaledDotProductSelfAttention: batch index mismatch"
-            assert (seq_coords[:, 1:].max(dim=0).values - seq_coords[:, 1:].min(dim=0).values < window_size).all(), \
-                    f"SparseWindowedScaledDotProductSelfAttention: window size exceeded"
+            seq_coords = qkv_coords[start : start + seq_lens[i]]
+            assert (
+                seq_coords[:, 0] == seq_batch_indices[i]
+            ).all(), (
+                f"SparseWindowedScaledDotProductSelfAttention: batch index mismatch"
+            )
+            assert (
+                seq_coords[:, 1:].max(dim=0).values
+                - seq_coords[:, 1:].min(dim=0).values
+                < window_size
+            ).all(), (
+                f"SparseWindowedScaledDotProductSelfAttention: window size exceeded"
+            )
             start += seq_lens[i]
 
     if all([seq_len == window_size for seq_len in seq_lens]):
         B = len(seq_lens)
         N = window_size
         qkv_feats = qkv_feats.reshape(B, N, 3, H, C)
-        if ATTN == 'xformers':
-            q, k, v = qkv_feats.unbind(dim=2)                       # [B, N, H, C]
-            out = xops.memory_efficient_attention(q, k, v)          # [B, N, H, C]
-        elif ATTN == 'flash_attn':
-            out = flash_attn.flash_attn_qkvpacked_func(qkv_feats)   # [B, N, H, C]
+        if ATTN == "xformers":
+            q, k, v = qkv_feats.unbind(dim=2)  # [B, N, H, C]
+            out = xops.memory_efficient_attention(q, k, v)  # [B, N, H, C]
+        elif ATTN == "flash_attn":
+            out = flash_attn.flash_attn_qkvpacked_func(qkv_feats)  # [B, N, H, C]
         else:
             raise ValueError(f"Unknown attention module: {ATTN}")
-        out = out.reshape(B * N, H, C)                              # [M, H, C]
+        out = out.reshape(B * N, H, C)  # [M, H, C]
     else:
-        if ATTN == 'xformers':
-            q, k, v = qkv_feats.unbind(dim=1)                       # [M, H, C]
-            q = q.unsqueeze(0)                                      # [1, M, H, C]
-            k = k.unsqueeze(0)                                      # [1, M, H, C]
-            v = v.unsqueeze(0)                                      # [1, M, H, C]
+        if ATTN == "xformers":
+            q, k, v = qkv_feats.unbind(dim=1)  # [M, H, C]
+            q = q.unsqueeze(0)  # [1, M, H, C]
+            k = k.unsqueeze(0)  # [1, M, H, C]
+            v = v.unsqueeze(0)  # [1, M, H, C]
             mask = xops.fmha.BlockDiagonalMask.from_seqlens(seq_lens)
-            out = xops.memory_efficient_attention(q, k, v, mask)[0] # [M, H, C]
-        elif ATTN == 'flash_attn':
-            cu_seqlens = torch.cat([torch.tensor([0]), torch.cumsum(torch.tensor(seq_lens), dim=0)], dim=0) \
-                        .to(qkv.device).int()
-            out = flash_attn.flash_attn_varlen_qkvpacked_func(qkv_feats, cu_seqlens, max(seq_lens)) # [M, H, C]
-
-    out = out[bwd_indices]      # [T, H, C]
+            out = xops.memory_efficient_attention(q, k, v, mask)[0]  # [M, H, C]
+        elif ATTN == "flash_attn":
+            cu_seqlens = (
+                torch.cat(
+                    [torch.tensor([0]), torch.cumsum(torch.tensor(seq_lens), dim=0)],
+                    dim=0,
+                )
+                .to(qkv.device)
+                .int()
+            )
+            out = flash_attn.flash_attn_varlen_qkvpacked_func(
+                qkv_feats, cu_seqlens, max(seq_lens)
+            )  # [M, H, C]
+
+    out = out[bwd_indices]  # [T, H, C]
 
     if DEBUG:
         qkv_coords = qkv_coords[bwd_indices]
-        assert torch.equal(qkv_coords, qkv.coords), "SparseWindowedScaledDotProductSelfAttention: coordinate mismatch"
+        assert torch.equal(
+            qkv_coords, qkv.coords
+        ), "SparseWindowedScaledDotProductSelfAttention: coordinate mismatch"
 
     return qkv.replace(out)

trellis/modules/sparse/basic.py

@@ -2,22 +2,23 @@ from typing import *
 import torch
 import torch.nn as nn
 from . import BACKEND, DEBUG
-SparseTensorData = None # Lazy import
+
+SparseTensorData = None  # Lazy import
 
 
 __all__ = [
-    'SparseTensor',
-    'sparse_batch_broadcast',
-    'sparse_batch_op',
-    'sparse_cat',
-    'sparse_unbind',
+    "SparseTensor",
+    "sparse_batch_broadcast",
+    "sparse_batch_op",
+    "sparse_cat",
+    "sparse_unbind",
 ]
 
 
 class SparseTensor:
     """
     Sparse tensor with support for both torchsparse and spconv backends.
-    
+
     Parameters:
     - feats (torch.Tensor): Features of the sparse tensor.
     - coords (torch.Tensor): Coordinates of the sparse tensor.
@@ -29,64 +30,89 @@ class SparseTensor:
     - Data corresponding to a same batch should be contiguous.
     - Coords should be in [0, 1023]
     """
+
     @overload
-    def __init__(self, feats: torch.Tensor, coords: torch.Tensor, shape: Optional[torch.Size] = None, layout: Optional[List[slice]] = None, **kwargs): ...
+    def __init__(
+        self,
+        feats: torch.Tensor,
+        coords: torch.Tensor,
+        shape: Optional[torch.Size] = None,
+        layout: Optional[List[slice]] = None,
+        **kwargs,
+    ):
+        ...
 
     @overload
-    def __init__(self, data, shape: Optional[torch.Size] = None, layout: Optional[List[slice]] = None, **kwargs): ...
+    def __init__(
+        self,
+        data,
+        shape: Optional[torch.Size] = None,
+        layout: Optional[List[slice]] = None,
+        **kwargs,
+    ):
+        ...
 
     def __init__(self, *args, **kwargs):
         # Lazy import of sparse tensor backend
         global SparseTensorData
         if SparseTensorData is None:
             import importlib
-            if BACKEND == 'torchsparse':
-                SparseTensorData = importlib.import_module('torchsparse').SparseTensor
-            elif BACKEND == 'spconv':
-                SparseTensorData = importlib.import_module('spconv.pytorch').SparseConvTensor
-                
+
+            if BACKEND == "torchsparse":
+                SparseTensorData = importlib.import_module("torchsparse").SparseTensor
+            elif BACKEND == "spconv":
+                SparseTensorData = importlib.import_module(
+                    "spconv.pytorch"
+                ).SparseConvTensor
+
         method_id = 0
         if len(args) != 0:
             method_id = 0 if isinstance(args[0], torch.Tensor) else 1
         else:
-            method_id = 1 if 'data' in kwargs else 0
+            method_id = 1 if "data" in kwargs else 0
 
         if method_id == 0:
             feats, coords, shape, layout = args + (None,) * (4 - len(args))
-            if 'feats' in kwargs:
-                feats = kwargs['feats']
-                del kwargs['feats']
-            if 'coords' in kwargs:
-                coords = kwargs['coords']
-                del kwargs['coords']
-            if 'shape' in kwargs:
-                shape = kwargs['shape']
-                del kwargs['shape']
-            if 'layout' in kwargs:
-                layout = kwargs['layout']
-                del kwargs['layout']
+            if "feats" in kwargs:
+                feats = kwargs["feats"]
+                del kwargs["feats"]
+            if "coords" in kwargs:
+                coords = kwargs["coords"]
+                del kwargs["coords"]
+            if "shape" in kwargs:
+                shape = kwargs["shape"]
+                del kwargs["shape"]
+            if "layout" in kwargs:
+                layout = kwargs["layout"]
+                del kwargs["layout"]
 
             if shape is None:
                 shape = self.__cal_shape(feats, coords)
             if layout is None:
                 layout = self.__cal_layout(coords, shape[0])
-            if BACKEND == 'torchsparse':
+            if BACKEND == "torchsparse":
                 self.data = SparseTensorData(feats, coords, **kwargs)
-            elif BACKEND == 'spconv':
+            elif BACKEND == "spconv":
                 spatial_shape = list(coords.max(0)[0] + 1)[1:]
-                self.data = SparseTensorData(feats.reshape(feats.shape[0], -1), coords, spatial_shape, shape[0], **kwargs)
+                self.data = SparseTensorData(
+                    feats.reshape(feats.shape[0], -1),
+                    coords,
+                    spatial_shape,
+                    shape[0],
+                    **kwargs,
+                )
                 self.data._features = feats
         elif method_id == 1:
             data, shape, layout = args + (None,) * (3 - len(args))
-            if 'data' in kwargs:
-                data = kwargs['data']
-                del kwargs['data']
-            if 'shape' in kwargs:
-                shape = kwargs['shape']
-                del kwargs['shape']
-            if 'layout' in kwargs:
-                layout = kwargs['layout']
-                del kwargs['layout']
+            if "data" in kwargs:
+                data = kwargs["data"]
+                del kwargs["data"]
+            if "shape" in kwargs:
+                shape = kwargs["shape"]
+                del kwargs["shape"]
+            if "layout" in kwargs:
+                layout = kwargs["layout"]
+                del kwargs["layout"]
 
             self.data = data
             if shape is None:
@@ -96,73 +122,84 @@ class SparseTensor:
 
         self._shape = shape
         self._layout = layout
-        self._scale = kwargs.get('scale', (1, 1, 1))
-        self._spatial_cache = kwargs.get('spatial_cache', {})
+        self._scale = kwargs.get("scale", (1, 1, 1))
+        self._spatial_cache = kwargs.get("spatial_cache", {})
 
         if DEBUG:
             try:
-                assert self.feats.shape[0] == self.coords.shape[0], f"Invalid feats shape: {self.feats.shape}, coords shape: {self.coords.shape}"
-                assert self.shape == self.__cal_shape(self.feats, self.coords), f"Invalid shape: {self.shape}"
-                assert self.layout == self.__cal_layout(self.coords, self.shape[0]), f"Invalid layout: {self.layout}"
+                assert (
+                    self.feats.shape[0] == self.coords.shape[0]
+                ), f"Invalid feats shape: {self.feats.shape}, coords shape: {self.coords.shape}"
+                assert self.shape == self.__cal_shape(
+                    self.feats, self.coords
+                ), f"Invalid shape: {self.shape}"
+                assert self.layout == self.__cal_layout(
+                    self.coords, self.shape[0]
+                ), f"Invalid layout: {self.layout}"
                 for i in range(self.shape[0]):
-                    assert torch.all(self.coords[self.layout[i], 0] == i), f"The data of batch {i} is not contiguous"
+                    assert torch.all(
+                        self.coords[self.layout[i], 0] == i
+                    ), f"The data of batch {i} is not contiguous"
             except Exception as e:
-                print('Debugging information:')
+                print("Debugging information:")
                 print(f"- Shape: {self.shape}")
                 print(f"- Layout: {self.layout}")
                 print(f"- Scale: {self._scale}")
                 print(f"- Coords: {self.coords}")
                 raise e
-        
+
     def __cal_shape(self, feats, coords):
         shape = []
         shape.append(coords[:, 0].max().item() + 1)
         shape.extend([*feats.shape[1:]])
         return torch.Size(shape)
-    
+
     def __cal_layout(self, coords, batch_size):
         seq_len = torch.bincount(coords[:, 0], minlength=batch_size)
-        offset = torch.cumsum(seq_len, dim=0) 
-        layout = [slice((offset[i] - seq_len[i]).item(), offset[i].item()) for i in range(batch_size)]
+        offset = torch.cumsum(seq_len, dim=0)
+        layout = [
+            slice((offset[i] - seq_len[i]).item(), offset[i].item())
+            for i in range(batch_size)
+        ]
         return layout
-    
+
     @property
     def shape(self) -> torch.Size:
         return self._shape
-    
+
     def dim(self) -> int:
         return len(self.shape)
-    
+
     @property
     def layout(self) -> List[slice]:
         return self._layout
 
     @property
     def feats(self) -> torch.Tensor:
-        if BACKEND == 'torchsparse':
+        if BACKEND == "torchsparse":
             return self.data.F
-        elif BACKEND == 'spconv':
+        elif BACKEND == "spconv":
             return self.data.features
-    
+
     @feats.setter
     def feats(self, value: torch.Tensor):
-        if BACKEND == 'torchsparse':
+        if BACKEND == "torchsparse":
             self.data.F = value
-        elif BACKEND == 'spconv':
+        elif BACKEND == "spconv":
             self.data.features = value
 
     @property
     def coords(self) -> torch.Tensor:
-        if BACKEND == 'torchsparse':
+        if BACKEND == "torchsparse":
             return self.data.C
-        elif BACKEND == 'spconv':
+        elif BACKEND == "spconv":
             return self.data.indices
-        
+
     @coords.setter
     def coords(self, value: torch.Tensor):
-        if BACKEND == 'torchsparse':
+        if BACKEND == "torchsparse":
             self.data.C = value
-        elif BACKEND == 'spconv':
+        elif BACKEND == "spconv":
             self.data.indices = value
 
     @property
@@ -174,12 +211,18 @@ class SparseTensor:
         return self.feats.device
 
     @overload
-    def to(self, dtype: torch.dtype) -> 'SparseTensor': ...
+    def to(self, dtype: torch.dtype) -> "SparseTensor":
+        ...
 
     @overload
-    def to(self, device: Optional[Union[str, torch.device]] = None, dtype: Optional[torch.dtype] = None) -> 'SparseTensor': ...
-
-    def to(self, *args, **kwargs) -> 'SparseTensor':
+    def to(
+        self,
+        device: Optional[Union[str, torch.device]] = None,
+        dtype: Optional[torch.dtype] = None,
+    ) -> "SparseTensor":
+        ...
+
+    def to(self, *args, **kwargs) -> "SparseTensor":
         device = None
         dtype = None
         if len(args) == 2:
@@ -189,13 +232,13 @@ class SparseTensor:
                 dtype = args[0]
             else:
                 device = args[0]
-        if 'dtype' in kwargs:
+        if "dtype" in kwargs:
             assert dtype is None, "to() received multiple values for argument 'dtype'"
-            dtype = kwargs['dtype']
-        if 'device' in kwargs:
+            dtype = kwargs["dtype"]
+        if "device" in kwargs:
             assert device is None, "to() received multiple values for argument 'device'"
-            device = kwargs['device']
-        
+            device = kwargs["device"]
+
         new_feats = self.feats.to(device=device, dtype=dtype)
         new_coords = self.coords.to(device=device)
         return self.replace(new_feats, new_coords)
@@ -204,46 +247,48 @@ class SparseTensor:
         new_feats = self.feats.type(dtype)
         return self.replace(new_feats)
 
-    def cpu(self) -> 'SparseTensor':
+    def cpu(self) -> "SparseTensor":
         new_feats = self.feats.cpu()
         new_coords = self.coords.cpu()
         return self.replace(new_feats, new_coords)
-    
-    def cuda(self) -> 'SparseTensor':
+
+    def cuda(self) -> "SparseTensor":
         new_feats = self.feats.cuda()
         new_coords = self.coords.cuda()
         return self.replace(new_feats, new_coords)
 
-    def half(self) -> 'SparseTensor':
+    def half(self) -> "SparseTensor":
         new_feats = self.feats.half()
         return self.replace(new_feats)
-    
-    def float(self) -> 'SparseTensor':
+
+    def float(self) -> "SparseTensor":
         new_feats = self.feats.float()
         return self.replace(new_feats)
-    
-    def detach(self) -> 'SparseTensor':
+
+    def detach(self) -> "SparseTensor":
         new_coords = self.coords.detach()
         new_feats = self.feats.detach()
         return self.replace(new_feats, new_coords)
 
     def dense(self) -> torch.Tensor:
-        if BACKEND == 'torchsparse':
+        if BACKEND == "torchsparse":
             return self.data.dense()
-        elif BACKEND == 'spconv':
+        elif BACKEND == "spconv":
             return self.data.dense()
 
-    def reshape(self, *shape) -> 'SparseTensor':
+    def reshape(self, *shape) -> "SparseTensor":
         new_feats = self.feats.reshape(self.feats.shape[0], *shape)
         return self.replace(new_feats)
-    
-    def unbind(self, dim: int) -> List['SparseTensor']:
+
+    def unbind(self, dim: int) -> List["SparseTensor"]:
         return sparse_unbind(self, dim)
 
-    def replace(self, feats: torch.Tensor, coords: Optional[torch.Tensor] = None) -> 'SparseTensor':
+    def replace(
+        self, feats: torch.Tensor, coords: Optional[torch.Tensor] = None
+    ) -> "SparseTensor":
         new_shape = [self.shape[0]]
         new_shape.extend(feats.shape[1:])
-        if BACKEND == 'torchsparse':
+        if BACKEND == "torchsparse":
             new_data = SparseTensorData(
                 feats=feats,
                 coords=self.data.coords if coords is None else coords,
@@ -251,7 +296,7 @@ class SparseTensor:
                 spatial_range=self.data.spatial_range,
             )
             new_data._caches = self.data._caches
-        elif BACKEND == 'spconv':
+        elif BACKEND == "spconv":
             new_data = SparseTensorData(
                 self.data.features.reshape(self.data.features.shape[0], -1),
                 self.data.indices,
@@ -259,7 +304,7 @@ class SparseTensor:
                 self.data.batch_size,
                 self.data.grid,
                 self.data.voxel_num,
-                self.data.indice_dict
+                self.data.indice_dict,
             )
             new_data._features = feats
             new_data.benchmark = self.data.benchmark
@@ -270,26 +315,39 @@ class SparseTensor:
             new_data.int8_scale = self.data.int8_scale
             if coords is not None:
                 new_data.indices = coords
-        new_tensor = SparseTensor(new_data, shape=torch.Size(new_shape), layout=self.layout, scale=self._scale, spatial_cache=self._spatial_cache)
+        new_tensor = SparseTensor(
+            new_data,
+            shape=torch.Size(new_shape),
+            layout=self.layout,
+            scale=self._scale,
+            spatial_cache=self._spatial_cache,
+        )
         return new_tensor
 
     @staticmethod
-    def full(aabb, dim, value, dtype=torch.float32, device=None) -> 'SparseTensor':
+    def full(aabb, dim, value, dtype=torch.float32, device=None) -> "SparseTensor":
         N, C = dim
         x = torch.arange(aabb[0], aabb[3] + 1)
         y = torch.arange(aabb[1], aabb[4] + 1)
         z = torch.arange(aabb[2], aabb[5] + 1)
-        coords = torch.stack(torch.meshgrid(x, y, z, indexing='ij'), dim=-1).reshape(-1, 3)
-        coords = torch.cat([
-            torch.arange(N).view(-1, 1).repeat(1, coords.shape[0]).view(-1, 1),
-            coords.repeat(N, 1),
-        ], dim=1).to(dtype=torch.int32, device=device)
+        coords = torch.stack(torch.meshgrid(x, y, z, indexing="ij"), dim=-1).reshape(
+            -1, 3
+        )
+        coords = torch.cat(
+            [
+                torch.arange(N).view(-1, 1).repeat(1, coords.shape[0]).view(-1, 1),
+                coords.repeat(N, 1),
+            ],
+            dim=1,
+        ).to(dtype=torch.int32, device=device)
         feats = torch.full((coords.shape[0], C), value, dtype=dtype, device=device)
         return SparseTensor(feats=feats, coords=coords)
 
-    def __merge_sparse_cache(self, other: 'SparseTensor') -> dict:
+    def __merge_sparse_cache(self, other: "SparseTensor") -> dict:
         new_cache = {}
-        for k in set(list(self._spatial_cache.keys()) + list(other._spatial_cache.keys())):
+        for k in set(
+            list(self._spatial_cache.keys()) + list(other._spatial_cache.keys())
+        ):
             if k in self._spatial_cache:
                 new_cache[k] = self._spatial_cache[k]
             if k in other._spatial_cache:
@@ -299,10 +357,12 @@ class SparseTensor:
                     new_cache[k].update(other._spatial_cache[k])
         return new_cache
 
-    def __neg__(self) -> 'SparseTensor':
+    def __neg__(self) -> "SparseTensor":
         return self.replace(-self.feats)
-    
-    def __elemwise__(self, other: Union[torch.Tensor, 'SparseTensor'], op: callable) -> 'SparseTensor':
+
+    def __elemwise__(
+        self, other: Union[torch.Tensor, "SparseTensor"], op: callable
+    ) -> "SparseTensor":
         if isinstance(other, torch.Tensor):
             try:
                 other = torch.broadcast_to(other, self.shape)
@@ -317,28 +377,44 @@ class SparseTensor:
             new_tensor._spatial_cache = self.__merge_sparse_cache(other)
         return new_tensor
 
-    def __add__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+    def __add__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, torch.add)
 
-    def __radd__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+    def __radd__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, torch.add)
-    
-    def __sub__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+
+    def __sub__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, torch.sub)
-    
-    def __rsub__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+
+    def __rsub__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, lambda x, y: torch.sub(y, x))
 
-    def __mul__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+    def __mul__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, torch.mul)
 
-    def __rmul__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+    def __rmul__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, torch.mul)
 
-    def __truediv__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+    def __truediv__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, torch.div)
 
-    def __rtruediv__(self, other: Union[torch.Tensor, 'SparseTensor', float]) -> 'SparseTensor':
+    def __rtruediv__(
+        self, other: Union[torch.Tensor, "SparseTensor", float]
+    ) -> "SparseTensor":
         return self.__elemwise__(other, lambda x, y: torch.div(y, x))
 
     def __getitem__(self, idx):
@@ -348,7 +424,9 @@ class SparseTensor:
             idx = range(*idx.indices(self.shape[0]))
         elif isinstance(idx, torch.Tensor):
             if idx.dtype == torch.bool:
-                assert idx.shape == (self.shape[0],), f"Invalid index shape: {idx.shape}"
+                assert idx.shape == (
+                    self.shape[0],
+                ), f"Invalid index shape: {idx.shape}"
                 idx = idx.nonzero().squeeze(1)
             elif idx.dtype in [torch.int32, torch.int64]:
                 assert len(idx.shape) == 1, f"Invalid index shape: {idx.shape}"
@@ -356,7 +434,7 @@ class SparseTensor:
                 raise ValueError(f"Unknown index type: {idx.dtype}")
         else:
             raise ValueError(f"Unknown index type: {type(idx)}")
-        
+
         coords = []
         feats = []
         for new_idx, old_idx in enumerate(idx):
@@ -392,7 +470,7 @@ class SparseTensor:
 def sparse_batch_broadcast(input: SparseTensor, other: torch.Tensor) -> torch.Tensor:
     """
     Broadcast a 1D tensor to a sparse tensor along the batch dimension then perform an operation.
-    
+
     Args:
         input (torch.Tensor): 1D tensor to broadcast.
         target (SparseTensor): Sparse tensor to broadcast to.
@@ -405,10 +483,12 @@ def sparse_batch_broadcast(input: SparseTensor, other: torch.Tensor) -> torch.Te
     return broadcasted
 
 
-def sparse_batch_op(input: SparseTensor, other: torch.Tensor, op: callable = torch.add) -> SparseTensor:
+def sparse_batch_op(
+    input: SparseTensor, other: torch.Tensor, op: callable = torch.add
+) -> SparseTensor:
     """
     Broadcast a 1D tensor to a sparse tensor along the batch dimension then perform an operation.
-    
+
     Args:
         input (torch.Tensor): 1D tensor to broadcast.
         target (SparseTensor): Sparse tensor to broadcast to.
@@ -420,7 +500,7 @@ def sparse_batch_op(input: SparseTensor, other: torch.Tensor, op: callable = tor
 def sparse_cat(inputs: List[SparseTensor], dim: int = 0) -> SparseTensor:
     """
     Concatenate a list of sparse tensors.
-    
+
     Args:
         inputs (List[SparseTensor]): List of sparse tensors to concatenate.
     """
@@ -447,7 +527,7 @@ def sparse_cat(inputs: List[SparseTensor], dim: int = 0) -> SparseTensor:
 def sparse_unbind(input: SparseTensor, dim: int) -> List[SparseTensor]:
     """
     Unbind a sparse tensor along a dimension.
-    
+
     Args:
         input (SparseTensor): Sparse tensor to unbind.
         dim (int): Dimension to unbind.

trellis/modules/sparse/conv/__init__.py

@@ -1,21 +1,26 @@
 from .. import BACKEND
 
 
-SPCONV_ALGO = 'auto'    # 'auto', 'implicit_gemm', 'native'
+SPCONV_ALGO = "auto"  # 'auto', 'implicit_gemm', 'native'
+
 
 def __from_env():
     import os
-        
+
     global SPCONV_ALGO
-    env_spconv_algo = os.environ.get('SPCONV_ALGO')
-    if env_spconv_algo is not None and env_spconv_algo in ['auto', 'implicit_gemm', 'native']:
+    env_spconv_algo = os.environ.get("SPCONV_ALGO")
+    if env_spconv_algo is not None and env_spconv_algo in [
+        "auto",
+        "implicit_gemm",
+        "native",
+    ]:
         SPCONV_ALGO = env_spconv_algo
     print(f"[SPARSE][CONV] spconv algo: {SPCONV_ALGO}")
-        
+
 
 __from_env()
 
-if BACKEND == 'torchsparse':
+if BACKEND == "torchsparse":
     from .conv_torchsparse import *
-elif BACKEND == 'spconv':
+elif BACKEND == "spconv":
     from .conv_spconv import *

trellis/modules/sparse/conv/conv_spconv.py

@@ -4,21 +4,54 @@ from .. import SparseTensor
 from .. import DEBUG
 from . import SPCONV_ALGO
 
+
 class SparseConv3d(nn.Module):
-    def __init__(self, in_channels, out_channels, kernel_size, stride=1, dilation=1, padding=None, bias=True, indice_key=None):
+    def __init__(
+        self,
+        in_channels,
+        out_channels,
+        kernel_size,
+        stride=1,
+        dilation=1,
+        padding=None,
+        bias=True,
+        indice_key=None,
+    ):
         super(SparseConv3d, self).__init__()
-        if 'spconv' not in globals():
+        if "spconv" not in globals():
             import spconv.pytorch as spconv
         algo = None
-        if SPCONV_ALGO == 'native':
+        if SPCONV_ALGO == "native":
             algo = spconv.ConvAlgo.Native
-        elif SPCONV_ALGO == 'implicit_gemm':
+        elif SPCONV_ALGO == "implicit_gemm":
             algo = spconv.ConvAlgo.MaskImplicitGemm
         if stride == 1 and (padding is None):
-            self.conv = spconv.SubMConv3d(in_channels, out_channels, kernel_size, dilation=dilation, bias=bias, indice_key=indice_key, algo=algo)
+            self.conv = spconv.SubMConv3d(
+                in_channels,
+                out_channels,
+                kernel_size,
+                dilation=dilation,
+                bias=bias,
+                indice_key=indice_key,
+                algo=algo,
+            )
         else:
-            self.conv = spconv.SparseConv3d(in_channels, out_channels, kernel_size, stride=stride, dilation=dilation, padding=padding, bias=bias, indice_key=indice_key, algo=algo)
-        self.stride = tuple(stride) if isinstance(stride, (list, tuple)) else (stride, stride, stride)
+            self.conv = spconv.SparseConv3d(
+                in_channels,
+                out_channels,
+                kernel_size,
+                stride=stride,
+                dilation=dilation,
+                padding=padding,
+                bias=bias,
+                indice_key=indice_key,
+                algo=algo,
+            )
+        self.stride = (
+            tuple(stride)
+            if isinstance(stride, (list, tuple))
+            else (stride, stride, stride)
+        )
         self.padding = padding
 
     def forward(self, x: SparseTensor) -> SparseTensor:
@@ -30,42 +63,65 @@ class SparseConv3d(nn.Module):
         if spatial_changed and (x.shape[0] != 1):
             # spconv was non-1 stride will break the contiguous of the output tensor, sort by the coords
             fwd = new_data.indices[:, 0].argsort()
-            bwd = torch.zeros_like(fwd).scatter_(0, fwd, torch.arange(fwd.shape[0], device=fwd.device))
+            bwd = torch.zeros_like(fwd).scatter_(
+                0, fwd, torch.arange(fwd.shape[0], device=fwd.device)
+            )
             sorted_feats = new_data.features[fwd]
             sorted_coords = new_data.indices[fwd]
             unsorted_data = new_data
             new_data = spconv.SparseConvTensor(sorted_feats, sorted_coords, unsorted_data.spatial_shape, unsorted_data.batch_size)  # type: ignore
 
         out = SparseTensor(
-            new_data, shape=torch.Size(new_shape), layout=new_layout,
+            new_data,
+            shape=torch.Size(new_shape),
+            layout=new_layout,
             scale=tuple([s * stride for s, stride in zip(x._scale, self.stride)]),
             spatial_cache=x._spatial_cache,
         )
 
         if spatial_changed and (x.shape[0] != 1):
-            out.register_spatial_cache(f'conv_{self.stride}_unsorted_data', unsorted_data)
-            out.register_spatial_cache(f'conv_{self.stride}_sort_bwd', bwd)
- 
+            out.register_spatial_cache(
+                f"conv_{self.stride}_unsorted_data", unsorted_data
+            )
+            out.register_spatial_cache(f"conv_{self.stride}_sort_bwd", bwd)
+
         return out
 
 
 class SparseInverseConv3d(nn.Module):
-    def __init__(self, in_channels, out_channels, kernel_size, stride=1, dilation=1, bias=True, indice_key=None):
+    def __init__(
+        self,
+        in_channels,
+        out_channels,
+        kernel_size,
+        stride=1,
+        dilation=1,
+        bias=True,
+        indice_key=None,
+    ):
         super(SparseInverseConv3d, self).__init__()
-        if 'spconv' not in globals():
+        if "spconv" not in globals():
             import spconv.pytorch as spconv
-        self.conv = spconv.SparseInverseConv3d(in_channels, out_channels, kernel_size, bias=bias, indice_key=indice_key)
-        self.stride = tuple(stride) if isinstance(stride, (list, tuple)) else (stride, stride, stride)
+        self.conv = spconv.SparseInverseConv3d(
+            in_channels, out_channels, kernel_size, bias=bias, indice_key=indice_key
+        )
+        self.stride = (
+            tuple(stride)
+            if isinstance(stride, (list, tuple))
+            else (stride, stride, stride)
+        )
 
     def forward(self, x: SparseTensor) -> SparseTensor:
         spatial_changed = any(s != 1 for s in self.stride)
         if spatial_changed:
             # recover the original spconv order
-            data = x.get_spatial_cache(f'conv_{self.stride}_unsorted_data')
-            bwd = x.get_spatial_cache(f'conv_{self.stride}_sort_bwd')
+            data = x.get_spatial_cache(f"conv_{self.stride}_unsorted_data")
+            bwd = x.get_spatial_cache(f"conv_{self.stride}_sort_bwd")
             data = data.replace_feature(x.feats[bwd])
             if DEBUG:
-                assert torch.equal(data.indices, x.coords[bwd]), 'Recover the original order failed'
+                assert torch.equal(
+                    data.indices, x.coords[bwd]
+                ), "Recover the original order failed"
         else:
             data = x.data
 
@@ -73,7 +129,9 @@ class SparseInverseConv3d(nn.Module):
         new_shape = [x.shape[0], self.conv.out_channels]
         new_layout = None if spatial_changed else x.layout
         out = SparseTensor(
-            new_data, shape=torch.Size(new_shape), layout=new_layout,
+            new_data,
+            shape=torch.Size(new_shape),
+            layout=new_layout,
             scale=tuple([s // stride for s, stride in zip(x._scale, self.stride)]),
             spatial_cache=x._spatial_cache,
         )

trellis/modules/sparse/conv/conv_torchsparse.py

@@ -4,35 +4,73 @@ from .. import SparseTensor
 
 
 class SparseConv3d(nn.Module):
-    def __init__(self, in_channels, out_channels, kernel_size, stride=1, dilation=1, bias=True, indice_key=None):
+    def __init__(
+        self,
+        in_channels,
+        out_channels,
+        kernel_size,
+        stride=1,
+        dilation=1,
+        bias=True,
+        indice_key=None,
+    ):
         super(SparseConv3d, self).__init__()
-        if 'torchsparse' not in globals():
+        if "torchsparse" not in globals():
             import torchsparse
-        self.conv = torchsparse.nn.Conv3d(in_channels, out_channels, kernel_size, stride, 0, dilation, bias)
+        self.conv = torchsparse.nn.Conv3d(
+            in_channels, out_channels, kernel_size, stride, 0, dilation, bias
+        )
 
     def forward(self, x: SparseTensor) -> SparseTensor:
         out = self.conv(x.data)
         new_shape = [x.shape[0], self.conv.out_channels]
-        out = SparseTensor(out, shape=torch.Size(new_shape), layout=x.layout if all(s == 1 for s in self.conv.stride) else None)
+        out = SparseTensor(
+            out,
+            shape=torch.Size(new_shape),
+            layout=x.layout if all(s == 1 for s in self.conv.stride) else None,
+        )
         out._spatial_cache = x._spatial_cache
-        out._scale = tuple([s * stride for s, stride in zip(x._scale, self.conv.stride)])
+        out._scale = tuple(
+            [s * stride for s, stride in zip(x._scale, self.conv.stride)]
+        )
         return out
 
 
 class SparseInverseConv3d(nn.Module):
-    def __init__(self, in_channels, out_channels, kernel_size, stride=1, dilation=1, bias=True, indice_key=None):
+    def __init__(
+        self,
+        in_channels,
+        out_channels,
+        kernel_size,
+        stride=1,
+        dilation=1,
+        bias=True,
+        indice_key=None,
+    ):
         super(SparseInverseConv3d, self).__init__()
-        if 'torchsparse' not in globals():
+        if "torchsparse" not in globals():
             import torchsparse
-        self.conv = torchsparse.nn.Conv3d(in_channels, out_channels, kernel_size, stride, 0, dilation, bias, transposed=True)
+        self.conv = torchsparse.nn.Conv3d(
+            in_channels,
+            out_channels,
+            kernel_size,
+            stride,
+            0,
+            dilation,
+            bias,
+            transposed=True,
+        )
 
     def forward(self, x: SparseTensor) -> SparseTensor:
-        out = self.conv(x.data)        
+        out = self.conv(x.data)
         new_shape = [x.shape[0], self.conv.out_channels]
-        out = SparseTensor(out, shape=torch.Size(new_shape), layout=x.layout if all(s == 1 for s in self.conv.stride) else None)
+        out = SparseTensor(
+            out,
+            shape=torch.Size(new_shape),
+            layout=x.layout if all(s == 1 for s in self.conv.stride) else None,
+        )
         out._spatial_cache = x._spatial_cache
-        out._scale = tuple([s // stride for s, stride in zip(x._scale, self.conv.stride)])
+        out._scale = tuple(
+            [s // stride for s, stride in zip(x._scale, self.conv.stride)]
+        )
         return out
-
-
-

trellis/modules/sparse/linear.py

@@ -2,9 +2,7 @@ import torch
 import torch.nn as nn
 from . import SparseTensor
 
-__all__ = [
-    'SparseLinear'
-]
+__all__ = ["SparseLinear"]
 
 
 class SparseLinear(nn.Linear):

trellis/modules/sparse/nonlinearity.py

@@ -2,18 +2,13 @@ import torch
 import torch.nn as nn
 from . import SparseTensor
 
-__all__ = [
-    'SparseReLU',
-    'SparseSiLU',
-    'SparseGELU',
-    'SparseActivation'
-]
+__all__ = ["SparseReLU", "SparseSiLU", "SparseGELU", "SparseActivation"]
 
 
 class SparseReLU(nn.ReLU):
     def forward(self, input: SparseTensor) -> SparseTensor:
         return input.replace(super().forward(input.feats))
-    
+
 
 class SparseSiLU(nn.SiLU):
     def forward(self, input: SparseTensor) -> SparseTensor:
@@ -32,4 +27,3 @@ class SparseActivation(nn.Module):
 
     def forward(self, input: SparseTensor) -> SparseTensor:
         return input.replace(self.activation(input.feats))
-    

trellis/modules/sparse/norm.py

@@ -4,10 +4,10 @@ from . import SparseTensor
 from . import DEBUG
 
 __all__ = [
-    'SparseGroupNorm',
-    'SparseLayerNorm',
-    'SparseGroupNorm32',
-    'SparseLayerNorm32',
+    "SparseGroupNorm",
+    "SparseLayerNorm",
+    "SparseGroupNorm32",
+    "SparseLayerNorm32",
 ]
 
 
@@ -19,7 +19,9 @@ class SparseGroupNorm(nn.GroupNorm):
         nfeats = torch.zeros_like(input.feats)
         for k in range(input.shape[0]):
             if DEBUG:
-                assert (input.coords[input.layout[k], 0] == k).all(), f"SparseGroupNorm: batch index mismatch"
+                assert (
+                    input.coords[input.layout[k], 0] == k
+                ).all(), f"SparseGroupNorm: batch index mismatch"
             bfeats = input.feats[input.layout[k]]
             bfeats = bfeats.permute(1, 0).reshape(1, input.shape[1], -1)
             bfeats = super().forward(bfeats)
@@ -47,12 +49,15 @@ class SparseGroupNorm32(SparseGroupNorm):
     """
     A GroupNorm layer that converts to float32 before the forward pass.
     """
+
     def forward(self, x: SparseTensor) -> SparseTensor:
         return super().forward(x.float()).type(x.dtype)
 
+
 class SparseLayerNorm32(SparseLayerNorm):
     """
     A LayerNorm layer that converts to float32 before the forward pass.
     """
+
     def forward(self, x: SparseTensor) -> SparseTensor:
         return super().forward(x.float()).type(x.dtype)

trellis/modules/sparse/spatial.py

@@ -3,11 +3,7 @@ import torch
 import torch.nn as nn
 from . import SparseTensor
 
-__all__ = [
-    'SparseDownsample',
-    'SparseUpsample',
-    'SparseSubdivide'
-]
+__all__ = ["SparseDownsample", "SparseUpsample", "SparseSubdivide"]
 
 
 class SparseDownsample(nn.Module):
@@ -15,6 +11,7 @@ class SparseDownsample(nn.Module):
     Downsample a sparse tensor by a factor of `factor`.
     Implemented as average pooling.
     """
+
     def __init__(self, factor: Union[int, Tuple[int, ...], List[int]]):
         super(SparseDownsample, self).__init__()
         self.factor = tuple(factor) if isinstance(factor, (list, tuple)) else factor
@@ -22,36 +19,47 @@ class SparseDownsample(nn.Module):
     def forward(self, input: SparseTensor) -> SparseTensor:
         DIM = input.coords.shape[-1] - 1
         factor = self.factor if isinstance(self.factor, tuple) else (self.factor,) * DIM
-        assert DIM == len(factor), 'Input coordinates must have the same dimension as the downsample factor.'
+        assert DIM == len(
+            factor
+        ), "Input coordinates must have the same dimension as the downsample factor."
 
         coord = list(input.coords.unbind(dim=-1))
         for i, f in enumerate(factor):
-            coord[i+1] = coord[i+1] // f
+            coord[i + 1] = coord[i + 1] // f
 
-        MAX = [coord[i+1].max().item() + 1 for i in range(DIM)]
+        MAX = [coord[i + 1].max().item() + 1 for i in range(DIM)]
         OFFSET = torch.cumprod(torch.tensor(MAX[::-1]), 0).tolist()[::-1] + [1]
         code = sum([c * o for c, o in zip(coord, OFFSET)])
         code, idx = code.unique(return_inverse=True)
 
         new_feats = torch.scatter_reduce(
-            torch.zeros(code.shape[0], input.feats.shape[1], device=input.feats.device, dtype=input.feats.dtype),
+            torch.zeros(
+                code.shape[0],
+                input.feats.shape[1],
+                device=input.feats.device,
+                dtype=input.feats.dtype,
+            ),
             dim=0,
             index=idx.unsqueeze(1).expand(-1, input.feats.shape[1]),
             src=input.feats,
-            reduce='mean'
+            reduce="mean",
         )
         new_coords = torch.stack(
-            [code // OFFSET[0]] +
-            [(code // OFFSET[i+1]) % MAX[i] for i in range(DIM)],
-            dim=-1
+            [code // OFFSET[0]]
+            + [(code // OFFSET[i + 1]) % MAX[i] for i in range(DIM)],
+            dim=-1,
+        )
+        out = SparseTensor(
+            new_feats,
+            new_coords,
+            input.shape,
         )
-        out = SparseTensor(new_feats, new_coords, input.shape,)
         out._scale = tuple([s // f for s, f in zip(input._scale, factor)])
         out._spatial_cache = input._spatial_cache
 
-        out.register_spatial_cache(f'upsample_{factor}_coords', input.coords)
-        out.register_spatial_cache(f'upsample_{factor}_layout', input.layout)
-        out.register_spatial_cache(f'upsample_{factor}_idx', idx)
+        out.register_spatial_cache(f"upsample_{factor}_coords", input.coords)
+        out.register_spatial_cache(f"upsample_{factor}_layout", input.layout)
+        out.register_spatial_cache(f"upsample_{factor}_idx", idx)
 
         return out
 
@@ -61,6 +69,7 @@ class SparseUpsample(nn.Module):
     Upsample a sparse tensor by a factor of `factor`.
     Implemented as nearest neighbor interpolation.
     """
+
     def __init__(self, factor: Union[int, Tuple[int, int, int], List[int]]):
         super(SparseUpsample, self).__init__()
         self.factor = tuple(factor) if isinstance(factor, (list, tuple)) else factor
@@ -68,24 +77,30 @@ class SparseUpsample(nn.Module):
     def forward(self, input: SparseTensor) -> SparseTensor:
         DIM = input.coords.shape[-1] - 1
         factor = self.factor if isinstance(self.factor, tuple) else (self.factor,) * DIM
-        assert DIM == len(factor), 'Input coordinates must have the same dimension as the upsample factor.'
+        assert DIM == len(
+            factor
+        ), "Input coordinates must have the same dimension as the upsample factor."
 
-        new_coords = input.get_spatial_cache(f'upsample_{factor}_coords')
-        new_layout = input.get_spatial_cache(f'upsample_{factor}_layout')
-        idx = input.get_spatial_cache(f'upsample_{factor}_idx')
+        new_coords = input.get_spatial_cache(f"upsample_{factor}_coords")
+        new_layout = input.get_spatial_cache(f"upsample_{factor}_layout")
+        idx = input.get_spatial_cache(f"upsample_{factor}_idx")
         if any([x is None for x in [new_coords, new_layout, idx]]):
-            raise ValueError('Upsample cache not found. SparseUpsample must be paired with SparseDownsample.')
+            raise ValueError(
+                "Upsample cache not found. SparseUpsample must be paired with SparseDownsample."
+            )
         new_feats = input.feats[idx]
         out = SparseTensor(new_feats, new_coords, input.shape, new_layout)
         out._scale = tuple([s * f for s, f in zip(input._scale, factor)])
         out._spatial_cache = input._spatial_cache
         return out
-    
+
+
 class SparseSubdivide(nn.Module):
     """
     Upsample a sparse tensor by a factor of `factor`.
     Implemented as nearest neighbor interpolation.
     """
+
     def __init__(self):
         super(SparseSubdivide, self).__init__()
 
@@ -96,15 +111,20 @@ class SparseSubdivide(nn.Module):
         n_coords = torch.nonzero(n_cube)
         n_coords = torch.cat([torch.zeros_like(n_coords[:, :1]), n_coords], dim=-1)
         factor = n_coords.shape[0]
-        assert factor == 2 ** DIM
+        assert factor == 2**DIM
         # print(n_coords.shape)
         new_coords = input.coords.clone()
         new_coords[:, 1:] *= 2
-        new_coords = new_coords.unsqueeze(1) + n_coords.unsqueeze(0).to(new_coords.dtype)
-        
-        new_feats = input.feats.unsqueeze(1).expand(input.feats.shape[0], factor, *input.feats.shape[1:])
-        out = SparseTensor(new_feats.flatten(0, 1), new_coords.flatten(0, 1), input.shape)
+        new_coords = new_coords.unsqueeze(1) + n_coords.unsqueeze(0).to(
+            new_coords.dtype
+        )
+
+        new_feats = input.feats.unsqueeze(1).expand(
+            input.feats.shape[0], factor, *input.feats.shape[1:]
+        )
+        out = SparseTensor(
+            new_feats.flatten(0, 1), new_coords.flatten(0, 1), input.shape
+        )
         out._scale = input._scale * 2
         out._spatial_cache = input._spatial_cache
         return out
-

trellis/modules/sparse/transformer/__init__.py

@@ -1,2 +1,2 @@
 from .blocks import *
-from .modulated import *
+from .modulated import *

trellis/modules/sparse/transformer/blocks.py

@@ -25,12 +25,15 @@ class SparseTransformerBlock(nn.Module):
     """
     Sparse Transformer block (MSA + FFN).
     """
+
     def __init__(
         self,
         channels: int,
         num_heads: int,
         mlp_ratio: float = 4.0,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "full",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "full",
         window_size: Optional[int] = None,
         shift_sequence: Optional[int] = None,
         shift_window: Optional[Tuple[int, int, int]] = None,
@@ -73,7 +76,9 @@ class SparseTransformerBlock(nn.Module):
 
     def forward(self, x: SparseTensor) -> SparseTensor:
         if self.use_checkpoint:
-            return torch.utils.checkpoint.checkpoint(self._forward, x, use_reentrant=False)
+            return torch.utils.checkpoint.checkpoint(
+                self._forward, x, use_reentrant=False
+            )
         else:
             return self._forward(x)
 
@@ -82,13 +87,16 @@ class SparseTransformerCrossBlock(nn.Module):
     """
     Sparse Transformer cross-attention block (MSA + MCA + FFN).
     """
+
     def __init__(
         self,
         channels: int,
         ctx_channels: int,
         num_heads: int,
         mlp_ratio: float = 4.0,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "full",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "full",
         window_size: Optional[int] = None,
         shift_sequence: Optional[int] = None,
         shift_window: Optional[Tuple[int, int, int]] = None,
@@ -146,6 +154,8 @@ class SparseTransformerCrossBlock(nn.Module):
 
     def forward(self, x: SparseTensor, context: torch.Tensor):
         if self.use_checkpoint:
-            return torch.utils.checkpoint.checkpoint(self._forward, x, context, use_reentrant=False)
+            return torch.utils.checkpoint.checkpoint(
+                self._forward, x, context, use_reentrant=False
+            )
         else:
             return self._forward(x, context)

trellis/modules/sparse/transformer/modulated.py

@@ -11,12 +11,15 @@ class ModulatedSparseTransformerBlock(nn.Module):
     """
     Sparse Transformer block (MSA + FFN) with adaptive layer norm conditioning.
     """
+
     def __init__(
         self,
         channels: int,
         num_heads: int,
         mlp_ratio: float = 4.0,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "full",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "full",
         window_size: Optional[int] = None,
         shift_sequence: Optional[int] = None,
         shift_window: Optional[Tuple[int, int, int]] = None,
@@ -50,15 +53,23 @@ class ModulatedSparseTransformerBlock(nn.Module):
         )
         if not share_mod:
             self.adaLN_modulation = nn.Sequential(
-                nn.SiLU(),
-                nn.Linear(channels, 6 * channels, bias=True)
+                nn.SiLU(), nn.Linear(channels, 6 * channels, bias=True)
             )
 
     def _forward(self, x: SparseTensor, mod: torch.Tensor) -> SparseTensor:
         if self.share_mod:
-            shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = mod.chunk(6, dim=1)
+            shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = mod.chunk(
+                6, dim=1
+            )
         else:
-            shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = self.adaLN_modulation(mod).chunk(6, dim=1)
+            (
+                shift_msa,
+                scale_msa,
+                gate_msa,
+                shift_mlp,
+                scale_mlp,
+                gate_mlp,
+            ) = self.adaLN_modulation(mod).chunk(6, dim=1)
         h = x.replace(self.norm1(x.feats))
         h = h * (1 + scale_msa) + shift_msa
         h = self.attn(h)
@@ -73,7 +84,9 @@ class ModulatedSparseTransformerBlock(nn.Module):
 
     def forward(self, x: SparseTensor, mod: torch.Tensor) -> SparseTensor:
         if self.use_checkpoint:
-            return torch.utils.checkpoint.checkpoint(self._forward, x, mod, use_reentrant=False)
+            return torch.utils.checkpoint.checkpoint(
+                self._forward, x, mod, use_reentrant=False
+            )
         else:
             return self._forward(x, mod)
 
@@ -82,13 +95,16 @@ class ModulatedSparseTransformerCrossBlock(nn.Module):
     """
     Sparse Transformer cross-attention block (MSA + MCA + FFN) with adaptive layer norm conditioning.
     """
+
     def __init__(
         self,
         channels: int,
         ctx_channels: int,
         num_heads: int,
         mlp_ratio: float = 4.0,
-        attn_mode: Literal["full", "shift_window", "shift_sequence", "shift_order", "swin"] = "full",
+        attn_mode: Literal[
+            "full", "shift_window", "shift_sequence", "shift_order", "swin"
+        ] = "full",
         window_size: Optional[int] = None,
         shift_sequence: Optional[int] = None,
         shift_window: Optional[Tuple[int, int, int]] = None,
@@ -99,7 +115,6 @@ class ModulatedSparseTransformerCrossBlock(nn.Module):
         qk_rms_norm_cross: bool = False,
         qkv_bias: bool = True,
         share_mod: bool = False,
-
     ):
         super().__init__()
         self.use_checkpoint = use_checkpoint
@@ -135,15 +150,25 @@ class ModulatedSparseTransformerCrossBlock(nn.Module):
         )
         if not share_mod:
             self.adaLN_modulation = nn.Sequential(
-                nn.SiLU(),
-                nn.Linear(channels, 6 * channels, bias=True)
+                nn.SiLU(), nn.Linear(channels, 6 * channels, bias=True)
             )
 
-    def _forward(self, x: SparseTensor, mod: torch.Tensor, context: torch.Tensor) -> SparseTensor:
+    def _forward(
+        self, x: SparseTensor, mod: torch.Tensor, context: torch.Tensor
+    ) -> SparseTensor:
         if self.share_mod:
-            shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = mod.chunk(6, dim=1)
+            shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = mod.chunk(
+                6, dim=1
+            )
         else:
-            shift_msa, scale_msa, gate_msa, shift_mlp, scale_mlp, gate_mlp = self.adaLN_modulation(mod).chunk(6, dim=1)
+            (
+                shift_msa,
+                scale_msa,
+                gate_msa,
+                shift_mlp,
+                scale_mlp,
+                gate_mlp,
+            ) = self.adaLN_modulation(mod).chunk(6, dim=1)
         h = x.replace(self.norm1(x.feats))
         h = h * (1 + scale_msa) + shift_msa
         h = self.self_attn(h)
@@ -159,8 +184,12 @@ class ModulatedSparseTransformerCrossBlock(nn.Module):
         x = x + h
         return x
 
-    def forward(self, x: SparseTensor, mod: torch.Tensor, context: torch.Tensor) -> SparseTensor:
+    def forward(
+        self, x: SparseTensor, mod: torch.Tensor, context: torch.Tensor
+    ) -> SparseTensor:
         if self.use_checkpoint:
-            return torch.utils.checkpoint.checkpoint(self._forward, x, mod, context, use_reentrant=False)
+            return torch.utils.checkpoint.checkpoint(
+                self._forward, x, mod, context, use_reentrant=False
+            )
         else:
             return self._forward(x, mod, context)