Fix CCCL 2.7.0-rc2 compile issue by removing reference from values.

[bdice]

We are planning to migrate to CCCL 2.7.0 in RAPIDS 25.02. I found a compilation error in libcugraph_etl that is fixed by removing a reference & from a device comparator, and passing the parameters by value instead.

This is a pretty small fix but I expect we will want to build 24.12 with newer CCCL versions for DLFW releases. To avoid needing a backport for DLFW, I am targeting 24.12 with this change. I can retarget to 25.02 if desired.

This is extracted from #4743, but that PR shouldn't be merged (it includes testing changes that are not intended for merge).

[bdice]

Full error output for posterity:

[1/3] Building CUDA object CMakeFiles/cugraph_etl.dir/src/renumbering.cu.o
FAILED: CMakeFiles/cugraph_etl.dir/src/renumbering.cu.o
/usr/bin/sccache /home/coder/.conda/envs/rapids/bin/nvcc -forward-unknown-to-host-compiler -ccbin=/home/coder/.conda/envs/rapids/bin/x86_64-conda-linux-gnu-c++ -DCUDA_API_PER_THREAD_DEFAULT_STREAM -DCUTLASS_NAMESPACE=raft_cutlass -DFMT_SHARED -DLIBCUDACXX_ENABLE_EXPERIMENTAL_MEMORY_RESOURCE -DRAFT_SYSTEM_LITTLE_ENDIAN=1 -DSPDLOG_ACTIVE_LEVEL=SPDLOG_LEVEL_INFO -DSPDLOG_FMT_EXTERN
AL -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA -DTHRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_CPP -Dcugraph_etl_EXPORTS -I/home/coder/cugraph/cpp/libcugraph_etl/src -I/home/coder/cugraph/cpp/libcugraph_etl/include -I/home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../.. -I/home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/l
ibcudacxx/lib/cmake/libcudacxx/../../../include -I/home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../.. -isystem /home/coder/.conda/envs/rapids/targets/x86_64-linux/include -isystem /home/coder/cugraph/cpp/include -isystem /home/coder/rmm/include -isystem /home/coder/rmm/build/conda/cuda-12.5/release/include -isystem /home/coder/rmm/build/conda
/cuda-12.5/release/_deps/nvtx3-src/c/include -isystem /home/coder/raft/cpp/include -isystem /home/coder/.conda/envs/rapids/include -isystem /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cuco-src/include -isystem /home/coder/raft/cpp/build/conda/cuda-12.5/release/_deps/nvidiacutlass-src/include -isystem /home/coder/raft/cpp/build/conda/cuda-12.5/release/_deps/nvidia
cutlass-build/include -isystem /home/coder/cudf/cpp/build/conda/cuda-12.5/release/_deps/dlpack-src/include -isystem /home/coder/cudf/cpp/build/conda/cuda-12.5/release/_deps/jitify-src -isystem /home/coder/cudf/cpp/include -isystem /home/coder/cudf/cpp/build/conda/cuda-12.5/release/include -t=1 -O3 -DNDEBUG -std=c++17 "--generate-code=arch=compute_70,code=[sm_70]" "--generate-cod
e=arch=compute_75,code=[sm_75]" "--generate-code=arch=compute_80,code=[sm_80]" "--generate-code=arch=compute_86,code=[sm_86]" "--generate-code=arch=compute_90,code=[compute_90,sm_90]" -Xcompiler=-fPIC --expt-extended-lambda --expt-relaxed-constexpr -Werror=cross-execution-space-call -Wno-deprecated-declarations -Xptxas=--disable-warnings -Xcompiler=-Wall,-Wno-error=sign-compare,
-Wno-error=unused-but-set-variable -Xfatbin=-compress-all -MD -MT CMakeFiles/cugraph_etl.dir/src/renumbering.cu.o -MF CMakeFiles/cugraph_etl.dir/src/renumbering.cu.o.d -x cu -c /home/coder/cugraph/cpp/libcugraph_etl/src/renumbering.cu -o CMakeFiles/cugraph_etl.dir/src/renumbering.cu.o
/home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/block/block_merge_sort.cuh(69): error: function "cugraph::etl::struct_sort_descending::operator()" cannot be called with the given argument list
            argument types are: (const cugraph::etl::str_hash_value, const cugraph::etl::str_hash_value)
            object type is: cugraph::etl::struct_sort_descending
      if (binary_pred(key2, key1))
          ^
/home/coder/cugraph/cpp/libcugraph_etl/src/renumbering.cu(733): note #3326-D: function "cugraph::etl::struct_sort_descending::operator()" does not match because argument #1 does not match parameter
                       bool operator()(str_hash_value& a, str_hash_value& b)
                            ^
          detected during:
            instantiation of "OffsetT cub::CUB_200700_700_750_800_860_900_NS::MergePath(KeyIt1, KeyIt2, OffsetT, OffsetT, OffsetT, BinaryPred) [with KeyIt1=cugraph::etl::str_hash_value *, KeyIt2=cugraph::etl::str_hash_value *, OffsetT=int, BinaryPred=cugraph::etl::struct_sort_descending]" at line 444
            instantiation of "void cub::CUB_200700_700_750_800_860_900_NS::BlockMergeSortStrategy<KeyT, ValueT, NUM_THREADS, ITEMS_PER_THREAD, SynchronizationPolicy>::Sort(KeyT (&)[ITEMS_PER_THREAD], ValueT (&)[ITEMS_PER_THREAD], CompareOp, int, KeyT) [with KeyT=cugraph::etl::str_hash_value, ValueT=cugraph::etl::accum_type, NUM_THREADS=256, ITEMS_PER_THREAD=5, SynchronizationPol
icy=cub::CUB_200700_700_750_800_860_900_NS::BlockMergeSort<cugraph::etl::str_hash_value, 256, 5, cugraph::etl::accum_type, 1, 1>, CompareOp=cugraph::etl::struct_sort_descending, IS_LAST_TILE=false]" at line 319
            instantiation of "void cub::CUB_200700_700_750_800_860_900_NS::BlockMergeSortStrategy<KeyT, ValueT, NUM_THREADS, ITEMS_PER_THREAD, SynchronizationPolicy>::Sort(KeyT (&)[ITEMS_PER_THREAD], ValueT (&)[ITEMS_PER_THREAD], CompareOp) [with KeyT=cugraph::etl::str_hash_value, ValueT=cugraph::etl::accum_type, NUM_THREADS=256, ITEMS_PER_THREAD=5, SynchronizationPolicy=cub::CU
B_200700_700_750_800_860_900_NS::BlockMergeSort<cugraph::etl::str_hash_value, 256, 5, cugraph::etl::accum_type, 1, 1>, CompareOp=cugraph::etl::struct_sort_descending]" at line 208 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/agent/agent_merge_sort.cuh
            instantiation of "void cub::CUB_200700_700_750_800_860_900_NS::AgentBlockSort<Policy, KeyInputIteratorT, ValueInputIteratorT, KeyIteratorT, ValueIteratorT, OffsetT, CompareOpT, KeyT, ValueT>::consume_tile<IS_LAST_TILE>(OffsetT, int) [with Policy=cub::CUB_200700_700_750_800_860_900_NS::AgentMergeSortPolicy<256, 5, cub::CUB_200700_700_750_800_860_900_NS::BLOCK_LOAD_WAR
P_TRANSPOSE, cub::CUB_200700_700_750_800_860_900_NS::LOAD_DEFAULT, cub::CUB_200700_700_750_800_860_900_NS::BLOCK_STORE_WARP_TRANSPOSE>, KeyInputIteratorT=cugraph::etl::str_hash_value *, ValueInputIteratorT=cugraph::etl::accum_type *, KeyIteratorT=cugraph::etl::str_hash_value *, ValueIteratorT=cugraph::etl::accum_type *, OffsetT=std::ptrdiff_t, CompareOpT=cugraph::etl::struct_sor
t_descending, KeyT=cugraph::etl::str_hash_value, ValueT=cugraph::etl::accum_type, IS_LAST_TILE=false]" at line 163 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/agent/agent_merge_sort.cuh
            instantiation of "void cub::CUB_200700_700_750_800_860_900_NS::AgentBlockSort<Policy, KeyInputIteratorT, ValueInputIteratorT, KeyIteratorT, ValueIteratorT, OffsetT, CompareOpT, KeyT, ValueT>::Process() [with Policy=cub::CUB_200700_700_750_800_860_900_NS::AgentMergeSortPolicy<256, 5, cub::CUB_200700_700_750_800_860_900_NS::BLOCK_LOAD_WARP_TRANSPOSE, cub::CUB_200700_70
0_750_800_860_900_NS::LOAD_DEFAULT, cub::CUB_200700_700_750_800_860_900_NS::BLOCK_STORE_WARP_TRANSPOSE>, KeyInputIteratorT=cugraph::etl::str_hash_value *, ValueInputIteratorT=cugraph::etl::accum_type *, KeyIteratorT=cugraph::etl::str_hash_value *, ValueIteratorT=cugraph::etl::accum_type *, OffsetT=std::ptrdiff_t, CompareOpT=cugraph::etl::struct_sort_descending, KeyT=cugraph::etl
::str_hash_value, ValueT=cugraph::etl::accum_type]" at line 209 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/device/dispatch/dispatch_merge_sort.cuh
            [ 8 instantiation contexts not shown ]
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::__merge_sort::merge_sort<SORT_ITEMS,STABLE,Derived,KeysIt,ItemsIt,CompareOp>(thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execution_policy<Derived> &, KeysIt, KeysIt, ItemsIt, CompareOp) [with SORT_ITEMS=thrust::THRUST_200700_700_750_800_860_900_NS::detail::true_type, STABLE=thr
ust::THRUST_200700_700_750_800_860_900_NS::detail::false_type, Derived=thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rmm::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execute_on_stream_base>, KeysIt=cugraph::etl::str_hash_value *, ItemsIt=cugraph::etl::accum_type *, CompareOp=cugraph::etl::struct_sort_descendi
ng]" at line 332 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../../thrust/system/cuda/detail/sort.h
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::__smart_sort::smart_sort<SORT_ITEMS,STABLE,Policy,KeysIt,ItemsIt,CompareOp,<unnamed>>(Policy &, KeysIt, KeysIt, ItemsIt, CompareOp) [with SORT_ITEMS=thrust::THRUST_200700_700_750_800_860_900_NS::detail::true_type, STABLE=thrust::THRUST_200700_700_750_800_860_900_NS::detail::false_type, Pol
icy=thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execution_policy<thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rmm::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execute_on_stream_base>>, KeysIt=cugraph::etl::str_hash_value *, ItemsIt=cugraph::etl::accum_type *, CompareOp=cugraph::etl::struct_sort_d
escending, <unnamed>=0]" at line 415 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../../thrust/system/cuda/detail/sort.h
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::sort_by_key(thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execution_policy<Derived> &, KeysIt, KeysIt, ValuesIt, CompareOp) [with Derived=thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rmm::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_80
0_860_900_NS::cuda_cub::execute_on_stream_base>, KeysIt=cugraph::etl::str_hash_value *, ValuesIt=cugraph::etl::accum_type *, CompareOp=cugraph::etl::struct_sort_descending]" at line 107 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../../thrust/detail/sort.inl
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::sort_by_key(const thrust::THRUST_200700_700_750_800_860_900_NS::detail::execution_policy_base<DerivedPolicy> &, RandomAccessIterator1, RandomAccessIterator1, RandomAccessIterator2, StrictWeakOrdering) [with DerivedPolicy=thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rm
m::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execute_on_stream_base>, RandomAccessIterator1=cugraph::etl::str_hash_value *, RandomAccessIterator2=cugraph::etl::accum_type *, StrictWeakOrdering=cugraph::etl::struct_sort_descending]" at line 865 of /home/coder/cugraph/cpp/libcugraph_etl/src/renumbering.cu
            instantiation of "std::tuple<std::unique_ptr<cudf::column, std::default_delete<cudf::column>>, std::unique_ptr<cudf::column, std::default_delete<cudf::column>>, std::unique_ptr<cudf::table, std::default_delete<cudf::table>>> cugraph::etl::renumber_functor::operator()<Dtype,<unnamed>>(const raft::handle_t &, const cudf::table_view &, const cudf::table_view &) [with Dt
ype=int8_t, <unnamed>=(void *)nullptr]" at line 461 of /home/coder/cudf/cpp/include/cudf/utilities/type_dispatcher.hpp

/home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/block/block_merge_sort.cuh(69): error: function "cugraph::etl::struct_sort_descending::operator()" cannot be called with the given argument list
            argument types are: (const cugraph::etl::str_hash_value, const cugraph::etl::str_hash_value)
            object type is: cugraph::etl::struct_sort_descending
      if (binary_pred(key2, key1))
          ^
/home/coder/cugraph/cpp/libcugraph_etl/src/renumbering.cu(733): note #3326-D: function "cugraph::etl::struct_sort_descending::operator()" does not match because argument #1 does not match parameter
                       bool operator()(str_hash_value& a, str_hash_value& b)
                            ^
          detected during:
            instantiation of "OffsetT cub::CUB_200700_700_750_800_860_900_NS::MergePath(KeyIt1, KeyIt2, OffsetT, OffsetT, OffsetT, BinaryPred) [with KeyIt1=cugraph::etl::str_hash_value *, KeyIt2=cugraph::etl::str_hash_value *, OffsetT=long, BinaryPred=cugraph::etl::struct_sort_descending]" at line 356 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cu
b/cmake/../../cub/agent/agent_merge_sort.cuh
            instantiation of "void cub::CUB_200700_700_750_800_860_900_NS::AgentPartition<KeyIteratorT, OffsetT, CompareOpT, KeyT>::Process() [with KeyIteratorT=cugraph::etl::str_hash_value *, OffsetT=std::ptrdiff_t, CompareOpT=cugraph::etl::struct_sort_descending, KeyT=cugraph::etl::str_hash_value]" at line 243 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl
-src/cub/cub/cmake/../../cub/device/dispatch/dispatch_merge_sort.cuh
            instantiation of "void cub::CUB_200700_700_750_800_860_900_NS::DeviceMergeSortPartitionKernel(__nv_bool, KeyIteratorT, KeyT *, OffsetT, OffsetT, OffsetT *, CompareOpT, OffsetT, int) [with KeyIteratorT=cugraph::etl::str_hash_value *, OffsetT=std::ptrdiff_t, CompareOpT=cugraph::etl::struct_sort_descending, KeyT=cugraph::etl::str_hash_value]" at line 632 of /home/coder/
cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/device/dispatch/dispatch_merge_sort.cuh
            instantiation of "cudaError_t cub::CUB_200700_700_750_800_860_900_NS::DispatchMergeSort<KeyInputIteratorT, ValueInputIteratorT, KeyIteratorT, ValueIteratorT, OffsetT, CompareOpT, SelectedPolicy>::Invoke<ActivePolicyT>() [with KeyInputIteratorT=cugraph::etl::str_hash_value *, ValueInputIteratorT=cugraph::etl::accum_type *, KeyIteratorT=cugraph::etl::str_hash_value *, 
ValueIteratorT=cugraph::etl::accum_type *, OffsetT=std::ptrdiff_t, CompareOpT=cugraph::etl::struct_sort_descending, SelectedPolicy=cub::CUB_200700_700_750_800_860_900_NS::DeviceMergeSortPolicy<cugraph::etl::str_hash_value *>, ActivePolicyT=cub::CUB_200700_700_750_800_860_900_NS::DeviceMergeSortPolicy<cugraph::etl::str_hash_value *>::Policy350]" at line 657 of /home/coder/cugraph
/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/util_device.cuh
            instantiation of "cudaError_t cub::CUB_200700_700_750_800_860_900_NS::ChainedPolicy<PTX_VERSION, PolicyT, PolicyT>::Invoke(int, FunctorT &) [with PTX_VERSION=350, PolicyT=cub::CUB_200700_700_750_800_860_900_NS::DeviceMergeSortPolicy<cugraph::etl::str_hash_value *>::Policy350, FunctorT=cub::CUB_200700_700_750_800_860_900_NS::DispatchMergeSort<cugraph::etl::str_hash_va
lue *, cugraph::etl::accum_type *, cugraph::etl::str_hash_value *, cugraph::etl::accum_type *, std::ptrdiff_t, cugraph::etl::struct_sort_descending, cub::CUB_200700_700_750_800_860_900_NS::DeviceMergeSortPolicy<cugraph::etl::str_hash_value *>>]" at line 640 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/cub/cub/cmake/../../cub/util_device.cuh
            [ 5 instantiation contexts not shown ]
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::__merge_sort::merge_sort<SORT_ITEMS,STABLE,Derived,KeysIt,ItemsIt,CompareOp>(thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execution_policy<Derived> &, KeysIt, KeysIt, ItemsIt, CompareOp) [with SORT_ITEMS=thrust::THRUST_200700_700_750_800_860_900_NS::detail::true_type, STABLE=thr
ust::THRUST_200700_700_750_800_860_900_NS::detail::false_type, Derived=thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rmm::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execute_on_stream_base>, KeysIt=cugraph::etl::str_hash_value *, ItemsIt=cugraph::etl::accum_type *, CompareOp=cugraph::etl::struct_sort_descendi
ng]" at line 332 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../../thrust/system/cuda/detail/sort.h
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::__smart_sort::smart_sort<SORT_ITEMS,STABLE,Policy,KeysIt,ItemsIt,CompareOp,<unnamed>>(Policy &, KeysIt, KeysIt, ItemsIt, CompareOp) [with SORT_ITEMS=thrust::THRUST_200700_700_750_800_860_900_NS::detail::true_type, STABLE=thrust::THRUST_200700_700_750_800_860_900_NS::detail::false_type, Pol
icy=thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execution_policy<thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rmm::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execute_on_stream_base>>, KeysIt=cugraph::etl::str_hash_value *, ItemsIt=cugraph::etl::accum_type *, CompareOp=cugraph::etl::struct_sort_d
escending, <unnamed>=0]" at line 415 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../../thrust/system/cuda/detail/sort.h
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::sort_by_key(thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execution_policy<Derived> &, KeysIt, KeysIt, ValuesIt, CompareOp) [with Derived=thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rmm::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_80
0_860_900_NS::cuda_cub::execute_on_stream_base>, KeysIt=cugraph::etl::str_hash_value *, ValuesIt=cugraph::etl::accum_type *, CompareOp=cugraph::etl::struct_sort_descending]" at line 107 of /home/coder/cugraph/cpp/build/conda/cuda-12.5/release/_deps/cccl-src/thrust/thrust/cmake/../../thrust/detail/sort.inl
            instantiation of "void thrust::THRUST_200700_700_750_800_860_900_NS::sort_by_key(const thrust::THRUST_200700_700_750_800_860_900_NS::detail::execution_policy_base<DerivedPolicy> &, RandomAccessIterator1, RandomAccessIterator1, RandomAccessIterator2, StrictWeakOrdering) [with DerivedPolicy=thrust::THRUST_200700_700_750_800_860_900_NS::detail::execute_with_allocator<rm
m::mr::thrust_allocator<char>, thrust::THRUST_200700_700_750_800_860_900_NS::cuda_cub::execute_on_stream_base>, RandomAccessIterator1=cugraph::etl::str_hash_value *, RandomAccessIterator2=cugraph::etl::accum_type *, StrictWeakOrdering=cugraph::etl::struct_sort_descending]" at line 865 of /home/coder/cugraph/cpp/libcugraph_etl/src/renumbering.cu
            instantiation of "std::tuple<std::unique_ptr<cudf::column, std::default_delete<cudf::column>>, std::unique_ptr<cudf::column, std::default_delete<cudf::column>>, std::unique_ptr<cudf::table, std::default_delete<cudf::table>>> cugraph::etl::renumber_functor::operator()<Dtype,<unnamed>>(const raft::handle_t &, const cudf::table_view &, const cudf::table_view &) [with Dt
ype=int8_t, <unnamed>=(void *)nullptr]" at line 461 of /home/coder/cudf/cpp/include/cudf/utilities/type_dispatcher.hpp

2 errors detected in the compilation of "/home/coder/cugraph/cpp/libcugraph_etl/src/renumbering.cu".
ninja: build stopped: subcommand failed.

[bdice]

/merge