diff --git a/cpp/benchmarks/groupby/group_sum_benchmark.cu b/cpp/benchmarks/groupby/group_sum_benchmark.cu
index f64022690d9..0e9f5061a1a 100644
--- a/cpp/benchmarks/groupby/group_sum_benchmark.cu
+++ b/cpp/benchmarks/groupby/group_sum_benchmark.cu
@@ -44,7 +44,6 @@ void BM_basic_sum(benchmark::State& state)
 {
   using wrapper = cudf::test::fixed_width_column_wrapper<int64_t>;
 
-  // const cudf::size_type num_columns{(cudf::size_type)state.range(0)};
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
   auto data_it = cudf::detail::make_counting_transform_iterator(
@@ -53,7 +52,7 @@ void BM_basic_sum(benchmark::State& state)
   wrapper keys(data_it, data_it + column_size);
   wrapper vals(data_it, data_it + column_size);
 
-  cudf::groupby::groupby gb_obj(cudf::table_view({keys}));
+  cudf::groupby::groupby gb_obj(cudf::table_view({keys, keys, keys}));
 
   std::vector<cudf::groupby::aggregation_request> requests;
   requests.emplace_back(cudf::groupby::aggregation_request());
@@ -73,7 +72,9 @@ BENCHMARK_REGISTER_F(Groupby, Basic)
   ->UseManualTime()
   ->Unit(benchmark::kMillisecond)
   ->Arg(10000)
-  ->Arg(10000000);
+  ->Arg(1000000)
+  ->Arg(10000000)
+  ->Arg(100000000);
 
 void BM_pre_sorted_sum(benchmark::State& state)
 {
diff --git a/cpp/benchmarks/hashing/hash_benchmark.cpp b/cpp/benchmarks/hashing/hash_benchmark.cpp
index 77b10399693..4ccb0bfad9d 100644
--- a/cpp/benchmarks/hashing/hash_benchmark.cpp
+++ b/cpp/benchmarks/hashing/hash_benchmark.cpp
@@ -25,10 +25,14 @@
 class HashBenchmark : public cudf::benchmark {
 };
 
-static void BM_hash(benchmark::State& state, cudf::hash_id hid)
+enum contains_nulls { no_nulls, nulls };
+
+static void BM_hash(benchmark::State& state, cudf::hash_id hid, contains_nulls has_nulls)
 {
   cudf::size_type const n_rows{(cudf::size_type)state.range(0)};
   auto const data = create_random_table({cudf::type_id::INT64}, 1, row_count{n_rows});
+  if (has_nulls == contains_nulls::no_nulls)
+    data->get_column(0).set_null_mask(rmm::device_buffer{}, 0);
 
   for (auto _ : state) {
     cuda_event_timer raii(state, true, rmm::cuda_stream_default);
@@ -36,16 +40,25 @@ static void BM_hash(benchmark::State& state, cudf::hash_id hid)
   }
 }
 
-#define HASH_BENCHMARK_DEFINE(name)                               \
-  BENCHMARK_DEFINE_F(HashBenchmark, name)                         \
-  (::benchmark::State & st) { BM_hash(st, cudf::hash_id::name); } \
-  BENCHMARK_REGISTER_F(HashBenchmark, name)                       \
-    ->RangeMultiplier(4)                                          \
-    ->Ranges({{1 << 14, 1 << 24}})                                \
-    ->UseManualTime()                                             \
+#define concat(a, b, c) a##b##c
+
+#define H_BENCHMARK_DEFINE(name, hid, n)                                            \
+  BENCHMARK_DEFINE_F(HashBenchmark, name)                                           \
+  (::benchmark::State & st) { BM_hash(st, cudf::hash_id::hid, contains_nulls::n); } \
+  BENCHMARK_REGISTER_F(HashBenchmark, name)                                         \
+    ->RangeMultiplier(4)                                                            \
+    ->Ranges({{1 << 14, 1 << 24}})                                                  \
+    ->UseManualTime()                                                               \
     ->Unit(benchmark::kMillisecond);
 
-HASH_BENCHMARK_DEFINE(HASH_MURMUR3)
-HASH_BENCHMARK_DEFINE(HASH_MD5)
-HASH_BENCHMARK_DEFINE(HASH_SERIAL_MURMUR3)
-HASH_BENCHMARK_DEFINE(HASH_SPARK_MURMUR3)
+#define HASH_BENCHMARK_DEFINE(hid, n) H_BENCHMARK_DEFINE(concat(hid, _, n), hid, n)
+
+HASH_BENCHMARK_DEFINE(HASH_MURMUR3, nulls)
+HASH_BENCHMARK_DEFINE(HASH_MD5, nulls)
+HASH_BENCHMARK_DEFINE(HASH_SERIAL_MURMUR3, nulls)
+HASH_BENCHMARK_DEFINE(HASH_SPARK_MURMUR3, nulls)
+
+HASH_BENCHMARK_DEFINE(HASH_MURMUR3, no_nulls)
+HASH_BENCHMARK_DEFINE(HASH_MD5, no_nulls)
+HASH_BENCHMARK_DEFINE(HASH_SERIAL_MURMUR3, no_nulls)
+HASH_BENCHMARK_DEFINE(HASH_SPARK_MURMUR3, no_nulls)
diff --git a/cpp/include/cudf/column/column_device_view.cuh b/cpp/include/cudf/column/column_device_view.cuh
index 6ecb0796283..a15f20ef52d 100644
--- a/cpp/include/cudf/column/column_device_view.cuh
+++ b/cpp/include/cudf/column/column_device_view.cuh
@@ -44,23 +44,30 @@
 namespace cudf {
 
 /**
- * @brief Policy for what assumptions the optional iterator has about null values
+ * @brief Indicates the presence of nulls at compile-time or runtime.
  *
- * - `YES` means that the column supports nulls and has null values, therefore
- *    the optional might not contain a value
+ * If used at compile-time, this indicator can tell the optimizer
+ * to include or exclude any null-checking clauses.
  *
- * - `NO` means that the column has no null values, therefore the optional will
- *    always have a value
- *
- * - `DYNAMIC` defers the assumption of nullability to runtime with the users stating
- *    on construction of the iterator if column has nulls.
  */
-struct contains_nulls {
-  struct YES {
+struct nullate {
+  struct YES : std::bool_constant<true> {
   };
-  struct NO {
+  struct NO : std::bool_constant<false> {
   };
   struct DYNAMIC {
+    DYNAMIC() = delete;
+    /**
+     * @brief Create a runtime nullate object.
+     *
+     * @see cudf::column_device_view::optional_begin for example usage
+     *
+     * @param b True if nulls are expected in the operation in which this
+     *          object is applied.
+     */
+    constexpr explicit DYNAMIC(bool b) noexcept : value{b} {}
+    constexpr operator bool() const noexcept { return value; }
+    bool value;  ///< True if nulls are expected
   };
 };
 
@@ -282,7 +289,7 @@ class alignas(16) column_device_view_base {
 // Forward declaration
 template <typename T>
 struct value_accessor;
-template <typename T, typename contains_nulls_mode>
+template <typename T, typename Nullate>
 struct optional_accessor;
 template <typename T, bool has_nulls>
 struct pair_accessor;
@@ -493,11 +500,11 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
   }
 
   /**
-   * @brief optional iterator for navigating this column
+   * @brief Optional iterator for navigating this column
    */
-  template <typename T, typename contains_nulls_mode>
+  template <typename T, typename Nullate>
   using const_optional_iterator =
-    thrust::transform_iterator<detail::optional_accessor<T, contains_nulls_mode>, count_it>;
+    thrust::transform_iterator<detail::optional_accessor<T, Nullate>, count_it>;
 
   /**
    * @brief Pair iterator for navigating this column
@@ -520,117 +527,57 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
    *
    * Dereferencing the returned iterator returns a `thrust::optional<T>`.
    *
-   * When the element of an iterator contextually converted to bool, the conversion returns true
+   * The element of this iterator contextually converts to bool. The conversion returns true
    * if the object contains a value and false if it does not contain a value.
    *
-   * optional_begin with mode `DYNAMIC` defers the assumption of nullability to
-   * runtime, with the user stating on construction of the iterator if column has nulls.
-   * `DYNAMIC` mode is nice when an algorithm is going to execute on multiple
-   * iterators and you don't want to compile all the combinations of iterator types
-   *
-   * Example:
+   * Calling this method with `nullate::DYNAMIC` defers the assumption of nullability to
+   * runtime with the caller indicating if the column has nulls. The `nullate::DYNAMIC` is
+   * useful when an algorithm is going to execute on multiple iterators and all the combinations of
+   * iterator types are not required at compile time.
    *
-   * \code{.cpp}
+   * @code{.cpp}
    * template<typename T>
    * void some_function(cudf::column_view<T> const& col_view){
    *    auto d_col = cudf::column_device_view::create(col_view);
    *    // Create a `DYNAMIC` optional iterator
-   *    auto optional_iterator = d_col->optional_begin<T>(cudf::contains_nulls::DYNAMIC{},
-   *                                                      col_view.has_nulls());
-   * }
-   * \endcode
-   *
-   * This function does not participate in overload resolution if
-   * `column_device_view::has_element_accessor<T>()` is false.
-   *
-   * @throws cudf::logic_error if the column is not nullable, and `DYNAMIC` mode used and
-   *         the user has stated nulls exist
-   * @throws cudf::logic_error if column datatype and Element type mismatch.
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  auto optional_begin(contains_nulls::DYNAMIC, bool has_nulls) const
-  {
-    return const_optional_iterator<T, contains_nulls::DYNAMIC>{
-      count_it{0}, detail::optional_accessor<T, contains_nulls::DYNAMIC>{*this, has_nulls}};
-  }
-
-  /**
-   * @brief Return an optional iterator to the first element of the column.
-   *
-   * Dereferencing the returned iterator returns a `thrust::optional<T>`.
-   *
-   * When the element of an iterator contextually converted to bool, the conversion returns true
-   * if the object contains a value and false if it does not contain a value.
-   *
-   * optional_begin with mode `YES` means that the column supports nulls and
-   * potentially has null values, therefore the optional might not contain a value
-   *
-   * Example:
-   *
-   * \code{.cpp}
-   * template<typename T, bool has_nulls>
-   * void some_function(cudf::column_view<T> const& col_view){
-   *    auto d_col = cudf::column_device_view::create(col_view);
-   *    if constexpr(has_nulls) {
-   *      auto optional_iterator = d_col->optional_begin<T>(cudf::contains_nulls::YES{});
-   *      //use optional_iterator
-   *    } else {
-   *      auto optional_iterator = d_col->optional_begin<T>(cudf::contains_nulls::NO{});
-   *      //use optional_iterator
-   *    }
+   *    auto optional_iterator =
+   *       d_col->optional_begin<T>(cudf::nullate::DYNAMIC{col_view.has_nulls()});
    * }
-   * \endcode
+   * @endcode
    *
-   * This function does not participate in overload resolution if
-   * `column_device_view::has_element_accessor<T>()` is false.
+   * Calling this method with `nullate::YES` means that the column supports nulls and
+   * the optional returned might not contain a value.
    *
-   * @throws cudf::logic_error if the column is not nullable, and `YES` mode used
-   * @throws cudf::logic_error if column datatype and Element type mismatch.
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  auto optional_begin(contains_nulls::YES) const
-  {
-    return const_optional_iterator<T, contains_nulls::YES>{
-      count_it{0}, detail::optional_accessor<T, contains_nulls::YES>{*this}};
-  }
-
-  /**
-   * @brief Return an optional iterator to the first element of the column.
+   * Calling this method with `nullate::NO` means that the column has no null values
+   * and the optional returned will always contain a value.
    *
-   * Dereferencing the returned iterator returns a `thrust::optional<T>`.
-   *
-   * When the element of an iterator contextually converted to bool, the conversion returns true
-   * if the object contains a value and false if it does not contain a value.
-   *
-   * optional_begin with mode `NO` means that the column has no null values,
-   * therefore the optional will always contain a value.
-   *
-   * Example:
-   *
-   * \code{.cpp}
+   * @code{.cpp}
    * template<typename T, bool has_nulls>
    * void some_function(cudf::column_view<T> const& col_view){
    *    auto d_col = cudf::column_device_view::create(col_view);
    *    if constexpr(has_nulls) {
-   *      auto optional_iterator = d_col->optional_begin<T>(cudf::contains_nulls::YES{});
+   *      auto optional_iterator = d_col->optional_begin<T>(cudf::nullate::YES{});
    *      //use optional_iterator
    *    } else {
-   *      auto optional_iterator = d_col->optional_begin<T>(cudf::contains_nulls::NO{});
+   *      auto optional_iterator = d_col->optional_begin<T>(cudf::nullate::NO{});
    *      //use optional_iterator
    *    }
    * }
-   * \endcode
+   * @endcode
    *
    * This function does not participate in overload resolution if
    * `column_device_view::has_element_accessor<T>()` is false.
    *
+   * @throws cudf::logic_error if the column is not nullable and `has_nulls` evaluates to true.
    * @throws cudf::logic_error if column datatype and Element type mismatch.
    */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  auto optional_begin(contains_nulls::NO) const
+  template <typename T,
+            typename Nullate,
+            CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
+  auto optional_begin(Nullate has_nulls) const
   {
-    return const_optional_iterator<T, contains_nulls::NO>{
-      count_it{0}, detail::optional_accessor<T, contains_nulls::NO>{*this}};
+    return const_optional_iterator<T, Nullate>{
+      count_it{0}, detail::optional_accessor<T, Nullate>{*this, has_nulls}};
   }
 
   /**
@@ -695,57 +642,21 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
    * @brief Return an optional iterator to the element following the last element of
    * the column.
    *
-   * Dereferencing the returned iterator returns a `thrust::optional<T>`.
+   * The returned iterator represents a `thrust::optional<T>` element.
    *
    * This function does not participate in overload resolution if
    * `column_device_view::has_element_accessor<T>()` is false.
    *
-   * @throws cudf::logic_error if the column is not nullable, and `DYNAMIC` mode used and
-   *         the user has stated nulls exist
+   * @throws cudf::logic_error if the column is not nullable and `has_nulls` is true
    * @throws cudf::logic_error if column datatype and Element type mismatch.
    */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  auto optional_end(contains_nulls::DYNAMIC, bool has_nulls) const
-  {
-    return const_optional_iterator<T, contains_nulls::DYNAMIC>{
-      count_it{size()}, detail::optional_accessor<T, contains_nulls::DYNAMIC>{*this, has_nulls}};
-  }
-
-  /**
-   * @brief Return an optional iterator to the element following the last element of
-   * the column.
-   *
-   * Dereferencing the returned iterator returns a `thrust::optional<T>`.
-   *
-   * This function does not participate in overload resolution if
-   * `column_device_view::has_element_accessor<T>()` is false.
-   *
-   * @throws cudf::logic_error if the column is not nullable, and `YES` mode used
-   * @throws cudf::logic_error if column datatype and Element type mismatch.
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  auto optional_end(contains_nulls::YES) const
-  {
-    return const_optional_iterator<T, contains_nulls::YES>{
-      count_it{size()}, detail::optional_accessor<T, contains_nulls::YES>{*this}};
-  }
-
-  /**
-   * @brief Return an optional iterator to the element following the last element of
-   * the column.
-   *
-   * Dereferencing the returned iterator returns a `thrust::optional<T>`.
-   *
-   * This function does not participate in overload resolution if
-   * `column_device_view::has_element_accessor<T>()` is false.
-   *
-   * @throws cudf::logic_error if column datatype and Element type mismatch.
-   */
-  template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
-  auto optional_end(contains_nulls::NO) const
+  template <typename T,
+            typename Nullate,
+            CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
+  auto optional_end(Nullate has_nulls) const
   {
-    return const_optional_iterator<T, contains_nulls::NO>{
-      count_it{size()}, detail::optional_accessor<T, contains_nulls::NO>{*this}};
+    return const_optional_iterator<T, Nullate>{
+      count_it{size()}, detail::optional_accessor<T, Nullate>{*this, has_nulls}};
   }
 
   /**
@@ -1201,77 +1112,56 @@ struct value_accessor {
  * @brief optional accessor of a column
  *
  *
- * The optional_accessor always returns a thrust::optional of column[i]. The validity
- * of the optional is determined by the contains_nulls_mode template parameter
- * which has the following modes:
+ * The optional_accessor always returns a `thrust::optional` of `column[i]`. The validity
+ * of the optional is determined by the `Nullate` parameter which may be one of the following:
  *
- * - `YES` means that the column supports nulls and has null values, therefore
- *    the optional might be valid or invalid
+ * - `nullate::YES` means that the column supports nulls and the optional returned
+ *    might be valid or invalid.
  *
- * - `NO` the user has attested that the column has no null values,
+ * - `nullate::NO` means the caller attests that the column has no null values,
  *    no checks will occur and `thrust::optional{column[i]}` will be
  *    return for each `i`.
  *
- * - `DYNAMIC` defers the assumption of nullability to runtime with the users stating
- *    on construction of the iterator if column has nulls.
- *    When `with_nulls=true` the return value validity will be determined if column[i]
- *    is not null.
- *    When `with_nulls=false` the return value will always be valid
+ * - `nullate::DYNAMIC` defers the assumption of nullability to runtime and the caller
+ *    specifies if the column has nulls at runtime.
+ *    For `DYNAMIC{true}` the return value will be `thrust::optional{column[i]}` if
+ *      element `i` is not null and `thrust::optional{}` if element `i` is null.
+ *    For `DYNAMIC{false}` the return value will always be `thrust::optional{column[i]}`.
  *
  * @throws cudf::logic_error if column datatype and template T type mismatch.
- * @throws cudf::logic_error if the column is not nullable, and `with_nulls=true`
- *
+ * @throws cudf::logic_error if the column is not nullable and `with_nulls` evaluates to true
  *
  * @tparam T The type of elements in the column
- * @tparam contains_nulls_mode Specifies if nulls are checked at runtime or compile time.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <typename T, typename contains_nulls_mode>
+template <typename T, typename Nullate>
 struct optional_accessor {
   column_device_view const col;  ///< column view of column in device
 
   /**
-   * @brief constructor
-   * @param[in] _col column device view of cudf column
+   * @brief Constructor
+   *
+   * @param col Column on which to iterator over its elements.
+   * @param with_nulls Indicates if the `col` should be checked for nulls.
    */
-  optional_accessor(column_device_view const& _col) : col{_col}
+  optional_accessor(column_device_view const& _col, Nullate with_nulls)
+    : col{_col}, has_nulls{with_nulls}
   {
     CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
+    if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
   }
 
   CUDA_DEVICE_CALLABLE
   thrust::optional<T> operator()(cudf::size_type i) const
   {
-    if constexpr (std::is_same_v<contains_nulls_mode, contains_nulls::YES>) {
+    if (has_nulls) {
       return (col.is_valid_nocheck(i)) ? thrust::optional<T>{col.element<T>(i)}
                                        : thrust::optional<T>{thrust::nullopt};
     }
     return thrust::optional<T>{col.element<T>(i)};
   }
-};
-
-template <typename T>
-struct optional_accessor<T, contains_nulls::DYNAMIC> {
-  column_device_view const col;  ///< column view of column in device
-  bool has_nulls;
-
-  /**
-   * @brief constructor
-   * @param[in] _col column device view of cudf column
-   * @param[in] with_nulls Indicates if @p _col has nulls
-   */
-  optional_accessor(column_device_view const& _col, bool with_nulls)
-    : col{_col}, has_nulls{with_nulls}
-  {
-    CUDF_EXPECTS(type_id_matches_device_storage_type<T>(col.type().id()), "the data type mismatch");
-    if (with_nulls) { CUDF_EXPECTS(_col.nullable(), "Unexpected non-nullable column."); }
-  }
 
-  CUDA_DEVICE_CALLABLE
-  thrust::optional<T> operator()(cudf::size_type i) const
-  {
-    return (has_nulls and col.is_null_nocheck(i)) ? thrust::optional<T>{thrust::nullopt}
-                                                  : thrust::optional<T>{col.element<T>(i)};
-  }
+  Nullate has_nulls{};
 };
 
 /**
diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh
index 3e789299716..01742384972 100644
--- a/cpp/include/cudf/detail/iterator.cuh
+++ b/cpp/include/cudf/detail/iterator.cuh
@@ -171,127 +171,61 @@ auto make_null_replacement_iterator(column_device_view const& column,
  *
  * Dereferencing the returned iterator returns a `thrust::optional<Element>`.
  *
- * When the element of an iterator contextually converted to bool, the conversion returns true
+ * The element of this iterator contextually converts to bool. The conversion returns true
  * if the object contains a value and false if it does not contain a value.
  *
- * make_optional_iterator with mode `DYNAMIC` defers the assumption of nullability to
- * runtime, with the user stating on construction of the iterator if column has nulls.
- * `DYNAMIC` mode is nice when an algorithm is going to execute on multiple
- * iterators and you don't want to compile all the combinations of iterator types
+ * Calling this function with `nullate::DYNAMIC` defers the assumption
+ * of nullability to runtime with the caller indicating if the column has nulls.
+ * This is useful when an algorithm is going to execute on multiple iterators and all
+ * the combinations of iterator types are not required at compile time.
  *
- * Example:
- *
- * \code{.cpp}
+ * @code{.cpp}
  * template<typename T>
  * void some_function(cudf::column_view<T> const& col_view){
  *    auto d_col = cudf::column_device_view::create(col_view);
  *    // Create a `DYNAMIC` optional iterator
- *    auto optional_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                cudf::contains_nulls::DYNAMIC{},
- *                                                col_view.has_nulls());
+ *    auto optional_iterator =
+ *      cudf::detail::make_optional_iterator<T>(
+ *        d_col, cudf::nullate::DYNAMIC{col_view.has_nulls()});
  * }
- * \endcode
- *
- * @throws cudf::logic_error if the column is not nullable, and `DYNAMIC` mode used and
- *         the user has stated nulls exist
- * @throws cudf::logic_error if column datatype and Element type mismatch.
- *
- * @tparam Element The type of elements in the column
- * @param column The column to iterate
- * @return Iterator that returns valid column elements and the validity of the
- * element in a thrust::optional
- */
-template <typename Element>
-auto make_optional_iterator(column_device_view const& column,
-                            contains_nulls::DYNAMIC,
-                            bool has_nulls)
-{
-  return column.optional_begin<Element>(contains_nulls::DYNAMIC{}, has_nulls);
-}
-
-/**
- * @brief Constructs an optional iterator over a column's values and its validity.
- *
- * Dereferencing the returned iterator returns a `thrust::optional<Element>`.
- *
- * When the element of an iterator contextually converted to bool, the conversion returns true
- * if the object contains a value and false if it does not contain a value.
- *
- * make_optional_iterator with mode `YES` means that the column supports nulls and
- * potentially has null values, therefore the optional might not contain a value
+ * @endcode
  *
- * Example:
+ * Calling this function with `nullate::YES` means that the column supports
+ * nulls and the optional returned might not contain a value.
+ * Calling this function with `nullate::NO` means that the column has no
+ * null values and the optional returned will always contain a value.
  *
- * \code{.cpp}
+ * @code{.cpp}
  * template<typename T, bool has_nulls>
  * void some_function(cudf::column_view<T> const& col_view){
  *    auto d_col = cudf::column_device_view::create(col_view);
  *    if constexpr(has_nulls) {
- *      auto optional_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                  cudf::contains_nulls::YES{});
+ *      auto optional_iterator =
+ *        cudf::detail::make_optional_iterator<T>(d_col, cudf::nullate::YES{});
  *      //use optional_iterator
  *    } else {
- *      auto optional_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                  cudf::contains_nulls::NO{});
+ *      auto optional_iterator =
+ *        cudf::detail::make_optional_iterator<T>(d_col, cudf::nullate::NO{});
  *      //use optional_iterator
  *    }
  * }
- * \endcode
+ * @endcode
  *
- * @throws cudf::logic_error if the column is not nullable, and `YES` mode used
+ * @throws cudf::logic_error if the column is not nullable and `has_nulls` is true.
  * @throws cudf::logic_error if column datatype and Element type mismatch.
  *
- * @tparam Element The type of elements in the column
- * @param column The column to iterate
- * @return Iterator that returns column elements and the validity of the
- * element as a thrust::optional
- */
-template <typename Element>
-auto make_optional_iterator(column_device_view const& column, contains_nulls::YES)
-{
-  return column.optional_begin<Element>(contains_nulls::YES{});
-}
-
-/**
- * @brief Constructs an optional iterator over a column's values and its validity.
- *
- * Dereferencing the returned iterator returns a `thrust::optional<Element>`.
- *
- * When the element of an iterator contextually converted to bool, the conversion returns true
- * if the object contains a value and false if it does not contain a value.
- *
- * make_optional_iterator with mode `NO` means that the column has no null values,
- * therefore the optional will always contain a value.
- *
- * Example:
- *
- * \code{.cpp}
- * template<typename T, bool has_nulls>
- * void some_function(cudf::column_view<T> const& col_view){
- *    auto d_col = cudf::column_device_view::create(col_view);
- *    if constexpr(has_nulls) {
- *      auto optional_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                  cudf::contains_nulls::YES{});
- *      //use optional_iterator
- *    } else {
- *      auto optional_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                  cudf::contains_nulls::NO{});
- *      //use optional_iterator
- *    }
- * }
- * \endcode
- *
- * @throws cudf::logic_error if column datatype and Element type mismatch.
+ * @tparam Element The type of elements in the column.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  *
- * @tparam Element The type of elements in the column
  * @param column The column to iterate
- * @return Iterator that returns column elements and the validity of the
- * element in a thrust::optional
+ * @param has_nulls Indicates whether `column` is checked for nulls.
+ * @return Iterator that returns valid column elements and the validity of the
+ * element in a `thrust::optional`
  */
-template <typename Element>
-auto make_optional_iterator(column_device_view const& column, contains_nulls::NO)
+template <typename Element, typename Nullate>
+auto make_optional_iterator(column_device_view const& column, Nullate has_nulls)
 {
-  return column.optional_begin<Element>(contains_nulls::NO{});
+  return column.optional_begin<Element, Nullate>(has_nulls);
 }
 
 /**
@@ -447,40 +381,38 @@ auto inline make_scalar_iterator(scalar const& scalar_value)
                                          scalar_value_accessor<Element>{scalar_value});
 }
 
-template <typename Element, typename contains_nulls_mode>
-struct scalar_optional_accessor;
-
 /**
- * @brief optional accessor of a maybe-nullable scalar
- *
- * The scalar_optional_accessor always returns a thrust::optional of the scalar.
- * The validity of the optional is determined by the contains_nulls_mode template parameter
- * which has the following modes:
+ * @brief Optional accessor for a scalar
  *
- * `DYNAMIC`: Defer nullability checks to runtime
+ * The `scalar_optional_accessor` always returns a `thrust::optional` of the scalar.
+ * The validity of the optional is determined by the `Nullate` parameter which may
+ * be one of the following:
  *
- *  - When `with_nulls=true` the return value will be a `thrust::optional{scalar}`
- *    when scalar is valid, and `thrust::optional{}` when the scalar is invalid.
+ * - `nullate::YES` means that the scalar may be valid or invalid and the optional returned
+ *    will contain a value only if the scalar is valid.
  *
- *  - When `with_nulls=false` the return value will always be `thrust::optional{scalar}`
+ * - `nullate::NO` means the caller attests that the scalar will always be valid,
+ *    no checks will occur and `thrust::optional{column[i]}` will return a value
+ *    for each `i`.
  *
- * `NO`: No null values will occur for this scalar, no checks will occur
- *  and `thrust::optional{scalar}` will always be returned.
- *
- * `YES`: null values will occur for this scalar,
- *  and `thrust::optional{scalar}` will always be returned.
+ * - `nullate::DYNAMIC` defers the assumption of nullability to runtime and the caller
+ *    specifies if the scalar may be valid or invalid.
+ *    For `DYNAMIC{true}` the return value will be a `thrust::optional{scalar}` when the
+ *      scalar is valid and a `thrust::optional{}` when the scalar is invalid.
+ *    For `DYNAMIC{false}` the return value will always be a `thrust::optional{scalar}`.
  *
  * @throws `cudf::logic_error` if scalar datatype and Element type mismatch.
  *
  * @tparam Element The type of return type of functor
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <typename Element, typename contains_nulls_mode>
+template <typename Element, typename Nullate>
 struct scalar_optional_accessor : public scalar_value_accessor<Element> {
   using super_t    = scalar_value_accessor<Element>;
   using value_type = thrust::optional<Element>;
 
-  scalar_optional_accessor(scalar const& scalar_value)
-    : scalar_value_accessor<Element>(scalar_value)
+  scalar_optional_accessor(scalar const& scalar_value, Nullate with_nulls)
+    : scalar_value_accessor<Element>(scalar_value), has_nulls{with_nulls}
   {
   }
 
@@ -494,32 +426,14 @@ struct scalar_optional_accessor : public scalar_value_accessor<Element> {
   CUDA_HOST_DEVICE_CALLABLE
   const value_type operator()(size_type) const
   {
-    if constexpr (std::is_same_v<contains_nulls_mode, contains_nulls::YES>) {
+    if (has_nulls) {
       return (super_t::dscalar.is_valid()) ? Element{super_t::dscalar.value()}
                                            : value_type{thrust::nullopt};
     }
     return Element{super_t::dscalar.value()};
   }
-};
 
-template <typename Element>
-struct scalar_optional_accessor<Element, cudf::contains_nulls::DYNAMIC>
-  : public scalar_value_accessor<Element> {
-  using super_t    = scalar_value_accessor<Element>;
-  using value_type = thrust::optional<Element>;
-  bool has_nulls;
-
-  scalar_optional_accessor(scalar const& scalar_value, bool with_nulls)
-    : scalar_value_accessor<Element>(scalar_value), has_nulls{with_nulls}
-  {
-  }
-
-  CUDA_HOST_DEVICE_CALLABLE
-  const value_type operator()(size_type) const
-  {
-    return (has_nulls and !super_t::dscalar.is_valid()) ? value_type{thrust::nullopt}
-                                                        : Element{super_t::dscalar.value()};
-  }
+  Nullate has_nulls{};
 };
 
 /**
@@ -622,156 +536,70 @@ struct scalar_representation_pair_accessor : public scalar_value_accessor<Elemen
  *
  * Dereferencing the returned iterator returns a `thrust::optional<Element>`.
  *
- * When the element of an iterator contextually converted to bool, the conversion returns true
+ * The element of this iterator contextually converts to bool. The conversion returns true
  * if the object contains a value and false if it does not contain a value.
  *
  * The iterator behavior is undefined if the scalar is destroyed before iterator dereferencing.
  *
- * make_optional_iterator with mode `DYNAMIC` defers the assumption of nullability to
- * runtime, with the user stating on construction of the iterator if scalar has nulls.
- *
- * Example:
+ * Calling this function with `nullate::DYNAMIC` defers the assumption
+ * of nullability to runtime with the caller indicating if the scalar is valid.
  *
- * \code{.cpp}
+ * @code{.cpp}
  * template<typename T>
  * void some_function(cudf::column_view<T> const& col_view,
  *                    scalar const& scalar_value,
  *                    bool col_has_nulls){
  *    auto d_col = cudf::column_device_view::create(col_view);
- *    auto column_iterator = cudf::detail::make_optional_iterator<T>(d_col,
-                                      cudf::contains_nulls::DYNAMIC{}, col_has_nulls);
- *    auto scalar_iterator = cudf::detail::make_optional_iterator<T>(scalar_value,
-                                      cudf::contains_nulls::DYNAMIC{}, scalar_value.is_valid());
+ *    auto column_iterator = cudf::detail::make_optional_iterator<T>(
+ *      d_col, cudf::nullate::DYNAMIC{col_has_nulls});
+ *    auto scalar_iterator = cudf::detail::make_optional_iterator<T>(
+ *      scalar_value, cudf::nullate::DYNAMIC{scalar_value.is_valid()});
  *    //use iterators
  * }
- * \endcode
- *
- * @throws cudf::logic_error if the scalar is not nullable, and `DYNAMIC` mode used and
- *         the user has stated nulls exist
- * @throws cudf::logic_error if scalar datatype and Element type mismatch.
- *
- * @tparam Element The type of elements in the scalar
- * @tparam has_nulls If the scalar value will have a null at runtime
- * @param scalar_value The scalar to iterate
- * @return Iterator that returns scalar elements and validity of the
- * element in a thrust::optional
- */
-template <typename Element>
-auto inline make_optional_iterator(scalar const& scalar_value,
-                                   contains_nulls::DYNAMIC,
-                                   bool has_nulls)
-{
-  CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
-               "the data type mismatch");
-  return thrust::make_transform_iterator(
-    thrust::make_constant_iterator<size_type>(0),
-    scalar_optional_accessor<Element, contains_nulls::DYNAMIC>{scalar_value, has_nulls});
-}
-
-/**
- * @brief Constructs an optional iterator over a scalar's values and its validity.
- *
- * Dereferencing the returned iterator returns a `thrust::optional<Element>`.
- *
- * When the element of an iterator contextually converted to bool, the conversion returns true
- * if the object contains a value and false if it does not contain a value.
- *
- * The iterator behavior is undefined if the scalar is destroyed before iterator dereferencing.
- *
- * make_optional_iterator ith mode `YES` means that the scalar supports nulls and
- * potentially has null values, therefore the optional might not contain a value
- * therefore the optional will always contain a value.
+ * @endcode
  *
- * Example:
+ * Calling this function with `nullate::YES` means that the scalar maybe invalid
+ * and the optional return might not contain a value.
+ * Calling this function with `nullate::NO` means that the scalar is valid
+ * and the optional returned will always contain a value.
  *
- * \code{.cpp}
+ * @code{.cpp}
  * template<typename T, bool any_nulls>
  * void some_function(cudf::column_view<T> const& col_view, scalar const& scalar_value){
  *    auto d_col = cudf::column_device_view::create(col_view);
  *    if constexpr(any_nulls) {
- *      auto column_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                cudf::contains_nulls::YES{});
- *      auto scalar_iterator = cudf::detail::make_optional_iterator<T>(scalar_value,
- *                                                cudf::contains_nulls::YES{});
+ *      auto column_iterator =
+ *        cudf::detail::make_optional_iterator<T>(d_col, cudf::nullate::YES{});
+ *      auto scalar_iterator =
+ *        cudf::detail::make_optional_iterator<T>(scalar_value, cudf::nullate::YES{});
  *      //use iterators
  *    } else {
- *      auto column_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                cudf::contains_nulls::NO{});
- *      auto scalar_iterator = cudf::detail::make_optional_iterator<T>(scalar_value,
- *                                                cudf::contains_nulls::NO{});
+ *      auto column_iterator =
+ *        cudf::detail::make_optional_iterator<T>(d_col, cudf::nullate::NO{});
+ *      auto scalar_iterator =
+ *        cudf::detail::make_optional_iterator<T>(scalar_value, cudf::nullate::NO{});
  *      //use iterators
  *    }
  * }
- * \endcode
+ * @endcode
  *
- * @throws cudf::logic_error if the scalar is not nullable, and `YES` mode used
  * @throws cudf::logic_error if scalar datatype and Element type mismatch.
  *
  * @tparam Element The type of elements in the scalar
- * @param scalar_value The scalar to iterate
- * @return Iterator that returns scalar elements and the validity of the
- * element in a thrust::optional
- */
-template <typename Element>
-auto inline make_optional_iterator(scalar const& scalar_value, contains_nulls::YES)
-{
-  CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
-               "the data type mismatch");
-  return thrust::make_transform_iterator(
-    thrust::make_constant_iterator<size_type>(0),
-    scalar_optional_accessor<Element, contains_nulls::YES>{scalar_value});
-}
-
-/**
- * @brief Constructs an optional iterator over a scalar's values and its validity.
- *
- * Dereferencing the returned iterator returns a `thrust::optional<Element>`.
- *
- * When the element of an iterator contextually converted to bool, the conversion returns true
- * if the object contains a value and false if it does not contain a value.
- *
- * The iterator behavior is undefined if the scalar is destroyed before iterator dereferencing.
- *
- * make_optional_iterator with mode `NO` means that the scalar has no null values,
- * therefore the optional will always contain a value.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  *
- * Example:
- *
- * \code{.cpp}
- * template<typename T, bool any_nulls>
- * void some_function(cudf::column_view<T> const& col_view, scalar const& scalar_value){
- *    auto d_col = cudf::column_device_view::create(col_view);
- *    if constexpr(any_nulls) {
- *      auto column_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                cudf::contains_nulls::YES{});
- *      auto scalar_iterator = cudf::detail::make_optional_iterator<T>(scalar_value,
- *                                                cudf::contains_nulls::YES{});
- *      //use iterators
- *    } else {
- *      auto column_iterator = cudf::detail::make_optional_iterator<T>(d_col,
- *                                                cudf::contains_nulls::NO{});
- *      auto scalar_iterator = cudf::detail::make_optional_iterator<T>(scalar_value,
- *                                                cudf::contains_nulls::NO{});
- *      //use iterators
- *    }
- * }
- * \endcode
- *
- * @throws cudf::logic_error if scalar datatype and Element type mismatch.
- *
- * @tparam Element The type of elements in the scalar
- * @param scalar_value The scalar to iterate
- * @return Iterator that returns scalar elements and the validity of the
- * element in a thrust::optional
+ * @param scalar_value The scalar to be returned by the iterator.
+ * @param has_nulls Indicates if the scalar value may be invalid.
+ * @return Iterator that returns scalar and the validity of the scalar in a thrust::optional
  */
-template <typename Element>
-auto inline make_optional_iterator(scalar const& scalar_value, contains_nulls::NO)
+template <typename Element, typename Nullate>
+auto inline make_optional_iterator(scalar const& scalar_value, Nullate has_nulls)
 {
   CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
                "the data type mismatch");
   return thrust::make_transform_iterator(
     thrust::make_constant_iterator<size_type>(0),
-    scalar_optional_accessor<Element, contains_nulls::NO>{scalar_value});
+    scalar_optional_accessor<Element, Nullate>{scalar_value, has_nulls});
 }
 
 /**
diff --git a/cpp/include/cudf/detail/merge.cuh b/cpp/include/cudf/detail/merge.cuh
index ec83e348e33..f141d9b5d59 100644
--- a/cpp/include/cudf/detail/merge.cuh
+++ b/cpp/include/cudf/detail/merge.cuh
@@ -90,8 +90,8 @@ struct tagged_element_relational_comparator {
 
     column_device_view const* ptr_right_dview{r_side == side::LEFT ? &lhs : &rhs};
 
-    auto erl_comparator =
-      element_relational_comparator<has_nulls>(*ptr_left_dview, *ptr_right_dview, null_precedence);
+    auto erl_comparator = element_relational_comparator(
+      nullate::DYNAMIC{has_nulls}, *ptr_left_dview, *ptr_right_dview, null_precedence);
 
     return cudf::type_dispatcher(lhs.type(), erl_comparator, l_indx, r_indx);
   }
diff --git a/cpp/include/cudf/table/row_operators.cuh b/cpp/include/cudf/table/row_operators.cuh
index 70ccac2f75d..0f3ca073380 100644
--- a/cpp/include/cudf/table/row_operators.cuh
+++ b/cpp/include/cudf/table/row_operators.cuh
@@ -50,9 +50,9 @@ namespace detail {
 /**
  * @brief Compare the elements ordering with respect to `lhs`.
  *
- * @param[in] lhs first element
- * @param[in] rhs second element
- * @return weak_ordering Indicates the relationship between the elements in
+ * @param lhs first element
+ * @param rhs second element
+ * @return Indicates the relationship between the elements in
  * the `lhs` and `rhs` columns.
  */
 template <typename Element>
@@ -69,14 +69,15 @@ __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`. Specialization is to
- * handle `nan` in the order shown below.
+ * to derive the order of the elements with respect to `lhs`.
+ *
+ * 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)`
  *
- * @param[in] lhs first element
- * @param[in] rhs second element
- * @return weak_ordering Indicates the relationship between the elements in
+ * @param lhs first element
+ * @param rhs second element
+ * @return Indicates the relationship between the elements in
  * the `lhs` and `rhs` columns.
  */
 template <typename Element, std::enable_if_t<std::is_floating_point<Element>::value>* = nullptr>
@@ -119,7 +120,7 @@ inline __device__ auto null_compare(bool lhs_is_null, bool rhs_is_null, null_ord
  *
  * @param[in] lhs first element
  * @param[in] rhs second element
- * @return weak_ordering Indicates the relationship between the elements in
+ * @return Indicates the relationship between the elements in
  * the `lhs` and `rhs` columns.
  */
 template <typename Element, std::enable_if_t<not std::is_floating_point<Element>::value>* = nullptr>
@@ -132,9 +133,9 @@ __device__ weak_ordering relational_compare(Element lhs, Element rhs)
  * @brief A specialization for floating-point `Element` type to check if
  * `lhs` is equivalent to `rhs`. `nan == nan`.
  *
- * @param[in] lhs first element
- * @param[in] rhs second element
- * @return bool `true` if `lhs` == `rhs` else `false`.
+ * @param lhs first element
+ * @param rhs second element
+ * @return `true` if `lhs` == `rhs` else `false`.
  */
 template <typename Element, std::enable_if_t<std::is_floating_point<Element>::value>* = nullptr>
 __device__ bool equality_compare(Element lhs, Element rhs)
@@ -147,9 +148,9 @@ __device__ bool equality_compare(Element lhs, Element rhs)
  * @brief A specialization for non-floating-point `Element` type to check if
  * `lhs` is equivalent to `rhs`.
  *
- * @param[in] lhs first element
- * @param[in] rhs second element
- * @return bool `true` if `lhs` == `rhs` else `false`.
+ * @param lhs first element
+ * @param rhs second element
+ * @return `true` if `lhs` == `rhs` else `false`.
  */
 template <typename Element, std::enable_if_t<not std::is_floating_point<Element>::value>* = nullptr>
 __device__ bool equality_compare(Element const lhs, Element const rhs)
@@ -160,9 +161,9 @@ __device__ bool equality_compare(Element const lhs, Element const rhs)
 /**
  * @brief Performs an equality comparison between two elements in two columns.
  *
- * @tparam has_nulls Indicates the potential for null values in either column.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <bool has_nulls = true>
+template <typename Nullate>
 class element_equality_comparator {
  public:
   /**
@@ -171,14 +172,17 @@ class element_equality_comparator {
    *
    * @note `lhs` and `rhs` may be the same.
    *
+   * @param has_nulls Indicates if either input column contains nulls.
    * @param lhs The column containing the first element
    * @param rhs The column containing the second element (may be the same as lhs)
    * @param nulls_are_equal Indicates if two null elements are treated as equivalent
    */
-  __host__ __device__ element_equality_comparator(column_device_view lhs,
-                                                  column_device_view rhs,
-                                                  bool nulls_are_equal = true)
-    : lhs{lhs}, rhs{rhs}, nulls_are_equal{nulls_are_equal}
+  __host__ __device__
+  element_equality_comparator(Nullate has_nulls,
+                              column_device_view lhs,
+                              column_device_view rhs,
+                              null_equality nulls_are_equal = null_equality::EQUAL)
+    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal}
   {
   }
 
@@ -188,18 +192,17 @@ class element_equality_comparator {
    * @param lhs_element_index The index of the first element
    * @param rhs_element_index The index of the second element
    * @return True if both lhs and rhs element are both nulls and `nulls_are_equal` is true, or equal
-   *
    */
   template <typename Element,
             std::enable_if_t<cudf::is_equality_comparable<Element, Element>()>* = nullptr>
   __device__ bool operator()(size_type lhs_element_index,
                              size_type rhs_element_index) const noexcept
   {
-    if (has_nulls) {
+    if (nulls) {
       bool const lhs_is_null{lhs.is_null(lhs_element_index)};
       bool const rhs_is_null{rhs.is_null(rhs_element_index)};
       if (lhs_is_null and rhs_is_null) {
-        return nulls_are_equal;
+        return nulls_are_equal == null_equality::EQUAL;
       } else if (lhs_is_null != rhs_is_null) {
         return false;
       }
@@ -220,14 +223,18 @@ class element_equality_comparator {
  private:
   column_device_view lhs;
   column_device_view rhs;
-  bool nulls_are_equal;
+  Nullate nulls;
+  null_equality nulls_are_equal;
 };
 
-template <bool has_nulls = true>
+template <typename Nullate>
 class row_equality_comparator {
  public:
-  row_equality_comparator(table_device_view lhs, table_device_view rhs, bool nulls_are_equal = true)
-    : lhs{lhs}, rhs{rhs}, nulls_are_equal{nulls_are_equal}
+  row_equality_comparator(Nullate has_nulls,
+                          table_device_view lhs,
+                          table_device_view rhs,
+                          null_equality nulls_are_equal = true)
+    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, nulls_are_equal{nulls_are_equal}
   {
     CUDF_EXPECTS(lhs.num_columns() == rhs.num_columns(), "Mismatched number of columns.");
   }
@@ -236,7 +243,7 @@ class row_equality_comparator {
   {
     auto equal_elements = [=](column_device_view l, column_device_view r) {
       return cudf::type_dispatcher(l.type(),
-                                   element_equality_comparator<has_nulls>{l, r, nulls_are_equal},
+                                   element_equality_comparator{nulls, l, r, nulls_are_equal},
                                    lhs_row_index,
                                    rhs_row_index);
     };
@@ -247,15 +254,16 @@ class row_equality_comparator {
  private:
   table_device_view lhs;
   table_device_view rhs;
-  bool nulls_are_equal;
+  Nullate nulls;
+  null_equality nulls_are_equal;
 };
 
 /**
  * @brief Performs a relational comparison between two elements in two columns.
  *
- * @tparam has_nulls Indicates the potential for null values in either column.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <bool has_nulls = true>
+template <typename Nullate>
 class element_relational_comparator {
  public:
   /**
@@ -266,13 +274,21 @@ class element_relational_comparator {
    *
    * @param lhs The column containing the first element
    * @param rhs The column containing the second element (may be the same as lhs)
-   * @param null_precedence Indicates how null values are ordered with other
-   * values
+   * @param has_nulls Indicates if either input column contains nulls.
+   * @param null_precedence Indicates how null values are ordered with other values
    */
-  __host__ __device__ element_relational_comparator(column_device_view lhs,
+  __host__ __device__ element_relational_comparator(Nullate has_nulls,
+                                                    column_device_view lhs,
                                                     column_device_view rhs,
                                                     null_order null_precedence)
-    : lhs{lhs}, rhs{rhs}, null_precedence{null_precedence}
+    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}, null_precedence{null_precedence}
+  {
+  }
+
+  __host__ __device__ element_relational_comparator(Nullate has_nulls,
+                                                    column_device_view lhs,
+                                                    column_device_view rhs)
+    : lhs{lhs}, rhs{rhs}, nulls{has_nulls}
   {
   }
 
@@ -281,7 +297,7 @@ class element_relational_comparator {
    *
    * @param lhs_element_index The index of the first element
    * @param rhs_element_index The index of the second element
-   * @return weak_ordering Indicates the relationship between the elements in
+   * @return Indicates the relationship between the elements in
    * the `lhs` and `rhs` columns.
    */
   template <typename Element,
@@ -289,11 +305,11 @@ class element_relational_comparator {
   __device__ weak_ordering operator()(size_type lhs_element_index,
                                       size_type rhs_element_index) const noexcept
   {
-    if (has_nulls) {
+    if (nulls) {
       bool const lhs_is_null{lhs.is_null(lhs_element_index)};
       bool const rhs_is_null{rhs.is_null(rhs_element_index)};
 
-      if (lhs_is_null or rhs_is_null) {  // atleast one is null
+      if (lhs_is_null or rhs_is_null) {  // at least one is null
         return null_compare(lhs_is_null, rhs_is_null, null_precedence);
       }
     }
@@ -313,7 +329,8 @@ class element_relational_comparator {
  private:
   column_device_view lhs;
   column_device_view rhs;
-  null_order null_precedence;
+  Nullate nulls;
+  null_order null_precedence{};
 };
 
 /**
@@ -329,9 +346,9 @@ class element_relational_comparator {
  * second letter in both words is the first non-equal letter, and `a < b`, thus
  * `aac < abb`.
  *
- * @tparam has_nulls Indicates the potential for null values in either row.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <bool has_nulls = true>
+template <typename Nullate>
 class row_lexicographic_comparator {
  public:
   /**
@@ -343,6 +360,7 @@ class row_lexicographic_comparator {
    *
    * @param lhs The first table
    * @param rhs The second table (may be the same table as `lhs`)
+   * @param has_nulls Indicates if either input table contains columns with nulls.
    * @param column_order Optional, device array the same length as a row that
    * indicates the desired ascending/descending order of each column in a row.
    * If `nullptr`, it is assumed all columns are sorted in ascending order.
@@ -351,11 +369,16 @@ class row_lexicographic_comparator {
    * it is nullptr, then null precedence would be `null_order::BEFORE` for all
    * columns.
    */
-  row_lexicographic_comparator(table_device_view lhs,
+  row_lexicographic_comparator(Nullate has_nulls,
+                               table_device_view lhs,
                                table_device_view rhs,
                                order const* column_order         = nullptr,
                                null_order const* null_precedence = nullptr)
-    : _lhs{lhs}, _rhs{rhs}, _column_order{column_order}, _null_precedence{null_precedence}
+    : _lhs{lhs},
+      _rhs{rhs},
+      _nulls{has_nulls},
+      _column_order{column_order},
+      _null_precedence{null_precedence}
   {
     CUDF_EXPECTS(_lhs.num_columns() == _rhs.num_columns(), "Mismatched number of columns.");
     CUDF_EXPECTS(detail::is_relationally_comparable(_lhs, _rhs),
@@ -376,14 +399,14 @@ class row_lexicographic_comparator {
     for (size_type i = 0; i < _lhs.num_columns(); ++i) {
       bool ascending = (_column_order == nullptr) or (_column_order[i] == order::ASCENDING);
 
-      weak_ordering state{weak_ordering::EQUIVALENT};
       null_order null_precedence =
         _null_precedence == nullptr ? null_order::BEFORE : _null_precedence[i];
 
       auto comparator =
-        element_relational_comparator<has_nulls>{_lhs.column(i), _rhs.column(i), null_precedence};
+        element_relational_comparator{_nulls, _lhs.column(i), _rhs.column(i), null_precedence};
 
-      state = cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index);
+      weak_ordering state =
+        cudf::type_dispatcher(_lhs.column(i).type(), comparator, lhs_index, rhs_index);
 
       if (state == weak_ordering::EQUIVALENT) { continue; }
 
@@ -395,6 +418,7 @@ class row_lexicographic_comparator {
  private:
   table_device_view _lhs;
   table_device_view _rhs;
+  Nullate _nulls{};
   null_order const* _null_precedence{};
   order const* _column_order{};
 };  // class row_lexicographic_comparator
@@ -403,9 +427,9 @@ class row_lexicographic_comparator {
  * @brief Computes the hash value of an element in the given column.
  *
  * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam has_nulls Indicates the potential for null values in the column.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <template <typename> class hash_function, bool has_nulls = true>
+template <template <typename> class hash_function, typename Nullate>
 class element_hasher {
  public:
   template <typename T, CUDF_ENABLE_IF(column_device_view::has_element_accessor<T>())>
@@ -421,15 +445,20 @@ class element_hasher {
     cudf_assert(false && "Unsupported type in hash.");
     return {};
   }
+
+  Nullate has_nulls;
 };
 
-template <template <typename> class hash_function, bool has_nulls = true>
+template <template <typename> class hash_function, typename Nullate>
 class element_hasher_with_seed {
  public:
-  element_hasher_with_seed() = default;
-  __device__ element_hasher_with_seed(uint32_t seed) : _seed{seed} {}
-  __device__ element_hasher_with_seed(uint32_t seed, hash_value_type null_hash)
-    : _seed{seed}, _null_hash(null_hash)
+  __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed)
+    : _seed{seed}, _has_nulls{has_nulls}
+  {
+  }
+
+  __device__ element_hasher_with_seed(Nullate has_nulls, uint32_t seed, hash_value_type null_hash)
+    : _seed{seed}, _null_hash{null_hash}, _has_nulls{has_nulls}
   {
   }
 
@@ -450,20 +479,24 @@ class element_hasher_with_seed {
  private:
   uint32_t _seed{DEFAULT_HASH_SEED};
   hash_value_type _null_hash{std::numeric_limits<hash_value_type>::max()};
+  Nullate _has_nulls;
 };
 
 /**
  * @brief Computes the hash value of a row in the given table.
  *
  * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam has_nulls Indicates the potential for null values in the table.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <template <typename> class hash_function, bool has_nulls = true>
+template <template <typename> class hash_function, typename Nullate>
 class row_hasher {
  public:
   row_hasher() = delete;
-  row_hasher(table_device_view t) : _table{t} {}
-  row_hasher(table_device_view t, uint32_t seed) : _table{t}, _seed(seed) {}
+  row_hasher(Nullate has_nulls, table_device_view t) : _table{t}, _has_nulls{has_nulls} {}
+  row_hasher(Nullate has_nulls, table_device_view t, uint32_t seed)
+    : _table{t}, _seed(seed), _has_nulls{has_nulls}
+  {
+  }
 
   __device__ auto operator()(size_type row_index) const
   {
@@ -476,7 +509,7 @@ class row_hasher {
       hash_combiner(hash_value_type{0},
                     type_dispatcher<dispatch_storage_type>(
                       _table.column(0).type(),
-                      element_hasher_with_seed<hash_function, has_nulls>{_seed},
+                      element_hasher_with_seed<hash_function, Nullate>{_has_nulls, _seed},
                       _table.column(0),
                       row_index));
 
@@ -484,7 +517,7 @@ class row_hasher {
     auto hasher = [=](size_type column_index) {
       return cudf::type_dispatcher<dispatch_storage_type>(
         _table.column(column_index).type(),
-        element_hasher<hash_function, has_nulls>{},
+        element_hasher<hash_function, Nullate>{_has_nulls},
         _table.column(column_index),
         row_index);
     };
@@ -502,6 +535,7 @@ class row_hasher {
 
  private:
   table_device_view _table;
+  Nullate _has_nulls;
   uint32_t _seed{DEFAULT_HASH_SEED};
 };
 
@@ -510,14 +544,14 @@ class row_hasher {
  * initial hash value for each column.
  *
  * @tparam hash_function Hash functor to use for hashing elements.
- * @tparam has_nulls Indicates the potential for null values in the table.
+ * @tparam Nullate A cudf::nullate type describing how to check for nulls.
  */
-template <template <typename> class hash_function, bool has_nulls = true>
+template <template <typename> class hash_function, typename Nullate>
 class row_hasher_initial_values {
  public:
   row_hasher_initial_values() = delete;
-  row_hasher_initial_values(table_device_view t, hash_value_type* initial_hash)
-    : _table{t}, _initial_hash(initial_hash)
+  row_hasher_initial_values(Nullate has_nulls, table_device_view t, hash_value_type* initial_hash)
+    : _table{t}, _initial_hash(initial_hash), _has_nulls{has_nulls}
   {
   }
 
@@ -529,11 +563,11 @@ class row_hasher_initial_values {
 
     // Hashes an element in a column and combines with an initial value
     auto hasher = [=](size_type column_index) {
-      auto hash_value =
-        cudf::type_dispatcher<dispatch_storage_type>(_table.column(column_index).type(),
-                                                     element_hasher<hash_function, has_nulls>{},
-                                                     _table.column(column_index),
-                                                     row_index);
+      auto hash_value = cudf::type_dispatcher<dispatch_storage_type>(
+        _table.column(column_index).type(),
+        element_hasher<hash_function, Nullate>{_has_nulls},
+        _table.column(column_index),
+        row_index);
 
       return hash_combiner(_initial_hash[column_index], hash_value);
     };
@@ -550,6 +584,7 @@ class row_hasher_initial_values {
  private:
   table_device_view _table;
   hash_value_type* _initial_hash;
+  Nullate _has_nulls;
 };
 
 }  // namespace cudf
diff --git a/cpp/src/copying/copy.cu b/cpp/src/copying/copy.cu
index f4d09c8e0be..55173fd409f 100644
--- a/cpp/src/copying/copy.cu
+++ b/cpp/src/copying/copy.cu
@@ -75,10 +75,8 @@ struct copy_if_else_functor_impl<T, std::enable_if_t<is_rep_layout_compatible<T>
     auto const& lhs = *p_lhs;
     auto const& rhs = *p_rhs;
 
-    auto lhs_iter =
-      cudf::detail::make_optional_iterator<T>(lhs, contains_nulls::DYNAMIC{}, left_nullable);
-    auto rhs_iter =
-      cudf::detail::make_optional_iterator<T>(rhs, contains_nulls::DYNAMIC{}, right_nullable);
+    auto lhs_iter = cudf::detail::make_optional_iterator<T>(lhs, nullate::DYNAMIC{left_nullable});
+    auto rhs_iter = cudf::detail::make_optional_iterator<T>(rhs, nullate::DYNAMIC{right_nullable});
     return detail::copy_if_else(left_nullable || right_nullable,
                                 lhs_iter,
                                 lhs_iter + size,
@@ -112,10 +110,8 @@ struct copy_if_else_functor_impl<string_view> {
     auto const& lhs = *p_lhs;
     auto const& rhs = *p_rhs;
 
-    auto lhs_iter =
-      cudf::detail::make_optional_iterator<T>(lhs, contains_nulls::DYNAMIC{}, left_nullable);
-    auto rhs_iter =
-      cudf::detail::make_optional_iterator<T>(rhs, contains_nulls::DYNAMIC{}, right_nullable);
+    auto lhs_iter = cudf::detail::make_optional_iterator<T>(lhs, nullate::DYNAMIC{left_nullable});
+    auto rhs_iter = cudf::detail::make_optional_iterator<T>(rhs, nullate::DYNAMIC{right_nullable});
     return strings::detail::copy_if_else(lhs_iter, lhs_iter + size, rhs_iter, filter, stream, mr);
   }
 };
diff --git a/cpp/src/copying/segmented_shift.cu b/cpp/src/copying/segmented_shift.cu
index 62f992012cd..b08eaa0862c 100644
--- a/cpp/src/copying/segmented_shift.cu
+++ b/cpp/src/copying/segmented_shift.cu
@@ -76,11 +76,10 @@ struct segmented_shift_functor<T, std::enable_if_t<is_rep_layout_compatible<T>()
   {
     auto values_device_view = column_device_view::create(segmented_values, stream);
     bool nullable           = not fill_value.is_valid(stream) or segmented_values.nullable();
-    auto input_iterator =
-      cudf::detail::make_optional_iterator<T>(
-        *values_device_view, contains_nulls::DYNAMIC{}, segmented_values.has_nulls()) -
-      offset;
-    auto fill_iterator = cudf::detail::make_optional_iterator<T>(fill_value, contains_nulls::YES{});
+    auto input_iterator     = cudf::detail::make_optional_iterator<T>(
+                            *values_device_view, nullate::DYNAMIC{segmented_values.has_nulls()}) -
+                          offset;
+    auto fill_iterator = cudf::detail::make_optional_iterator<T>(fill_value, nullate::YES{});
     return copy_if_else(nullable,
                         input_iterator,
                         input_iterator + segmented_values.size(),
@@ -105,12 +104,10 @@ struct segmented_shift_functor<string_view> {
                                      rmm::mr::device_memory_resource* mr)
   {
     auto values_device_view = column_device_view::create(segmented_values, stream);
-    auto input_iterator =
-      make_optional_iterator<cudf::string_view>(
-        *values_device_view, contains_nulls::DYNAMIC{}, segmented_values.has_nulls()) -
-      offset;
-    auto fill_iterator =
-      make_optional_iterator<cudf::string_view>(fill_value, contains_nulls::YES{});
+    auto input_iterator     = make_optional_iterator<cudf::string_view>(
+                            *values_device_view, nullate::DYNAMIC{segmented_values.has_nulls()}) -
+                          offset;
+    auto fill_iterator = make_optional_iterator<cudf::string_view>(fill_value, nullate::YES{});
     return strings::detail::copy_if_else(input_iterator,
                                          input_iterator + segmented_values.size(),
                                          fill_iterator,
diff --git a/cpp/src/groupby/hash/groupby.cu b/cpp/src/groupby/hash/groupby.cu
index e35fa36a289..0c3e79ea36c 100644
--- a/cpp/src/groupby/hash/groupby.cu
+++ b/cpp/src/groupby/hash/groupby.cu
@@ -416,8 +416,8 @@ void sparse_to_dense_results(table_view const& keys,
  * @brief Construct hash map that uses row comparator and row hasher on
  * `d_keys` table and stores indices
  */
-template <bool keys_have_nulls>
 auto create_hash_map(table_device_view const& d_keys,
+                     bool keys_have_nulls,
                      null_policy include_null_keys,
                      rmm::cuda_stream_view stream)
 {
@@ -426,15 +426,17 @@ auto create_hash_map(table_device_view const& d_keys,
 
   using map_type = concurrent_unordered_map<size_type,
                                             size_type,
-                                            row_hasher<default_hash, keys_have_nulls>,
-                                            row_equality_comparator<keys_have_nulls>>;
+                                            row_hasher<default_hash, nullate::DYNAMIC>,
+                                            row_equality_comparator<nullate::DYNAMIC>>;
 
   using allocator_type = typename map_type::allocator_type;
 
-  bool const null_keys_are_equal{include_null_keys == null_policy::INCLUDE};
+  auto const null_keys_are_equal =
+    include_null_keys == null_policy::INCLUDE ? null_equality::EQUAL : null_equality::UNEQUAL;
 
-  row_hasher<default_hash, keys_have_nulls> hasher{d_keys};
-  row_equality_comparator<keys_have_nulls> rows_equal{d_keys, d_keys, null_keys_are_equal};
+  row_hasher<default_hash, nullate::DYNAMIC> hasher{nullate::DYNAMIC{keys_have_nulls}, d_keys};
+  row_equality_comparator rows_equal{
+    nullate::DYNAMIC{keys_have_nulls}, d_keys, d_keys, null_keys_are_equal};
 
   return map_type::create(compute_hash_table_size(d_keys.num_rows()),
                           stream,
@@ -481,14 +483,13 @@ auto create_sparse_results_table(table_view const& flattened_values,
 /**
  * @brief Computes all aggregations from `requests` that require a single pass
  * over the data and stores the results in `sparse_results`
- *
- * @see groupby_null_templated()
  */
-template <bool keys_have_nulls, typename Map>
+template <typename Map>
 void compute_single_pass_aggs(table_view const& keys,
                               host_span<aggregation_request const> requests,
                               cudf::detail::result_cache* sparse_results,
                               Map& map,
+                              bool keys_have_nulls,
                               null_policy include_null_keys,
                               rmm::cuda_stream_view stream)
 {
@@ -578,24 +579,24 @@ rmm::device_uvector<size_type> extract_populated_keys(Map map,
  * results using the aforementioned index vector. Dense results are stored into
  * the in/out parameter `cache`.
  */
-template <bool keys_have_nulls>
-std::unique_ptr<table> groupby_null_templated(table_view const& keys,
-                                              host_span<aggregation_request const> requests,
-                                              cudf::detail::result_cache* cache,
-                                              null_policy include_null_keys,
-                                              rmm::cuda_stream_view stream,
-                                              rmm::mr::device_memory_resource* mr)
+std::unique_ptr<table> groupby(table_view const& keys,
+                               host_span<aggregation_request const> requests,
+                               cudf::detail::result_cache* cache,
+                               bool keys_have_nulls,
+                               null_policy include_null_keys,
+                               rmm::cuda_stream_view stream,
+                               rmm::mr::device_memory_resource* mr)
 {
   auto d_keys_ptr = table_device_view::create(keys, stream);
-  auto map        = create_hash_map<keys_have_nulls>(*d_keys_ptr, include_null_keys, stream);
+  auto map        = create_hash_map(*d_keys_ptr, keys_have_nulls, include_null_keys, stream);
 
   // Cache of sparse results where the location of aggregate value in each
   // column is indexed by the hash map
   cudf::detail::result_cache sparse_results(requests.size());
 
   // Compute all single pass aggs first
-  compute_single_pass_aggs<keys_have_nulls>(
-    keys, requests, &sparse_results, *map, include_null_keys, stream);
+  compute_single_pass_aggs(
+    keys, requests, &sparse_results, *map, keys_have_nulls, include_null_keys, stream);
 
   // Extract the populated indices from the hash map and create a gather map.
   // Gathering using this map from sparse results will give dense results.
@@ -664,14 +665,8 @@ std::pair<std::unique_ptr<table>, std::vector<aggregation_result>> groupby(
 {
   cudf::detail::result_cache cache(requests.size());
 
-  std::unique_ptr<table> unique_keys;
-  if (has_nulls(keys)) {
-    unique_keys =
-      groupby_null_templated<true>(keys, requests, &cache, include_null_keys, stream, mr);
-  } else {
-    unique_keys =
-      groupby_null_templated<false>(keys, requests, &cache, include_null_keys, stream, mr);
-  }
+  std::unique_ptr<table> unique_keys =
+    groupby(keys, requests, &cache, has_nulls(keys), include_null_keys, stream, mr);
 
   return std::make_pair(std::move(unique_keys), extract_results(requests, cache));
 }
diff --git a/cpp/src/groupby/sort/group_nunique.cu b/cpp/src/groupby/sort/group_nunique.cu
index b8c4571961d..5154c867095 100644
--- a/cpp/src/groupby/sort/group_nunique.cu
+++ b/cpp/src/groupby/sort/group_nunique.cu
@@ -49,7 +49,7 @@ struct nunique_functor {
 
     auto values_view = column_device_view::create(values, stream);
     if (values.has_nulls()) {
-      auto equal              = element_equality_comparator<true>{*values_view, *values_view};
+      auto equal = element_equality_comparator{nullate::YES{}, *values_view, *values_view};
       auto is_unique_iterator = thrust::make_transform_iterator(
         thrust::make_counting_iterator<size_type>(0),
         [v = *values_view,
@@ -72,7 +72,7 @@ struct nunique_functor {
                             thrust::make_discard_iterator(),
                             result->mutable_view().begin<size_type>());
     } else {
-      auto equal              = element_equality_comparator<false>{*values_view, *values_view};
+      auto equal = element_equality_comparator{nullate::NO{}, *values_view, *values_view};
       auto is_unique_iterator = thrust::make_transform_iterator(
         thrust::make_counting_iterator<size_type>(0),
         [v = *values_view,
diff --git a/cpp/src/groupby/sort/group_rank_scan.cu b/cpp/src/groupby/sort/group_rank_scan.cu
index f36bdc0a660..62aa3df8e5c 100644
--- a/cpp/src/groupby/sort/group_rank_scan.cu
+++ b/cpp/src/groupby/sort/group_rank_scan.cu
@@ -31,7 +31,6 @@ namespace {
 /**
  * @brief generate grouped row ranks or dense ranks using a row comparison then scan the results
  *
- * @tparam has_nulls if the order_by column has nulls
  * @tparam value_resolver flag value resolver function with boolean first and row number arguments
  * @tparam scan_operator scan function ran on the flag values
  * @param order_by input column to generate ranks for
@@ -39,23 +38,26 @@ namespace {
  * @param group_offsets group index offsets with group ID indices
  * @param resolver flag value resolver
  * @param scan_op scan operation ran on the flag results
+ * @param has_nulls true if nulls are included in the `order_by` column
  * @param stream CUDA stream used for device memory operations and kernel launches
  * @param mr Device memory resource used to allocate the returned column's device memory
  * @return std::unique_ptr<column> rank values
  */
-template <bool has_nulls, typename value_resolver, typename scan_operator>
+template <typename value_resolver, typename scan_operator>
 std::unique_ptr<column> rank_generator(column_view const& order_by,
                                        device_span<size_type const> group_labels,
                                        device_span<size_type const> group_offsets,
                                        value_resolver resolver,
                                        scan_operator scan_op,
+                                       bool has_nulls,
                                        rmm::cuda_stream_view stream,
                                        rmm::mr::device_memory_resource* mr)
 {
   auto const flattened = cudf::structs::detail::flatten_nested_columns(
     table_view{{order_by}}, {}, {}, structs::detail::column_nullability::MATCH_INCOMING);
   auto const d_flat_order = table_device_view::create(flattened, stream);
-  row_equality_comparator<has_nulls> comparator(*d_flat_order, *d_flat_order, true);
+  row_equality_comparator comparator(
+    nullate::DYNAMIC{has_nulls}, *d_flat_order, *d_flat_order, null_equality::EQUAL);
   auto ranks         = make_fixed_width_column(data_type{type_to_id<size_type>()},
                                        flattened.flattened_columns().num_rows(),
                                        mask_state::UNALLOCATED,
@@ -92,22 +94,13 @@ std::unique_ptr<column> rank_scan(column_view const& order_by,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr)
 {
-  if (has_nested_nulls(table_view{{order_by}})) {
-    return rank_generator<true>(
-      order_by,
-      group_labels,
-      group_offsets,
-      [] __device__(bool equality, auto row_index) { return equality ? row_index + 1 : 0; },
-      DeviceMax{},
-      stream,
-      mr);
-  }
-  return rank_generator<false>(
+  return rank_generator(
     order_by,
     group_labels,
     group_offsets,
     [] __device__(bool equality, auto row_index) { return equality ? row_index + 1 : 0; },
     DeviceMax{},
+    has_nested_nulls(table_view{{order_by}}),
     stream,
     mr);
 }
@@ -118,22 +111,13 @@ std::unique_ptr<column> dense_rank_scan(column_view const& order_by,
                                         rmm::cuda_stream_view stream,
                                         rmm::mr::device_memory_resource* mr)
 {
-  if (has_nested_nulls(table_view{{order_by}})) {
-    return rank_generator<true>(
-      order_by,
-      group_labels,
-      group_offsets,
-      [] __device__(bool equality, auto row_index) { return equality; },
-      DeviceSum{},
-      stream,
-      mr);
-  }
-  return rank_generator<false>(
+  return rank_generator(
     order_by,
     group_labels,
     group_offsets,
     [] __device__(bool equality, auto row_index) { return equality; },
     DeviceSum{},
+    has_nested_nulls(table_view{{order_by}}),
     stream,
     mr);
 }
diff --git a/cpp/src/groupby/sort/group_scan_util.cuh b/cpp/src/groupby/sort/group_scan_util.cuh
index e25fdd6fc27..08f65536466 100644
--- a/cpp/src/groupby/sort/group_scan_util.cuh
+++ b/cpp/src/groupby/sort/group_scan_util.cuh
@@ -220,21 +220,12 @@ struct group_scan_functor<K,
 
     // Find the indices of the prefix min/max elements within each group.
     auto const count_iter = thrust::make_counting_iterator<size_type>(0);
-    if (values.has_nulls()) {
-      auto const binop =
-        cudf::reduction::detail::row_arg_minmax_fn<true>(values.size(),
-                                                         *d_flattened_values_ptr,
-                                                         flattened_null_precedences.data(),
-                                                         K == aggregation::MIN);
-      do_scan(count_iter, map_begin, binop);
-    } else {
-      auto const binop =
-        cudf::reduction::detail::row_arg_minmax_fn<false>(values.size(),
-                                                          *d_flattened_values_ptr,
-                                                          flattened_null_precedences.data(),
-                                                          K == aggregation::MIN);
-      do_scan(count_iter, map_begin, binop);
-    }
+    auto const binop      = cudf::reduction::detail::row_arg_minmax_fn(values.size(),
+                                                                  *d_flattened_values_ptr,
+                                                                  values.has_nulls(),
+                                                                  flattened_null_precedences.data(),
+                                                                  K == aggregation::MIN);
+    do_scan(count_iter, map_begin, binop);
 
     auto gather_map_view =
       column_view(data_type{type_to_id<offset_type>()}, gather_map.size(), gather_map.data());
diff --git a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
index 95a36f40e57..4fde825c0e0 100644
--- a/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
+++ b/cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
@@ -270,14 +270,14 @@ struct group_reduction_functor<
 
     auto const count_iter   = thrust::make_counting_iterator<ResultType>(0);
     auto const result_begin = result->mutable_view().template begin<ResultType>();
-    if (values.has_nulls()) {
-      auto const binop =
-        cudf::reduction::detail::row_arg_minmax_fn<true>(values.size(),
-                                                         *d_flattened_values_ptr,
-                                                         flattened_null_precedences.data(),
-                                                         K == aggregation::ARGMIN);
-      do_reduction(count_iter, result_begin, binop);
+    auto const binop        = cudf::reduction::detail::row_arg_minmax_fn(values.size(),
+                                                                  *d_flattened_values_ptr,
+                                                                  values.has_nulls(),
+                                                                  flattened_null_precedences.data(),
+                                                                  K == aggregation::ARGMIN);
+    do_reduction(count_iter, result_begin, binop);
 
+    if (values.has_nulls()) {
       // Generate bitmask for the output by segmented reduction of the input bitmask.
       auto const d_values_ptr = column_device_view::create(values, stream);
       auto validity           = rmm::device_uvector<bool>(num_groups, stream);
@@ -288,13 +288,6 @@ struct group_reduction_functor<
       auto [null_mask, null_count] =
         cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr);
       result->set_null_mask(std::move(null_mask), null_count);
-    } else {
-      auto const binop =
-        cudf::reduction::detail::row_arg_minmax_fn<false>(values.size(),
-                                                          *d_flattened_values_ptr,
-                                                          flattened_null_precedences.data(),
-                                                          K == aggregation::ARGMIN);
-      do_reduction(count_iter, result_begin, binop);
     }
 
     return result;
diff --git a/cpp/src/groupby/sort/sort_helper.cu b/cpp/src/groupby/sort/sort_helper.cu
index 7adb4ccec76..8d09728b771 100644
--- a/cpp/src/groupby/sort/sort_helper.cu
+++ b/cpp/src/groupby/sort/sort_helper.cu
@@ -50,9 +50,8 @@ namespace {
  * @brief Compares two `table` rows for equality as if the table were
  * ordered according to a specified permutation map.
  */
-template <bool nullable = true>
 struct permuted_row_equality_comparator {
-  cudf::row_equality_comparator<nullable> _comparator;
+  cudf::row_equality_comparator<cudf::nullate::DYNAMIC> _comparator;
   cudf::size_type const* _map;
 
   /**
@@ -60,10 +59,12 @@ struct permuted_row_equality_comparator {
    *
    * @param t The `table` whose rows will be compared
    * @param map The permutation map that specifies the effective ordering of
-   *`t`. Must be the same size as `t.num_rows()`
+   * `t`. Must be the same size as `t.num_rows()`
    */
-  permuted_row_equality_comparator(cudf::table_device_view const& t, cudf::size_type const* map)
-    : _comparator(t, t, true), _map{map}
+  permuted_row_equality_comparator(cudf::table_device_view const& t,
+                                   cudf::size_type const* map,
+                                   bool nullable = true)
+    : _comparator(cudf::nullate::DYNAMIC{nullable}, t, t, cudf::null_equality::EQUAL), _map{map}
   {
   }
 
@@ -76,7 +77,7 @@ struct permuted_row_equality_comparator {
    *
    * @param lhs The index of the first row
    * @param rhs The index of the second row
-   * @returns if the two specified rows in the permuted order are equivalent
+   * @returns true if the two specified rows in the permuted order are equivalent
    */
   CUDA_DEVICE_CALLABLE
   bool operator()(cudf::size_type lhs, cudf::size_type rhs)
@@ -196,21 +197,12 @@ sort_groupby_helper::index_vector const& sort_groupby_helper::group_offsets(
   auto sorted_order       = key_sort_order(stream).data<size_type>();
   decltype(_group_offsets->begin()) result_end;
 
-  if (has_nulls(_keys)) {
-    result_end = thrust::unique_copy(
-      rmm::exec_policy(stream),
-      thrust::make_counting_iterator<size_type>(0),
-      thrust::make_counting_iterator<size_type>(num_keys(stream)),
-      _group_offsets->begin(),
-      permuted_row_equality_comparator<true>(*device_input_table, sorted_order));
-  } else {
-    result_end = thrust::unique_copy(
-      rmm::exec_policy(stream),
-      thrust::make_counting_iterator<size_type>(0),
-      thrust::make_counting_iterator<size_type>(num_keys(stream)),
-      _group_offsets->begin(),
-      permuted_row_equality_comparator<false>(*device_input_table, sorted_order));
-  }
+  result_end = thrust::unique_copy(
+    rmm::exec_policy(stream),
+    thrust::make_counting_iterator<size_type>(0),
+    thrust::make_counting_iterator<size_type>(num_keys(stream)),
+    _group_offsets->begin(),
+    permuted_row_equality_comparator(*device_input_table, sorted_order, has_nulls(_keys)));
 
   size_type num_groups = thrust::distance(_group_offsets->begin(), result_end);
   _group_offsets->set_element(num_groups, num_keys(stream), stream);
diff --git a/cpp/src/hash/hashing.cu b/cpp/src/hash/hashing.cu
index a882b33bcdf..039652e0012 100644
--- a/cpp/src/hash/hashing.cu
+++ b/cpp/src/hash/hashing.cu
@@ -64,43 +64,24 @@ std::unique_ptr<column> serial_murmur_hash3_32(table_view const& input,
   auto const device_input = table_device_view::create(leaf_table, stream);
   auto output_view        = output->mutable_view();
 
-  if (has_nulls(leaf_table)) {
-    thrust::tabulate(rmm::exec_policy(stream),
-                     output_view.begin<int32_t>(),
-                     output_view.end<int32_t>(),
-                     [device_input = *device_input, seed] __device__(auto row_index) {
-                       return thrust::reduce(
-                         thrust::seq,
-                         device_input.begin(),
-                         device_input.end(),
-                         seed,
-                         [rindex = row_index] __device__(auto hash, auto column) {
-                           return cudf::type_dispatcher(
-                             column.type(),
-                             element_hasher_with_seed<hash_function, true>{hash, hash},
-                             column,
-                             rindex);
-                         });
-                     });
-  } else {
-    thrust::tabulate(rmm::exec_policy(stream),
-                     output_view.begin<int32_t>(),
-                     output_view.end<int32_t>(),
-                     [device_input = *device_input, seed] __device__(auto row_index) {
-                       return thrust::reduce(
-                         thrust::seq,
-                         device_input.begin(),
-                         device_input.end(),
-                         seed,
-                         [rindex = row_index] __device__(auto hash, auto column) {
-                           return cudf::type_dispatcher(
-                             column.type(),
-                             element_hasher_with_seed<hash_function, false>{hash, hash},
-                             column,
-                             rindex);
-                         });
-                     });
-  }
+  thrust::tabulate(
+    rmm::exec_policy(stream),
+    output_view.begin<int32_t>(),
+    output_view.end<int32_t>(),
+    [device_input = *device_input, nulls = has_nulls(leaf_table), seed] __device__(auto row_index) {
+      return thrust::reduce(thrust::seq,
+                            device_input.begin(),
+                            device_input.end(),
+                            seed,
+                            [rindex = row_index, nulls] __device__(auto hash, auto column) {
+                              return cudf::type_dispatcher(
+                                column.type(),
+                                element_hasher_with_seed<hash_function, nullate::DYNAMIC>{
+                                  nullate::DYNAMIC{nulls}, hash, hash},
+                                column,
+                                rindex);
+                            });
+    });
 
   return output;
 }
diff --git a/cpp/src/hash/murmur_hash.cu b/cpp/src/hash/murmur_hash.cu
index 81be4d0eabe..a761d058180 100644
--- a/cpp/src/hash/murmur_hash.cu
+++ b/cpp/src/hash/murmur_hash.cu
@@ -50,31 +50,17 @@ std::unique_ptr<column> murmur_hash3_32(table_view const& input,
                  "Expected same size of initial hash values as number of columns");
     auto device_initial_hash = make_device_uvector_async(initial_hash, stream);
 
-    if (nullable) {
-      thrust::tabulate(
-        rmm::exec_policy(stream),
-        output_view.begin<int32_t>(),
-        output_view.end<int32_t>(),
-        row_hasher_initial_values<MurmurHash3_32, true>(*device_input, device_initial_hash.data()));
-    } else {
-      thrust::tabulate(rmm::exec_policy(stream),
-                       output_view.begin<int32_t>(),
-                       output_view.end<int32_t>(),
-                       row_hasher_initial_values<MurmurHash3_32, false>(
-                         *device_input, device_initial_hash.data()));
-    }
+    thrust::tabulate(rmm::exec_policy(stream),
+                     output_view.begin<int32_t>(),
+                     output_view.end<int32_t>(),
+                     row_hasher_initial_values<MurmurHash3_32, nullate::DYNAMIC>(
+                       nullate::DYNAMIC{nullable}, *device_input, device_initial_hash.data()));
   } else {
-    if (nullable) {
-      thrust::tabulate(rmm::exec_policy(stream),
-                       output_view.begin<int32_t>(),
-                       output_view.end<int32_t>(),
-                       row_hasher<MurmurHash3_32, true>(*device_input));
-    } else {
-      thrust::tabulate(rmm::exec_policy(stream),
-                       output_view.begin<int32_t>(),
-                       output_view.end<int32_t>(),
-                       row_hasher<MurmurHash3_32, false>(*device_input));
-    }
+    thrust::tabulate(
+      rmm::exec_policy(stream),
+      output_view.begin<int32_t>(),
+      output_view.end<int32_t>(),
+      row_hasher<MurmurHash3_32, nullate::DYNAMIC>(nullate::DYNAMIC{nullable}, *device_input));
   }
 
   return output;
diff --git a/cpp/src/join/hash_join.cu b/cpp/src/join/hash_join.cu
index c5b680f129e..eee0a8cc6f0 100644
--- a/cpp/src/join/hash_join.cu
+++ b/cpp/src/join/hash_join.cu
@@ -81,7 +81,7 @@ void build_join_hash_table(cudf::table_view const& build,
   CUDF_EXPECTS(0 != build_table_ptr->num_columns(), "Selected build dataset is empty");
   CUDF_EXPECTS(0 != build_table_ptr->num_rows(), "Build side table has no rows");
 
-  row_hash hash_build{*build_table_ptr};
+  row_hash hash_build{nullate::YES{}, *build_table_ptr};
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
   make_pair_function pair_func{hash_build, empty_key_sentinel};
 
@@ -147,9 +147,9 @@ probe_join_hash_table(cudf::table_device_view build_table,
   auto left_indices  = std::make_unique<rmm::device_uvector<size_type>>(join_size, stream, mr);
   auto right_indices = std::make_unique<rmm::device_uvector<size_type>>(join_size, stream, mr);
 
-  pair_equality equality{probe_table, build_table, compare_nulls == null_equality::EQUAL};
+  pair_equality equality{probe_table, build_table, compare_nulls};
 
-  row_hash hash_probe{probe_table};
+  row_hash hash_probe{nullate::YES{}, probe_table};
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
   make_pair_function pair_func{hash_probe, empty_key_sentinel};
 
@@ -212,9 +212,9 @@ std::size_t get_full_join_size(cudf::table_device_view build_table,
   auto left_indices  = std::make_unique<rmm::device_uvector<size_type>>(join_size, stream, mr);
   auto right_indices = std::make_unique<rmm::device_uvector<size_type>>(join_size, stream, mr);
 
-  pair_equality equality{probe_table, build_table, compare_nulls == null_equality::EQUAL};
+  pair_equality equality{probe_table, build_table, compare_nulls};
 
-  row_hash hash_probe{probe_table};
+  row_hash hash_probe{nullate::YES{}, probe_table};
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
   make_pair_function pair_func{hash_probe, empty_key_sentinel};
 
diff --git a/cpp/src/join/hash_join.cuh b/cpp/src/join/hash_join.cuh
index 10b0e420ef6..976b0c81ead 100644
--- a/cpp/src/join/hash_join.cuh
+++ b/cpp/src/join/hash_join.cuh
@@ -117,9 +117,9 @@ std::size_t compute_join_output_size(table_device_view build_table,
     }
   }
 
-  pair_equality equality{probe_table, build_table, compare_nulls == null_equality::EQUAL};
+  pair_equality equality{probe_table, build_table, compare_nulls};
 
-  row_hash hash_probe{probe_table};
+  row_hash hash_probe{nullate::YES{}, probe_table};
   auto const empty_key_sentinel = hash_table.get_empty_key_sentinel();
   make_pair_function pair_func{hash_probe, empty_key_sentinel};
 
diff --git a/cpp/src/join/join_common_utils.cuh b/cpp/src/join/join_common_utils.cuh
index cec633765c7..4b33772dd69 100644
--- a/cpp/src/join/join_common_utils.cuh
+++ b/cpp/src/join/join_common_utils.cuh
@@ -32,8 +32,10 @@ namespace detail {
  */
 class pair_equality {
  public:
-  pair_equality(table_device_view lhs, table_device_view rhs, bool nulls_are_equal = true)
-    : _check_row_equality{lhs, rhs, nulls_are_equal}
+  pair_equality(table_device_view lhs,
+                table_device_view rhs,
+                null_equality nulls_are_equal = null_equality::EQUAL)
+    : _check_row_equality{cudf::nullate::YES{}, lhs, rhs, nulls_are_equal}
   {
   }
 
diff --git a/cpp/src/join/join_common_utils.hpp b/cpp/src/join/join_common_utils.hpp
index d6eb5e93a98..84506daf2f1 100644
--- a/cpp/src/join/join_common_utils.hpp
+++ b/cpp/src/join/join_common_utils.hpp
@@ -45,9 +45,9 @@ using multimap_type =
                         default_allocator<char>,
                         cuco::double_hashing<DEFAULT_JOIN_CG_SIZE, hash_type, hash_type>>;
 
-using row_hash = cudf::row_hasher<default_hash>;
+using row_hash = cudf::row_hasher<default_hash, cudf::nullate::YES>;
 
-using row_equality = cudf::row_equality_comparator<true>;
+using row_equality = cudf::row_equality_comparator<cudf::nullate::YES>;
 
 enum class join_kind { INNER_JOIN, LEFT_JOIN, FULL_JOIN, LEFT_SEMI_JOIN, LEFT_ANTI_JOIN };
 
diff --git a/cpp/src/join/semi_join.cu b/cpp/src/join/semi_join.cu
index 5b5dd418a97..3d27c5740f4 100644
--- a/cpp/src/join/semi_join.cu
+++ b/cpp/src/join/semi_join.cu
@@ -77,13 +77,13 @@ std::unique_ptr<rmm::device_uvector<cudf::size_type>> left_semi_anti_join(
   // Create hash table containing all keys found in right table
   auto right_rows_d            = table_device_view::create(right_flattened_keys, stream);
   size_t const hash_table_size = compute_hash_table_size(right_num_rows);
-  row_hash hash_build{*right_rows_d};
-  row_equality equality_build{*right_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL};
+  row_hash hash_build{cudf::nullate::YES{}, *right_rows_d};
+  row_equality equality_build{cudf::nullate::YES{}, *right_rows_d, *right_rows_d, compare_nulls};
 
   // Going to join it with left table
   auto left_rows_d = table_device_view::create(left_flattened_keys, stream);
-  row_hash hash_probe{*left_rows_d};
-  row_equality equality_probe{*left_rows_d, *right_rows_d, compare_nulls == null_equality::EQUAL};
+  row_hash hash_probe{cudf::nullate::YES{}, *left_rows_d};
+  row_equality equality_probe{cudf::nullate::YES{}, *left_rows_d, *right_rows_d, compare_nulls};
 
   auto hash_table_ptr = hash_table_type::create(hash_table_size,
                                                 stream,
diff --git a/cpp/src/partitioning/partitioning.cu b/cpp/src/partitioning/partitioning.cu
index e8c56cdafd8..7b3b7b0f3fd 100644
--- a/cpp/src/partitioning/partitioning.cu
+++ b/cpp/src/partitioning/partitioning.cu
@@ -486,7 +486,8 @@ std::pair<std::unique_ptr<table>, std::vector<size_type>> hash_partition_table(
     cudf::detail::make_zeroed_device_uvector_async<size_type>(num_rows, stream);
 
   auto const device_input = table_device_view::create(table_to_hash, stream);
-  auto const hasher       = row_hasher<hash_function, hash_has_nulls>(*device_input, seed);
+  auto const hasher       = row_hasher<hash_function, nullate::DYNAMIC>(
+    nullate::DYNAMIC{hash_has_nulls}, *device_input, seed);
 
   // If the number of partitions is a power of two, we can compute the partition
   // number of each row more efficiently with bitwise operations
diff --git a/cpp/src/reductions/arg_minmax_util.cuh b/cpp/src/reductions/arg_minmax_util.cuh
index 40df23bcd8e..5694d0ed0fa 100644
--- a/cpp/src/reductions/arg_minmax_util.cuh
+++ b/cpp/src/reductions/arg_minmax_util.cuh
@@ -24,21 +24,20 @@ namespace detail {
 
 /**
  * @brief Binary operator ArgMin/ArgMax with index values into the input table.
- *
- * @tparam T Type of the underlying data. This is the fallback for the cases when T does not support
- * '<' operator.
  */
-template <bool has_nulls>
 struct row_arg_minmax_fn {
   size_type const num_rows;
-  row_lexicographic_comparator<has_nulls> const comp;
+  row_lexicographic_comparator<nullate::DYNAMIC> const comp;
   bool const arg_min;
 
-  row_arg_minmax_fn(size_type const num_rows_,
-                    table_device_view const& table_,
-                    null_order const* null_precedence_,
-                    bool const arg_min_)
-    : num_rows(num_rows_), comp(table_, table_, nullptr, null_precedence_), arg_min(arg_min_)
+  row_arg_minmax_fn(size_type const num_rows,
+                    table_device_view const& table,
+                    bool has_nulls,
+                    null_order const* null_precedence,
+                    bool const arg_min)
+    : num_rows(num_rows),
+      comp(nullate::DYNAMIC{has_nulls}, table, table, nullptr, null_precedence),
+      arg_min(arg_min)
   {
   }
 
diff --git a/cpp/src/reductions/scan/rank_scan.cu b/cpp/src/reductions/scan/rank_scan.cu
index e7f1e867a41..9ac4db3a34b 100644
--- a/cpp/src/reductions/scan/rank_scan.cu
+++ b/cpp/src/reductions/scan/rank_scan.cu
@@ -33,18 +33,19 @@ namespace {
 /**
  * @brief generate row ranks or dense ranks using a row comparison then scan the results
  *
- * @tparam has_nulls if the order_by column has nulls
  * @tparam value_resolver flag value resolver with boolean first and row number arguments
  * @tparam scan_operator scan function ran on the flag values
  * @param order_by input column to generate ranks for
+ * @param has_nulls if the order_by column has nested nulls
  * @param resolver flag value resolver
  * @param scan_op scan operation ran on the flag results
  * @param stream CUDA stream used for device memory operations and kernel launches
  * @param mr Device memory resource used to allocate the returned column's device memory
  * @return std::unique_ptr<column> rank values
  */
-template <bool has_nulls, typename value_resolver, typename scan_operator>
+template <typename value_resolver, typename scan_operator>
 std::unique_ptr<column> rank_generator(column_view const& order_by,
+                                       bool has_nulls,
                                        value_resolver resolver,
                                        scan_operator scan_op,
                                        rmm::cuda_stream_view stream,
@@ -53,7 +54,8 @@ std::unique_ptr<column> rank_generator(column_view const& order_by,
   auto const flattened = cudf::structs::detail::flatten_nested_columns(
     table_view{{order_by}}, {}, {}, structs::detail::column_nullability::MATCH_INCOMING);
   auto const d_flat_order = table_device_view::create(flattened, stream);
-  row_equality_comparator<has_nulls> comparator(*d_flat_order, *d_flat_order, true);
+  row_equality_comparator comparator(
+    nullate::DYNAMIC{has_nulls}, *d_flat_order, *d_flat_order, null_equality::EQUAL);
   auto ranks         = make_fixed_width_column(data_type{type_to_id<size_type>()},
                                        flattened.flattened_columns().num_rows(),
                                        mask_state::UNALLOCATED,
@@ -85,16 +87,9 @@ std::unique_ptr<column> inclusive_dense_rank_scan(column_view const& order_by,
 {
   CUDF_EXPECTS(!cudf::structs::detail::is_or_has_nested_lists(order_by),
                "Unsupported list type in dense_rank scan.");
-  if (has_nested_nulls(table_view{{order_by}})) {
-    return rank_generator<true>(
-      order_by,
-      [] __device__(bool equality, auto row_index) { return equality; },
-      DeviceSum{},
-      stream,
-      mr);
-  }
-  return rank_generator<false>(
+  return rank_generator(
     order_by,
+    has_nested_nulls(table_view{{order_by}}),
     [] __device__(bool equality, auto row_index) { return equality; },
     DeviceSum{},
     stream,
@@ -107,16 +102,9 @@ std::unique_ptr<column> inclusive_rank_scan(column_view const& order_by,
 {
   CUDF_EXPECTS(!cudf::structs::detail::is_or_has_nested_lists(order_by),
                "Unsupported list type in rank scan.");
-  if (has_nested_nulls(table_view{{order_by}})) {
-    return rank_generator<true>(
-      order_by,
-      [] __device__(bool equality, auto row_index) { return equality ? row_index + 1 : 0; },
-      DeviceMax{},
-      stream,
-      mr);
-  }
-  return rank_generator<false>(
+  return rank_generator(
     order_by,
+    has_nested_nulls(table_view{{order_by}}),
     [] __device__(bool equality, auto row_index) { return equality ? row_index + 1 : 0; },
     DeviceMax{},
     stream,
diff --git a/cpp/src/reductions/scan/scan_inclusive.cu b/cpp/src/reductions/scan/scan_inclusive.cu
index b0e761c4c3b..2b1ac8aa704 100644
--- a/cpp/src/reductions/scan/scan_inclusive.cu
+++ b/cpp/src/reductions/scan/scan_inclusive.cu
@@ -182,15 +182,12 @@ struct scan_functor<Op, cudf::struct_view> {
       is_min_op ? cudf::detail::make_device_uvector_async(flattened_input.null_orders(), stream)
                 : rmm::device_uvector<cudf::null_order>(0, stream);
 
-    if (input.has_nulls()) {
-      auto const binop = cudf::reduction::detail::row_arg_minmax_fn<true>(
-        input.size(), *d_flattened_input_ptr, flattened_null_precedences.data(), is_min_op);
-      do_scan(binop);
-    } else {
-      auto const binop = cudf::reduction::detail::row_arg_minmax_fn<false>(
-        input.size(), *d_flattened_input_ptr, flattened_null_precedences.data(), is_min_op);
-      do_scan(binop);
-    }
+    auto const binop = cudf::reduction::detail::row_arg_minmax_fn(input.size(),
+                                                                  *d_flattened_input_ptr,
+                                                                  input.has_nulls(),
+                                                                  flattened_null_precedences.data(),
+                                                                  is_min_op);
+    do_scan(binop);
 
     // Gather the children columns of the input column. Must use `get_sliced_child` to properly
     // handle input in case it is a sliced view.
diff --git a/cpp/src/reductions/simple.cuh b/cpp/src/reductions/simple.cuh
index 7dd54e9250a..e5633341ffa 100644
--- a/cpp/src/reductions/simple.cuh
+++ b/cpp/src/reductions/simple.cuh
@@ -317,15 +317,13 @@ struct same_element_type_dispatcher {
     };
 
     auto const minmax_idx = [&] {
-      if (input.has_nulls()) {
-        auto const binop = cudf::reduction::detail::row_arg_minmax_fn<true>(
-          input.size(), *d_flattened_input_ptr, flattened_null_precedences.data(), is_min_op);
-        return do_reduction(binop);
-      } else {
-        auto const binop = cudf::reduction::detail::row_arg_minmax_fn<false>(
-          input.size(), *d_flattened_input_ptr, flattened_null_precedences.data(), is_min_op);
-        return do_reduction(binop);
-      }
+      auto const binop =
+        cudf::reduction::detail::row_arg_minmax_fn(input.size(),
+                                                   *d_flattened_input_ptr,
+                                                   input.has_nulls(),
+                                                   flattened_null_precedences.data(),
+                                                   is_min_op);
+      return do_reduction(binop);
     }();
 
     return cudf::detail::get_element(input, minmax_idx, stream, mr);
diff --git a/cpp/src/replace/clamp.cu b/cpp/src/replace/clamp.cu
index 713b3b27a2b..fae02805620 100644
--- a/cpp/src/replace/clamp.cu
+++ b/cpp/src/replace/clamp.cu
@@ -180,7 +180,7 @@ std::enable_if_t<cudf::is_fixed_width<T>(), std::unique_ptr<cudf::column>> clamp
   };
 
   auto input_pair_iterator =
-    make_optional_iterator<T>(*input_device_view, contains_nulls::DYNAMIC{}, input.has_nulls());
+    make_optional_iterator<T>(*input_device_view, nullate::DYNAMIC{input.has_nulls()});
   thrust::transform(rmm::exec_policy(stream),
                     input_pair_iterator,
                     input_pair_iterator + input.size(),
@@ -232,10 +232,10 @@ struct dispatch_clamp {
   {
     CUDF_EXPECTS(lo.type() == input.type(), "mismatching types of scalar and input");
 
-    auto lo_itr         = make_optional_iterator<T>(lo, contains_nulls::YES{});
-    auto hi_itr         = make_optional_iterator<T>(hi, contains_nulls::YES{});
-    auto lo_replace_itr = make_optional_iterator<T>(lo_replace, contains_nulls::NO{});
-    auto hi_replace_itr = make_optional_iterator<T>(hi_replace, contains_nulls::NO{});
+    auto lo_itr         = make_optional_iterator<T>(lo, nullate::YES{});
+    auto hi_itr         = make_optional_iterator<T>(hi, nullate::YES{});
+    auto lo_replace_itr = make_optional_iterator<T>(lo_replace, nullate::NO{});
+    auto hi_replace_itr = make_optional_iterator<T>(hi_replace, nullate::NO{});
 
     return clamp<T>(input, lo_itr, lo_replace_itr, hi_itr, hi_replace_itr, stream, mr);
   }
diff --git a/cpp/src/replace/nans.cu b/cpp/src/replace/nans.cu
index c1c26573692..7b47f8df28d 100644
--- a/cpp/src/replace/nans.cu
+++ b/cpp/src/replace/nans.cu
@@ -56,9 +56,9 @@ struct replace_nans_functor {
     };
 
     auto input_iterator =
-      make_optional_iterator<T>(*input_device_view, contains_nulls::DYNAMIC{}, input.has_nulls());
+      make_optional_iterator<T>(*input_device_view, nullate::DYNAMIC{input.has_nulls()});
     auto replacement_iterator =
-      make_optional_iterator<T>(replacement, contains_nulls::DYNAMIC{}, replacement_nullable);
+      make_optional_iterator<T>(replacement, nullate::DYNAMIC{replacement_nullable});
     return copy_if_else(input.has_nulls() or replacement_nullable,
                         input_iterator,
                         input_iterator + size,
diff --git a/cpp/src/search/search.cu b/cpp/src/search/search.cu
index 462d0678eab..9a677d7907a 100644
--- a/cpp/src/search/search.cu
+++ b/cpp/src/search/search.cu
@@ -124,17 +124,13 @@ std::unique_ptr<column> search_ordered(table_view const& t,
     detail::make_device_uvector_async(null_precedence_flattened, stream);
 
   auto const count_it = thrust::make_counting_iterator<size_type>(0);
-  if (has_null_elements) {
-    auto const comp = row_lexicographic_comparator<true>(
-      lhs, rhs, column_order_dv.data(), null_precedence_dv.data());
-    launch_search(
-      count_it, count_it, t.num_rows(), values.num_rows(), result_out, comp, find_first, stream);
-  } else {
-    auto const comp = row_lexicographic_comparator<false>(
-      lhs, rhs, column_order_dv.data(), null_precedence_dv.data());
-    launch_search(
-      count_it, count_it, t.num_rows(), values.num_rows(), result_out, comp, find_first, stream);
-  }
+  auto const comp     = row_lexicographic_comparator(nullate::DYNAMIC{has_null_elements},
+                                                 lhs,
+                                                 rhs,
+                                                 column_order_dv.data(),
+                                                 null_precedence_dv.data());
+  launch_search(
+    count_it, count_it, t.num_rows(), values.num_rows(), result_out, comp, find_first, stream);
 
   return result;
 }
diff --git a/cpp/src/sort/is_sorted.cu b/cpp/src/sort/is_sorted.cu
index b08baaa0261..a8820204c22 100644
--- a/cpp/src/sort/is_sorted.cu
+++ b/cpp/src/sort/is_sorted.cu
@@ -27,12 +27,14 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <thrust/sort.h>
+
 namespace cudf {
 namespace detail {
 
-template <bool has_nulls>
 auto is_sorted(cudf::table_view const& in,
                std::vector<order> const& column_order,
+               bool has_nulls,
                std::vector<null_order> const& null_precedence,
                rmm::cuda_stream_view stream)
 {
@@ -45,8 +47,11 @@ auto is_sorted(cudf::table_view const& in,
                                    ? make_device_uvector_async(flattened.null_orders(), stream)
                                    : rmm::device_uvector<null_order>(0, stream);
 
-  auto comparator = row_lexicographic_comparator<has_nulls>(
-    *d_input, *d_input, d_column_order.data(), d_null_precedence.data());
+  auto comparator = row_lexicographic_comparator(nullate::DYNAMIC{has_nulls},
+                                                 *d_input,
+                                                 *d_input,
+                                                 d_column_order.data(),
+                                                 d_null_precedence.data());
 
   auto sorted = thrust::is_sorted(rmm::exec_policy(stream),
                                   thrust::make_counting_iterator(0),
@@ -76,11 +81,8 @@ bool is_sorted(cudf::table_view const& in,
       "Number of columns in the table doesn't match the vector null_precedence's size .\n");
   }
 
-  if (has_nulls(in)) {
-    return detail::is_sorted<true>(in, column_order, null_precedence, rmm::cuda_stream_default);
-  } else {
-    return detail::is_sorted<false>(in, column_order, null_precedence, rmm::cuda_stream_default);
-  }
+  return detail::is_sorted(
+    in, column_order, has_nulls(in), null_precedence, rmm::cuda_stream_default);
 }
 
 }  // namespace cudf
diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu
index de0a44e3234..e17a18997e8 100644
--- a/cpp/src/sort/rank.cu
+++ b/cpp/src/sort/rank.cu
@@ -37,10 +37,11 @@ namespace cudf {
 namespace detail {
 namespace {
 // Functor to identify unique elements in a sorted order table/column
-template <bool has_nulls, typename ReturnType, typename Iterator>
+template <typename ReturnType, typename Iterator>
 struct unique_comparator {
-  unique_comparator(table_device_view device_table, Iterator const sorted_order)
-    : comparator(device_table, device_table, true), permute(sorted_order)
+  unique_comparator(table_device_view device_table, Iterator const sorted_order, bool has_nulls)
+    : comparator(nullate::DYNAMIC{has_nulls}, device_table, device_table, null_equality::EQUAL),
+      permute(sorted_order)
   {
   }
   __device__ ReturnType operator()(size_type index) const noexcept
@@ -49,7 +50,7 @@ struct unique_comparator {
   };
 
  private:
-  row_equality_comparator<has_nulls> comparator;
+  row_equality_comparator<nullate::DYNAMIC> comparator;
   Iterator const permute;
 };
 
@@ -63,21 +64,12 @@ rmm::device_uvector<size_type> sorted_dense_rank(column_view input_col,
   rmm::device_uvector<size_type> dense_rank_sorted(input_size, stream);
   auto sorted_index_order = thrust::make_permutation_iterator(
     sorted_order_view.begin<size_type>(), thrust::make_counting_iterator<size_type>(0));
-  if (input_col.has_nulls()) {
-    auto conv = unique_comparator<true, size_type, decltype(sorted_index_order)>(
-      *device_table, sorted_index_order);
-    auto unique_it = cudf::detail::make_counting_transform_iterator(0, conv);
+  auto conv = unique_comparator<size_type, decltype(sorted_index_order)>(
+    *device_table, sorted_index_order, input_col.has_nulls());
+  auto unique_it = cudf::detail::make_counting_transform_iterator(0, conv);
 
-    thrust::inclusive_scan(
-      rmm::exec_policy(stream), unique_it, unique_it + input_size, dense_rank_sorted.data());
-  } else {
-    auto conv = unique_comparator<false, size_type, decltype(sorted_index_order)>(
-      *device_table, sorted_index_order);
-    auto unique_it = cudf::detail::make_counting_transform_iterator(0, conv);
-
-    thrust::inclusive_scan(
-      rmm::exec_policy(stream), unique_it, unique_it + input_size, dense_rank_sorted.data());
-  }
+  thrust::inclusive_scan(
+    rmm::exec_policy(stream), unique_it, unique_it + input_size, dense_rank_sorted.data());
   return dense_rank_sorted;
 }
 
diff --git a/cpp/src/sort/sort_impl.cuh b/cpp/src/sort/sort_impl.cuh
index 25f0815e645..881503a49e3 100644
--- a/cpp/src/sort/sort_impl.cuh
+++ b/cpp/src/sort/sort_impl.cuh
@@ -127,40 +127,25 @@ std::unique_ptr<column> sorted_order(table_view input,
   auto device_table = table_device_view::create(flattened, stream);
   auto const d_column_order = make_device_uvector_async(flattened.orders(), stream);
 
-  if (has_nulls(flattened)) {
-    auto const d_null_precedence = make_device_uvector_async(flattened.null_orders(), stream);
-    auto const comparator        = row_lexicographic_comparator<true>(
-      *device_table, *device_table, d_column_order.data(), d_null_precedence.data());
-    if (stable) {
-      thrust::stable_sort(rmm::exec_policy(stream),
-                          mutable_indices_view.begin<size_type>(),
-                          mutable_indices_view.end<size_type>(),
-                          comparator);
-    } else {
-      thrust::sort(rmm::exec_policy(stream),
-                   mutable_indices_view.begin<size_type>(),
-                   mutable_indices_view.end<size_type>(),
-                   comparator);
-    }
-    // protection for temporary d_column_order and d_null_precedence
-    stream.synchronize();
+  auto const d_null_precedence = make_device_uvector_async(flattened.null_orders(), stream);
+  auto const comparator = row_lexicographic_comparator(nullate::DYNAMIC{has_nulls(flattened)},
+                                                       *device_table,
+                                                       *device_table,
+                                                       d_column_order.data(),
+                                                       d_null_precedence.data());
+  if (stable) {
+    thrust::stable_sort(rmm::exec_policy(stream),
+                        mutable_indices_view.begin<size_type>(),
+                        mutable_indices_view.end<size_type>(),
+                        comparator);
   } else {
-    auto const comparator =
-      row_lexicographic_comparator<false>(*device_table, *device_table, d_column_order.data());
-    if (stable) {
-      thrust::stable_sort(rmm::exec_policy(stream),
-                          mutable_indices_view.begin<size_type>(),
-                          mutable_indices_view.end<size_type>(),
-                          comparator);
-    } else {
-      thrust::sort(rmm::exec_policy(stream),
-                   mutable_indices_view.begin<size_type>(),
-                   mutable_indices_view.end<size_type>(),
-                   comparator);
-    }
-    // protection for temporary d_column_order
-    stream.synchronize();
+    thrust::sort(rmm::exec_policy(stream),
+                 mutable_indices_view.begin<size_type>(),
+                 mutable_indices_view.end<size_type>(),
+                 comparator);
   }
+  // protection for temporary d_column_order and d_null_precedence
+  stream.synchronize();
 
   return sorted_indices;
 }
diff --git a/cpp/src/stream_compaction/distinct_count.cu b/cpp/src/stream_compaction/distinct_count.cu
index 4e142204a29..5c695f8a16f 100644
--- a/cpp/src/stream_compaction/distinct_count.cu
+++ b/cpp/src/stream_compaction/distinct_count.cu
@@ -50,27 +50,15 @@ cudf::size_type distinct_count(table_view const& keys,
   auto sorted_row_index   = sorted_indices->view().data<cudf::size_type>();
   auto device_input_table = cudf::table_device_view::create(keys, stream);
 
-  if (cudf::has_nulls(keys)) {
-    row_equality_comparator<true> comp(
-      *device_input_table, *device_input_table, nulls_equal == null_equality::EQUAL);
-    return thrust::count_if(
-      rmm::exec_policy(stream),
-      thrust::counting_iterator<cudf::size_type>(0),
-      thrust::counting_iterator<cudf::size_type>(keys.num_rows()),
-      [sorted_row_index, comp] __device__(cudf::size_type i) {
-        return (i == 0 || not comp(sorted_row_index[i], sorted_row_index[i - 1]));
-      });
-  } else {
-    row_equality_comparator<false> comp(
-      *device_input_table, *device_input_table, nulls_equal == null_equality::EQUAL);
-    return thrust::count_if(
-      rmm::exec_policy(stream),
-      thrust::counting_iterator<cudf::size_type>(0),
-      thrust::counting_iterator<cudf::size_type>(keys.num_rows()),
-      [sorted_row_index, comp] __device__(cudf::size_type i) {
-        return (i == 0 || not comp(sorted_row_index[i], sorted_row_index[i - 1]));
-      });
-  }
+  row_equality_comparator comp(
+    nullate::DYNAMIC{cudf::has_nulls(keys)}, *device_input_table, *device_input_table, nulls_equal);
+  return thrust::count_if(
+    rmm::exec_policy(stream),
+    thrust::counting_iterator<cudf::size_type>(0),
+    thrust::counting_iterator<cudf::size_type>(keys.num_rows()),
+    [sorted_row_index, comp] __device__(cudf::size_type i) {
+      return (i == 0 || not comp(sorted_row_index[i], sorted_row_index[i - 1]));
+    });
 }
 
 /**
diff --git a/cpp/src/stream_compaction/drop_duplicates.cu b/cpp/src/stream_compaction/drop_duplicates.cu
index f236e6a5f53..abc34663aee 100644
--- a/cpp/src/stream_compaction/drop_duplicates.cu
+++ b/cpp/src/stream_compaction/drop_duplicates.cu
@@ -83,35 +83,18 @@ column_view get_unique_ordered_indices(cudf::table_view const& keys,
   // extract unique indices
   auto device_input_table = cudf::table_device_view::create(keys, stream);
 
-  if (cudf::has_nulls(keys)) {
-    auto comp = row_equality_comparator<true>(
-      *device_input_table, *device_input_table, nulls_equal == null_equality::EQUAL);
-    auto result_end = unique_copy(sorted_indices->view().begin<cudf::size_type>(),
-                                  sorted_indices->view().end<cudf::size_type>(),
-                                  unique_indices.begin<cudf::size_type>(),
-                                  comp,
-                                  keep,
-                                  stream);
-
-    return cudf::detail::slice(
-      column_view(unique_indices),
-      0,
-      thrust::distance(unique_indices.begin<cudf::size_type>(), result_end));
-  } else {
-    auto comp = row_equality_comparator<false>(
-      *device_input_table, *device_input_table, nulls_equal == null_equality::EQUAL);
-    auto result_end = unique_copy(sorted_indices->view().begin<cudf::size_type>(),
-                                  sorted_indices->view().end<cudf::size_type>(),
-                                  unique_indices.begin<cudf::size_type>(),
-                                  comp,
-                                  keep,
-                                  stream);
-
-    return cudf::detail::slice(
-      column_view(unique_indices),
-      0,
-      thrust::distance(unique_indices.begin<cudf::size_type>(), result_end));
-  }
+  auto comp = row_equality_comparator(
+    nullate::DYNAMIC{cudf::has_nulls(keys)}, *device_input_table, *device_input_table, nulls_equal);
+  auto result_end = unique_copy(sorted_indices->view().begin<cudf::size_type>(),
+                                sorted_indices->view().end<cudf::size_type>(),
+                                unique_indices.begin<cudf::size_type>(),
+                                comp,
+                                keep,
+                                stream);
+
+  return cudf::detail::slice(column_view(unique_indices),
+                             0,
+                             thrust::distance(unique_indices.begin<cudf::size_type>(), result_end));
 }
 }  // namespace
 
diff --git a/cpp/src/transform/one_hot_encode.cu b/cpp/src/transform/one_hot_encode.cu
index 5b5c35df551..16aee349bb5 100644
--- a/cpp/src/transform/one_hot_encode.cu
+++ b/cpp/src/transform/one_hot_encode.cu
@@ -37,10 +37,11 @@ namespace detail {
 
 namespace {
 
-template <typename InputType, bool has_nulls>
+template <typename InputType>
 struct one_hot_encode_functor {
-  one_hot_encode_functor(column_device_view input, column_device_view category)
-    : _equality_comparator{input, category}, _input_size{input.size()}
+  one_hot_encode_functor(column_device_view input, column_device_view category, bool nulls)
+    : _equality_comparator{nullate::DYNAMIC{nulls}, input, category, null_equality::EQUAL},
+      _input_size{input.size()}
   {
   }
 
@@ -52,13 +53,12 @@ struct one_hot_encode_functor {
   }
 
  private:
-  element_equality_comparator<has_nulls> const _equality_comparator;
+  element_equality_comparator<nullate::DYNAMIC> const _equality_comparator;
   size_type const _input_size;
 };
 
 }  // anonymous namespace
 
-template <bool has_nulls>
 struct one_hot_encode_launcher {
   template <typename InputType, CUDF_ENABLE_IF(is_equality_comparable<InputType, InputType>())>
   std::pair<std::unique_ptr<column>, table_view> operator()(column_view const& input_column,
@@ -72,8 +72,8 @@ struct one_hot_encode_launcher {
 
     auto d_input_column    = column_device_view::create(input_column, stream);
     auto d_category_column = column_device_view::create(categories, stream);
-    one_hot_encode_functor<InputType, has_nulls> one_hot_encoding_compute_f(*d_input_column,
-                                                                            *d_category_column);
+    one_hot_encode_functor<InputType> one_hot_encoding_compute_f(
+      *d_input_column, *d_category_column, input_column.nullable() || categories.nullable());
 
     thrust::transform(rmm::exec_policy(stream),
                       thrust::make_counting_iterator(0),
@@ -118,11 +118,7 @@ std::pair<std::unique_ptr<column>, table_view> one_hot_encode(column_view const&
     return std::make_pair(std::move(empty_data), table_view{views});
   }
 
-  return (!input.nullable() && !categories.nullable())
-           ? type_dispatcher(
-               input.type(), one_hot_encode_launcher<false>{}, input, categories, stream, mr)
-           : type_dispatcher(
-               input.type(), one_hot_encode_launcher<true>{}, input, categories, stream, mr);
+  return type_dispatcher(input.type(), one_hot_encode_launcher{}, input, categories, stream, mr);
 }
 
 }  // namespace detail
diff --git a/cpp/tests/iterator/optional_iterator_test.cuh b/cpp/tests/iterator/optional_iterator_test.cuh
index 6b1c2b360eb..d19c9e49ad9 100644
--- a/cpp/tests/iterator/optional_iterator_test.cuh
+++ b/cpp/tests/iterator/optional_iterator_test.cuh
@@ -36,11 +36,11 @@ void nonull_optional_iterator(IteratorTest<T>& testFixture)
   // GPU test
   testFixture.iterator_test_thrust(
     replaced_array,
-    cudf::detail::make_optional_iterator<T>(*d_col, cudf::contains_nulls::DYNAMIC{}, false),
+    cudf::detail::make_optional_iterator<T>(*d_col, cudf::nullate::DYNAMIC{false}),
     host_values.size());
   testFixture.iterator_test_thrust(
     replaced_array,
-    cudf::detail::make_optional_iterator<T>(*d_col, cudf::contains_nulls::NO{}),
+    cudf::detail::make_optional_iterator<T>(*d_col, cudf::nullate::NO{}),
     host_values.size());
 }
 
@@ -72,22 +72,20 @@ void null_optional_iterator(IteratorTest<T>& testFixture)
                  [](auto s, bool b) { return thrust::optional<T>{s}; });
 
   // GPU test for correct null mapping
-  testFixture.iterator_test_thrust(optional_values,
-                                   d_col->optional_begin<T>(cudf::contains_nulls::DYNAMIC{}, true),
-                                   host_values.size());
+  testFixture.iterator_test_thrust(
+    optional_values, d_col->optional_begin<T>(cudf::nullate::DYNAMIC{true}), host_values.size());
 
   testFixture.iterator_test_thrust(
-    optional_values, d_col->optional_begin<T>(cudf::contains_nulls::YES{}), host_values.size());
+    optional_values, d_col->optional_begin<T>(cudf::nullate::YES{}), host_values.size());
   testFixture.iterator_test_thrust(
-    optional_values, d_col->optional_begin<T>(cudf::contains_nulls::YES{}), host_values.size());
+    optional_values, d_col->optional_begin<T>(cudf::nullate::YES{}), host_values.size());
 
   // GPU test for ignoring null mapping
-  testFixture.iterator_test_thrust(value_all_valid,
-                                   d_col->optional_begin<T>(cudf::contains_nulls::DYNAMIC{}, false),
-                                   host_values.size());
+  testFixture.iterator_test_thrust(
+    value_all_valid, d_col->optional_begin<T>(cudf::nullate::DYNAMIC{false}), host_values.size());
 
   testFixture.iterator_test_thrust(
-    value_all_valid, d_col->optional_begin<T>(cudf::contains_nulls::NO{}), host_values.size());
+    value_all_valid, d_col->optional_begin<T>(cudf::nullate::NO{}), host_values.size());
   testFixture.iterator_test_thrust(
-    value_all_valid, d_col->optional_begin<T>(cudf::contains_nulls::NO{}), host_values.size());
+    value_all_valid, d_col->optional_begin<T>(cudf::nullate::NO{}), host_values.size());
 }
diff --git a/cpp/tests/iterator/optional_iterator_test_numeric.cu b/cpp/tests/iterator/optional_iterator_test_numeric.cu
index a8c135a726f..1d23d73bf36 100644
--- a/cpp/tests/iterator/optional_iterator_test_numeric.cu
+++ b/cpp/tests/iterator/optional_iterator_test_numeric.cu
@@ -99,7 +99,7 @@ TYPED_TEST(NumericOptionalIteratorTest, mean_var_output)
     replaced_array.begin(), replaced_array.end(), T{0}, [](T acc, T i) { return acc + i * i; });
 
   // GPU test
-  auto it_dev         = d_col->optional_begin<T>(cudf::contains_nulls::YES{});
+  auto it_dev         = d_col->optional_begin<T>(cudf::nullate::YES{});
   auto it_dev_squared = thrust::make_transform_iterator(it_dev, transformer);
 
   // this can be computed with a single reduce and without a temporary output vector
diff --git a/cpp/tests/table/table_view_tests.cu b/cpp/tests/table/table_view_tests.cu
index 3a792573108..c94963525a0 100644
--- a/cpp/tests/table/table_view_tests.cu
+++ b/cpp/tests/table/table_view_tests.cu
@@ -45,8 +45,8 @@ void row_comparison(cudf::table_view input1,
   auto device_table_2 = cudf::table_device_view::create(input2, stream);
   auto d_column_order = cudf::detail::make_device_uvector_sync(column_order);
 
-  auto comparator = cudf::row_lexicographic_comparator<false>(
-    *device_table_1, *device_table_2, d_column_order.data());
+  auto comparator = cudf::row_lexicographic_comparator(
+    cudf::nullate::NO{}, *device_table_1, *device_table_2, d_column_order.data());
 
   thrust::transform(rmm::exec_policy(stream),
                     thrust::make_counting_iterator(0),
diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu
index 8341425e9e7..2d9fe7775d5 100644
--- a/cpp/tests/utilities/column_utilities.cu
+++ b/cpp/tests/utilities/column_utilities.cu
@@ -18,6 +18,7 @@
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/get_value.cuh>
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/lists/lists_column_view.hpp>
 #include <cudf/strings/convert/convert_datetime.hpp>
@@ -35,16 +36,17 @@
 
 #include <jit/type.hpp>
 
+#include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
 #include <thrust/equal.h>
+#include <thrust/execution_policy.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/logical.h>
+#include <thrust/sequence.h>
 
 #include <numeric>
 #include <sstream>
-#include "cudf/detail/utilities/vector_factories.hpp"
-#include "rmm/cuda_stream_view.hpp"
 
 namespace cudf {
 
@@ -325,11 +327,13 @@ class corresponding_rows_unequal {
                              column_device_view lhs_row_indices_,
                              column_device_view rhs_row_indices_,
                              size_type /*fp_ulps*/)
-    : comp(d_lhs, d_rhs), lhs_row_indices(lhs_row_indices_), rhs_row_indices(rhs_row_indices_)
+    : comp(cudf::nullate::YES{}, d_lhs, d_rhs, cudf::null_equality::EQUAL),
+      lhs_row_indices(lhs_row_indices_),
+      rhs_row_indices(rhs_row_indices_)
   {
   }
 
-  cudf::row_equality_comparator<true> comp;
+  cudf::row_equality_comparator<cudf::nullate::YES> comp;
 
   __device__ bool operator()(size_type index)
   {
@@ -358,7 +362,7 @@ class corresponding_rows_not_equivalent {
                                     size_type fp_ulps_)
     : d_lhs(d_lhs),
       d_rhs(d_rhs),
-      comp(d_lhs, d_rhs),
+      comp(cudf::nullate::YES{}, d_lhs, d_rhs, null_equality::EQUAL),
       lhs_row_indices(lhs_row_indices_),
       rhs_row_indices(rhs_row_indices_),
       fp_ulps(fp_ulps_)
@@ -404,7 +408,7 @@ class corresponding_rows_not_equivalent {
     }
   };
 
-  cudf::row_equality_comparator<true> comp;
+  cudf::row_equality_comparator<cudf::nullate::YES> comp;
 
   __device__ bool operator()(size_type index)
   {