Windows/Blackwell support, plus fix DINOv3 feature extraction on Transformers 5.x

[rwfsmith]

Summary

Adds Windows + Blackwell (sm_120) support and fixes two correctness bugs found while getting the pipeline running there. Tested end-to-end on Windows 11 / Python 3.13 / CUDA 13.0 / RTX 5090, without conda.

The two correctness fixes are not Windows-specific and are the most valuable part of this PR — they are cleanly separable if you'd prefer to take them alone.

---

1. DINOv3 feature extraction is broken on Transformers 5.x (cross-platform)

DinoV3FeatureExtractor.extract_features iterates self.model.layer. In current Transformers, DINOv3ViTModel has no .layer — the blocks moved into the encoder submodule at DINOv3ViTModel.model.layer. Verified on transformers==5.15.1:

transformers 5.15.1
has .layer      : False
has .model.layer: True
type(m.model)   : DINOv3ViTEncoder
has .norm       : True LayerNorm
norm affine     : True

So main raises AttributeError: 'DINOv3ViTModel' object has no attribute 'layer' on any recent Transformers.

The tempting one-line fix — returning self.model(...).last_hidden_state — is wrong, and fails quietly. DINOv3ViTModel.forward applies the backbone's trained final LayerNorm (elementwise_affine=True, confirmed above), whereas this code is built on a parameter-free F.layer_norm of the pre-norm states. That learned affine shifts the conditioning distribution enough that the sparse-structure flow sampler collapses z_s to ≈0 and decodes an all-empty voxel grid on roughly 1 seed in 6. It surfaces ~300 lines downstream as:

RuntimeError: max(): Expected reduction dim to be specified for input.numel() == 0

I hit exactly this and spent a while chasing it, which is why the fix is written defensively.

This PR adds _encoder_layers(), which resolves the block list across Transformers versions, then keeps the original manual iteration and parameter-free F.layer_norm semantics unchanged.

Measured on a fixed 40-seed sparse-structure sweep:

	before	after
empty grids	7/40	0/40
voxel count	139–2169	3299–4338
z_s std	0.11–0.41	0.58–0.65

The whole distribution tightens, so this improves output fidelity on every seed, not just the ones that crashed.

Applied to both copies of DinoV3FeatureExtractor (inference and trainer); the two extract_features/_encoder_layers implementations are verified AST-identical.

How this was isolated: running the same 40 seeds under both flash_attn and sdpa produced the identical failing seed set, which ruled out attention-backend numerics and pointed upstream to the conditioning.

2. Invalid C++ literal suffixes in o-voxel (portability)

flexible_dual_grid.cpp uses 1e-6d and 0.0d. d is not a valid floating-point suffix in C++ (it's Java/D); in C++11 it parses as a user-defined literal operator""d, which isn't defined here. MSVC rejects it. Changed to 1e-6 / 0.0 — these were already double contexts, so no behavior change.

Also torch::from_blob(..., {svo.size()}, ...): size() returns size_t, and narrowing inside a braced-init-list is ill-formed. Added an explicit static_cast<int64_t>.

---

3. Windows support (additive)

• requirements.txt — pip-only install; no conda required
• setup_windows.ps1 / .bat / .py — three entry points (PowerShell, cmd.exe, pure Python); locate VS2022 build tools, set CUDA_HOME / TORCH_CUDA_ARCH_LIST, then pip install
• Attention backend defaults: on Windows, probe whether flash_attn actually runs on the GPU and fall back to sdpa if not. Linux behavior is unchanged — _default_backend() returns 'flash_attn' unconditionally off Windows.
• Added missing 'sdpa' to the SPARSE_ATTN_BACKEND env allowlist and to set_attn_backend's type hint; the sdpa path existed but was unreachable via env var.
• sample_sparse_structure() now raises an actionable error on an empty grid instead of the cryptic max() failure 300 lines later.
• README section for the Windows path; .gitignore entries for tmp/ and example.py output.

---

⚠️ Needs discussion before merge

requirements.txt pins two dependencies to my forks:

cumesh     @ git+https://github.com/rwfsmith/CuMesh.git@windows-py313-cuda13
flex_gemm  @ git+https://github.com/rwfsmith/FlexGEMM.git@windows-py313-cuda13

These carry Windows/CUDA-13 build fixes, and the CuMesh branch also fixes an out-of-bounds GPU read in the narrow-band dual contouring remesh kernel that made to_glb(..., remesh=True) non-deterministically produce shredded meshes — same inputs, different result run to run. That one is a real correctness bug worth upstreaming to CuMesh independently.

I don't expect you'd want to ship pins to a personal fork. Happy to split this PR so the two correctness fixes land on their own, or to rework the packaging however you prefer.

Verification

• example.py → sample.mp4 (4.0 MB) + sample.glb (44.1 MB); mesh checks out at 808k verts / 973k faces, 0 NaNs, 6 degenerate faces of 973k, sane extents
• app.py end-to-end through the Gradio UI, including GLB extraction
• Seed 1 + assets/example_image/T.png, a combination that previously crashed, now produces a clean mesh

Scope / limitations

• Validated on a single configuration (Windows 11, Python 3.13, CUDA 13.0, RTX 5090). Not tested on other Windows GPUs or CUDA versions.
• I did not regression-test on Linux. The Windows-specific changes are guarded by platform.system() == 'Windows', and the DINOv3 and C++ changes preserve existing semantics, but a Linux CI run would be worth doing.
• The DINOv3 fix is validated by distribution behavior and by restoring the original intent of the code, not by a bit-for-bit comparison against a known-good older-Transformers reference.

[rwfsmith]

Following up on the "Needs discussion before merge" section — I've now filed the upstream PRs so the fork pins in requirements.txt can be removed rather than merged:

CuMesh

• JeffreyXiang/CuMesh#39 — out-of-bounds UDF read in the dual contouring kernel. This one is not Windows-specific: get_vertex_val() uses a hashmap lookup result as an array index without checking the miss sentinel, so it reads ~16 GB past the base pointer on a miss. It's UB on every platform and causes silent, intermittent mesh corruption; it just happened to be reproducible on my setup. Worth taking regardless of anything in this PR.
• JeffreyXiang/CuMesh#40 — C++20 for the existing Windows MSVC flags (PyTorch header requirement).

FlexGEMM

• JeffreyXiang/FlexGEMM#31 — template disambiguator on four dependent data_ptr<T>() calls that MSVC rejects, plus the same C++20 flag bump.

Once those land, I'll update requirements.txt here to point back at the official JeffreyXiang/* repos, which removes the only part of this PR that pins to a personal fork.

Happy to also split the two correctness fixes in this PR (the DINOv3 Transformers 5.x fix and the o-voxel invalid C++ literal suffixes) into a separate, smaller PR if you'd prefer to evaluate those independently of the Windows support work — the DINOv3 one in particular fixes a hard AttributeError crash on current Transformers that affects Linux users too.

[rwfsmith]

@microsoft-github-policy-service agree

[rwfsmith]

Update: while testing further I found and fixed a third correctness bug, this one also cross-platform, in CuMesh::fill_holes. It's now in the pinned rwfsmith/CuMesh@windows-py313-cuda13 branch and upstreamed separately as JeffreyXiang/CuMesh#41.

Symptom: every generation after the first one in a process produced a mesh smeared into a cylindrical column of sliver triangles. Because the first run in a fresh process was almost always fine, single-shot repros never showed it — but the Gradio app is long-lived, so in practice the 2nd image onward was corrupt.

Cause: fill_holes averages hole-cap vertices with cub::DeviceSegmentedReduce::Sum over Vec3f. CUB builds the reduction identity on the host as InitT{}. Vec3f has a user-provided default constructor, so Vec3f{} calls it rather than zero-filling — and that constructor is __device__-only, so it isn't callable from host code. The identity was therefore raw host stack memory, folded into every segment sum.

The corrupted coordinates decoded to exactly G / n for small integers n (G ≈ 9.52e9), which is only possible if a single constant is added to each segment before dividing by segment size — that pinned it to the identity. Fix is 6 lines: make Vec3f's constructors __host__ __device__.

Validation (N generations in one process, same input):

	before	after
run 0	ok	ok
runs 1–4	broken — max edge 4.76e9, 3910 long edges, 52.3% mask coverage	ok — max edge 0.005, 0 long edges, 30.0% coverage

0/5 clean before, 5/5 after, confirmed visually as well.

No change to this PR's diff is needed — requirements.txt already pins that branch. Flagging it here since it's the third cross-platform correctness fix uncovered by this work, alongside CuMesh#39, CuMesh#40, and FlexGEMM#31.