CuACD

Read the paper

Context

• Basically same search-based algorithm as CoACD except fully on GPU, more candidate planes give quality improvement while GPU gives speed improvement
• Problem before is that cut + hull + concavity scoring runs hundreds of thousands of times per mesh, only on CPU. VisACD was previous attempt to move some ops onto GPU

GPU defs

• SIMT: 32 lanes share one instruction pointer and step together / warp (threads/blocks are virtual and get regrouped into warps by the runtime)
• Divergence: lanes in a warp take different branches so the warp runs each branch serially with the other lanes idle (e.g. if else)
• Warp intrinsics run in one instruction with no memory traffic: __ballot_sync (each lane votes a bool → 32-bit mask), __popc (count bits), __shfl_sync (read another lane’s register)
• Shared memory is small fast per-block scratch. Registers hold each warp’s state & more state per warp means fewer warps fit at once (occupancy)
• Little’s law: throughput = in-flight work / latency, so you need enough concurrent work to hide latency

Per-phase CUDA decomposition

• Usually we split the algorithm into steps, run each as a separate kernel over the whole GPU, one after another (~7 steps per cut for the clipper)
• Problem is that near the end of a step only a few threads are still working, but the next step can’t start until all finish so most of the GPU idles. Also each step’s output size (e.g. # triangles the plane crossed) is only known after it runs, so the CPU has to check it and set up memory before launching the next step × hundreds of thousands of cuts. Note that CUDA Graphs (pre-recorded launches) don’t help since sizes change every cut
• CuACD runs all steps inside one warp with the device-side heap handling memory

Warps

Idea: design the algorithm around warps (groups of 32 lanes) instead of individual threads, Hong et al. 2011 (graph algorithms)

There are two levels of parallelism:

• Across warps: each warp gets its own independent job (e.g. one piece of the mesh). A GPU only needs ~100 busy warps to be full, and ACD always has 100–1000 pieces to work on, so the GPU stays busy
• Inside a warp: the 32 lanes split up the work of that one job (e.g. each lane handles a different set of triangles) Since a whole job stays inside one warp, moving to the next step is just the next line of code, not a new kernel $⟹$ no waiting on other warps, no CPU check-ins

Everything is built from 4 basic patterns (ordered from most to least GPU-friendly):

1. Go through an array together: lane 0 takes items 0, 32, 64…, lane 1 takes 1, 33, 65…. Used to transform every item, add things up, or keep only items that pass a test
2. Sorting, removing duplicates, checking if something is in a set (all built on a warp sort)
3. Paired work: split the warp into 16 pairs of lanes, each pair handles two related things (e.g. both ends of an edge)
4. Divide and conquer: each lane takes a sub-problem, results get merged. Least preferred bc lanes may branch differently/diverge

Note: small bookkeeping (e.g. allocating memory) is just done by lane 0 alone

Warp sort: quicksort done by 1 warp. Splitting around the pivot is pattern 1 (keep items < pivot, keep items > pivot). Tiny pieces (≤32 items) get sorted directly. Uses almost no shared memory, so it handles big inputs where NVIDIA’s library sort (CUB) runs out of memory

Device-side heap allocator*

CUDA has a GPU-side malloc, but every warp goes through one shared lock, so warps wait in line and bottlenecks

We grab 70% of GPU memory once at startup and hand out pieces.

Memory is divided into 64 separate pools, each warp uses pool (warp id mod 64). Free blocks are sorted into bins by size & a 64-bit number marks which bins have something free $\to$ finding a big-enough block is one instruction. Freed blocks merge with free neighbors so memory doesn’t get chopped into useless tiny pieces. Only lane 0 of a warp talks to the allocator.

Result: 5× faster than CUDA’s malloc, 10–11× when many warps share a pool.

Pipeline

Note: same overall loop, concavity metric, and hull-merge step as CoACD

• More candidate planes: on top of CoACD’s axis-aligned planes, adds planes through sharp inward creases (“concave edges”, from Thul et al. 2018). 94 candidates vs CoACD’s 60. Ablation shows that this causes the quality gain
• Simple look-ahead search instead of MCTS: try each candidate, then look 2 cuts ahead (5 options each), pick the best path. Many pieces and many candidates are evaluated at once
• Connected-component splitting done with a version of union-find that needs no locks
• Cutting a mesh with a plane (clipper): find which triangles cross the plane, split them, then fill in the hole on the cut face. Filling the hole works by repeatedly clipping off triangle “ears”, with all 32 lanes checking at once whether an ear is valid
• Convex hull: split the points into 16 groups, each lane pair hulls one group, then merge pairwise. For big inputs, first throw away points that are obviously inside (~70% of them). >300× faster than the CoACD CPU hull
• Hausdorff distance (used in the concavity score): one warp per piece, ~29–50× faster than CoACD’s CPU version

Baseline

CoACD, NavACD, VisACD (thresholds swept to match mean concavity 0.05)

Results

Table 1. Concavity, number of parts, and decomposition time across three benchmarks, averaged over the per-mesh results. All baselines are operated at the threshold sweep that targets CoACD’s default mean concavity (0.05); residual concavity differences reflect the granularity of each system’s threshold knob rather than a quality gap. Bold marks the best entry per column; lower is better on all three axes.

Table 2. Convex hull construction time on batches of 256 independent point sets, in milliseconds. Within each row, bold marks the faster GPU method and underline the second.

Future directions

• Hand-designed heuristics, no learned proposals like RL-ACD (unfort no code released)
  • TODO reproduce results of RL-ACD & link
  • RL-ACD’s training bottleneck is exact cuts for the replay buffer & CuACD’s GPU components make that cheap
  • Could use a Q-net to prune/rank the candidate pool before the warp look-ahead -ablations show the pool matters and search depth doesn’t
• Heap never compacts or returns memory, only 25–50% occupancy due to register pressure
  • Authors’ fix: combine the warp allocator with modern memory management + register-aware scheduling
• Speed (~0.2s/mesh) opens up interactive use: ACD inside authoring tools, or at runtime (e.g. dynamic fracture, cf. Muller et al. 2013)
• Components (heap, warp sort, clipper, hull, Hausdorff) are drop-in CUDA modules, so other search-based geometry pipelines could reuse them