Fix thrust failure when transfering data from device_vector to host_vector with vectors of size 1

cpp/include/cudf_test/column_utilities.hpp

@@ -216,6 +216,21 @@ inline std::pair<thrust::host_vector<std::string>, std::vector<bitmask_type>> to
   return {host_data, bitmask_to_host(c)};
 }
 
+/**
+ * @brief Print the different values between two columns
+ *
+ * @param differences contains indices of the different values
+ * @param lhs and rhs are the two input columns
+ * @param all_differences set to true will print all the different values and false will print only
+ * the first value
+ * @param depth the depth level of column data
+ */
+std::string differences_message(thrust::device_vector<int> const& differences,
+                                column_view const& lhs,
+                                column_view const& rhs,
+                                bool all_differences,
+                                int depth);
+
 }  // namespace test
 }  // namespace cudf
 

cpp/include/cudf_test/column_wrapper.hpp

@@ -94,7 +94,7 @@ struct fixed_width_type_converter {
   template <typename FromT                                                        = From,
             typename ToT                                                          = To,
             typename std::enable_if<std::is_same<FromT, ToT>::value, void>::type* = nullptr>
-  ToT operator()(FromT element) const
+  constexpr ToT operator()(FromT element) const
   {
     return element;
   }
@@ -106,7 +106,7 @@ struct fixed_width_type_converter {
                                       (cudf::is_convertible<FromT, ToT>::value ||
                                        std::is_constructible<ToT, FromT>::value),
                                     void>::type* = nullptr>
-  ToT operator()(FromT element) const
+  constexpr ToT operator()(FromT element) const
   {
     return static_cast<ToT>(element);
   }
@@ -117,7 +117,7 @@ struct fixed_width_type_converter {
     typename ToT                         = To,
     typename std::enable_if<std::is_integral<FromT>::value && cudf::is_timestamp_t<ToT>::value,
                             void>::type* = nullptr>
-  ToT operator()(FromT element) const
+  constexpr ToT operator()(FromT element) const
   {
     return ToT{typename ToT::duration{element}};
   }

cpp/tests/reshape/byte_cast_tests.cpp

@@ -26,6 +26,23 @@ using namespace cudf::test;
 class ByteCastTest : public cudf::test::BaseFixture {
 };
 
+// Previously, there was a bug that causes thrust to crash
+// This test is to make sure that bug does not pop up again
+// If this test fails, the program will crash
+TEST_F(ByteCastTest, thurstTest)
+{
+  cudf::test::fixed_width_column_wrapper<int> a{0, 1, 2, 3};
+  cudf::test::fixed_width_column_wrapper<int> b{0, 1, 2, 4};
+  cudf::test::fixed_width_column_wrapper<int> c{10, 11, 12, 13};
+
+  thrust::device_vector<int> diff_ab = std::vector<int>{3};
+  thrust::device_vector<int> diff_ac = std::vector<int>{0, 1, 2, 3};
+
+  // Prevent the compiler to optimize away the calls to differences_message
+  volatile auto str1 = cudf::test::differences_message(diff_ab, a, b, true, 0);
+  volatile auto str2 = cudf::test::differences_message(diff_ac, a, c, true, 0);
+}
+
 TEST_F(ByteCastTest, int16ValuesWithSplit)
 {
   using limits = std::numeric_limits<int16_t>;

cpp/tests/utilities/column_utilities.cu

@@ -160,45 +160,6 @@ class corresponding_rows_not_equivalent {
   }
 };
 
-std::string differences_message(thrust::device_vector<int> const& differences,
-                                column_view const& lhs,
-                                column_view const& rhs,
-                                bool all_differences,
-                                int depth)
-{
-  CUDF_EXPECTS(not differences.empty(), "Shouldn't enter this function if `differences` is empty");
-
-  std::string const depth_str = depth > 0 ? "depth " + std::to_string(depth) + '\n' : "";
-
-  if (all_differences) {
-    std::ostringstream buffer;
-    buffer << depth_str << "differences:" << std::endl;
-
-    auto source_table = cudf::table_view({lhs, rhs});
-    auto diff_column  = fixed_width_column_wrapper<int32_t>(differences.begin(), differences.end());
-    auto diff_table   = cudf::gather(source_table, diff_column);
-
-    //  Need to pull back the differences
-    auto const h_left_strings  = to_strings(diff_table->get_column(0));
-    auto const h_right_strings = to_strings(diff_table->get_column(1));
-
-    for (size_t i = 0; i < differences.size(); ++i)
-      buffer << depth_str << "lhs[" << differences[i] << "] = " << h_left_strings[i] << ", rhs["
-             << differences[i] << "] = " << h_right_strings[i] << std::endl;
-
-    return buffer.str();
-  } else {
-    int index = differences[0];  // only stringify first difference
-
-    auto diff_lhs = cudf::detail::slice(lhs, index, index + 1);
-    auto diff_rhs = cudf::detail::slice(rhs, index, index + 1);
-
-    return depth_str + "first difference: " + "lhs[" + std::to_string(index) +
-           "] = " + to_string(diff_lhs, "") + ", rhs[" + std::to_string(index) +
-           "] = " + to_string(diff_rhs, "");
-  }
-}
-
 // non-nested column types
 template <typename T, bool check_exact_equality>
 struct column_comparator_impl {
@@ -373,6 +334,45 @@ struct column_comparator {
 
 }  // namespace
 
+/**
+ * @copydoc cudf::test::differences_message
+ */
+std::string differences_message(thrust::device_vector<int> const& differences,
+                                column_view const& lhs,
+                                column_view const& rhs,
+                                bool all_differences,
+                                int depth)
+{
+  CUDF_EXPECTS(not differences.empty(), "Shouldn't enter this function if `differences` is empty");
+  std::string const depth_str = depth > 0 ? "depth " + std::to_string(depth) + '\n' : "";
+  if (all_differences) {
+    std::ostringstream buffer;
+    buffer << depth_str << "differences:" << std::endl;
+
+    // thrust may crash if a device_vector is passed to fixed_width_column_wrapper,
+    // thus we construct fixed_width_column_wrapper from a host_vector instead
+    thrust::host_vector<int> h_differences(differences);
+    auto source_table = cudf::table_view({lhs, rhs});
+    auto diff_column =
+      fixed_width_column_wrapper<int32_t>(h_differences.begin(), h_differences.end());
+    auto diff_table = cudf::gather(source_table, diff_column);
+    //  Need to pull back the differences
+    auto const h_left_strings  = to_strings(diff_table->get_column(0));
+    auto const h_right_strings = to_strings(diff_table->get_column(1));
+    for (size_t i = 0; i < h_differences.size(); ++i)
+      buffer << depth_str << "lhs[" << h_differences[i] << "] = " << h_left_strings[i] << ", rhs["
+             << h_differences[i] << "] = " << h_right_strings[i] << std::endl;
+    return buffer.str();
+  } else {
+    int index     = differences[0];  // only stringify first difference
+    auto diff_lhs = cudf::detail::slice(lhs, index, index + 1);
+    auto diff_rhs = cudf::detail::slice(rhs, index, index + 1);
+    return depth_str + "first difference: " + "lhs[" + std::to_string(index) +
+           "] = " + to_string(diff_lhs, "") + ", rhs[" + std::to_string(index) +
+           "] = " + to_string(diff_rhs, "");
+  }
+}
+
 /**
  * @copydoc cudf::test::expect_column_properties_equal
  */