29ec6a5 · proc-nodes · pear-git-v3

Changed files (6) +310 -17

demo/bake_bridged_graph.gd +4 -1

…	…	@@ -31,7 +31,10 @@ func _init() -> void:
31	31	print("cage: %d verts, %d triangles" % [verts.size(), idx.size() / 3])
32	32
33	33	var host := ProcGeometryNode.new()
34		host.use_gpu = false
	34	# GPU on/off from the command line. It was hardcoded false, so this harness
	35	# could only ever measure the CPU path -- and measuring the CPU path to
	36	# explain why a node ran on the CPU proves nothing.
	37	host.use_gpu = a.size() > 3 and a[3] == "gpu"
35	38	host.set_input_mesh(cage_mesh)
36	39	host.set_graph(g)
37	40	get_root().add_child(host)

docs/performance-audit-vs-blender.md +115 -0

…	…	@@ -0,0 +1,115 @@
	1	# Why proc-nodes is 32x slower than Blender on the same graph
	2
	3	Measured 2026-09-08 on the bridged voronoi mountain (`mountain_voronoi1.blend`,
	4	`VoronoiMountain`, 2,582,018 output vertices), 16-core machine, RTX 2070.
	5
	6	## The numbers
	7
	8	proc-nodes, full evaluation 74.0 s
	9	Blender, forced full re-evaluation 2.29 s
	10	ratio 32x
	11
	12	⚠ Blender is not "milliseconds". It caches: `to_mesh()` on an unchanged
	13	graph returns in 0.0 ms because nothing was recomputed. Forcing a real
	14	re-evaluation by changing a socket the graph reads gives 2287.7 ms. The
	15	millisecond impression is the cache, and that cache is itself finding #4.
	16
	17	proc-nodes, by node
	18	MeshToSDFGrid 53.7 s 73%
	19	SetPosition 17.8 s 24% (the whole value cone is charged here,
	20	because its fields are lazy and get
	21	materialised when Set Position asks)
	22	SubdivideMesh 1.7 s 2%
	23	GridToMesh 0.5 s <1%
	24
	25	## 1. Nothing is threaded except one file
	26
	27	`grep -rln "std::thread\\|parallel_for\\|WorkerThreadPool" src/nodes/ src/graph/`
	28	returns exactly one file: `volume_grid_nodes.cpp`. Every value node -- the noise,
	29	voronoi, math and vector-math nodes that make up the 17.8 s cone -- runs on one
	30	core.
	31
	32	Blender parallelises every field and multi-function evaluation:
	33
	34	FN_multi_function.hh:111 int64_t min_grain_size = 10000;
	35	intern/multi_function.cc:136 threading::parallel_for_aligned(...)
	36	intern/field.cc:477 threading::parallel_for(mask.index_range(), 2048, ...)
	37
	38	On this machine that is up to 16x on its own, and it is the single largest
	39	factor outside the SDF node.
	40
	41	## 2. Every node materialises its whole output
	42
	43	`VoronoiTextureNode::execute` allocates `PackedFloat32Array dist_a`,
	44	`PackedColorArray col_a`, `PackedVector3Array pos_a`, `PackedFloat32Array w_a`
	45	and a radius array -- five outputs, each `n` elements. At 2.58M elements the
	46	colour array alone is 41 MB and the position array 31 MB, for ONE node. A value
	47	cone of ~25 nodes writes and re-reads hundreds of megabytes that are never
	48	looked at again.
	49
	50	Blender evaluates fields over an `IndexMask` in chunks of 2048 (`field.cc:477`),
	51	so a chunk's intermediates stay in cache and the whole chain runs while the data
	52	is hot. Same arithmetic, a different memory profile.
	53
	54	The per-element accessors compound it: `SocketValue::float_at` /`vec_at` branch
	55	on `is_field` and then test `float_array.size() > idx`, `int_array.size() > idx`
	56	in sequence -- a type dispatch per element, where Blender devirtualises the span
	57	once per chunk.
	58
	59	## 3. The SDF grid is dense and exact; Blender's is sparse and banded
	60
	61	`MeshToSDFGridNode::execute` allocates a DENSE grid --
	62	`vol->density_grid.resize(res[0] * res[1] * res[2])` -- and runs
	63	`sdf.evaluate(...)`, an exact closest-point BVH query, for EVERY voxel. For this
	64	mountain that is 251 x 251 x 162 = 10.2 million queries against a 5.16M
	65	triangle BVH. It is threaded, and still 53.7 s.
	66
	67	Blender's `Mesh to SDF Grid` is
	68	`openvdb::tools::meshToLevelSet<openvdb::FloatGrid>`
	69	(`geometry/intern/mesh_to_volume.cc:199`), reached from
	70	`node_geo_mesh_to_sdf_grid.cc`. That is a SCAN CONVERSION into a sparse tree:
	71	work is proportional to surface area / voxel², not to volume / voxel³, and only
	72	voxels within the band are stored or computed at all.
	73
	74	⛔ **Our node takes the same "Band Width" input and uses it for the opposite
	75	thing.** Blender's band is the width of the computed shell. Ours only PADS the
	76	bounds -- `const float pad = voxel * band;` -- and then computes every voxel
	77	inside those padded bounds regardless. Same parameter name, inverted meaning.
	78
	79	Two independent costs follow:
	80
	81	- count: a band of 3 voxels over this surface is on the order of 3e5 voxels
	82	against 1.02e7 computed -- roughly 34x more work than the result needs.
	83	- per query: 53.7 s x 16 threads / 10.2e6 = **84 us of core time per
	84	query**. A BVH closest-point descent over 5M triangles should be single-digit
	85	microseconds. So the query itself is ~20x slower than it should be, on top of
	86	there being ~34x too many of them.
	87
	88	## 4. There is no cache, and the bridge threw Blender's answer away
	89
	90	Blender re-evaluates only what changed. proc-nodes re-evaluates the whole graph
	91	on every incoming message, and `blender_bridge.cpp` additionally called
	92	`evaluate()` and then `mesh_inst->set_mesh(Ref<Mesh>())` on import -- discarding
	93	the mesh Blender had just computed and shipped in the same message, to spend
	94	74 s reproducing it.
	95
	96	## What to fix, in order of measured value
	97
	98	\| fix \| expected \|
	99	\|---\|---\|
	100	\| SDF narrow band (compute the shell, not the volume) \| 53.7 s -> ~2 s \|
	101	\| Thread the value-node loops the way `volume_grid_nodes` already does \| 17.8 s -> ~1.5 s \|
	102	\| Do not re-bake on import; show what Blender sent \| removes 74 s from the bridge loop entirely \|
	103	\| Dirty-tracking so unchanged graphs are not re-run \| removes it again on every subsequent message \|
	104	\| Chunked field evaluation (the deep fix for #2) \| large, and the most invasive \|
	105
	106	The GPU batch is a fifth path and is NOT measurable headlessly: a batch needs a
	107	`RenderingDevice`, and `--headless` has none, so `gpu_available` is false and
	108	every harness run reports CPU regardless. `SwitchNode` gained a `compile_glsl`
	109	in this session because its absence was refusing the batch that would have taken
	110	Set Position:
	111
	112	GPU batch REFUSED: node 31 (CombineXYZ) socket 2 is connected from outside
	113	the batch and has no push-constant slot
	114
	115	Whether that now lands has to be measured in the editor.

src/graph/graph_executor.cpp +11 -0

…	…	@@ -1019,10 +1019,21 @@ static bool _batch_is_representable(ProcEvalContext &p_ctx, const GpuBatch &p_ba
1019	1019	// Otherwise the only channel is a push constant, and only for a
1020	1020	// value that is ONE value -- never a field.
1021	1021	if (!pc_map.has(si)) {
	1022	// ⛔ SAY WHY. A refused batch falls back to the CPU silently,
	1023	// and on a 2.58M-vertex mountain that is 17.8 s of Set Position
	1024	// with nothing in the log but "CPU node [SetPosition]". The
	1025	// whole point of the batch is that it does not happen.
	1026	UtilityFunctions::print(" GPU batch REFUSED: node ", node_idx, " (",
	1027	node->get_node_type(), ") socket ", si, " is connected from outside "
	1028	"the batch and has no push-constant slot");
1022	1029	return false;
1023	1030	}
1024	1031	const SocketValue raw = p_ctx._read_input_raw(node_idx, si);
1025	1032	if (raw.is_field \|\| raw.is_lazy) {
	1033	UtilityFunctions::print(" GPU batch REFUSED: node ", node_idx, " (",
	1034	node->get_node_type(), ") socket ", si,
	1035	raw.is_lazy ? " is a LAZY field" : " is a FIELD",
	1036	" arriving from outside the batch -- a push constant is one value");
1026	1037	return false;
1027	1038	}
1028	1039	}

src/nodes/selection_nodes.cpp +10 -0

…	…	@@ -1,4 +1,5 @@
1	1	#include "selection_nodes.h"
	2	#include "../gpu/glsl_emit_context.h"
2	3	#include "../graph/graph_executor.h"
3	4	#include "../shared/blender_noise.h"
4	5	#include <godot_cpp/variant/utility_functions.hpp>
…	…	@@ -374,6 +375,15 @@ void SwitchNode::set_input_type(int p_type) {
374	375	output_defs.ptrw()[0].type = input_type;
375	376	}
376	377
	378	String SwitchNode::compile_glsl(const GlslEmitContext &p_ctx, int p_node_idx) const {
	379	// execute()'s scalar path is `cond ? t : f`, and its element-wise path is the
	380	// same choice per element -- which is what a per-vertex shader already is.
	381	// Socket 1 is False, socket 2 is True.
	382	return p_ctx.declare(p_node_idx, 0) + String(" = ") + p_ctx.in(p_node_idx, 0) +
	383	String(" ? ") + p_ctx.in(p_node_idx, 2) + String(" : ") +
	384	p_ctx.in(p_node_idx, 1) + String(";");
	385	}
	386
377	387	void SwitchNode::execute(ProcEvalContext &p_ctx, int p_node_idx) {
378	388	if (input_type == SocketType::Geometry) {
379	389	bool condition = p_ctx.read_bool(p_node_idx, 0);

src/nodes/selection_nodes.h +17 -0

…	…	@@ -113,6 +113,23 @@ public:
113	113	void set_input_type(int p_type);
114	114	int get_input_type_prop() const { return (int)input_type; }
115	115
	116	// ⛔ GEOMETRY SWITCHING IS NOT A VALUE, so only the value forms emit. A
	117	// geometry Switch picks a BRANCH of the graph, which a per-vertex shader
	118	// cannot express, and claiming GPU support for it would make
	119	// compile_to_function refuse the whole graph.
	120	//
	121	// Why this matters beyond the density path: a batch is built backwards over
	122	// the value cone feeding Set Position, and it stops at the first node that
	123	// cannot emit. Measured on the bridged mountain --
	124	//
	125	// GPU batch REFUSED: node 31 (CombineXYZ) socket 2 is connected from
	126	// outside the batch and has no push-constant slot
	127	//
	128	// -- because CombineXYZ's Z came from this node. One missing compile_glsl
	129	// put Set Position on the CPU for 2.58M vertices: 17.8 s an evaluation.
	130	bool supports_gpu() const override { return input_type != SocketType::Geometry; }
	131	String compile_glsl(const GlslEmitContext &p_ctx, int p_node_idx) const override;
	132
116	133	void execute(ProcEvalContext &p_ctx, int p_node_idx) override;
117	134
118	135	SwitchNode();

src/nodes/volume_grid_nodes.cpp +153 -16

…	…	@@ -193,32 +193,117 @@ void MeshToSDFGridNode::execute(ProcEvalContext &p_ctx, int p_node_idx) {
193	193	vol->grid_size = Vector3(res[0] * voxel, res[1] * voxel, res[2] * voxel);
194	194	vol->density_grid.resize(res[0] * res[1] * res[2]);
195	195	float *w = vol->density_grid.ptrw();
196		// ⛔ THREADED, because this is the whole cost of the node. Every sample is a
197		// closest-point query against the BVH, and the grid is capped at 256 per
198		// axis -- 16.7 million of them for the bridged mountain, which took 75.7 s
199		// serially and ran TWICE per evaluation, which is what "the mountain did not
200		// come across" looked like from the outside.
	196
	197	const int64_t nx = res[0], ny = res[1], nz = res[2];
	198	const int64_t total = nx * ny * nz;
	199	auto at = [nx, ny](int64_t x, int64_t y, int64_t z) { return (z * ny + y) * nx + x; };
	200
	201	// ── NARROW BAND ─────────────────────────────────────────────────────────
	202	//
	203	// Every voxel used to get an exact closest-point BVH query. For the bridged
	204	// voronoi mountain that is 251 x 251 x 162 = 10.2 MILLION queries against a
	205	// 5.16M triangle BVH -- 53.7 s, threaded, 73% of the whole evaluation.
	206	//
	207	// Dual contouring only reads values NEAR the zero crossing: it finds sign
	208	// changes on edges and interpolates where the magnitude is small. Every sign
	209	// change lies inside a shell around the triangles, so the exact field is
	210	// only needed there. Blender does the same thing and has always done --
	211	// openvdb::tools::meshToLevelSet scan-converts into a sparse tree, and the
	212	// node's "Band Width" IS that shell's width.
201	213	//
202		// libdc parallelises the equivalent loops with `#pragma omp parallel for`,
203		// but proc-nodes builds without -fopenmp, so those pragmas are silently
204		// ignored here and every libdc loop this extension compiles is serial.
205		// std::thread rather than adding the flag: it changes one loop instead of
206		// every libdc loop's behaviour inside this library, and pulls in no runtime.
	214	// ⛔ OURS TOOK THE SAME PARAMETER AND USED IT FOR THE OPPOSITE THING: `band`
	215	// only padded the bounds (`pad` above) and then every voxel inside those
	216	// padded bounds was computed anyway. Same name, inverted meaning.
	217	//
	218	// Requires a closed mesh, because the outside is found by flooding from the
	219	// grid boundary. `signed_volume` already says whether it is closed, and an
	220	// open one falls back to computing every voxel, which is what this has
	221	// always done.
	222	const bool closed = Math::abs(signed_volume) >= 1e-9;
	223	const int bandv = MAX(1, (int)band);
	224
	225	std::vector<uint8_t> in_band;
	226	if (closed) {
	227	in_band.assign((size_t)total, 0);
	228	// 1. Stamp each triangle's voxel AABB. No expansion here: dilating
	229	// afterwards is O(grid) once instead of O(triangles x band^3), and
	230	// these triangles are far smaller than a voxel -- expanding each one
	231	// by 3 would be 5.16M x 343 stamps.
	232	for (size_t t = 0; t + 2 < idx.size(); t += 3) {
	233	const Vector3 &a = mesh->positions[idx[t]];
	234	const Vector3 &b = mesh->positions[idx[t + 1]];
	235	const Vector3 &c = mesh->positions[idx[t + 2]];
	236	const Vector3 lo = a.min(b).min(c);
	237	const Vector3 hi = a.max(b).max(c);
	238	const int x0 = CLAMP((int)Math::floor((lo.x - mn[0]) / voxel), 0, (int)nx - 1);
	239	const int x1 = CLAMP((int)Math::floor((hi.x - mn[0]) / voxel), 0, (int)nx - 1);
	240	const int y0 = CLAMP((int)Math::floor((lo.y - mn[1]) / voxel), 0, (int)ny - 1);
	241	const int y1 = CLAMP((int)Math::floor((hi.y - mn[1]) / voxel), 0, (int)ny - 1);
	242	const int z0 = CLAMP((int)Math::floor((lo.z - mn[2]) / voxel), 0, (int)nz - 1);
	243	const int z1 = CLAMP((int)Math::floor((hi.z - mn[2]) / voxel), 0, (int)nz - 1);
	244	for (int z = z0; z <= z1; z++) {
	245	for (int y = y0; y <= y1; y++) {
	246	uint8_t *row = &in_band[(size_t)at(0, y, z)];
	247	for (int x = x0; x <= x1; x++) {
	248	row[x] = 1;
	249	}
	250	}
	251	}
	252	}
	253	// 2. Dilate by `band` voxels, separably: three 1D max-filters rather
	254	// than one 3D neighbourhood, so the cost is 3 x O(grid) x band.
	255	std::vector<uint8_t> tmp(in_band.size());
	256	auto dilate_axis = [&](int axis) {
	257	tmp = in_band;
	258	for (int64_t z = 0; z < nz; z++) {
	259	for (int64_t y = 0; y < ny; y++) {
	260	for (int64_t x = 0; x < nx; x++) {
	261	if (tmp[(size_t)at(x, y, z)] == 0) {
	262	continue;
	263	}
	264	for (int d = -bandv; d <= bandv; d++) {
	265	int64_t xx = x, yy = y, zz = z;
	266	if (axis == 0) xx += d; else if (axis == 1) yy += d; else zz += d;
	267	if (xx < 0 \|\| xx >= nx \|\| yy < 0 \|\| yy >= ny \|\| zz < 0 \|\| zz >= nz) {
	268	continue;
	269	}
	270	in_band[(size_t)at(xx, yy, zz)] = 1;
	271	}
	272	}
	273	}
	274	}
	275	};
	276	dilate_axis(0);
	277	dilate_axis(1);
	278	dilate_axis(2);
	279	}
	280
	281	// ⛔ THREADED, because this is the whole cost of the node. libdc parallelises
	282	// its equivalent loops with `#pragma omp parallel for`, but proc-nodes builds
	283	// without -fopenmp so those pragmas are silently ignored and every libdc loop
	284	// this extension compiles is serial. std::thread rather than adding the flag:
	285	// it changes one loop instead of every libdc loop's behaviour in this
	286	// library, and pulls in no runtime.
207	287	//
208	288	// Safe to share `sdf`: evaluate() is non-const by declaration only -- it
209	289	// calls closest_distance_sq/query_bvh, both const, and the BVH is immutable
210	290	// after load_arrays.
211	291	{
212		const int nz = res[2];
	292	const int nzi = res[2];
213	293	unsigned hw = std::thread::hardware_concurrency();
214	294	int nthreads = (int)(hw > 0 ? hw : 1u);
215		if (nthreads > nz) nthreads = nz;
	295	if (nthreads > nzi) nthreads = nzi;
216	296	if (nthreads < 1) nthreads = 1;
	297	const bool banded = closed && !in_band.empty();
217	298	auto slab = [&](int z0, int z1) {
218	299	for (int z = z0; z < z1; z++) {
219	300	for (int y = 0; y < res[1]; y++) {
220	301	for (int x = 0; x < res[0]; x++) {
221		w[(z * res[1] + y) * res[0] + x] = sdf_sign * sdf.evaluate(
	302	const int64_t o = at(x, y, z);
	303	if (banded && !in_band[(size_t)o]) {
	304	continue; // filled by the flood below
	305	}
	306	w[o] = sdf_sign * sdf.evaluate(
222	307	mn[0] + x * voxel, mn[1] + y * voxel,
223	308	mn[2] + z * voxel, voxel);
224	309	}
…	…	@@ -226,20 +311,72 @@ void MeshToSDFGridNode::execute(ProcEvalContext &p_ctx, int p_node_idx) {
226	311	}
227	312	};
228	313	if (nthreads == 1) {
229		slab(0, nz);
	314	slab(0, nzi);
230	315	} else {
231	316	std::vector<std::thread> pool;
232	317	pool.reserve(nthreads);
233		const int per = (nz + nthreads - 1) / nthreads;
	318	const int per = (nzi + nthreads - 1) / nthreads;
234	319	for (int t = 0; t < nthreads; t++) {
235	320	const int z0 = t * per;
236		const int z1 = MIN(z0 + per, nz);
	321	const int z1 = MIN(z0 + per, nzi);
237	322	if (z0 >= z1) break;
238	323	pool.emplace_back(slab, z0, z1);
239	324	}
240	325	for (std::thread &th : pool) th.join();
241	326	}
242	327	}
	328
	329	// ── OUTSIDE BY FLOOD, INSIDE BY WHAT IS LEFT ────────────────────────────
	330	//
	331	// A voxel outside the band cannot contain a sign change, so it needs only
	332	// the right SIGN and a magnitude bigger than anything in the band, so no
	333	// edge between the two reads as a crossing. The bounds were padded by
	334	// `pad` above, so the whole grid boundary is outside the mesh and a flood
	335	// from it reaches exactly the outside.
	336	if (closed && !in_band.empty()) {
	337	// positive = solid after sdf_sign (see the winding note above)
	338	const float far_v = voxel * (float)(bandv + 1);
	339	std::vector<uint8_t> outside((size_t)total, 0);
	340	std::vector<int64_t> stack;
	341	stack.reserve(4096);
	342	auto push_if_open = [&](int64_t x, int64_t y, int64_t z) {
	343	const int64_t o = at(x, y, z);
	344	if (in_band[(size_t)o] \|\| outside[(size_t)o]) {
	345	return;
	346	}
	347	outside[(size_t)o] = 1;
	348	stack.push_back(o);
	349	};
	350	for (int64_t z = 0; z < nz; z++) {
	351	for (int64_t y = 0; y < ny; y++) {
	352	for (int64_t x = 0; x < nx; x++) {
	353	if (x == 0 \|\| y == 0 \|\| z == 0 \|\| x == nx - 1 \|\| y == ny - 1 \|\| z == nz - 1) {
	354	push_if_open(x, y, z);
	355	}
	356	}
	357	}
	358	}
	359	while (!stack.empty()) {
	360	const int64_t o = stack.back();
	361	stack.pop_back();
	362	const int64_t x = o % nx;
	363	const int64_t y = (o / nx) % ny;
	364	const int64_t z = o / (nx * ny);
	365	if (x > 0) push_if_open(x - 1, y, z);
	366	if (x < nx - 1) push_if_open(x + 1, y, z);
	367	if (y > 0) push_if_open(x, y - 1, z);
	368	if (y < ny - 1) push_if_open(x, y + 1, z);
	369	if (z > 0) push_if_open(x, y, z - 1);
	370	if (z < nz - 1) push_if_open(x, y, z + 1);
	371	}
	372	for (int64_t o = 0; o < total; o++) {
	373	if (in_band[(size_t)o]) {
	374	continue;
	375	}
	376	w[o] = outside[(size_t)o] ? -far_v : far_v;
	377	}
	378	}
	379
243	380	out_geo->set_volume(vol);
244	381	p_ctx.write_geometry(p_node_idx, 0, out_geo);
245	382	}

proc-nodes commits / 29ec6a5

Changed files (6) +310 -17