GPU bug fix (#7)

* [bug] fix typos

* [bug] fix compile error when using single-precision floating-point arithmetic

Author: WYK96

src/scalable_ccd/cuda/ccd.cu

@@ -77,7 +77,7 @@ namespace {
     }
 } // namespace
 
-double
+Scalar
 ccd(const Eigen::MatrixXd& vertices_t0,
     const Eigen::MatrixXd& vertices_t1,
     const Eigen::MatrixXi& edges,

src/scalable_ccd/cuda/ipc_ccd_strategy.hpp

@@ -1,6 +1,7 @@
 #pragma once
 
 #include <Eigen/Core>
+#include <scalable_ccd/scalar.hpp>
 
 namespace scalable_ccd::cuda {
 
@@ -13,13 +14,13 @@ namespace scalable_ccd::cuda {
 /// @param max_iter Maximum number of iterations
 /// @param tolerance Tolerance for the CCD algorithm
 /// @return Time of impact
-double ipc_ccd_strategy(
+Scalar ipc_ccd_strategy(
     const Eigen::MatrixXd& V0,
     const Eigen::MatrixXd& V1,
     const Eigen::MatrixXi& E,
     const Eigen::MatrixXi& F,
-    const double min_distance,
+    const Scalar min_distance,
     const int max_iter,
-    const double tolerance);
+    const Scalar tolerance);
 
 } // namespace scalable_ccd::cuda

src/scalable_ccd/cuda/scalar.cuh

@@ -28,6 +28,8 @@ make_Scalar3(const Scalar& a, const Scalar& b, const Scalar& c)
 
 #else
 
+__device__ constexpr float SCALAR_MAX = scalable_ccd::SCALAR_MAX;
+
 using Scalar2 = float2;
 using Scalar3 = float3;
 

src/scalable_ccd/cuda/utils/device_matrix.cuh

@@ -34,14 +34,23 @@ public:
         thrust::copy(mat.data(), mat.data() + mat.size(), m_data.begin());
     }
 
-#ifndef SCALABLE_CCD_USE_DOUBLE
-    DeviceMatrix(const Eigen::MatrixXd& mat) : DeviceMatrix(mat.cast<T>()) { }
+    template <typename U = T, typename std::enable_if<!std::is_same<U, double>::value>::type* = nullptr>
+    DeviceMatrix(const Eigen::MatrixXd& mat)
+        : m_rows(mat.rows())
+        , m_cols(mat.cols())
+        , m_data(mat.size())
+    {
+        thrust::copy(mat.data(), mat.data() + mat.size(), m_data.begin());
+    }
 
+    template <typename U = T, typename std::enable_if<!std::is_same<U, double>::value>::type* = nullptr>
     void operator=(const Eigen::MatrixXd& mat)
     {
-        return operator=(mat.cast<T>());
+        m_rows = mat.rows();
+        m_cols = mat.cols();
+        m_data.resize(mat.size());
+        thrust::copy(mat.data(), mat.data() + mat.size(), m_data.begin());
     }
-#endif
 
     // TODO: Add function to retrieve the matrix from the device.