Repository navigation
Conversation
… parameter groups
In Caspar kernels, factors write to multiple distinct output parameter groups
(e.g., residual 'res', camera pose, camera calibration, and 3D point).
Between these parameter groups, the kernel emits barrier synchronizations
and shared memory reduction flushes (FlushSumShared).
Previously, Common Subexpression Elimination (CSE) in make_dabseg and
do_cpse operated globally across the entire factor DAG. When a subexpression
(such as -(1 + k*r^2)) computed during the camera pose Jacobian was reused
hundreds of lines later in the calibration Jacobian across 12 FlushSumShared
barriers, the variable remained live throughout all intermediate reductions.
Under __launch_bounds__(1024, 1) (which limits threads to 64 registers), this
extreme register pressure caused registers to be spilled to local memory and
corrupted in kernels with high occupancy.
This change scopes subexpression elimination per output parameter group:
1. Accessor accepts an optional group: str | None parameter, and Factor.make_kernels
explicitly tags output accessors with their parameter group ('res' or parameter name).
2. CallData includes group in its hash and equality, preventing Dabseg.data_to_id
from deduplicating identical subexpressions across group boundaries.
3. make_dabseg resets its expression-to-value map across output group transitions.
4. do_cpse partitions calls by call.group so partial associative/commutative
reductions do not cross group boundaries.
5. Function fixers preserve call.group across DAG transformations.
This eliminates cross-barrier live ranges between parameter groups, preventing
stack spilling and register corruption across reduction flushes.
sarlinpe
added a commit
to colmap/colmap
that referenced
this pull request
Oct 2, 2026
Replaces manual kernel workaround with freshly generated CUDA code from SymForce branch incorporating barrier-scoped common subexpression elimination (symforce-org/symforce#475). Cross-barrier register reuse between output parameter groups (e.g. pose -> calib) is eliminated at code generation time.
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 Caspar kernels, factors write to multiple distinct output parameter groups (e.g., residual
out_res, camera poseout_pose_*, camera calibrationout_calib_*, and 3D pointout_point_*). Between these parameter groups, the kernel emits barrier synchronizations and shared memory reduction flushes (FlushSumShared).Previously, Common Subexpression Elimination (CSE) in
make_dabseganddo_cpseoperated globally across the entire factor DAG. When a subexpression (such as-(1 + k*r^2)) computed during the camera pose Jacobian was reused hundreds of lines later in the calibration Jacobian across 12FlushSumSharedbarriers, the variable remained live throughout all intermediate reductions. Under__launch_bounds__(1024, 1)(which limits threads to 64 registers), this extreme register pressure (>1.8 KB stack spill) caused registers to be spilled to local memory and corrupted in kernels with high occupancy.In COLMAP's Caspar bundle adjustment integration (see colmap/colmap#4799), this register corruption caused double-precision calibration Jacobians to evaluate to wildly inaccurate values (~10^4 instead of fractional derivatives).
Proposed Solution
Rather than relying on string heuristics, this PR scopes subexpression elimination by output parameter group:
Accessor: Addedgroup: str | None = Noneparameter.Factor.make_kernels: Explicitly tags output accessors with their parameter group (group=kfor node variables andgroup="res"for residual outputs).CallData&Dabseg: IncludedgroupinCallData's hash and equality, preventingDabseg.data_to_idfrom deduplicating expressions across group boundaries.make_dabseg: Resetsexpr_mapacrossout.grouptransitions during DAG construction.do_cpse: Partitions associative/commutative calls bycall.groupso partial CSE reductions only combine terms within the same group.function_fixers: Preservescall.groupacross DAG transformations.Verification
test_generate_solverinsymforce_caspar_codegen_test.pypasses, confirming that cross-barrier live ranges are eliminated in the test Rosenbrock factor as well (reducing register pressure from 9 to 8 registers inkernel_rosenbrock_res_jac_first.cu).bundle_adjustment_caspar_testpass cleanly in bothf32andf64without needing any manual kernel workarounds.