Skip to content

texture: add CUDA orthomap generation backend (cherry-pick of 3181f4a)#1269

Open
cdcseacave wants to merge 4 commits into
cuda-texturefrom
cuda-ortho
Open

texture: add CUDA orthomap generation backend (cherry-pick of 3181f4a)#1269
cdcseacave wants to merge 4 commits into
cuda-texturefrom
cuda-ortho

Conversation

@cdcseacave

@cdcseacave cdcseacave commented May 10, 2026

Copy link
Copy Markdown
Owner

Summary

Cherry-picks 3181f4a from justin/justin/cuda-ortho (Justin Gabriel's fork) on top of the cuda-texture branch (#1268): adds a CUDA-accelerated orthographic-map generator that runs without full mesh texturing. Useful for top-down map generation on geo-referenced reconstructions.

The new path is gated behind #ifdef _USE_CUDA AND a runtime --ortho-only flag, so default behavior is unchanged.

PR base is cuda-texture (not develop) so the diff is just the ortho work. When #1268 merges, this PR auto-rebases to develop.

New CLI flags on TextureMesh (require _USE_CUDA)

  • --ortho-only — generate only the orthomap, skip mesh texturing
  • --tile-size — world-space tile size (0 = auto binary search, default)
  • --max-tile-res — max per-tile pixel resolution (default 8192)
  • --resume-orthomap — skip re-rasterization when tile PNGs already exist on disk

New API

Scene::GenerateOrthoMap(tileSize, maxImgRes, maxTileRes, outputName, bResume) in libs/MVS/Scene.h, implemented in the new libs/MVS/SceneOrthoCUDA.inl. Pipeline:

  1. Per-view mesh projection (face/visibility/score maps)
  2. Auto tile-size binary search (or fixed value)
  3. Per-tile multi-view rasterization → multi-view stack on GPU
  4. Per-tile blending (best-views + outlier removal + Laplacian-pyramid recombination)
  5. Final tile assembly with overlap feathering, written to outputName (resolved against WORKING_FOLDER if relative; absolute paths also accepted)

Build fix (CCCL 13 / Thrust)

Two make_transform_iterator lambdas in curast.cu were declared __device__-only. CCCL 13 (CUDA 13.x) instantiates cuda::std::invoke_result_t<lambda, ...> from host code, which requires the lambda to be __host__ __device__ (or wrapped via cuda::proclaim_return_type, or a named functor). Promoted both lambdas to __host__ __device__ and gave them an explicit -> uint trailing return type — see commit e2f8ba2.

Path-handling cleanup

Replaced manual WORKING_FOLDER + name concatenations in SceneOrthoCUDA.inl with the existing MAKE_PATH / MAKE_PATH_SAFE macros from libs/Common/Common.h. The user-provided output filename now uses MAKE_PATH_SAFE, so both relative and absolute -o paths work correctly (previously an absolute path silently failed cv::imwrite because of the literal string concatenation).

Documentation

Added docs/CUDA_TEXTURE_AND_ORTHOMAP.md — a short user-facing README covering the goals, methods, example invocations, and observed behavior of both cuda-texture and cuda-ortho.

Validation

In-tree Tests 2 1 0 (regression check on cuda-texture path), 3 consecutive runs:

Run CPU score CUDA score
1 46.86 45.40
2 47.15 45.44
3 47.33 45.00

(One earlier run flaked with CUDA <45 — known pre-existing PatchMatch threading non-determinism, not caused by this PR.)

Tanks-and-Temples Truck (251 images, mesh decimated to 30% = 900k faces):

Metric Value
--ortho-only wall time 10.75 min
Auto tile-size convergence ~30 binary-search iterations → 8.21 m world-space
Tile count 2 (0,0) + (0,1)
Final orthomap size 0.97 MB PNG
Peak RSS 5.1 GB
Peak pagefile 10.4 GB

Credit

Original author: JuSiGab (Justin Gabriel) — the cuda-ortho commit was authored by him on justin/justin/cuda-ortho. Cherry-picked here with author re-attribution to facilitate review on this fork.

Test plan

  • Verify build with OpenMVS_USE_CUDA=ON on Windows (MSVC, CUDA 13.x) and Linux
  • Verify default TextureMesh (no --ortho-only) behavior is unchanged
  • Run Tests 2 1 0 — expect cuda-texture path to still pass (CPU + CUDA scores ≥ 45)
  • Run TextureMesh --ortho-only -w <scene_dir> -o ortho.png on a geo-referenced scene — expect <scene_dir>/ortho.png
  • Run with absolute -o /abs/path/ortho.png — should also work after the path-handling cleanup

🤖 Generated with Claude Code

Copilot AI review requested due to automatic review settings May 10, 2026 20:40

Copilot AI left a comment

Copy link
Copy Markdown

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Pull request overview

Adds a CUDA-backed orthographic map (“orthomap”) generation path to the MVS scene texturing pipeline, intended to generate top-down maps on geo-referenced reconstructions without running full mesh texturing. The new functionality is gated behind _USE_CUDA and a new TextureMesh --ortho-only runtime flag.

Changes:

  • Introduces Scene::GenerateOrthoMap(...) and implements the CUDA orthomap pipeline in a new SceneOrthoCUDA.inl.
  • Extends CURAST with orthographic projection/rasterization + a substantial GPU-side collapse/blending pipeline for per-tile orthomap synthesis.
  • Adds new TextureMesh CLI flags (--ortho-only, --tile-size, --max-tile-res, --resume-orthomap) to drive the orthomap generation mode.

Reviewed changes

Copilot reviewed 7 out of 7 changed files in this pull request and generated 10 comments.

Show a summary per file
File Description
libs/MVS/SceneTextureCUDA.inl Includes the new orthomap implementation into the CUDA scene texturing compilation unit.
libs/MVS/SceneOrthoCUDA.inl New orthomap generation implementation: view projection, tile-size search, per-tile rasterization, tile blending/assembly.
libs/MVS/Scene.h Adds the CUDA-gated Scene::GenerateOrthoMap(...) API.
libs/MVS/CURAST/curast.h Extends CURAST/OrthoCURAST interfaces and adds new CUDA-side structs/workspace declarations for orthomap collapse.
libs/MVS/CURAST/curast.cu Adds orthographic projection/rasterization kernels and a large orthomap stack-collapse implementation.
libs/MVS/CURAST/curast.cpp Implements OrthoCURAST methods (projection, rasterization, normal map generation, collapse orchestration).
apps/TextureMesh/TextureMesh.cpp Adds CLI flags and an early-exit --ortho-only mode that calls Scene::GenerateOrthoMap(...).

Comment thread libs/MVS/CURAST/curast.cu
Comment thread libs/MVS/CURAST/curast.cu
Comment thread libs/MVS/CURAST/curast.cu Outdated
Comment thread libs/MVS/CURAST/curast.cu Outdated
Comment thread libs/MVS/SceneOrthoCUDA.inl
Comment thread libs/MVS/SceneOrthoCUDA.inl Outdated
Comment thread libs/MVS/SceneOrthoCUDA.inl Outdated
Comment thread apps/TextureMesh/TextureMesh.cpp Outdated
Comment thread libs/MVS/CURAST/curast.h Outdated
Comment thread libs/MVS/CURAST/curast.h Outdated
- curast.h: strip UTF-8 BOM; fix "fort"->"for" typo in SetFacesForView doc
- curast.cu UpdateBestViewsOrtho: apply computed edgeMap to visibilityMap
  via MultiplyKernel (matches perspective UpdateBestViews; previously the
  JFA distance-to-edge map was computed and freed unused)
- curast.cu OrthoRasterizeKernelBis: add per-tile depth buffer gated by
  atomicMaxFloat(pW.z()) so only the topmost world-Z surface contributes
  samples per pixel (eliminates the inter-triangle race that allowed far
  surfaces to overwrite near ones)
- curast.cu OrthoRasterize: drop cudaDeviceSynchronize and switch the
  pointer-table + new depth buffer to cudaMallocAsync/cudaFreeAsync on
  the stream, keeping the per-tile pipeline fully stream-scoped
- curast.cu GetNormalOrthoMap: switch per-call depth-buffer malloc/free
  to async on the stream and move thrust fills onto thrust::cuda::par.on
  to avoid implicit host syncs
- SceneOrthoCUDA.inl auto tile-size search: when no candidate satisfies
  max-tile-res / GPU memory, clamp tileRes to maxTileRes with a warning
  instead of allocating beyond the user/GPU budget
- SceneOrthoCUDA.inl final crop: guard the no-non-black-pixel case so we
  fall back to writing the full orthoMap rather than building a
  degenerate ROI
- TextureMesh --ortho-only: clarify --output-file semantics in this mode
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants