Hessian block assembly + BSR

CMakeLists.txt

@@ -137,16 +137,15 @@ if(BUILD_TESTING)
   add_executable(
     nls_tests
     tests/utils.cpp
+    tests/bsr_expansion.cu
     tests/factor_batch_test.cpp
     tests/problem_test.cpp
     tests/prior_vector_prior_factor_test.cpp
     tests/residual_batch_test.cpp
     tests/loss_function_test.cpp
-    tests/jacobian_ops_test.cpp
     tests/state_batch_ops_test.cpp
     tests/sparse_matrix_test.cpp
-    tests/cusparse_matrix_multiplier_test.cpp
-    tests/fast_matrix_multiplier_test.cpp
+    tests/block_hessian_assembler_test.cpp
     tests/sparse_linear_solver_test.cpp
     tests/dense_linear_solver_test.cpp
     tests/dense_cholesky_solver_test.cpp

LICENSE

@@ -200,4 +200,3 @@
    WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
    See the License for the specific language governing permissions and
    limitations under the License.
-

NOTICE

@@ -40,4 +40,4 @@ This product includes third-party software components:
    https://github.com/NVIDIA/warp
 
 Full license texts for third-party components are provided in the
-third_party/LICENSES directory.
+third_party/LICENSES directory.

cunls/common/cublas_helper.cpp

@@ -15,9 +15,10 @@
  * limitations under the License.
  */
 
+#include "cunls/common/cublas_helper.h"
+
 #include <cublas_v2.h>
 
-#include "cunls/common/cublas_helper.h"
 #include "cunls/common/log.h"
 
 namespace cunls {
@@ -77,4 +78,4 @@ void *cuBLASHandle::GetHandle(cudaStream_t stream) {
   return handle_;
 }
 
-} // namespace cunls
+}  // namespace cunls

cunls/common/cublas_helper.h

@@ -38,17 +38,15 @@ const char *cublasGetErrorString(int status);
  * If the cuBLAS status indicates an error, throws an exception with
  * a descriptive error message.
  */
-#define THROW_ON_CUBLAS_ERROR(status)                                          \
-  CHECK_CUDA_ERROR(status, cublasGetErrorString, true)
+#define THROW_ON_CUBLAS_ERROR(status) CHECK_CUDA_ERROR(status, cublasGetErrorString, true)
 
 /**
  * @brief Macro to log a warning on cuBLAS errors.
  *
  * If the cuBLAS status indicates an error, logs a warning message
  * but does not throw an exception.
  */
-#define WARN_ON_CUBLAS_ERROR(status)                                           \
-  CHECK_CUDA_ERROR(status, cublasGetErrorString, false)
+#define WARN_ON_CUBLAS_ERROR(status) CHECK_CUDA_ERROR(status, cublasGetErrorString, false)
 
 /**
  * @brief RAII wrapper for cuBLAS handle management.
@@ -61,7 +59,7 @@ const char *cublasGetErrorString(int status);
  * Non-copyable: Prevents accidental handle duplication.
  */
 class cuBLASHandle {
-public:
+ public:
   cuBLASHandle() = default;
 
   cuBLASHandle(const cuBLASHandle &) = delete;
@@ -89,9 +87,9 @@ class cuBLASHandle {
    */
   void *GetHandle(cudaStream_t stream);
 
-private:
-  cudaStream_t stream_ = nullptr; ///< Currently associated CUDA stream.
-  void *handle_ = nullptr;        ///< The cuBLAS handle.
+ private:
+  cudaStream_t stream_ = nullptr;  ///< Currently associated CUDA stream.
+  void *handle_ = nullptr;         ///< The cuBLAS handle.
 };
 
-} // namespace cunls
+}  // namespace cunls

cunls/common/cuda_stream.cpp

@@ -30,8 +30,7 @@ namespace cunls {
  * @param sync_on_destroy If true, the stream will be synchronized before
  *                        destruction.
  */
-CudaStream::CudaStream(bool sync_on_destroy)
-    : sync_on_destroy_(sync_on_destroy) {
+CudaStream::CudaStream(bool sync_on_destroy) : sync_on_destroy_(sync_on_destroy) {
   THROW_ON_CUDA_ERROR(cudaStreamCreate(&stream));
 }
 
@@ -51,4 +50,4 @@ CudaStream::~CudaStream() {
 
 /** @brief Returns a reference to the underlying cudaStream_t handle. */
 cudaStream_t &CudaStream::GetStream() { return stream; }
-} // namespace cunls
+}  // namespace cunls

cunls/common/cuda_stream.h

@@ -31,7 +31,7 @@ namespace cunls {
  * Non-copyable to prevent accidental sharing of stream ownership.
  */
 class CudaStream {
-public:
+ public:
   /**
    * @brief Constructs a new CUDA stream.
    *
@@ -55,8 +55,8 @@ class CudaStream {
    */
   cudaStream_t &GetStream();
 
-private:
-  cudaStream_t stream;   ///< The underlying CUDA stream handle.
-  bool sync_on_destroy_; ///< Whether to synchronize on destruction.
+ private:
+  cudaStream_t stream;    ///< The underlying CUDA stream handle.
+  bool sync_on_destroy_;  ///< Whether to synchronize on destruction.
 };
-} // namespace cunls
+}  // namespace cunls

cunls/common/cudss_helper.h

@@ -17,9 +17,9 @@
 
 #pragma once
 
-#include <cstddef>
 #include <cuda_runtime.h>
 
+#include <cstddef>
 #include <mutex>
 #include <vector>
 
@@ -42,17 +42,15 @@ const char *cudssGetErrorString(int status);
  * If the status indicates an error, this macro will throw an exception with
  * a descriptive error message.
  */
-#define THROW_ON_CUDSS_ERROR(status)                                           \
-  CHECK_CUDA_ERROR(status, cudssGetErrorString, true)
+#define THROW_ON_CUDSS_ERROR(status) CHECK_CUDA_ERROR(status, cudssGetErrorString, true)
 
 /**
  * @brief Macro to check cuDSS status and log a warning on error.
  *
  * If the status indicates an error, this macro will log a warning but will
  * not throw an exception.
  */
-#define WARN_ON_CUDSS_ERROR(status)                                            \
-  CHECK_CUDA_ERROR(status, cudssGetErrorString, false)
+#define WARN_ON_CUDSS_ERROR(status) CHECK_CUDA_ERROR(status, cudssGetErrorString, false)
 
 /**
  * @brief Reusable device memory pool used by cuDSS callbacks.
@@ -62,7 +60,7 @@ const char *cudssGetErrorString(int status);
  * the retained capacity of an available block.
  */
 class cuDSSDeviceMemPool {
-public:
+ public:
   cuDSSDeviceMemPool() = default;
 
   cuDSSDeviceMemPool(const cuDSSDeviceMemPool &) = delete;
@@ -91,7 +89,7 @@ class cuDSSDeviceMemPool {
    */
   int Dealloc(void *ptr, size_t size, cudaStream_t stream);
 
-private:
+ private:
   struct Block {
     void *ptr = nullptr;
     size_t capacity = 0;
@@ -105,14 +103,12 @@ class cuDSSDeviceMemPool {
 /**
  * @brief C callback wrapper for cuDSS device allocation.
  */
-int cuDSSDeviceMemPoolAlloc(void *ctx, void **ptr, size_t size,
-                            cudaStream_t stream);
+int cuDSSDeviceMemPoolAlloc(void *ctx, void **ptr, size_t size, cudaStream_t stream);
 
 /**
  * @brief C callback wrapper for cuDSS device deallocation.
  */
-int cuDSSDeviceMemPoolDealloc(void *ctx, void *ptr, size_t size,
-                              cudaStream_t stream);
+int cuDSSDeviceMemPoolDealloc(void *ctx, void *ptr, size_t size, cudaStream_t stream);
 
 /**
  * @brief Installs a cuDSS memory handler backed by a custom pool.
@@ -139,7 +135,7 @@ void DetachcuDSSDeviceMemHandler(void *handle);
  * and automatically destroyed in the destructor.
  */
 class cuDSSHandle {
-public:
+ public:
   cuDSSHandle() = default;
 
   cuDSSHandle(const cuDSSHandle &) = delete;
@@ -161,9 +157,9 @@ class cuDSSHandle {
    */
   void *GetHandle(cudaStream_t stream);
 
-private:
-  cudaStream_t stream_ = nullptr; ///< Currently associated CUDA stream.
-  void *handle_ = nullptr;        ///< The cuDSS handle.
+ private:
+  cudaStream_t stream_ = nullptr;  ///< Currently associated CUDA stream.
+  void *handle_ = nullptr;         ///< The cuDSS handle.
 };
 
 /**
@@ -173,7 +169,7 @@ class cuDSSHandle {
  * or a dense vector and manages its lifecycle.
  */
 class cuDSSDescription {
-public:
+ public:
   /**
    * @brief Constructs a cuDSS matrix descriptor from a CSR sparse matrix.
    *
@@ -200,8 +196,8 @@ class cuDSSDescription {
    */
   void *GetDescription() { return matrix_; }
 
-private:
-  void *matrix_; ///< The cuDSS matrix descriptor.
+ private:
+  void *matrix_;  ///< The cuDSS matrix descriptor.
 };
 
 /**
@@ -211,7 +207,7 @@ class cuDSSDescription {
  * parameters and options.
  */
 class cuDSSConfig {
-public:
+ public:
   /** @brief Constructor that creates a cuDSS configuration object. */
   cuDSSConfig(int reordering_algorithm = 0, int nthreads = 1);
 
@@ -224,8 +220,8 @@ class cuDSSConfig {
    */
   void *GetData() const { return config_; }
 
-private:
-  void *config_ = nullptr; ///< The cuDSS configuration handle.
+ private:
+  void *config_ = nullptr;  ///< The cuDSS configuration handle.
 };
 
 /**
@@ -235,7 +231,7 @@ class cuDSSConfig {
  * and working memory during the factorization and solve phases.
  */
 class cuDSSData {
-public:
+ public:
   cuDSSData() = default;
 
   /** @brief Destructor that releases the cuDSS data object. */
@@ -253,9 +249,9 @@ class cuDSSData {
    */
   void *GetData(void *handle);
 
-private:
-  void *handle_ = nullptr; ///< Associated cuDSS handle.
-  void *data_ = nullptr;   ///< The cuDSS data handle.
+ private:
+  void *handle_ = nullptr;  ///< Associated cuDSS handle.
+  void *data_ = nullptr;    ///< The cuDSS data handle.
 };
 
-} // namespace cunls
+}  // namespace cunls

cunls/common/cusolver_helper.cpp

@@ -15,9 +15,10 @@
  * limitations under the License.
  */
 
+#include "cunls/common/cusolver_helper.h"
+
 #include <cusolverDn.h>
 
-#include "cunls/common/cusolver_helper.h"
 #include "cunls/common/log.h"
 
 namespace cunls {
@@ -85,8 +86,7 @@ cuSolverHandle::cuSolverHandle() {
 }
 
 cuSolverHandle::~cuSolverHandle() {
-  WARN_ON_CUSOLVER_ERROR(
-      cusolverDnDestroy(static_cast<cusolverDnHandle_t>(handle_)));
+  WARN_ON_CUSOLVER_ERROR(cusolverDnDestroy(static_cast<cusolverDnHandle_t>(handle_)));
 }
 
 void *cuSolverHandle::GetHandle(cudaStream_t stream) {
@@ -101,8 +101,7 @@ void *cuSolverHandle::GetHandle(cudaStream_t stream) {
   }
 
   if (handle_ != nullptr) {
-    THROW_ON_CUSOLVER_ERROR(
-        cusolverDnDestroy(static_cast<cusolverDnHandle_t>(handle_)));
+    THROW_ON_CUSOLVER_ERROR(cusolverDnDestroy(static_cast<cusolverDnHandle_t>(handle_)));
   }
 
   stream_ = stream;
@@ -120,8 +119,7 @@ cuSolverInfo::cuSolverInfo() {
 }
 
 cuSolverInfo::~cuSolverInfo() {
-  WARN_ON_CUSOLVER_ERROR(
-      cusolverDnDestroySyevjInfo(static_cast<syevjInfo_t>(info_)));
+  WARN_ON_CUSOLVER_ERROR(cusolverDnDestroySyevjInfo(static_cast<syevjInfo_t>(info_)));
 }
 
-} // namespace cunls
+}  // namespace cunls

cunls/common/cusolver_helper.h

@@ -38,17 +38,15 @@ const char *cusolverGetErrorString(int status);
  * Checks the cuSolver status and throws std::runtime_error with a descriptive
  * message if the status indicates an error.
  */
-#define THROW_ON_CUSOLVER_ERROR(status)                                        \
-  CHECK_CUDA_ERROR(status, cusolverGetErrorString, true)
+#define THROW_ON_CUSOLVER_ERROR(status) CHECK_CUDA_ERROR(status, cusolverGetErrorString, true)
 
 /**
  * @brief Macro that logs a warning if cuSolver operation fails.
  *
  * Checks the cuSolver status and logs a warning message if the status indicates
  * an error, but does not throw an exception.
  */
-#define WARN_ON_CUSOLVER_ERROR(status)                                         \
-  CHECK_CUDA_ERROR(status, cusolverGetErrorString, false)
+#define WARN_ON_CUSOLVER_ERROR(status) CHECK_CUDA_ERROR(status, cusolverGetErrorString, false)
 
 /**
  * @brief RAII wrapper for cuSolver handle management.
@@ -58,7 +56,7 @@ const char *cusolverGetErrorString(int status);
  * a specific CUDA stream when GetHandle is called.
  */
 class cuSolverHandle {
-public:
+ public:
   /// Constructs a cuSolver handle (handle is created lazily on first GetHandle
   /// call)
   cuSolverHandle();
@@ -88,9 +86,9 @@ class cuSolverHandle {
    */
   void *GetHandle(cudaStream_t stream);
 
-private:
-  cudaStream_t stream_ = nullptr; ///< Currently associated CUDA stream
-  void *handle_ = nullptr;        ///< cuSolver handle
+ private:
+  cudaStream_t stream_ = nullptr;  ///< Currently associated CUDA stream
+  void *handle_ = nullptr;         ///< cuSolver handle
 };
 
 /**
@@ -100,7 +98,7 @@ class cuSolverHandle {
  * symmetric eigenvalue decomposition.
  */
 class cuSolverInfo {
-public:
+ public:
   /// Constructs a cuSolver info object
   cuSolverInfo();
   /// Destroys the cuSolver info object
@@ -121,8 +119,8 @@ class cuSolverInfo {
    */
   void *GetInfo() const { return info_; }
 
-private:
-  void *info_ = nullptr; ///< cuSolver eigenvalue solver info handle
+ private:
+  void *info_ = nullptr;  ///< cuSolver eigenvalue solver info handle
 };
 
-} // namespace cunls
+}  // namespace cunls