Hessian block assembly + BSR - #20
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
💤 Files with no reviewable changes (1)
📝 WalkthroughWalkthroughThe minimizer now assembles normal equations directly from dense per-factor Jacobians. It supports automatic CSR or BSR storage, storage-aware PCG solving, BSR matrix utilities, revised scaling and Python APIs, and expanded CUDA test coverage. ChangesNormal-equation pipeline
Estimated code review effort: 5 (Critical) | ~120 minutes Sequence Diagram(s)sequenceDiagram
participant Problem
participant GaussNewtonMinimizer
participant BlockHessianAssembler
participant NormalEquations
participant BlockSparsePCGSolver
participant State
Problem->>GaussNewtonMinimizer: provide factor graph and state batches
GaussNewtonMinimizer->>BlockHessianAssembler: evaluate dense factor Jacobians
BlockHessianAssembler->>NormalEquations: assemble H and rhs
NormalEquations->>NormalEquations: select CSR or BSR storage
NormalEquations->>BlockSparsePCGSolver: initialize active sparse solver
GaussNewtonMinimizer->>NormalEquations: solve normal equations
NormalEquations->>BlockSparsePCGSolver: run PCG
BlockSparsePCGSolver->>GaussNewtonMinimizer: return tangent step
GaussNewtonMinimizer->>State: apply tangent step
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✨ Finishing Touches 💡 1📝 Generate docstrings 💡
🧪 Generate unit tests (beta)
Comment |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 17
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
python/src/bind_types.cpp (1)
76-96: 🗄️ Data Integrity & Integration | 🟡 Minor | ⚡ Quick winRemove
jacobian_column_normfrom the Python API documentation.The stub, package initializer, examples, and benchmarks contain no removed API names.
docs/sphinx/api/minimizer.rst:495still listsjacobian_column_normas a validColumnScalingvalue.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@python/src/bind_types.cpp` around lines 76 - 96, Remove the stale jacobian_column_norm entry from the ColumnScaling value documentation in the minimizer API reference, while preserving the currently exposed values none and hessian_diagonal in the binding around MinimizerOptions.cunls/minimizer/gauss_newton_minimizer.h (1)
308-321: 📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick winStale
BuildSystemdocumentation in both minimizer headers. TheBuildSystemsignature lost itslhsandrhsoutput parameters and the CSR-only assembly path was replaced byNormalEquations, but both doc blocks still describe the previous interface.
cunls/minimizer/gauss_newton_minimizer.h#L308-L321: remove the "converts to CSR format" wording and the reference to the removedhessian_member; describe assembly throughnormal_equations_and the CSR-or-BSR storage choice.cunls/minimizer/levenberg_marquardt_minimizer.h#L143-L158: delete the@param[out] lhsand@param[out] rhslines, which document parameters the declaration no longer takes.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/gauss_newton_minimizer.h` around lines 308 - 321, Update the BuildSystem documentation in cunls/minimizer/gauss_newton_minimizer.h lines 308-321 to remove CSR-only and hessian_ references, and describe assembly through normal_equations_ with CSR-or-BSR storage. In cunls/minimizer/levenberg_marquardt_minimizer.h lines 143-158, remove the obsolete `@param`[out] lhs and `@param`[out] rhs entries; no other interface changes are needed.
🧹 Nitpick comments (14)
tests/sparse_matrix_test.cpp (1)
146-158: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueClear the column and value vectors before generating.
Line 146 resets
csr_row_offsetswithassign, butcsr_col_idxandcsr_valuesare only appended to.SetUpcalls this method once on a fresh fixture, so the current tests pass.GenerateRandomCSRMatrixis public, so a second call on the same instance produces an inconsistent matrix: the row offsets restart at zero while the column indices retain the previous content. Clear all three vectors so the method is self-contained.♻️ Proposed fix
csr_row_offsets.assign(1, 0); + csr_col_idx.clear(); + csr_values.clear(); for (int row = 0; row < rows; ++row) {🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/sparse_matrix_test.cpp` around lines 146 - 158, Update GenerateRandomCSRMatrix to clear csr_row_offsets, csr_col_idx, and csr_values before generating rows, preserving the existing initialization of row offsets and ensuring repeated calls produce a self-contained consistent matrix.tests/block_hessian_assembler_test.cpp (1)
1125-1127: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueThe comment describes a case the test does not construct.
Line 1125 states that "a coprime mix must fall back to scalar storage".
MakeMixedDimbuilds tangent dimensions 3 and 6, whose gcd is 3, and Line 1127 asserts 3. No coprime mix exists in this test, and no fallback is exercised. Either correct the comment, or add a state-batch pair with coprime tangent dimensions and assert thatChooseHessianBlockSizereturns 1.📝 Proposed comment fix
- // gcd(3, 6) = 3 stays, but a coprime mix must fall back to scalar storage. + // Two batches with tangent dims 3 and 6 tile at gcd(3, 6) = 3. auto mixed = MakeMixedDim(8); EXPECT_EQ(3, ChooseHessianBlockSize(mixed->problem));🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/block_hessian_assembler_test.cpp` around lines 1125 - 1127, Update the test around MakeMixedDim and ChooseHessianBlockSize so the scenario matches its description: either revise the comment to describe the existing 3-and-6 dimensions and expected block size 3, or construct a mixed state-batch pair with coprime tangent dimensions and assert that ChooseHessianBlockSize returns 1. Prefer adding the coprime case if the fallback behavior is intended to be covered.cunls/minimizer/sparse_matrix.h (1)
139-144: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueDocument the second
ComputeWeightedSquaredStepoverload.The Doxygen block above Line 139 documents only the first overload. The second overload at Line 142-144 adds
num_rows,num_cols, andnum_nonzerosand has no documentation. Every other declaration in this header carries a Doxygen block. Add one that explains when a caller supplies precomputed dimensions instead of letting the function derive them.📝 Proposed documentation addition
float ComputeWeightedSquaredStep(cudaStream_t stream, void *handle, const CSRSparseMatrix &matrix, const dvector<float> &step, dvector<uint8_t> &buffer); +/** + * `@brief` Same as above, with precomputed matrix dimensions. + * + * Use this overload on the solver hot path to skip the device query that + * derives the dimensions from `@p` matrix. + * + * `@param` num_rows Row count of `@p` matrix. + * `@param` num_cols Column count of `@p` matrix. + * `@param` num_nonzeros Stored entry count of `@p` matrix. + */ float ComputeWeightedSquaredStep(cudaStream_t stream, void *handle, const CSRSparseMatrix &matrix, int num_rows, int num_cols, int num_nonzeros, const dvector<float> &step, dvector<uint8_t> &buffer);🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/sparse_matrix.h` around lines 139 - 144, Add a Doxygen block for the second ComputeWeightedSquaredStep overload, documenting that callers provide the precomputed num_rows, num_cols, and num_nonzeros values instead of the function deriving them. Keep the existing parameter and return documentation style consistent with the surrounding declarations.cunls/linear_solver/block_sparse_pcg_solver.cu (2)
360-406: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low valueAlign the specialized block sizes between the two dispatch paths.
The thread-per-row path specializes 2, 3, 4, 6, 7, 15 and 16. The warp-per-block-row path specializes 2 through 8. A pose-graph problem with
block_size == 5or8therefore takes the runtime-bkernelBsrMultiplyRowKernel<0>and loses the compile-time unrolling, while the same block size is specialized on the other path. The result stays correct, so this is a performance detail only.♻️ Proposed change to the row-path specializations
LAUNCH_ROW(2); LAUNCH_ROW(3); LAUNCH_ROW(4); + LAUNCH_ROW(5); LAUNCH_ROW(6); LAUNCH_ROW(7); + LAUNCH_ROW(8); LAUNCH_ROW(15); LAUNCH_ROW(16);🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/linear_solver/block_sparse_pcg_solver.cu` around lines 360 - 406, Align the block-size specializations in the thread-per-row dispatch with the warp-per-block-row dispatch by adding cases for block sizes 5 and 8 to the LAUNCH_ROW list in BsrMultiply dispatch. Keep the existing specialized cases and generic fallback unchanged.
1232-1252: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low value
InitializeCommondoes not usestream.The parameter
streamis unused in the body (Lines 1253-1310). The CSR overload passes the stream separately toInitializeCsrSpMV. Either drop the parameter or mark it as intentionally unused, so the signature matches the work the function performs.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/linear_solver/block_sparse_pcg_solver.cu` around lines 1232 - 1252, Update BlockSparsePCGSolver::InitializeCommon to remove its unused stream parameter, then adjust both Initialize overloads to call it without stream while preserving the CSR-specific InitializeCsrSpMV(stream, spd_matrix) call.cunls/minimizer/normal_equations.cu (1)
46-55: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low valueRefresh only values after the first assembly. Preserve one full structural copy after each
Initialize, then copy onlyvalueson later iterations. Keep the generic copy helpers as full-copy operations because reused solves can have different sparsity patterns.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/normal_equations.cu` around lines 46 - 55, Update NormalEquations::Assemble to perform a full bsr_hessian_/csr_hessian_ to bsr_lhs_/csr_lhs_ copy immediately after Initialize, then copy only the matrix values on subsequent assemblies. Track initialization state separately for each storage path, and leave CopyBSRSparseMatrix and CopyCSRSparseMatrix as full-copy helpers.cunls/minimizer/block_hessian_assembler.h (1)
96-101: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueThe
problemparameter ofAssembleis unused.Both
Assembledefinitions incunls/minimizer/block_hessian_assembler.cu(lines 309-325) ignoreproblem. The parameter documents a contract but produces an unused-parameter warning and suggests a data dependency that does not exist. Either drop the parameter or omit the name in the definitions.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/block_hessian_assembler.h` around lines 96 - 101, Remove the unused problem parameter from both Assemble overloads, or omit its name in their definitions while preserving the existing declaration contract. Update the CSR and BSR Assemble definitions consistently so they no longer trigger an unused-parameter warning.cunls/linear_solver/csr_sparse_linear_solver.h (1)
98-109: 📐 Maintainability & Code Quality | 🔵 Trivial | 💤 Low valueOmit parameter names in the default BSR bodies.
Both default implementations ignore every argument. Named parameters can trigger
-Wunused-parameterin builds that enable it. Comment out or drop the names to keep the intent explicit.♻️ Proposed change
- virtual bool Initialize(cudaStream_t stream, const Problem &problem, - const BSRSparseMatrix &spd_matrix, const dvector<float> &rhs, - dvector<float> &result) { + virtual bool Initialize(cudaStream_t /*stream*/, const Problem & /*problem*/, + const BSRSparseMatrix & /*spd_matrix*/, const dvector<float> & /*rhs*/, + dvector<float> & /*result*/) { return false; } /** `@brief` BSR counterpart of Solve; only called when supported. */ - virtual bool Solve(cudaStream_t stream, const BSRSparseMatrix &spd_matrix, - const dvector<float> &rhs, dvector<float> &result) { + virtual bool Solve(cudaStream_t /*stream*/, const BSRSparseMatrix & /*spd_matrix*/, + const dvector<float> & /*rhs*/, dvector<float> & /*result*/) { return false; }🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/linear_solver/csr_sparse_linear_solver.h` around lines 98 - 109, Remove or comment out the parameter names in the default BSR overloads of Initialize and Solve, while preserving their signatures and false return values. Keep the parameter types and virtual methods unchanged so derived implementations remain compatible and unused-parameter warnings are avoided.cunls/minimizer/block_hessian_assembler.cu (2)
334-344: 🩺 Stability & Availability | 🔵 Trivial | ⚡ Quick winGuard
LaunchAssemblyagainst empty scatter maps.Line 338 calls
WriteOffsets().data()unconditionally.HessianStructureBuilder::WriteOffsets()returns an empty vector whenBuildruns withwant_scatter_maps == false, andDeviceVector::data()then returnsnullptr. The kernel would dereference a null pointer at line 122. The twoInitializeoverloads always passtrue, so the path is currently unreachable. Add an explicit check so a future caller cannot reach the kernel with a null map.🛡️ Proposed guard
const int *factor_cols = structure_builder_.FactorCols().data(); const int *write_offsets = structure_builder_.WriteOffsets().data(); + if (write_offsets == nullptr && !plans_.empty()) { + throw std::runtime_error( + "BlockHessianAssembler: scatter maps are missing; Initialize() was not called"); + }🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/block_hessian_assembler.cu` around lines 334 - 344, Guard BlockHessianAssembler::LaunchAssembly against an empty WriteOffsets() result before launching any kernels. Validate that the scatter map is present and return or otherwise fail safely when WriteOffsets().data() is null, ensuring no kernel receives a null write-offset map while preserving normal assembly for initialized scatter maps.
203-222: 🩺 Stability & Availability | 🔵 Trivial | ⚖️ Poor tradeoffQuery the device shared-memory limit instead of hardcoding 48 KiB.
kMaxSharedBytesfixes the budget at 48 KiB. Modern architectures allow a larger dynamic shared-memory block aftercudaFuncSetAttribute(..., cudaFuncAttributeMaxDynamicSharedMemorySizeOptin, ...). With the fixed cap, a factor whose staging arena needs more than 48 KiB throws at line 218 even when the device could run it, and problems with a largeresidual_dim * tangent_dimfail instead of degrading.Read
cudaDevAttrMaxSharedMemoryPerBlockOptinonce and use it as the budget, or document the 48 KiB cap as a deliberate portability floor.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/block_hessian_assembler.cu` around lines 203 - 222, Update PickWarpsPerBlock to obtain the device’s cudaDevAttrMaxSharedMemoryPerBlockOptin limit once and use it instead of the fixed kMaxSharedBytes budget, while retaining the existing block and warp-byte calculations and minimum-warp validation. Ensure the selected dynamic shared-memory size remains compatible with the corresponding kernel opt-in configuration.cunls/minimizer/bsr_matrix.cu (2)
155-165: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick win
FillRowOfTileKerneluses a thread-per-block-row schedule on a skewed workload.The file argues at lines 174-179 that bundle-adjustment Hessians are extremely skewed: a pose block row can hold thousands of tiles while a landmark row holds a handful. This kernel assigns one thread per block row and loops serially over that row's tiles. The threads on pose rows then serialize for the length of the whole kernel, which is the exact pattern the SpMV avoids.
ScaleSymmetricandConvertBSRToCSRboth call this kernel, andScaleSymmetricruns on the solver's per-iteration path.Consider one thread per tile with a
thrust::upper_boundoverrow_offsets, or a warp per block row.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/bsr_matrix.cu` around lines 155 - 165, Rework FillRowOfTileKernel’s workload mapping so tile writes are parallelized instead of assigning one thread to each block row and serializing across its tile range. Use a one-thread-per-tile approach with row identification via thrust::upper_bound, or an equivalent warp-per-row schedule, while preserving the row_of_tile values required by ScaleSymmetric and ConvertBSRToCSR.
333-376: 🚀 Performance & Scalability | 🔵 Trivial | 💤 Low valueThe specialized tile edges differ between the two schedules.
The row path instantiates 2, 3, 4, 6, 7, 15, 16. The warp path instantiates 2, 3, 4, 5, 6, 7, 8. A matrix with
block_size == 5or8falls to the runtime-edge row kernel, and a matrix withblock_size == 15or16falls to the generic warp kernel. Both fallbacks are correct but slower.Align the two lists, or add a comment that explains why the sets differ.
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/bsr_matrix.cu` around lines 333 - 376, Align the specialized block-size cases in the row and warp dispatch switches so both schedules cover the same supported tile sizes, including 5, 8, 15, and 16 as appropriate. Update the LAUNCH_ROW and LAUNCH_WARP lists while preserving their existing kernel signatures and fallback behavior; if the sets must remain different, document the scheduling-specific reason directly beside the switches.cunls/minimizer/hessian_structure.cu (1)
644-651: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick winReplace the device-pointer dereference with a stream-ordered readback.
Line 648 dereferences the iterator returned by
thrust::max_element. That dereference performs a blockingcudaMemcpyon the legacy default stream, not onstream. The surrounding work is enqueued onstreamwithpar_nosync. The pattern relies on Thrust's internal synchronization for correctness, and it adds a hidden host block on the setup path.The function already keeps
pinned_buf_with two slots. Use one slot and copy both scalars with stream-orderedcudaMemcpyAsync.♻️ Proposed change
if (num_block_rows > 0) { thrust::device_ptr<int> counts(block_row_counts.data()); - // Peak row length before the scan destroys it; the SpMV picks its schedule - // from this. - output.max_tiles_per_row = *thrust::max_element(stream_policy, counts, counts + num_block_rows); + // Peak row length before the scan destroys it; the SpMV picks its schedule + // from this. + auto max_it = thrust::max_element(stream_policy, counts, counts + num_block_rows); + if (pinned_buf_.size() < 2) { + pinned_buf_.resize(2); + } + THROW_ON_CUDA_ERROR(cudaMemcpyAsync(pinned_buf_.data() + 1, thrust::raw_pointer_cast(max_it), + sizeof(int), cudaMemcpyDeviceToHost, stream)); + THROW_ON_CUDA_ERROR(cudaStreamSynchronize(stream)); + output.max_tiles_per_row = pinned_buf_[1]; thrust::device_ptr<int> offsets(output.row_offsets.data()); thrust::inclusive_scan(stream_policy, counts, counts + num_block_rows, offsets + 1); }🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/hessian_structure.cu` around lines 644 - 651, In the block handling num_block_rows in the surrounding setup function, replace the blocking dereference of thrust::max_element with a stream-ordered cudaMemcpyAsync readback using one slot of pinned_buf_. Copy both required scalar values into that pinned slot on stream, then consume the results only after the stream-ordered work is complete; preserve the existing max_tiles_per_row assignment and row-offset scan behavior.cunls/minimizer/bsr_matrix.h (1)
47-50: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winUse
kMaxHessianBlockSizefor the default argument.The literal
16appears as the default formax_block_sizeon line 47, askMaxHessianBlockSizeon line 50, and again askMaxBlockSizeinsideBsrMultiplyWarpGenericKernelincunls/minimizer/bsr_matrix.cu(line 243). The generic kernel sizes a register array from its copy. If one value changes and the others do not, the generic kernel overrunsacc.Move the constant above the declaration and reference it in all three places.
♻️ Proposed change
+/** `@brief` Largest tile edge the block SpMV supports. */ +constexpr int kMaxHessianBlockSize = 16; + -int ChooseHessianBlockSize(const Problem &problem, int max_block_size = 16); - -/** `@brief` Largest tile edge the block SpMV supports. */ -constexpr int kMaxHessianBlockSize = 16; +int ChooseHessianBlockSize(const Problem &problem, int max_block_size = kMaxHessianBlockSize);🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/bsr_matrix.h` around lines 47 - 50, Move kMaxHessianBlockSize above ChooseHessianBlockSize, use it as the max_block_size default argument, and update BsrMultiplyWarpGenericKernel’s kMaxBlockSize to reference the same constant instead of duplicating 16.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cunls/linear_solver/block_sparse_pcg_solver.cu`:
- Around line 634-643: Update ExtractScalarJacobi so its diagonal flooring
preserves the sign of d, matching ExtractAndFactorBlockDiagonalsKernel,
ExtractAndFactorGenericKernel, and ExtractAndFactorPerThreadKernel. Replace the
absolute-value-only behavior with the existing sign-preserving floor logic while
retaining pivot_floor handling for zero and small-magnitude diagonals.
- Around line 267-306: Add a host-side validation guard in LaunchBsrMultiply to
reject block_size values greater than 16 before dispatching
BsrMultiplyWarpGenericKernel, while preserving existing handling for supported
sizes and other dispatch paths.
In `@cunls/linear_solver/block_sparse_pcg_solver.h`:
- Around line 231-243: Move the existing CSR Solve documentation block from
before SupportsBlockStorage() to immediately precede the CSRSparseMatrix
overload of Solve in the block solver declaration. Leave the
SupportsBlockStorage(), Initialize(), and BSRSparseMatrix Solve declarations
without that PCG-loop documentation.
In `@cunls/minimizer/bsr_matrix.cu`:
- Around line 238-276: Prevent BsrMultiplyWarpGenericKernel from indexing acc
beyond its fixed capacity by explicitly rejecting block_size values greater than
kMaxBlockSize before launching the kernel. Apply the same validation in the
LaunchBsrMultiply path so generic-kernel selection cannot proceed with oversized
blocks, and derive the local array bound from kMaxBlockSize rather than
duplicating the literal 16.
- Around line 497-523: Update ConvertBSRToCSR so ExpandTilesToCSRKernel is
launched only when input.values.size() is nonzero; preserve row-offset
construction and the existing empty-matrix handling, while still producing valid
empty CSR output when the BSR matrix has rows but no tiles.
- Around line 41-54: Update FillExpandedRowOffsetsKernel so the terminating row
(row == num_rows) does not read block_row_offsets[block_row + 1]; handle that
boundary case using the valid final block-row offset while preserving the
existing calculation for rows before num_rows.
In `@cunls/minimizer/gauss_newton_minimizer.h`:
- Around line 356-368: Remove the obsolete “After the unscaled Hessian copy is
in lhs and rhs holds b = -J^T r,” fragment from the
ApplyColumnScalingToNormalEquations documentation, leaving the current
Assemble()-based description intact.
In `@cunls/minimizer/hessian_structure.cu`:
- Around line 501-521: Initialize write_offsets_ to -1 before either early
return in both HessianStructureBuilder::Build overloads, including the overload
near the second cited block. After resizing write_offsets_, issue an
asynchronous device memset on the provided stream when the buffer is non-empty,
preserving the existing normal-path behavior while ensuring num_cols == 0 or
total_pairs_ == 0 cannot leave entries uninitialized.
- Around line 547-562: Add an overflow guard immediately after reading total_nnz
in the Hessian construction flow, using the existing CUDA/error-checking
mechanism to reject a negative or otherwise invalid 32-bit count before resizing
output.col_ids and output.values. Keep the inclusive scan and readback
unchanged, and ensure the failure clearly identifies CSR non-zero count
overflow.
In `@cunls/minimizer/levenberg_marquardt_minimizer.cpp`:
- Around line 70-87: Update both LevenbergMarquardtMinimizer convergence
overrides, CheckConvergence and EvaluateAndCheckConvergence, to use their step
parameter for every squared-step reduction and related size calculation instead
of the step_ member. Preserve the existing reduction flow while ensuring all
computations reflect the vector supplied by the caller.
- Around line 83-84: Update both convergence checks in the Levenberg–Marquardt
minimizer to use matching damping coordinates: pair the original Hessian
diagonal with the scaled step dx, or pair diagonal_ representing diag(S H S)
with the unscaled step z. Ensure ComputeWeightedSquaredStepAsync and the
corresponding convergence calculation use the same consistent pairing.
In `@cunls/minimizer/llms.txt`:
- Around line 17-23: Update the `normal_equations.h` documentation to state that
`NormalEquations` centralizes storage-layout selection and owns the layout,
rather than claiming it is the only place CSR/BSR is named or branched. Also
mention that storage-specific assembly overloads in `block_hessian_assembler.h`
and storage-aware solver branches remain outside `NormalEquations`.
In `@cunls/minimizer/sparse_matrix.cu`:
- Around line 53-67: Update ScaleSymmetricCSR, ExtractDiagonal, and
AddScaledDiagonal to detect zero-row CSR metadata and return before launching
any zero-work CUDA kernels; AddScaledDiagonal must still copy the input matrix
before returning. Update ExtractMatrixMetadata’s contract for the valid empty
CSR case (row_offsets = {0}), and add tests covering normal-equation scaling,
diagonal extraction, and damping.
In `@cunls/state/state_batch_ops.cu`:
- Around line 127-130: Replace the incorrect `@copydoc` target above
ComputeStateBlockColumnOffsets with the appropriate documentation reference for
this free function, or remove the directive if no matching documented symbol
exists; do not reference StateBatchOps::InitUpdatesVector.
In `@docs/sphinx/api/minimizer.rst`:
- Around line 149-151: Remove the remaining JacobianColumnNorm references in the
column-scaling theory note around S_{jj} and the pycunls.ColumnScaling enum
documentation. Update both sections to list only the supported HessianDiagonal
behavior and values, matching the bindings exposed by bind_types.cpp.
In `@tests/block_hessian_assembler_test.cpp`:
- Around line 977-980: Replace the exact LastIterations equality assertion in
the CSR/BSR solver comparison with an assertion that their iteration counts
differ by no more than an appropriate bounded tolerance, while preserving the
existing recurrence comparison and solver behavior checks.
- Around line 1229-1250: Update MatchesReferenceWithColumnScaling so the
reference SystemBuilder uses scalar storage while the block builder retains
block storage, then assert each builder’s UsesBlockStorage() result to verify
the intended layouts. Keep column scaling enabled and compare the block output
against the scalar reference, ensuring the test exercises distinct assembly
paths rather than comparing identical configurations.
---
Outside diff comments:
In `@cunls/minimizer/gauss_newton_minimizer.h`:
- Around line 308-321: Update the BuildSystem documentation in
cunls/minimizer/gauss_newton_minimizer.h lines 308-321 to remove CSR-only and
hessian_ references, and describe assembly through normal_equations_ with
CSR-or-BSR storage. In cunls/minimizer/levenberg_marquardt_minimizer.h lines
143-158, remove the obsolete `@param`[out] lhs and `@param`[out] rhs entries; no
other interface changes are needed.
In `@python/src/bind_types.cpp`:
- Around line 76-96: Remove the stale jacobian_column_norm entry from the
ColumnScaling value documentation in the minimizer API reference, while
preserving the currently exposed values none and hessian_diagonal in the binding
around MinimizerOptions.
---
Nitpick comments:
In `@cunls/linear_solver/block_sparse_pcg_solver.cu`:
- Around line 360-406: Align the block-size specializations in the
thread-per-row dispatch with the warp-per-block-row dispatch by adding cases for
block sizes 5 and 8 to the LAUNCH_ROW list in BsrMultiply dispatch. Keep the
existing specialized cases and generic fallback unchanged.
- Around line 1232-1252: Update BlockSparsePCGSolver::InitializeCommon to remove
its unused stream parameter, then adjust both Initialize overloads to call it
without stream while preserving the CSR-specific InitializeCsrSpMV(stream,
spd_matrix) call.
In `@cunls/linear_solver/csr_sparse_linear_solver.h`:
- Around line 98-109: Remove or comment out the parameter names in the default
BSR overloads of Initialize and Solve, while preserving their signatures and
false return values. Keep the parameter types and virtual methods unchanged so
derived implementations remain compatible and unused-parameter warnings are
avoided.
In `@cunls/minimizer/block_hessian_assembler.cu`:
- Around line 334-344: Guard BlockHessianAssembler::LaunchAssembly against an
empty WriteOffsets() result before launching any kernels. Validate that the
scatter map is present and return or otherwise fail safely when
WriteOffsets().data() is null, ensuring no kernel receives a null write-offset
map while preserving normal assembly for initialized scatter maps.
- Around line 203-222: Update PickWarpsPerBlock to obtain the device’s
cudaDevAttrMaxSharedMemoryPerBlockOptin limit once and use it instead of the
fixed kMaxSharedBytes budget, while retaining the existing block and warp-byte
calculations and minimum-warp validation. Ensure the selected dynamic
shared-memory size remains compatible with the corresponding kernel opt-in
configuration.
In `@cunls/minimizer/block_hessian_assembler.h`:
- Around line 96-101: Remove the unused problem parameter from both Assemble
overloads, or omit its name in their definitions while preserving the existing
declaration contract. Update the CSR and BSR Assemble definitions consistently
so they no longer trigger an unused-parameter warning.
In `@cunls/minimizer/bsr_matrix.cu`:
- Around line 155-165: Rework FillRowOfTileKernel’s workload mapping so tile
writes are parallelized instead of assigning one thread to each block row and
serializing across its tile range. Use a one-thread-per-tile approach with row
identification via thrust::upper_bound, or an equivalent warp-per-row schedule,
while preserving the row_of_tile values required by ScaleSymmetric and
ConvertBSRToCSR.
- Around line 333-376: Align the specialized block-size cases in the row and
warp dispatch switches so both schedules cover the same supported tile sizes,
including 5, 8, 15, and 16 as appropriate. Update the LAUNCH_ROW and LAUNCH_WARP
lists while preserving their existing kernel signatures and fallback behavior;
if the sets must remain different, document the scheduling-specific reason
directly beside the switches.
In `@cunls/minimizer/bsr_matrix.h`:
- Around line 47-50: Move kMaxHessianBlockSize above ChooseHessianBlockSize, use
it as the max_block_size default argument, and update
BsrMultiplyWarpGenericKernel’s kMaxBlockSize to reference the same constant
instead of duplicating 16.
In `@cunls/minimizer/hessian_structure.cu`:
- Around line 644-651: In the block handling num_block_rows in the surrounding
setup function, replace the blocking dereference of thrust::max_element with a
stream-ordered cudaMemcpyAsync readback using one slot of pinned_buf_. Copy both
required scalar values into that pinned slot on stream, then consume the results
only after the stream-ordered work is complete; preserve the existing
max_tiles_per_row assignment and row-offset scan behavior.
In `@cunls/minimizer/normal_equations.cu`:
- Around line 46-55: Update NormalEquations::Assemble to perform a full
bsr_hessian_/csr_hessian_ to bsr_lhs_/csr_lhs_ copy immediately after
Initialize, then copy only the matrix values on subsequent assemblies. Track
initialization state separately for each storage path, and leave
CopyBSRSparseMatrix and CopyCSRSparseMatrix as full-copy helpers.
In `@cunls/minimizer/sparse_matrix.h`:
- Around line 139-144: Add a Doxygen block for the second
ComputeWeightedSquaredStep overload, documenting that callers provide the
precomputed num_rows, num_cols, and num_nonzeros values instead of the function
deriving them. Keep the existing parameter and return documentation style
consistent with the surrounding declarations.
In `@tests/block_hessian_assembler_test.cpp`:
- Around line 1125-1127: Update the test around MakeMixedDim and
ChooseHessianBlockSize so the scenario matches its description: either revise
the comment to describe the existing 3-and-6 dimensions and expected block size
3, or construct a mixed state-batch pair with coprime tangent dimensions and
assert that ChooseHessianBlockSize returns 1. Prefer adding the coprime case if
the fallback behavior is intended to be covered.
In `@tests/sparse_matrix_test.cpp`:
- Around line 146-158: Update GenerateRandomCSRMatrix to clear csr_row_offsets,
csr_col_idx, and csr_values before generating rows, preserving the existing
initialization of row offsets and ensuring repeated calls produce a
self-contained consistent matrix.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Enterprise
Run ID: 1ee2a0e7-504c-4df1-9cc7-6ad7b330d9ec
📒 Files selected for processing (51)
CMakeLists.txtcunls/common/llms.txtcunls/common/types.hcunls/linear_solver/block_sparse_pcg_solver.cucunls/linear_solver/block_sparse_pcg_solver.hcunls/linear_solver/csr_sparse_linear_solver.hcunls/linear_solver/cudss_sparse_linear_solver.hcunls/linear_solver/dense_cholesky_solver.hcunls/linear_solver/dense_linear_solver.hcunls/linear_solver/dense_qr_solver.hcunls/minimizer/CMakeLists.txtcunls/minimizer/block_hessian_assembler.cucunls/minimizer/block_hessian_assembler.hcunls/minimizer/bsr_matrix.cucunls/minimizer/bsr_matrix.hcunls/minimizer/cusparse_matrix_multiplier.cppcunls/minimizer/cusparse_matrix_multiplier.hcunls/minimizer/fast_matrix_multiplier.cucunls/minimizer/fast_matrix_multiplier.hcunls/minimizer/gauss_newton_minimizer.cucunls/minimizer/gauss_newton_minimizer.hcunls/minimizer/hessian_structure.cucunls/minimizer/hessian_structure.hcunls/minimizer/jacobian_ops.cucunls/minimizer/levenberg_marquardt_minimizer.cppcunls/minimizer/levenberg_marquardt_minimizer.hcunls/minimizer/llms.txtcunls/minimizer/minimizer_state.cucunls/minimizer/minimizer_state.hcunls/minimizer/normal_equations.cucunls/minimizer/normal_equations.hcunls/minimizer/residual_batch.cucunls/minimizer/sparse_matrix.cucunls/minimizer/sparse_matrix.hcunls/minimizer/sparse_matrix_multiplier.cppcunls/minimizer/sparse_matrix_multiplier.hcunls/state/state_batch_ops.cucunls/state/state_batch_ops.hdocs/sphinx/api/common.rstdocs/sphinx/api/minimizer.rstpython/pycunls/__init__.pypython/pycunls/_pycunls_core.pyipython/src/bind_types.cpppython/tests/test_minimizer.pytests/block_hessian_assembler_test.cpptests/cusparse_matrix_multiplier_test.cpptests/fast_matrix_multiplier_test.cpptests/gauss_newton_test.cpptests/jacobian_ops_test.cpptests/sparse_matrix_test.cpptests/utils.h
💤 Files with no reviewable changes (13)
- python/pycunls/_pycunls_core.pyi
- cunls/minimizer/sparse_matrix_multiplier.cpp
- cunls/minimizer/sparse_matrix_multiplier.h
- cunls/minimizer/fast_matrix_multiplier.h
- cunls/minimizer/jacobian_ops.cu
- tests/cusparse_matrix_multiplier_test.cpp
- cunls/minimizer/fast_matrix_multiplier.cu
- cunls/minimizer/cusparse_matrix_multiplier.h
- tests/jacobian_ops_test.cpp
- tests/fast_matrix_multiplier_test.cpp
- python/pycunls/init.py
- cunls/minimizer/cusparse_matrix_multiplier.cpp
- tests/gauss_newton_test.cpp
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (1)
cunls/minimizer/block_hessian_assembler.h (1)
44-49: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winUpdate the storage documentation for the BSR API.
Line 44 through Line 49 state that storage remains plain CSR and that no block-format Hessian is needed. Line 76 through Line 94 add BSR initialization and assembly overloads. Describe both CSR and BSR output contracts in the class comment.
Proposed documentation update
- * Storage stays plain CSR: the pattern produced by ComputeHessianStructure + * Output storage can be scalar CSR or uniform block BSR. The pattern produced + * by ComputeHessianStructure🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/minimizer/block_hessian_assembler.h` around lines 44 - 49, Update the class-level storage documentation near the existing “plain CSR” description to document both output contracts: the existing CSR layout and the BSR layout supported by the BSR initialization and assembly overloads. Clarify that block-format Hessian storage is available through the BSR API while preserving the existing CSR behavior.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@cunls/minimizer/sparse_matrix.cu`:
- Around line 416-424: Replace the static thread_local spmv_result in
ComputeWeightedSquaredStepAsync with a caller-owned dvector<float>& scratch
parameter, update its declaration in sparse_matrix.h, and pass the existing
NormalEquations-owned scratch buffer from the normal_equations.cu call site,
matching the BSR overload’s ownership pattern.
---
Nitpick comments:
In `@cunls/minimizer/block_hessian_assembler.h`:
- Around line 44-49: Update the class-level storage documentation near the
existing “plain CSR” description to document both output contracts: the existing
CSR layout and the BSR layout supported by the BSR initialization and assembly
overloads. Clarify that block-format Hessian storage is available through the
BSR API while preserving the existing CSR behavior.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Enterprise
Run ID: b723a21c-2cc8-4685-817f-d093ac5f7b88
📒 Files selected for processing (60)
CMakeLists.txtLICENSENOTICEcunls/common/cublas_helper.cppcunls/common/cublas_helper.hcunls/common/cuda_stream.cppcunls/common/cuda_stream.hcunls/common/cudss_helper.hcunls/common/cusolver_helper.cppcunls/common/cusolver_helper.hcunls/common/cusparse_helper.cppcunls/common/cusparse_helper.hcunls/common/device_vector.hcunls/common/helper.hcunls/common/log.cppcunls/common/log.hcunls/common/pinned_vector.hcunls/common/profiler.cppcunls/common/profiler.hcunls/common/type_traits.hcunls/common/types.hcunls/common/utils.cppcunls/common/utils.hcunls/factor/factor_batch.hcunls/factor/information_factor_batch.cppcunls/factor/information_factor_batch.hcunls/factor/pnp_factor_batch.cucunls/factor/pnp_factor_batch.hcunls/factor/point_to_plane_factor_batch.cucunls/linear_solver/block_sparse_pcg_solver.cucunls/linear_solver/block_sparse_pcg_solver.hcunls/linear_solver/csr_sparse_linear_solver.hcunls/linear_solver/cudss_sparse_linear_solver.hcunls/linear_solver/dense_cholesky_solver.hcunls/linear_solver/dense_linear_solver.hcunls/linear_solver/dense_qr_solver.hcunls/minimizer/block_hessian_assembler.cucunls/minimizer/block_hessian_assembler.hcunls/minimizer/bsr_matrix.cucunls/minimizer/bsr_matrix.hcunls/minimizer/gauss_newton_minimizer.cucunls/minimizer/gauss_newton_minimizer.hcunls/minimizer/hessian_structure.cucunls/minimizer/hessian_structure.hcunls/minimizer/levenberg_marquardt_minimizer.cppcunls/minimizer/levenberg_marquardt_minimizer.hcunls/minimizer/minimizer_state.hcunls/minimizer/normal_equations.cucunls/minimizer/normal_equations.hcunls/minimizer/residual_batch.cucunls/minimizer/sparse_matrix.cucunls/minimizer/sparse_matrix.hcunls/state/state_batch_ops.cucunls/state/state_batch_ops.hpython/src/bind_types.cpptests/block_hessian_assembler_test.cpptests/bsr_expansion.cutests/bsr_expansion.htests/sparse_matrix_test.cpptests/utils.h
💤 Files with no reviewable changes (3)
- LICENSE
- cunls/minimizer/levenberg_marquardt_minimizer.h
- cunls/minimizer/gauss_newton_minimizer.cu
🚧 Files skipped from review as they are similar to previous changes (22)
- tests/utils.h
- cunls/state/state_batch_ops.h
- cunls/linear_solver/dense_cholesky_solver.h
- CMakeLists.txt
- cunls/linear_solver/dense_linear_solver.h
- cunls/linear_solver/csr_sparse_linear_solver.h
- cunls/minimizer/minimizer_state.h
- cunls/state/state_batch_ops.cu
- cunls/minimizer/hessian_structure.h
- cunls/minimizer/normal_equations.h
- python/src/bind_types.cpp
- cunls/linear_solver/cudss_sparse_linear_solver.h
- cunls/minimizer/block_hessian_assembler.cu
- cunls/minimizer/normal_equations.cu
- cunls/linear_solver/block_sparse_pcg_solver.h
- cunls/linear_solver/dense_qr_solver.h
- tests/block_hessian_assembler_test.cpp
- cunls/minimizer/hessian_structure.cu
- cunls/minimizer/residual_batch.cu
- cunls/minimizer/gauss_newton_minimizer.h
- cunls/common/types.h
- cunls/linear_solver/block_sparse_pcg_solver.cu
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
tests/block_hessian_assembler_test.cpp (1)
1291-1309: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winAssert the system is actually empty, not just zero-valued.
AllConstantStatesProduceZeroSystemmakes every state block constant, sonum_colsis 0 andsnapshot.valuesandsnapshot.rhsare both expected to be empty. Theforloops at Lines 1303 and 1306 iterate zero times over empty vectors, soEXPECT_EQ(v, 0.f)never executes. The test currently passes even if assembly produced a corrupted but empty-looking result, because the loops check nothing when the containers are empty.Add explicit size assertions so the test verifies the "zero system" claim in its name.
🔧 Proposed fix
SystemSnapshot snapshot = Snapshot(block, stream.GetStream()); + EXPECT_TRUE(snapshot.values.empty()); + EXPECT_TRUE(snapshot.rhs.empty()); for (float v : snapshot.values) { EXPECT_EQ(v, 0.f); } for (float v : snapshot.rhs) { EXPECT_EQ(v, 0.f); }🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tests/block_hessian_assembler_test.cpp` around lines 1291 - 1309, Update AllConstantStatesProduceZeroSystem to explicitly assert that snapshot.values and snapshot.rhs are empty before or alongside the existing value checks. Use ASSERT/EXPECT on both container sizes so the test verifies the system has zero entries rather than relying on loops that may execute zero times.
♻️ Duplicate comments (1)
cunls/linear_solver/block_sparse_pcg_solver.cu (1)
644-656: 🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winPivot-sign flooring is still inconsistent between the scalar and block Jacobi paths.
ExtractScalarJacobiat Line 655 storesfmaxf(fabsf(d), pivot_floor), which discards the sign of a negative diagonal entry. Every block-path factorization kernel (ExtractAndFactorBlockDiagonalsKernelat Lines 481-483,ExtractAndFactorGenericKernelat Lines 546-548,ExtractAndFactorPerThreadKernelat Lines 619-621) applies a sign-preserving floor instead:d = (d >= 0.f) ? pivot_floor : -pivot_floor;.For a strictly SPD Hessian the two agree. For an indefinite or damped-then-scaled Hessian, the scalar (
B == 1) preconditioner flips the sign of that diagonal component, changing the PCG search direction relative to theB > 1layout. This was flagged in a previous review round and remains unresolved in this revision.🔧 Proposed fix
float d = 0.f; GatherTileRow(row_off, col_idx, values, block_storage, block_size, row_start + idx, 1, 0, &d); - factors[factor_offset + idx] = fmaxf(fabsf(d), pivot_floor); + if (fabsf(d) < pivot_floor) { + d = (d >= 0.f) ? pivot_floor : -pivot_floor; + } + factors[factor_offset + idx] = d;🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@cunls/linear_solver/block_sparse_pcg_solver.cu` around lines 644 - 656, Update ExtractScalarJacobi to apply the same sign-preserving pivot floor as the block factorization kernels: preserve d when its magnitude exceeds pivot_floor, otherwise use pivot_floor with d’s original sign. Keep the factor write and scalar Jacobi flow unchanged.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@tests/block_hessian_assembler_test.cpp`:
- Around line 1291-1309: Update AllConstantStatesProduceZeroSystem to explicitly
assert that snapshot.values and snapshot.rhs are empty before or alongside the
existing value checks. Use ASSERT/EXPECT on both container sizes so the test
verifies the system has zero entries rather than relying on loops that may
execute zero times.
---
Duplicate comments:
In `@cunls/linear_solver/block_sparse_pcg_solver.cu`:
- Around line 644-656: Update ExtractScalarJacobi to apply the same
sign-preserving pivot floor as the block factorization kernels: preserve d when
its magnitude exceeds pivot_floor, otherwise use pivot_floor with d’s original
sign. Keep the factor write and scalar Jacobi flow unchanged.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Enterprise
Run ID: f672318c-313c-48b3-ac56-7da4343b90cb
📒 Files selected for processing (8)
cunls/linear_solver/block_sparse_pcg_solver.cucunls/minimizer/bsr_matrix.cucunls/minimizer/normal_equations.cucunls/minimizer/normal_equations.hcunls/minimizer/sparse_matrix.cucunls/minimizer/sparse_matrix.htests/block_hessian_assembler_test.cpptests/sparse_matrix_test.cpp
🚧 Files skipped from review as they are similar to previous changes (5)
- cunls/minimizer/normal_equations.h
- cunls/minimizer/bsr_matrix.cu
- cunls/minimizer/normal_equations.cu
- cunls/minimizer/sparse_matrix.h
- cunls/minimizer/sparse_matrix.cu
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (1)
docs/sphinx/api/minimizer.rst (1)
386-402: 🗄️ Data Integrity & Integration | 🟡 Minor | ⚡ Quick winScope the BSR and conversion claims.
The text first states that the Hessian is stored as BSR, but Lines 395-401 document an automatic CSR fallback. Also, Line 402 says that no conversion occurs on the solver path, while Lines 137-141 and 454-457 state that dense backends convert CSR to dense.
State that BSR is used only when block storage is selected, and limit the no-conversion claim to BSR-to-CSR conversion.
Proposed documentation fix
-The Hessian of a factor graph is block structured, so it is stored as BSR — one -column index per dense tile instead of one per scalar entry. That is bandwidth -the iterative solver's SpMV no longer has to move. +When block storage is selected, the Hessian is stored as BSR — one column index +per dense tile instead of one per scalar entry. This reduces the bandwidth that +the iterative solver's SpMV must move. ... -No conversion is ever performed on the solver path. +No BSR-to-CSR conversion is performed on the solver path. Dense backends may +still convert CSR to dense as documented above.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@docs/sphinx/api/minimizer.rst` around lines 386 - 402, Update the documentation paragraph to state that the Hessian uses BSR only when block storage is selected, otherwise it remains scalar CSR. Revise the no-conversion statement so it specifically claims that the solver path performs no BSR-to-CSR conversion, while preserving the documented dense-backend CSR-to-dense conversions elsewhere.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Outside diff comments:
In `@docs/sphinx/api/minimizer.rst`:
- Around line 386-402: Update the documentation paragraph to state that the
Hessian uses BSR only when block storage is selected, otherwise it remains
scalar CSR. Revise the no-conversion statement so it specifically claims that
the solver path performs no BSR-to-CSR conversion, while preserving the
documented dense-backend CSR-to-dense conversions elsewhere.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yml
Review profile: CHILL
Plan: Enterprise
Run ID: 2df04fbb-9c9e-4172-88f3-4a125b2f759e
📒 Files selected for processing (4)
cunls/linear_solver/block_sparse_pcg_solver.cucunls/linear_solver/block_sparse_pcg_solver.hcunls/state/state_batch_ops.cudocs/sphinx/api/minimizer.rst
🚧 Files skipped from review as they are similar to previous changes (3)
- cunls/state/state_batch_ops.cu
- cunls/linear_solver/block_sparse_pcg_solver.h
- cunls/linear_solver/block_sparse_pcg_solver.cu
1 participant
Summary by CodeRabbit