Fixes to #88, #89

nvdiffrast/common/cudaraster/impl/BinRaster.inl

@@ -212,6 +212,7 @@ __device__ __inline__ void binRasterImpl(const CRParams p)
                 hiy = add_clamp_0_x((v0y + max_max(d01y, 0, d02y)) >> binLog, 0, p.heightBins - 1);
 
                 U32 bit = 1 << threadIdx.x;
+#if __CUDA_ARCH__ >= 700
                 bool multi = (hix != lox || hiy != loy);
                 if (!__any_sync(hasTriMask, multi))
                 {
@@ -220,6 +221,7 @@ __device__ __inline__ void binRasterImpl(const CRParams p)
                     s_outMask[threadIdx.y][binIdx] = mask;
                     __syncwarp(hasTriMask);
                 } else
+#endif
                 {
                     bool complex = (hix > lox+1 || hiy > loy+1);
                     if (!__any_sync(hasTriMask, complex))

nvdiffrast/common/cudaraster/impl/Buffer.cpp

@@ -23,7 +23,7 @@ Buffer::Buffer(void)
 Buffer::~Buffer(void)
 {
     if (m_gpuPtr)
-        NVDR_CHECK_CUDA_ERROR(cudaFree(m_gpuPtr));
+        cudaFree(m_gpuPtr); // Don't throw an exception.
 }
 
 //------------------------------------------------------------------------

nvdiffrast/common/cudaraster/impl/TriangleSetup.inl

@@ -80,6 +80,12 @@ __device__ __inline__ bool prepareTriangle(
         S32 e0 = t0.x * t1.y - t0.y * t1.x;
         S32 e1 = t1.x * t2.y - t1.y * t2.x;
         S32 e2 = t2.x * t0.y - t2.y * t0.x;
+        if (area < 0)
+        {
+            e0 = -e0;
+            e1 = -e1;
+            e2 = -e2;
+        }
 
         if (e0 < 0 || e1 < 0 || e2 < 0)
         {
@@ -92,6 +98,12 @@ __device__ __inline__ bool prepareTriangle(
             e0 = t0.x * t1.y - t0.y * t1.x;
             e1 = t1.x * t2.y - t1.y * t2.x;
             e2 = t2.x * t0.y - t2.y * t0.x;
+            if (area < 0)
+            {
+                e0 = -e0;
+                e1 = -e1;
+                e2 = -e2;
+            }
 
             if (e0 < 0 || e1 < 0 || e2 < 0)
                 return false; // Between pixels.