         __inline__ __device__
         scalar warpReduceSum(scalar val)
         {
-            #if (__CUDA_ARCH__ < 700) || (CUDART_VERSION < 9000)
+            #if (CUDART_VERSION >= 9000)
+            for (int offset = warpSize/2; offset > 0; offset /= 2)
+                val += __shfl_down_sync(0xffffffff, val, offset);
+            #else
             for (int offset = warpSize/2; offset > 0; offset /= 2)
                 val += __shfl_down(val, offset);
-            #else
-            for (int offset = warpSize/2; offset > 0; offset /= 2)
-                val += __shfl_down_sync(0xffffffff, val, offset);
             #endif
             return val;
         }

        __inline__ __device__
        scalar warpReduceMin(scalar val)
        {
            #if (__CUDA_ARCH__ < 700) || (CUDART_VERSION < 9000)
            for (int offset = warpSize/2; offset > 0; offset /= 2)
                val  = min(val, __shfl_down(val, offset));
            #else
            for (int offset = warpSize/2; offset > 0; offset /= 2)
                val  = min(val, __shfl_down_sync(0xffffffff, val, offset));
            #endif
            return val;
        }
        __inline__ __device__
        scalar warpReduceMax(scalar val)
        {
            #if (__CUDA_ARCH__ < 700) || (CUDART_VERSION < 9000)
            for (int offset = warpSize/2; offset > 0; offset /= 2)
                val = max(val, __shfl_down(val, offset));
            #else
            for (int offset = warpSize/2; offset > 0; offset /= 2)
                val = max(val, __shfl_down_sync(0xffffffff, val, offset));
            #endif
            return val;
        }

        __inline__ __device__
        Vector3 warpReduceSum(Vector3 val)
        {
            #if (CUDART_VERSION >= 9000)
            for (int offset = warpSize/2; offset > 0; offset /= 2)
            {
                val[0] += __shfl_down_sync(0xffffffff, val[0], offset);
                val[1] += __shfl_down_sync(0xffffffff, val[1], offset);
                val[2] += __shfl_down_sync(0xffffffff, val[2], offset);
            }
            #else
            for (int offset = warpSize/2; offset > 0; offset /= 2)
            {
                val[0] += __shfl_down(val[0], offset);
                val[1] += __shfl_down(val[1], offset);
                val[2] += __shfl_down(val[2], offset);
            }
            #endif
            return val;
