Add workaround for crash in libcudf debug build using output_indexalator in thrust::lower_bound

cpp/src/dictionary/detail/concatenate.cu

@@ -153,19 +153,38 @@ struct dispatch_compute_indices {
         }));
 
     auto new_keys_view = column_device_view::create(new_keys, stream);
+
+    auto begin = new_keys_view->begin<Element>();
+    auto end   = new_keys_view->end<Element>();
+
     // create the indices output column
     auto result = make_numeric_column(
       all_indices.type(), all_indices.size(), mask_state::UNALLOCATED, stream, mr);
     auto result_itr =
       cudf::detail::indexalator_factory::make_output_iterator(result->mutable_view());
     // new indices values are computed by matching the concatenated keys to the new key set
+
+#ifdef NDEBUG
     thrust::lower_bound(rmm::exec_policy(stream),
-                        new_keys_view->begin<Element>(),
-                        new_keys_view->end<Element>(),
+                        begin,
+                        end,
                         all_itr,
                         all_itr + all_indices.size(),
                         result_itr,
                         thrust::less<Element>());
+#else
+    // There is a problem with thrust::lower_bound and the output_indexalator.
+    // https://github.com/NVIDIA/thrust/issues/1452; thrust team created nvbug 3322776
+    // This is a workaround.
+    thrust::transform(rmm::exec_policy(stream),
+                      all_itr,
+                      all_itr + all_indices.size(),
+                      result_itr,
+                      [begin, end] __device__(auto key) {
+                        auto itr = thrust::lower_bound(thrust::seq, begin, end, key);
+                        return static_cast<size_type>(thrust::distance(begin, itr));
+                      });
+#endif
     return result;
   }
 

cpp/src/dictionary/set_keys.cu

@@ -23,6 +23,7 @@
 #include <cudf/detail/stream_compaction.hpp>
 #include <cudf/detail/valid_if.cuh>
 #include <cudf/dictionary/detail/encode.hpp>
+#include <cudf/dictionary/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 #include <cudf/dictionary/dictionary_factories.hpp>
 #include <cudf/stream_compaction.hpp>
@@ -38,6 +39,7 @@ namespace cudf {
 namespace dictionary {
 namespace detail {
 namespace {
+
 /**
  * @brief Type-dispatch functor for remapping the old indices to new values based on the new
  * key-set.
@@ -55,16 +57,11 @@ struct dispatch_compute_indices {
              rmm::mr::device_memory_resource* mr)
   {
     auto dictionary_view = column_device_view::create(input.parent(), stream);
-    auto d_dictionary    = *dictionary_view;
-    auto keys_view       = column_device_view::create(input.keys(), stream);
-    auto dictionary_itr  = thrust::make_permutation_iterator(
-      keys_view->begin<Element>(),
-      thrust::make_transform_iterator(
-        thrust::make_counting_iterator<size_type>(0), [d_dictionary] __device__(size_type idx) {
-          if (d_dictionary.is_null(idx)) return 0;
-          return static_cast<size_type>(d_dictionary.element<dictionary32>(idx));
-        }));
-    auto new_keys_view = column_device_view::create(new_keys, stream);
+    auto dictionary_itr  = make_dictionary_iterator<Element>(*dictionary_view);
+    auto new_keys_view   = column_device_view::create(new_keys, stream);
+
+    auto begin = new_keys_view->begin<Element>();
+    auto end   = new_keys_view->end<Element>();
 
     // create output indices column
     auto result = make_numeric_column(get_indices_type_for_size(new_keys.size()),
@@ -74,14 +71,30 @@ struct dispatch_compute_indices {
                                       mr);
     auto result_itr =
       cudf::detail::indexalator_factory::make_output_iterator(result->mutable_view());
+
+#ifdef NDEBUG
     thrust::lower_bound(rmm::exec_policy(stream),
-                        new_keys_view->begin<Element>(),
-                        new_keys_view->end<Element>(),
+                        begin,
+                        end,
                         dictionary_itr,
                         dictionary_itr + input.size(),
                         result_itr,
                         thrust::less<Element>());
+#else
+    // There is a problem with thrust::lower_bound and the output_indexalator
+    // https://github.com/NVIDIA/thrust/issues/1452; thrust team created nvbug 3322776
+    // This is a workaround.
+    thrust::transform(rmm::exec_policy(stream),
+                      dictionary_itr,
+                      dictionary_itr + input.size(),
+                      result_itr,
+                      [begin, end] __device__(auto key) {
+                        auto itr = thrust::lower_bound(thrust::seq, begin, end, key);
+                        return static_cast<size_type>(thrust::distance(begin, itr));
+                      });
+#endif
     result->set_null_count(0);
+
     return result;
   }