From 33173e433ca264c24a12871a691032be83875cf8 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 3 May 2022 14:09:16 -0500
Subject: [PATCH 1/2] Small changes from #9452.

---
 cpp/CMakeLists.txt                            |  2 +-
 cpp/include/cudf/detail/structs/utilities.hpp | 16 +++++++++-
 cpp/include/cudf/table/row_operators.cuh      |  2 +-
 cpp/include/cudf/utilities/traits.hpp         | 31 +++++++++++++++++++
 cpp/src/binaryop/compiled/ATan2.cu            |  4 +--
 cpp/src/binaryop/compiled/Add.cu              |  4 +--
 cpp/src/binaryop/compiled/BitwiseAnd.cu       |  4 +--
 cpp/src/binaryop/compiled/BitwiseOr.cu        |  4 +--
 cpp/src/binaryop/compiled/BitwiseXor.cu       |  4 +--
 cpp/src/binaryop/compiled/Div.cu              |  4 +--
 cpp/src/binaryop/compiled/FloorDiv.cu         |  4 +--
 cpp/src/binaryop/compiled/Greater.cu          |  4 +--
 cpp/src/binaryop/compiled/GreaterEqual.cu     |  4 +--
 cpp/src/binaryop/compiled/Less.cu             |  4 +--
 cpp/src/binaryop/compiled/LessEqual.cu        |  4 +--
 cpp/src/binaryop/compiled/LogBase.cu          |  4 +--
 cpp/src/binaryop/compiled/LogicalAnd.cu       |  4 +--
 cpp/src/binaryop/compiled/LogicalOr.cu        |  4 +--
 cpp/src/binaryop/compiled/Mod.cu              |  4 +--
 cpp/src/binaryop/compiled/Mul.cu              |  4 +--
 cpp/src/binaryop/compiled/NullEquals.cu       |  4 +--
 cpp/src/binaryop/compiled/NullLogicalAnd.cu   |  2 +-
 cpp/src/binaryop/compiled/NullLogicalOr.cu    |  2 +-
 cpp/src/binaryop/compiled/NullMax.cu          |  4 +--
 cpp/src/binaryop/compiled/NullMin.cu          |  4 +--
 cpp/src/binaryop/compiled/PMod.cu             |  4 +--
 cpp/src/binaryop/compiled/Pow.cu              |  4 +--
 cpp/src/binaryop/compiled/PyMod.cu            |  4 +--
 cpp/src/binaryop/compiled/ShiftLeft.cu        |  4 +--
 cpp/src/binaryop/compiled/ShiftRight.cu       |  4 +--
 .../binaryop/compiled/ShiftRightUnsigned.cu   |  4 +--
 cpp/src/binaryop/compiled/Sub.cu              |  4 +--
 cpp/src/binaryop/compiled/TrueDiv.cu          |  4 +--
 cpp/src/binaryop/compiled/binary_ops.cu       |  4 +--
 cpp/src/binaryop/compiled/binary_ops.cuh      |  4 +--
 cpp/src/binaryop/compiled/binary_ops.hpp      |  8 ++---
 cpp/src/binaryop/compiled/equality_ops.cu     |  4 +--
 cpp/src/binaryop/compiled/util.cpp            |  2 +-
 cpp/src/structs/utilities.cpp                 |  6 ++++
 39 files changed, 121 insertions(+), 70 deletions(-)

diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index cbe2811afe4..7870366b714 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -189,7 +189,6 @@ add_library(
   src/ast/expression_parser.cpp
   src/ast/expressions.cpp
   src/binaryop/binaryop.cpp
-  src/binaryop/compiled/binary_ops.cu
   src/binaryop/compiled/Add.cu
   src/binaryop/compiled/ATan2.cu
   src/binaryop/compiled/BitwiseAnd.cu
@@ -220,6 +219,7 @@ add_library(
   src/binaryop/compiled/ShiftRightUnsigned.cu
   src/binaryop/compiled/Sub.cu
   src/binaryop/compiled/TrueDiv.cu
+  src/binaryop/compiled/binary_ops.cu
   src/binaryop/compiled/util.cpp
   src/labeling/label_bins.cu
   src/bitmask/null_mask.cu
diff --git a/cpp/include/cudf/detail/structs/utilities.hpp b/cpp/include/cudf/detail/structs/utilities.hpp
index 751b7c00e8a..45d4c3b5ae4 100644
--- a/cpp/include/cudf/detail/structs/utilities.hpp
+++ b/cpp/include/cudf/detail/structs/utilities.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -245,6 +245,20 @@ std::tuple<cudf::table_view, std::vector<rmm::device_buffer>> superimpose_parent
   table_view const& table,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+/**
+ * @brief Checks if a column or any of its children is a struct column with structs that are null.
+ *
+ * This function searches for structs that are null -- differentiating between structs that are null
+ * and structs containing null values. Null structs add a column to the result of the flatten column
+ * utility and necessitates column_nullability::FORCE when flattening the column for comparison
+ * operations.
+ *
+ * @param col Column to check for null structs
+ * @return A boolean indicating if the column is or contains a struct column that contains a null
+ * struct.
+ */
+bool contains_null_structs(column_view const& col);
 }  // namespace detail
 }  // namespace structs
 }  // namespace cudf
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
index 4eca03a800c..4d503cd53b8 100644
--- a/cpp/include/cudf/table/row_operators.cuh
+++ b/cpp/include/cudf/table/row_operators.cuh
@@ -74,7 +74,7 @@ __device__ weak_ordering compare_elements(Element lhs, Element rhs)
  * @brief A specialization for floating-point `Element` type relational comparison
  * to derive the order of the elements with respect to `lhs`.
  *
- * This Specialization handles `nan` in the following order:
+ * This specialization handles `nan` in the following order:
  * `[-Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN, null] (for null_order::AFTER)`
  * `[null, -Inf, -ve, 0, -0, +ve, +Inf, NaN, NaN] (for null_order::BEFORE)`
  *
diff --git a/cpp/include/cudf/utilities/traits.hpp b/cpp/include/cudf/utilities/traits.hpp
index ed24517f55b..d8fa7bff0b8 100644
--- a/cpp/include/cudf/utilities/traits.hpp
+++ b/cpp/include/cudf/utilities/traits.hpp
@@ -699,6 +699,37 @@ constexpr inline bool is_nested(data_type type)
   return cudf::type_dispatcher(type, is_nested_impl{});
 }
 
+/**
+ * @brief Indicates whether `T` is a struct type.
+ *
+ * @param T The type to verify
+ * @return A boolean indicating if T is a struct type
+ */
+template <typename T>
+constexpr inline bool is_struct()
+{
+  return std::is_same_v<T, cudf::struct_view>;
+}
+
+struct is_struct_impl {
+  template <typename T>
+  constexpr bool operator()()
+  {
+    return is_struct<T>();
+  }
+};
+
+/**
+ * @brief Indicates whether `type` is a struct type.
+ *
+ * @param type The `data_type` to verify
+ * @return A boolean indicating if `type` is a struct type
+ */
+constexpr inline bool is_struct(data_type type)
+{
+  return cudf::type_dispatcher(type, is_struct_impl{});
+}
+
 template <typename FromType>
 struct is_bit_castable_to_impl {
   template <typename ToType, std::enable_if_t<is_compound<ToType>()>* = nullptr>
diff --git a/cpp/src/binaryop/compiled/ATan2.cu b/cpp/src/binaryop/compiled/ATan2.cu
index 8e5cbf57f55..90c2bead8a7 100644
--- a/cpp/src/binaryop/compiled/ATan2.cu
+++ b/cpp/src/binaryop/compiled/ATan2.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ATan2>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Add.cu b/cpp/src/binaryop/compiled/Add.cu
index 4cd2ced66f4..18375cc0680 100644
--- a/cpp/src/binaryop/compiled/Add.cu
+++ b/cpp/src/binaryop/compiled/Add.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Add>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/BitwiseAnd.cu b/cpp/src/binaryop/compiled/BitwiseAnd.cu
index 6abac2bd197..77374683c4f 100644
--- a/cpp/src/binaryop/compiled/BitwiseAnd.cu
+++ b/cpp/src/binaryop/compiled/BitwiseAnd.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::BitwiseAnd>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/BitwiseOr.cu b/cpp/src/binaryop/compiled/BitwiseOr.cu
index 6d523cbf1d1..7eda9800598 100644
--- a/cpp/src/binaryop/compiled/BitwiseOr.cu
+++ b/cpp/src/binaryop/compiled/BitwiseOr.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::BitwiseOr>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/BitwiseXor.cu b/cpp/src/binaryop/compiled/BitwiseXor.cu
index 45175681574..5d6bb9d2cb3 100644
--- a/cpp/src/binaryop/compiled/BitwiseXor.cu
+++ b/cpp/src/binaryop/compiled/BitwiseXor.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::BitwiseXor>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Div.cu b/cpp/src/binaryop/compiled/Div.cu
index 7cc895ecd06..ce8a7519da0 100644
--- a/cpp/src/binaryop/compiled/Div.cu
+++ b/cpp/src/binaryop/compiled/Div.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Div>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/FloorDiv.cu b/cpp/src/binaryop/compiled/FloorDiv.cu
index 99ea2706b86..5b68defa598 100644
--- a/cpp/src/binaryop/compiled/FloorDiv.cu
+++ b/cpp/src/binaryop/compiled/FloorDiv.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::FloorDiv>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Greater.cu b/cpp/src/binaryop/compiled/Greater.cu
index 679e029b5fc..00195237486 100644
--- a/cpp/src/binaryop/compiled/Greater.cu
+++ b/cpp/src/binaryop/compiled/Greater.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Greater>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/GreaterEqual.cu b/cpp/src/binaryop/compiled/GreaterEqual.cu
index 23b0c6aaa0d..d1cd0e484a5 100644
--- a/cpp/src/binaryop/compiled/GreaterEqual.cu
+++ b/cpp/src/binaryop/compiled/GreaterEqual.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::GreaterEqual>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Less.cu b/cpp/src/binaryop/compiled/Less.cu
index 7ab5dfe3478..8ca237f564a 100644
--- a/cpp/src/binaryop/compiled/Less.cu
+++ b/cpp/src/binaryop/compiled/Less.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Less>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LessEqual.cu b/cpp/src/binaryop/compiled/LessEqual.cu
index 983c50c9575..61031155dc6 100644
--- a/cpp/src/binaryop/compiled/LessEqual.cu
+++ b/cpp/src/binaryop/compiled/LessEqual.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LessEqual>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LogBase.cu b/cpp/src/binaryop/compiled/LogBase.cu
index bdc709b86bf..b75d1723299 100644
--- a/cpp/src/binaryop/compiled/LogBase.cu
+++ b/cpp/src/binaryop/compiled/LogBase.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LogBase>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LogicalAnd.cu b/cpp/src/binaryop/compiled/LogicalAnd.cu
index 08112fadfff..ad4cea633fb 100644
--- a/cpp/src/binaryop/compiled/LogicalAnd.cu
+++ b/cpp/src/binaryop/compiled/LogicalAnd.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LogicalAnd>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LogicalOr.cu b/cpp/src/binaryop/compiled/LogicalOr.cu
index bc400afd4cd..1bb35ed3510 100644
--- a/cpp/src/binaryop/compiled/LogicalOr.cu
+++ b/cpp/src/binaryop/compiled/LogicalOr.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LogicalOr>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Mod.cu b/cpp/src/binaryop/compiled/Mod.cu
index 0b82c09c8a6..e15e33ad664 100644
--- a/cpp/src/binaryop/compiled/Mod.cu
+++ b/cpp/src/binaryop/compiled/Mod.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Mod>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Mul.cu b/cpp/src/binaryop/compiled/Mul.cu
index 15394245259..e6b15942c4f 100644
--- a/cpp/src/binaryop/compiled/Mul.cu
+++ b/cpp/src/binaryop/compiled/Mul.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Mul>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullEquals.cu b/cpp/src/binaryop/compiled/NullEquals.cu
index 3fc76e804f7..3d1207b1a6c 100644
--- a/cpp/src/binaryop/compiled/NullEquals.cu
+++ b/cpp/src/binaryop/compiled/NullEquals.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullEquals>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullLogicalAnd.cu b/cpp/src/binaryop/compiled/NullLogicalAnd.cu
index 48ae125bc93..747a6175b33 100644
--- a/cpp/src/binaryop/compiled/NullLogicalAnd.cu
+++ b/cpp/src/binaryop/compiled/NullLogicalAnd.cu
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullLogicalAnd>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullLogicalOr.cu b/cpp/src/binaryop/compiled/NullLogicalOr.cu
index e0ea95ac3ee..6203cf899eb 100644
--- a/cpp/src/binaryop/compiled/NullLogicalOr.cu
+++ b/cpp/src/binaryop/compiled/NullLogicalOr.cu
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullLogicalOr>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullMax.cu b/cpp/src/binaryop/compiled/NullMax.cu
index 78a44041cba..2624cb86945 100644
--- a/cpp/src/binaryop/compiled/NullMax.cu
+++ b/cpp/src/binaryop/compiled/NullMax.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullMax>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullMin.cu b/cpp/src/binaryop/compiled/NullMin.cu
index 629ab600fd7..09e4781610b 100644
--- a/cpp/src/binaryop/compiled/NullMin.cu
+++ b/cpp/src/binaryop/compiled/NullMin.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullMin>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/PMod.cu b/cpp/src/binaryop/compiled/PMod.cu
index 36902c0ed10..3ad6f61b66c 100644
--- a/cpp/src/binaryop/compiled/PMod.cu
+++ b/cpp/src/binaryop/compiled/PMod.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::PMod>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Pow.cu b/cpp/src/binaryop/compiled/Pow.cu
index c6f897ee18d..3fc8dcee206 100644
--- a/cpp/src/binaryop/compiled/Pow.cu
+++ b/cpp/src/binaryop/compiled/Pow.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Pow>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/PyMod.cu b/cpp/src/binaryop/compiled/PyMod.cu
index b05dcd8e7bc..1d59a6aac4d 100644
--- a/cpp/src/binaryop/compiled/PyMod.cu
+++ b/cpp/src/binaryop/compiled/PyMod.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::PyMod>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/ShiftLeft.cu b/cpp/src/binaryop/compiled/ShiftLeft.cu
index 6cc950b2d50..050a7632d20 100644
--- a/cpp/src/binaryop/compiled/ShiftLeft.cu
+++ b/cpp/src/binaryop/compiled/ShiftLeft.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ShiftLeft>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/ShiftRight.cu b/cpp/src/binaryop/compiled/ShiftRight.cu
index 1ddd7100a73..cf17a5af549 100644
--- a/cpp/src/binaryop/compiled/ShiftRight.cu
+++ b/cpp/src/binaryop/compiled/ShiftRight.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ShiftRight>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu b/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu
index a87b4b9f9ac..a3bfe9e2d15 100644
--- a/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu
+++ b/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ShiftRightUnsigned>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Sub.cu b/cpp/src/binaryop/compiled/Sub.cu
index e0cf47c1310..327a7a7c7de 100644
--- a/cpp/src/binaryop/compiled/Sub.cu
+++ b/cpp/src/binaryop/compiled/Sub.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Sub>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/TrueDiv.cu b/cpp/src/binaryop/compiled/TrueDiv.cu
index d8f1d956340..bf01226f96a 100644
--- a/cpp/src/binaryop/compiled/TrueDiv.cu
+++ b/cpp/src/binaryop/compiled/TrueDiv.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::TrueDiv>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu
index c01359b80d0..a50be6379e6 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cu
+++ b/cpp/src/binaryop/compiled/binary_ops.cu
@@ -14,8 +14,8 @@
  * limitations under the License.
  */
 
-#include "binary_ops.hpp"
-#include "operation.cuh"
+#include <binaryop/compiled/binary_ops.hpp>
+#include <binaryop/compiled/operation.cuh>
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
diff --git a/cpp/src/binaryop/compiled/binary_ops.cuh b/cpp/src/binaryop/compiled/binary_ops.cuh
index ec41fbb8883..be3d1ab775a 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cuh
+++ b/cpp/src/binaryop/compiled/binary_ops.cuh
@@ -16,8 +16,8 @@
 
 #pragma once
 
-#include "binary_ops.hpp"
-#include "operation.cuh"
+#include <binaryop/compiled/binary_ops.hpp>
+#include <binaryop/compiled/operation.cuh>
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/detail/utilities/integer_utils.hpp>
diff --git a/cpp/src/binaryop/compiled/binary_ops.hpp b/cpp/src/binaryop/compiled/binary_ops.hpp
index 26a0f26b59c..85ef254c5a1 100644
--- a/cpp/src/binaryop/compiled/binary_ops.hpp
+++ b/cpp/src/binaryop/compiled/binary_ops.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2021, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -167,7 +167,7 @@ std::optional<data_type> get_common_type(data_type out, data_type lhs, data_type
  * @param out output type of the binary operation
  * @param lhs first operand type of the binary operation
  * @param rhs second operand type of the binary operation
- * @param op binary operator enum.
+ * @param op binary operator enum
  * @return true if given binary operator supports given input and output types.
  */
 bool is_supported_operation(data_type out, data_type lhs, data_type rhs, binary_operator op);
@@ -195,8 +195,8 @@ void apply_binary_op(mutable_column_device_view&,
                      bool is_rhs_scalar,
                      rmm::cuda_stream_view stream);
 /**
- * @brief Deploys single type or double type dispatcher that runs equality operation on each element
- * of @p lhsd and @p rhsd columns.
+ * @brief Deploys single type or double type dispatcher that runs equality operation on each
+ * element of @p lhsd and @p rhsd columns.
  *
  * Comparison operators are EQUAL, NOT_EQUAL, NULL_EQUALS.
  * @p outd type is boolean.
diff --git a/cpp/src/binaryop/compiled/equality_ops.cu b/cpp/src/binaryop/compiled/equality_ops.cu
index 03c3e373476..fe34f3ef373 100644
--- a/cpp/src/binaryop/compiled/equality_ops.cu
+++ b/cpp/src/binaryop/compiled/equality_ops.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021, NVIDIA CORPORATION.
+ * Copyright (c) 2021-2022, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "binary_ops.cuh"
+#include <binaryop/compiled/binary_ops.cuh>
 
 namespace cudf::binops::compiled {
 void dispatch_equality_op(mutable_column_device_view& outd,
diff --git a/cpp/src/binaryop/compiled/util.cpp b/cpp/src/binaryop/compiled/util.cpp
index d8f1eb03a16..b79bc52fc94 100644
--- a/cpp/src/binaryop/compiled/util.cpp
+++ b/cpp/src/binaryop/compiled/util.cpp
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include "operation.cuh"
+#include <binaryop/compiled/operation.cuh>
 
 #include <cudf/binaryop.hpp>
 #include <cudf/column/column_device_view.cuh>
diff --git a/cpp/src/structs/utilities.cpp b/cpp/src/structs/utilities.cpp
index a2c173cae5f..5baab0f09a2 100644
--- a/cpp/src/structs/utilities.cpp
+++ b/cpp/src/structs/utilities.cpp
@@ -441,6 +441,12 @@ std::tuple<cudf::table_view, std::vector<rmm::device_buffer>> superimpose_parent
   return {table_view{superimposed_columns}, std::move(superimposed_nullmasks)};
 }
 
+bool contains_null_structs(column_view const& col)
+{
+  return (is_struct(col) && col.has_nulls()) ||
+         std::any_of(col.child_begin(), col.child_end(), contains_null_structs);
+}
+
 }  // namespace detail
 }  // namespace structs
 }  // namespace cudf

From 8ebe02b6c40cb1bbc3a4408ca8676078bc4df298 Mon Sep 17 00:00:00 2001
From: Bradley Dice <bdice@bradleydice.com>
Date: Tue, 3 May 2022 16:49:16 -0500
Subject: [PATCH 2/2] Revert include changes.

---
 cpp/src/binaryop/compiled/ATan2.cu              | 4 ++--
 cpp/src/binaryop/compiled/Add.cu                | 4 ++--
 cpp/src/binaryop/compiled/BitwiseAnd.cu         | 4 ++--
 cpp/src/binaryop/compiled/BitwiseOr.cu          | 4 ++--
 cpp/src/binaryop/compiled/BitwiseXor.cu         | 4 ++--
 cpp/src/binaryop/compiled/Div.cu                | 4 ++--
 cpp/src/binaryop/compiled/FloorDiv.cu           | 4 ++--
 cpp/src/binaryop/compiled/Greater.cu            | 4 ++--
 cpp/src/binaryop/compiled/GreaterEqual.cu       | 4 ++--
 cpp/src/binaryop/compiled/Less.cu               | 4 ++--
 cpp/src/binaryop/compiled/LessEqual.cu          | 4 ++--
 cpp/src/binaryop/compiled/LogBase.cu            | 4 ++--
 cpp/src/binaryop/compiled/LogicalAnd.cu         | 4 ++--
 cpp/src/binaryop/compiled/LogicalOr.cu          | 4 ++--
 cpp/src/binaryop/compiled/Mod.cu                | 4 ++--
 cpp/src/binaryop/compiled/Mul.cu                | 4 ++--
 cpp/src/binaryop/compiled/NullEquals.cu         | 4 ++--
 cpp/src/binaryop/compiled/NullLogicalAnd.cu     | 2 +-
 cpp/src/binaryop/compiled/NullLogicalOr.cu      | 2 +-
 cpp/src/binaryop/compiled/NullMax.cu            | 4 ++--
 cpp/src/binaryop/compiled/NullMin.cu            | 4 ++--
 cpp/src/binaryop/compiled/PMod.cu               | 4 ++--
 cpp/src/binaryop/compiled/Pow.cu                | 4 ++--
 cpp/src/binaryop/compiled/PyMod.cu              | 4 ++--
 cpp/src/binaryop/compiled/ShiftLeft.cu          | 4 ++--
 cpp/src/binaryop/compiled/ShiftRight.cu         | 4 ++--
 cpp/src/binaryop/compiled/ShiftRightUnsigned.cu | 4 ++--
 cpp/src/binaryop/compiled/Sub.cu                | 4 ++--
 cpp/src/binaryop/compiled/TrueDiv.cu            | 4 ++--
 cpp/src/binaryop/compiled/binary_ops.cu         | 4 ++--
 cpp/src/binaryop/compiled/binary_ops.cuh        | 4 ++--
 cpp/src/binaryop/compiled/binary_ops.hpp        | 8 ++++----
 cpp/src/binaryop/compiled/equality_ops.cu       | 4 ++--
 cpp/src/binaryop/compiled/util.cpp              | 2 +-
 34 files changed, 67 insertions(+), 67 deletions(-)

diff --git a/cpp/src/binaryop/compiled/ATan2.cu b/cpp/src/binaryop/compiled/ATan2.cu
index 90c2bead8a7..8e5cbf57f55 100644
--- a/cpp/src/binaryop/compiled/ATan2.cu
+++ b/cpp/src/binaryop/compiled/ATan2.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ATan2>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Add.cu b/cpp/src/binaryop/compiled/Add.cu
index 18375cc0680..4cd2ced66f4 100644
--- a/cpp/src/binaryop/compiled/Add.cu
+++ b/cpp/src/binaryop/compiled/Add.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Add>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/BitwiseAnd.cu b/cpp/src/binaryop/compiled/BitwiseAnd.cu
index 77374683c4f..6abac2bd197 100644
--- a/cpp/src/binaryop/compiled/BitwiseAnd.cu
+++ b/cpp/src/binaryop/compiled/BitwiseAnd.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::BitwiseAnd>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/BitwiseOr.cu b/cpp/src/binaryop/compiled/BitwiseOr.cu
index 7eda9800598..6d523cbf1d1 100644
--- a/cpp/src/binaryop/compiled/BitwiseOr.cu
+++ b/cpp/src/binaryop/compiled/BitwiseOr.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::BitwiseOr>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/BitwiseXor.cu b/cpp/src/binaryop/compiled/BitwiseXor.cu
index 5d6bb9d2cb3..45175681574 100644
--- a/cpp/src/binaryop/compiled/BitwiseXor.cu
+++ b/cpp/src/binaryop/compiled/BitwiseXor.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::BitwiseXor>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Div.cu b/cpp/src/binaryop/compiled/Div.cu
index ce8a7519da0..7cc895ecd06 100644
--- a/cpp/src/binaryop/compiled/Div.cu
+++ b/cpp/src/binaryop/compiled/Div.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Div>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/FloorDiv.cu b/cpp/src/binaryop/compiled/FloorDiv.cu
index 5b68defa598..99ea2706b86 100644
--- a/cpp/src/binaryop/compiled/FloorDiv.cu
+++ b/cpp/src/binaryop/compiled/FloorDiv.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::FloorDiv>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Greater.cu b/cpp/src/binaryop/compiled/Greater.cu
index 00195237486..679e029b5fc 100644
--- a/cpp/src/binaryop/compiled/Greater.cu
+++ b/cpp/src/binaryop/compiled/Greater.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Greater>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/GreaterEqual.cu b/cpp/src/binaryop/compiled/GreaterEqual.cu
index d1cd0e484a5..23b0c6aaa0d 100644
--- a/cpp/src/binaryop/compiled/GreaterEqual.cu
+++ b/cpp/src/binaryop/compiled/GreaterEqual.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::GreaterEqual>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Less.cu b/cpp/src/binaryop/compiled/Less.cu
index 8ca237f564a..7ab5dfe3478 100644
--- a/cpp/src/binaryop/compiled/Less.cu
+++ b/cpp/src/binaryop/compiled/Less.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Less>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LessEqual.cu b/cpp/src/binaryop/compiled/LessEqual.cu
index 61031155dc6..983c50c9575 100644
--- a/cpp/src/binaryop/compiled/LessEqual.cu
+++ b/cpp/src/binaryop/compiled/LessEqual.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LessEqual>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LogBase.cu b/cpp/src/binaryop/compiled/LogBase.cu
index b75d1723299..bdc709b86bf 100644
--- a/cpp/src/binaryop/compiled/LogBase.cu
+++ b/cpp/src/binaryop/compiled/LogBase.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LogBase>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LogicalAnd.cu b/cpp/src/binaryop/compiled/LogicalAnd.cu
index ad4cea633fb..08112fadfff 100644
--- a/cpp/src/binaryop/compiled/LogicalAnd.cu
+++ b/cpp/src/binaryop/compiled/LogicalAnd.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LogicalAnd>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/LogicalOr.cu b/cpp/src/binaryop/compiled/LogicalOr.cu
index 1bb35ed3510..bc400afd4cd 100644
--- a/cpp/src/binaryop/compiled/LogicalOr.cu
+++ b/cpp/src/binaryop/compiled/LogicalOr.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::LogicalOr>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Mod.cu b/cpp/src/binaryop/compiled/Mod.cu
index e15e33ad664..0b82c09c8a6 100644
--- a/cpp/src/binaryop/compiled/Mod.cu
+++ b/cpp/src/binaryop/compiled/Mod.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Mod>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Mul.cu b/cpp/src/binaryop/compiled/Mul.cu
index e6b15942c4f..15394245259 100644
--- a/cpp/src/binaryop/compiled/Mul.cu
+++ b/cpp/src/binaryop/compiled/Mul.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Mul>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullEquals.cu b/cpp/src/binaryop/compiled/NullEquals.cu
index 3d1207b1a6c..3fc76e804f7 100644
--- a/cpp/src/binaryop/compiled/NullEquals.cu
+++ b/cpp/src/binaryop/compiled/NullEquals.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullEquals>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullLogicalAnd.cu b/cpp/src/binaryop/compiled/NullLogicalAnd.cu
index 747a6175b33..48ae125bc93 100644
--- a/cpp/src/binaryop/compiled/NullLogicalAnd.cu
+++ b/cpp/src/binaryop/compiled/NullLogicalAnd.cu
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullLogicalAnd>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullLogicalOr.cu b/cpp/src/binaryop/compiled/NullLogicalOr.cu
index 6203cf899eb..e0ea95ac3ee 100644
--- a/cpp/src/binaryop/compiled/NullLogicalOr.cu
+++ b/cpp/src/binaryop/compiled/NullLogicalOr.cu
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullLogicalOr>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullMax.cu b/cpp/src/binaryop/compiled/NullMax.cu
index 2624cb86945..78a44041cba 100644
--- a/cpp/src/binaryop/compiled/NullMax.cu
+++ b/cpp/src/binaryop/compiled/NullMax.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullMax>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/NullMin.cu b/cpp/src/binaryop/compiled/NullMin.cu
index 09e4781610b..629ab600fd7 100644
--- a/cpp/src/binaryop/compiled/NullMin.cu
+++ b/cpp/src/binaryop/compiled/NullMin.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::NullMin>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/PMod.cu b/cpp/src/binaryop/compiled/PMod.cu
index 3ad6f61b66c..36902c0ed10 100644
--- a/cpp/src/binaryop/compiled/PMod.cu
+++ b/cpp/src/binaryop/compiled/PMod.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::PMod>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Pow.cu b/cpp/src/binaryop/compiled/Pow.cu
index 3fc8dcee206..c6f897ee18d 100644
--- a/cpp/src/binaryop/compiled/Pow.cu
+++ b/cpp/src/binaryop/compiled/Pow.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Pow>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/PyMod.cu b/cpp/src/binaryop/compiled/PyMod.cu
index 1d59a6aac4d..b05dcd8e7bc 100644
--- a/cpp/src/binaryop/compiled/PyMod.cu
+++ b/cpp/src/binaryop/compiled/PyMod.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::PyMod>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/ShiftLeft.cu b/cpp/src/binaryop/compiled/ShiftLeft.cu
index 050a7632d20..6cc950b2d50 100644
--- a/cpp/src/binaryop/compiled/ShiftLeft.cu
+++ b/cpp/src/binaryop/compiled/ShiftLeft.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ShiftLeft>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/ShiftRight.cu b/cpp/src/binaryop/compiled/ShiftRight.cu
index cf17a5af549..1ddd7100a73 100644
--- a/cpp/src/binaryop/compiled/ShiftRight.cu
+++ b/cpp/src/binaryop/compiled/ShiftRight.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ShiftRight>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu b/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu
index a3bfe9e2d15..a87b4b9f9ac 100644
--- a/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu
+++ b/cpp/src/binaryop/compiled/ShiftRightUnsigned.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::ShiftRightUnsigned>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/Sub.cu b/cpp/src/binaryop/compiled/Sub.cu
index 327a7a7c7de..e0cf47c1310 100644
--- a/cpp/src/binaryop/compiled/Sub.cu
+++ b/cpp/src/binaryop/compiled/Sub.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::Sub>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/TrueDiv.cu b/cpp/src/binaryop/compiled/TrueDiv.cu
index bf01226f96a..d8f1d956340 100644
--- a/cpp/src/binaryop/compiled/TrueDiv.cu
+++ b/cpp/src/binaryop/compiled/TrueDiv.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 template void apply_binary_op<ops::TrueDiv>(mutable_column_device_view&,
diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu
index a50be6379e6..c01359b80d0 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cu
+++ b/cpp/src/binaryop/compiled/binary_ops.cu
@@ -14,8 +14,8 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.hpp>
-#include <binaryop/compiled/operation.cuh>
+#include "binary_ops.hpp"
+#include "operation.cuh"
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
diff --git a/cpp/src/binaryop/compiled/binary_ops.cuh b/cpp/src/binaryop/compiled/binary_ops.cuh
index be3d1ab775a..ec41fbb8883 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cuh
+++ b/cpp/src/binaryop/compiled/binary_ops.cuh
@@ -16,8 +16,8 @@
 
 #pragma once
 
-#include <binaryop/compiled/binary_ops.hpp>
-#include <binaryop/compiled/operation.cuh>
+#include "binary_ops.hpp"
+#include "operation.cuh"
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/detail/utilities/integer_utils.hpp>
diff --git a/cpp/src/binaryop/compiled/binary_ops.hpp b/cpp/src/binaryop/compiled/binary_ops.hpp
index 85ef254c5a1..26a0f26b59c 100644
--- a/cpp/src/binaryop/compiled/binary_ops.hpp
+++ b/cpp/src/binaryop/compiled/binary_ops.hpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2018-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -167,7 +167,7 @@ std::optional<data_type> get_common_type(data_type out, data_type lhs, data_type
  * @param out output type of the binary operation
  * @param lhs first operand type of the binary operation
  * @param rhs second operand type of the binary operation
- * @param op binary operator enum
+ * @param op binary operator enum.
  * @return true if given binary operator supports given input and output types.
  */
 bool is_supported_operation(data_type out, data_type lhs, data_type rhs, binary_operator op);
@@ -195,8 +195,8 @@ void apply_binary_op(mutable_column_device_view&,
                      bool is_rhs_scalar,
                      rmm::cuda_stream_view stream);
 /**
- * @brief Deploys single type or double type dispatcher that runs equality operation on each
- * element of @p lhsd and @p rhsd columns.
+ * @brief Deploys single type or double type dispatcher that runs equality operation on each element
+ * of @p lhsd and @p rhsd columns.
  *
  * Comparison operators are EQUAL, NOT_EQUAL, NULL_EQUALS.
  * @p outd type is boolean.
diff --git a/cpp/src/binaryop/compiled/equality_ops.cu b/cpp/src/binaryop/compiled/equality_ops.cu
index fe34f3ef373..03c3e373476 100644
--- a/cpp/src/binaryop/compiled/equality_ops.cu
+++ b/cpp/src/binaryop/compiled/equality_ops.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2021-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/binary_ops.cuh>
+#include "binary_ops.cuh"
 
 namespace cudf::binops::compiled {
 void dispatch_equality_op(mutable_column_device_view& outd,
diff --git a/cpp/src/binaryop/compiled/util.cpp b/cpp/src/binaryop/compiled/util.cpp
index b79bc52fc94..d8f1eb03a16 100644
--- a/cpp/src/binaryop/compiled/util.cpp
+++ b/cpp/src/binaryop/compiled/util.cpp
@@ -14,7 +14,7 @@
  * limitations under the License.
  */
 
-#include <binaryop/compiled/operation.cuh>
+#include "operation.cuh"
 
 #include <cudf/binaryop.hpp>
 #include <cudf/column/column_device_view.cuh>