Fix MinkowskiInterpolation zero/under-weighted results at sparse boundaries (renormalize over existing neighbors)

[alpsaur]

Fixes the long-standing boundary bug in MinkowskiInterpolation / SparseTensor.features_at_coordinates reported upstream as NVIDIA/MinkowskiEngine#477 (duplicates #363, #383).

Root cause

The backend (interpolation_map_weight in src/coordinate_map_cpu.hpp / src/coordinate_map_gpu.cu) computes multilinear weights prod_d (1 - |q_d - c_d| / stride_d) over the 2^D lattice corners surrounding each query point and silently drops corners that are not occupied in the sparse tensor (CPU: hash miss skipped; GPU: thrust::remove_if on the unused key). The surviving weights were never renormalized, so any query with a partially missing neighborhood was under-weighted by the factor sum(existing weights) < 1 — down to exactly zero for queries adjacent to valid data.

Before / after

3D lattice: fully occupied 2x2x2 cube at (0..1)^3 with features 1..8, plus an isolated voxel at (5,5,5) with feature 9. Identical results on CPU and CUDA:

query	corners occupied	before	after	expected (renormalized)
(0,0,0) exact voxel	on-voxel	1.000000	1.000000 (bit-identical)	1.0
(5,5,5) exact voxel	on-voxel	9.000000	9.000000 (bit-identical)	9.0
(0.25,0.5,0.75) interior	8/8	3.750000	3.750000 (bit-identical)	3.75
(5.25,5.25,5.25) next to isolated voxel	1/8	3.796875	9.000000	9.0
(0.5,0.5,1.75) just outside cube face	4/8	1.250000	5.000000	5.0
(20.5,20.5,20.5) empty region	0/8	0.0	0.0 (unchanged)	0.0

Fix

Python-side, in MinkowskiInterpolationFunction.forward: scatter-add the backend-returned weights per output row; for rows whose found-corner count is below 2^D (and weight sum > 0), divide both the output features and the weights by the surviving weight sum. The renormalized weights are saved for backward and returned via return_weights=True, so gradients and reported weights stay exactly consistent with the produced features — the existing backward kernels are untouched.

Why Python-side rather than in the kernels:

• one implementation covers the CPU kernel, the GPU coo_spmm path, and all dtypes;
• backward correctness falls out of saving the renormalized weights (no changes to InterpolationBackwardCPU/GPU);
• far smaller blast radius than modifying interpolation_map_weight in both backends plus the spmm accumulation.

Backward compatibility

• Exact-on-voxel and full-neighborhood queries are bit-identical to the previous output: their rows are multiplied by exactly 1.0, and when a batch has no partial row at all the backend tensors are returned untouched (fast path).
• Queries with no occupied corner still return zeros (documented behavior of features_at_coordinates).
• TensorField.splat() consumes raw weights through the coordinate-manager API and is intentionally unaffected — splatting is a scatter, not an average, so un-renormalized weights are correct there.
• fp16/bf16 (autocast) supported; the weight sum accumulates in fp32.
• The fix is on by default: it is a correctness fix, and no test or example depends on the un-renormalized values.

Validation

• New tests/python/interpolation_renorm.py (GPU cases conftest-gated like the rest of the suite): exact-voxel bit-equality, full-neighborhood bit-equality vs the raw backend output, boundary queries vs a renormalized fp64 reference, per-query weight sums of 1, fp64 torch.autograd.gradcheck over partial neighborhoods (CPU+CUDA), fp16/bf16 autocast tolerance variants. 9/9 pass locally on an RTX 5090 (CUDA 12.8, torch 2.9.1).
• tests/python/interpolation.py green except the pre-existing conftest-skipped entry (test_gpu's 10M-iteration leak loop validated via a trimmed 200-iteration run incl. gradcheck).
• Core subset green: convolution.py (minus TestPCD), pool.py, half_precision.py — 37 passed, 3 pre-existing skips.
• No kernel changes; no rebuild required.

🤖 Generated with Claude Code