floor_divide now rounds signed integers toward negative infinity

- Resolves #1247

Author: ndgrigorian

dpctl/tensor/libtensor/include/kernels/elementwise_functions/floor_divide.hpp

@@ -56,14 +56,25 @@ struct FloorDivideFunctor
     using supports_sg_loadstore =
         std::negation<std::disjunction<tu_ns::is_complex<argT1>,
                                        tu_ns::is_complex<argT2>>>; // TRUE
-    using supports_vec = std::negation<
-        std::disjunction<tu_ns::is_complex<argT1>, tu_ns::is_complex<argT2>>>;
+    using supports_vec = std::negation<std::disjunction<
+        tu_ns::is_complex<argT1>,
+        tu_ns::is_complex<argT2>,
+        std::conjunction<std::is_integral<argT1>, std::is_signed<argT1>>,
+        std::conjunction<std::is_integral<argT2>, std::is_signed<argT2>>>>;
+    // no vec overload for signed integers to avoid loop
 
     resT operator()(const argT1 &in1, const argT2 &in2)
     {
         auto tmp = in1 / in2;
         if constexpr (std::is_integral_v<decltype(tmp)>) {
-            return tmp;
+            if constexpr (std::is_unsigned_v<decltype(tmp)>) {
+                return tmp;
+            }
+            else {
+                auto rem = in1 % in2;
+                auto corr = (rem != 0 && ((rem < 0) != (in2 < 0)));
+                return (tmp - corr);
+            }
         }
         else {
             return sycl::floor(tmp);

dpctl/tests/elementwise/test_floor_divide.py

@@ -186,3 +186,28 @@ def __sycl_usm_array_interface__(self):
     c = Canary()
     with pytest.raises(ValueError):
         dpt.floor_divide(a, c)
+
+
+def test_floor_divide_gh_1247():
+    get_queue_or_skip()
+
+    x = dpt.ones(1, dtype="i4")
+    res = dpt.floor_divide(x, -2)
+    np.testing.assert_array_equal(
+        dpt.asnumpy(res), np.full(res.shape, -1, dtype=res.dtype)
+    )
+
+    x = dpt.full(1, -1, dtype="i4")
+    res = dpt.floor_divide(x, 2)
+    np.testing.assert_array_equal(
+        dpt.asnumpy(res), np.full(res.shape, -1, dtype=res.dtype)
+    )
+
+    x = dpt.arange(-5, 6, 1, dtype="i4")
+    np.testing.assert_array_equal(
+        dpt.asnumpy(dpt.floor_divide(x, 3)), np.floor_divide(dpt.asnumpy(x), 3)
+    )
+    np.testing.assert_array_equal(
+        dpt.asnumpy(dpt.floor_divide(x, -3)),
+        np.floor_divide(dpt.asnumpy(x), -3),
+    )