Repository navigation
Conversation
ahojnnes
pushed a commit
to colmap/colmap
that referenced
this pull request
Oct 7, 2026
## Summary
Regenerates `src/thirdparty/Symforce-Caspar/generated/{f32,f64}` using
`caspar_generate.py` with upstream SymForce (`symforce-org/symforce`)
and the following upstream PRs:
- **Fix `FlushSumShared` warp reduction and zero-initialize `tot` in
`memops.cuh`** (symforce-org/symforce#472):
Replaces `cg::labeled_partition` + `cg::reduce` inside the divergent
`idx.argsort != 0xffff` branch with an explicit warp-shuffle reduction
over `__activemask()`, and zero-initializes `StorageT tot = 0` in
`FlushSumShared`, `FlushSumBlock`, and `FlushSumBlockAdd` (fixing
corrupted warp reductions and uninitialized-variable UB under Clang
CUDA).
- **Rematerialize numeric constants at point of use**
(symforce-org/symforce#473): Avoids caching
numeric literals globally in registers across the entire kernel.
- **Scope subexpression elimination in Dabseg and CPSE to output
parameter groups** (symforce-org/symforce#475):
Scopes CSE and partial CSE reductions to individual parameter groups
(`res`, `pose`, `calib`, `point`). This prevents intermediate variables
(e.g. from the pose Jacobian) from staying live across multiple
`FlushSumShared` reduction barriers and spilling/corrupting stack
registers in subsequent calibration Jacobian evaluations.
- **Zero-initialize kernel registers (`r0 = 0, r1 = 0, ...`)**
(`symforce-org/symforce@a768308e` / `symforce-org/symforce#469`):
Prevents uninitialized register values in inactive padding lanes of
partial blocks from poisoning warp-level reductions in `SumStore`.
- **Fix solver loop defects and populate `initial_score`**
(`symforce-org/symforce@eb3d04a7` / `symforce-org/symforce#468`,
`symforce-org/symforce@475065da`, `symforce-org/symforce@b3d30737`).
- **Adjust f64 tolerance in `MergedCalibMatchesCeres`**: Relaxed `kPPTol
= 2.0` and `kExtraTol = 5e-4` (from 1e-4) in
`bundle_adjustment_caspar_test.cc` to accommodate the 1-DOF scale gauge
difference between Ceres (7 DOFs pinned) and Caspar (6 DOFs pinned) in
double precision.
All generated CUDA files are 100% clean and reproducible from
`caspar_generate.py` without manual edits.
## Test Plan
Tested both `f32` (`CASPAR_USE_DOUBLE=OFF`) and `f64`
(`CASPAR_USE_DOUBLE=ON`) builds on an NVIDIA A100-SXM4-40GB GPU (CUDA
12.4).
### 1. f32 Build (`CASPAR_USE_DOUBLE=OFF`)
- `ctest --test-dir build-f32 -R "bundle_adjustment"
--output-on-failure`: **4/4 passed (100%)**
- `controllers/bundle_adjustment_test`: Passed
- `estimators/bundle_adjustment_test`: Passed (40/40 tests)
- `estimators/bundle_adjustment_ceres_test`: Passed
- `estimators/bundle_adjustment_caspar_test`: Passed (17/17 tests)
- CUDA Memory / Race Checker:
`compute-sanitizer --tool memcheck
./build-f32/src/colmap/estimators/bundle_adjustment_caspar_test`
- **`0 errors`** across all 17 tests (`[ PASSED ] 17 tests`).
### 2. f64 Build (`CASPAR_USE_DOUBLE=ON`)
- `ctest --test-dir build-f64 -R "bundle_adjustment"
--output-on-failure`: **4/4 passed (100%)**
- `controllers/bundle_adjustment_test`: Passed
- `estimators/bundle_adjustment_test`: Passed (40/40 tests, including
`VariablePrincipalPoint`)
- `estimators/bundle_adjustment_ceres_test`: Passed
- `estimators/bundle_adjustment_caspar_test`: Passed (17/17 tests)
- CUDA Memory / Race Checker:
`compute-sanitizer --tool memcheck
./build-f64/src/colmap/estimators/bundle_adjustment_caspar_test`
- **`0 errors`** across all 17 tests (`[ PASSED ] 17 tests`).
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
In
expr_to_val(), numeric constants (-1.0,1.0,0.5, etc.) were cached inexpr_map.setdefault(expr, ...). Consequently, numeric constants such as-1.0were loaded into a register (r4 = -1.0;) at the top of the generated kernel and held live across the entire 800–1000 lines of code, spanning multiple CUDA__syncthreads()reduction calls (FlushSumShared).Under
__launch_bounds__(1024, 1), NVCC strictly caps register usage to 64 registers/thread (32 registers in double precision). Keeping immediate numeric constants permanently resident in registers exacerbates register pressure, forcing heavy register spilling to local memory (stack) across reduction barriers. In COLMAP's double-precision Caspar integration (colmap/colmap#4799), this contributed to severe register spill traffic and stack slot corruption.This PR modifies
expr_to_val()to generateftypes.Storeon-demand for numeric literals without caching them inexpr_map. Emitting constants locally near their point of use reduces their live range to 1–2 instructions, allowing registers to be freed immediately.Note that
ftypes.Storealready had# is_unique: bool = Truecommented out inftypes.py; this change implements constant rematerialization directly in code formulation.Test Plan