static_cast to avoid bug on windows (#571)

Signed-off-by: Clement Fuji Tsang <cfujitsang@nvidia.com>
NVIDIAGameWorks · May 26, 2022 · eeab9c8 · eeab9c8
1 parent b52da8f
commit eeab9c8
Show file tree

Hide file tree

Showing 3 changed files with 8 additions and 5 deletions.
diff --git a/kaolin/csrc/render/mesh/deftet_cuda.cu b/kaolin/csrc/render/mesh/deftet_cuda.cu
@@ -139,7 +139,8 @@ __global__ void deftet_sparse_render_forward_cuda_kernel(
 	      scalar_t _w1 = c_edge_x * a_edge_y - c_edge_y * a_edge_x;
 	      scalar_t _w2 = a_edge_x * b_edge_y - a_edge_y * b_edge_x;
 	      scalar_t norm = _w0 + _w1 + _w2;
-	      scalar_t norm_eps = copysignf(eps, norm);
+	      scalar_t norm_eps = copysignf(static_cast<double>(eps),
+                                      static_cast<double>(norm));
 	      w0 = _w0 / (norm + norm_eps);
 	      w1 = _w1 / (norm + norm_eps);
 	      w2 = _w2 / (norm + norm_eps);
@@ -323,7 +324,8 @@ __global__ void deftet_sparse_render_backward_cuda_kernel(
       const scalar_t k1 = s * q - n * t;
       const scalar_t k2 = m * t - s * p;
       scalar_t k3 = m * q - n * p;
-      k3 += copysign(eps, k3);
+      // Need to explicitly cast because there is a bug on windows.
+      k3 += copysign(static_cast<double>(eps), static_cast<double>(k3));
 
       const scalar_t dk1dm = 0;
       const scalar_t dk1dn = -t;

diff --git a/kaolin/csrc/render/mesh/rasterization_cuda.cu b/kaolin/csrc/render/mesh/rasterization_cuda.cu
@@ -138,7 +138,8 @@ __global__ void packed_rasterize_forward_cuda_kernel(
           scalar_t w1 = c_edge_x * a_edge_y - c_edge_y * a_edge_x;
           scalar_t w2 = a_edge_x * b_edge_y - a_edge_y * b_edge_x;
           scalar_t norm = w0 + w1 + w2;
-          norm += copysign(eps, norm);
+          norm += copysign(static_cast<double>(eps),
+                           static_cast<double>(norm));
           w0 /= norm;
           w1 /= norm;
           w2 /= norm;
@@ -318,7 +319,7 @@ __global__ void rasterize_backward_cuda_kernel(
       const scalar_t k1 = s * q - n * t;
       const scalar_t k2 = m * t - s * p;
       scalar_t k3 = m * q - n * p;
-      k3 += copysign(eps, k3);
+      k3 += copysign(static_cast<double>(eps), static_cast<double>(k3));
 
       const scalar_t dk1dm = 0;
       const scalar_t dk1dn = -t;

diff --git a/setup.py b/setup.py
@@ -14,7 +14,7 @@
 import warnings
 
 TORCH_MIN_VER = '1.5.0'
-TORCH_MAX_VER = '1.10.2'
+TORCH_MAX_VER = '1.11.0'
 CYTHON_MIN_VER = '0.29.20'
 INCLUDE_EXPERIMENTAL = os.getenv('KAOLIN_INSTALL_EXPERIMENTAL') is not None
 IGNORE_TORCH_VER = os.getenv('IGNORE_TORCH_VER') is not None