Warp at the Geometry BoundaryNative NVIDIA Warp Kernel Execution in Houdini, and Where Simulation Time Actually Goes
GPU kernel languages have made the compute half of procedural simulation close to free. We built a compiled Houdini operator that executes NVIDIA Warp kernels directly on geometry, then measured where the time goes. Across every workload we tested, the kernel was never the cost. The boundary between the host application and the compute layer was, by an order of magnitude or more, and it is the part almost nobody instruments.
Context
A shift is underway in how simulation code gets written. Kernel languages like NVIDIA Warp compile ordinary Python into native CPU and CUDA code, which means the arithmetic that used to demand C++ and a build system is now something a technical artist can write in an afternoon. Warp is also differentiable, which is why it turns up in machine-learning-adjacent simulation research rather than only in graphics.
That is a genuinely new capability, and it is arriving fastest in research environments: notebooks, standalone Python, training loops. It is arriving slowest in the places where film and commercial visual effects are actually authored. Houdini is where the industry builds procedural simulation, and the path from a Warp kernel to a Houdini scene has until now run through a Python SOP, a scripted node that reimports your modules on every cook, carries Python startup cost, and leaves you marshalling geometry by hand at both ends.
We are interested in that gap for a specific reason. A technology is only as useful as the number of people who can reach it, and the people who most need fast custom simulation are the ones least likely to write an HDK plugin to get it. So we built the bridge, published it, and more usefully, measured it.
The measurement is the part we think is worth reading. We expected to be reporting on kernel speed. We ended up reporting on something else entirely.
Why This Matters
If you don't write simulation code, here is the short version.
Most interesting behaviour in a shot is arithmetic repeated across hundreds of thousands of points: cloth settling, a surface reacting, a pattern spreading across a mesh. Houdini is excellent at this when the operation is one it already ships. When it isn't, an artist either builds the effect out of nodes designed for something else, or drops into Python and waits most of a minute per frame.
The new kernel languages remove that wait in principle. What we found is that removing it in practice depends far less on the maths than on the plumbing around it, on how geometry gets handed back and forth between the application and the compute layer. Get that wrong and a fast kernel still gives you a slow frame.
This is the sort of thing worth publishing rather than keeping. The tool does not decide what the effect should look like; it shortens the distance between an artist deciding and seeing, which means more attempts inside the same schedule. The judgment stays with the artists and supervisors making it. We would rather the whole industry got the faster loop than that we alone did.
What We Built
A compiled Houdini SOP that executes Warp kernels against the geometry flowing through it.
The design rests on three decisions, each of which we can now defend with numbers rather than taste.
It is a real HDK operator, not a scripted node. Every parameter is a PRM_Template. This is not an aesthetic preference: a scripted node pays Python startup on the cook path and acquires hou.session dependencies that make it fragile to ship. A compiled operator pays neither.
The geometry is already in scope. The kernel author gets geo (read-only input), npoints, and point_attrib_array(name), a float32 numpy view of a point attribute, without writing any marshalling code. This is the decision the results turned out to be about.
The authoring surface mirrors Warp's own model rather than inventing one. Kernel code and host code are separate fields, because that is the separation Warp itself draws:
# Kernel: device code
@wp.kernel
def measure(p: wp.array(dtype=wp.vec3), out: wp.array(dtype=float)):
i = wp.tid()
out[i] = wp.length(p[i])
# Host: runs once per cook, geometry already in scope
P = point_attrib_array("P")
with wp.ScopedDevice(device):
points = wp.array(P, dtype=wp.vec3)
result = wp.zeros(npoints, dtype=float)
wp.launch(measure, dim=npoints, inputs=[points], outputs=[result])
out = result
One deliberately unforgiving detail: leaving out unassigned is an error, not a no-op. A node that cooks cleanly and silently writes nothing is a failure that survives all the way to a render, and we would rather it surface at the point of the typo.
Method
The benchmark workload is a Gray-Scott reaction-diffusion step, chosen because it is a real simulation with a real neighbour dependency rather than a trivially parallel per-point operation that would flatter any GPU path.
Measurements were taken on a 300x300 grid: 90,000 points, 89,401 primitives, one diffusion step per cook. A second set of measurements on a 150x150 grid isolates the topology cost separately from the simulation cost.
The comparison is internally controlled. Every configuration ran the same arithmetic, on the same geometry, on the same machine, in the same session, and only the route between Houdini's geometry and the kernel differed. The result reported here is therefore the ratio between routes, which is the quantity that survives a change of hardware. Treat the absolute millisecond figures as single-machine and indicative; they are there to give the ratios a sense of scale, not to rank hardware. All figures are CPU-side, for the reason given under Limitations and Validity.
Results
Per-cook time on the 300x300 grid:
| Route | Cook time |
|---|---|
| Python SOP performing identical work | 155.6 ms |
Compiled node, geometry via houdiniGeo | 145.9 ms |
Compiled node, geometry via geo and out | 96.0 ms |
The headline comparison, 155.6 ms against 96.0 ms, is a 1.6x improvement, and on its own it is the least interesting number here.
The interesting number is the 49.9 ms separating the node's own two routes. Both run the same kernel. Both produce the same frame. The only difference is that one copies the input geometry so it can be written through the full HOM API, and the other reads attribute buffers directly and returns arrays. A third of the frame time was being spent on the trip, not the work.
Two further measurements make the same point more sharply.
Topology reconstruction. On the smaller 150x150 grid, building a neighbour graph in Python cost approximately 618 ms per frame against a 24 ms simulation, a 25:1 ratio of bookkeeping to arithmetic. Caching that graph across cooks, keyed on topology-derived values so it rebuilds when the geometry actually changes, brought the frame to roughly 7 ms. That is an improvement of about 88x in which the kernel was never modified.
Attribute access granularity. Moving whole buffers (point_attrib_array, setPointFloatAttribValuesFromString) against the equivalent per-element Python list calls is roughly an 18x difference for identical results.
Finding
Across every configuration we measured, the compute was not the constraint. The host boundary was.
This is worth stating plainly because it inverts the intuition the current wave of GPU tooling encourages. The pitch for kernel languages is speed of arithmetic, and that pitch is honest, because Warp delivers it. But once arithmetic is cheap, it stops being the term that dominates. What dominates is everything the arithmetic is wrapped in: the copy, the graph rebuild, the per-element access, the crossing.
Three practical consequences follow, and they generalise well beyond this node.
Make the zero-copy path the default one. If the fast route and the convenient route differ, authors will take the convenient one, and a third of the frame goes with it. Our node makes the copy lazy, so it happens only if writable geometry is actually referenced, and the cheap path is what you get by not asking for anything.
Give the tool somewhere to keep things. The single largest win we measured, by a wide margin, came from persisting topology-derived data across cooks. Any tool in this class needs an explicit place for that, or every author reinvents it badly.
Instrument the boundary, not the kernel. A profiler pointed at the kernel will report that everything is fine while the frame takes 618 ms.
None of this argues against fast kernels. It argues that reporting kernel speed alone, which is how most of this technology is currently benchmarked, measures the term that already stopped mattering.
Limitations and Validity
Stated plainly, because a result is only as good as the conditions around it.
- The CUDA path is implemented but unvalidated. The development machine's Warp build reports
CUDA not enabled, so every figure above is CPU-side. We expect the boundary cost to matter more on GPU rather than less, since device transfer joins the crossing, but we have not measured it and we are not going to claim it. - One workload family. Reaction-diffusion is representative of neighbour-dependent grid simulation. It is not representative of sparse, adaptive, or heavily branching workloads.
- Single-machine figures. The ratios are internally controlled, as described under Method, but they have not been reproduced across a range of hardware.
- One Houdini version. Only the
houdini-20.0variant is built and tested. Houdini 19.5 ships Python 3.9, which cannot resolve a Python 3.10warp_lang. - Structural constraints of the current node. A single input, and
outwrites float and vector point attributes only, with other classes and types going through the writable-geometry route. - A hard boundary, not a design choice.
hou.pwd().geometry()does not work here, or in any node's own parameters. It asks the node to cook while it is already cooking, which Houdini reports as Infinite recursion in evaluation. A Python SOP is the exception only because Houdini hands it working geometry through a layer that is not in the public HDK. This is worth knowing before you design against it.
Open Questions
The measurements point at more work than they close.
Does the boundary finding hold on GPU? We believe it strengthens. We need the hardware to say so.
What does differentiability change here? Warp's gradients are the reason this technology matters beyond raw throughput. A node that can carry gradients back through a Houdini cook is a materially different proposition to a fast one, and the boundary design is what would decide whether it is feasible.
Where should topology caching live? We solved it with an explicit per-node store, which puts the burden on the author. Whether the correct answer is invalidation the framework understands natively is an open design question.
Availability
Open source under MIT, for production use, modification, and redistribution.
GitHub: github.com/plattipus/houdini_nvidia_warp, including the Gray-Scott benchmark scene used above, so the numbers can be contested rather than taken on trust.
| Houdini | 20.0 (Python 3.10) |
| warp-lang | >=1.17,<2 |
| License | MIT |
Prior Art and Credits
This work sits on top of NVIDIA Warp, by NVIDIA, under the Apache 2.0 licence, and the Houdini Development Kit, by SideFX. It is an independent integration, not affiliated with or endorsed by either.
"NVIDIA" and "Warp" are trademarks of NVIDIA Corporation. "Houdini" is a trademark of Side Effects Software Inc. Both are used here only to describe what this work interoperates with.
From the Lab
This is one output of our open research practice. We publish the method and the measurements, including the ones that undercut the pitch, because a result nobody can check is marketing.
It sits alongside Crucible, which removes the bake between a procedural asset and the final render, and houdini_usd_gsplat, which brings Gaussian splats into Solaris as first-class USD primitives. The through-line across all three is the same as the finding above: the expensive part of a pipeline is rarely the computation. It is the crossings, the conversions, and the waiting, and those are what we go after, so the visual effects and animation and film and commercial production teams get more attempts inside the same schedule.