[BUG]: The latest `cub::DeviceRunLengthEncode` fail to compile with thrust transform iterator

[trivialfis]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Compile-time Error

Component

CUB

Describe the bug

Ref: dmlc/xgboost#11784

The compilation log is provided below:
rle-nvcc.log

The error happens only with the latest CCCL, but it works fine with older versions.

How to Reproduce

I haven't been able to extract an MRE yet. Creating an independent example somehow couldn't reproduce the error.

Expected behavior

The compilation should succeed. RLE should support a transform iterator with a lambda function.

Reproduction link

No response

Operating System

Ubuntu 24.04, Windows 11

nvidia-smi output

No response

NVCC version

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Tue_May_27_02:21:03_PDT_2025
Cuda compilation tools, release 12.9, V12.9.86
Build cuda_12.9.r12.9/compiler.36037853_0

[github-actions]

Hi @trivialfis!

Thanks for submitting this issue - the CCCL team has been notified and we'll get back to you as soon as we can!
In the mean time, feel free to add any relevant information to this issue.

[miscco]

@elstehle this seems like a bug that we rely on the streaming context to be default constructible.

my suggestion is to pass an optional<streaming_context_t> instead of a plain streaming context, especially given that its an optional argument?

[bernhardmgruber]

streaming_context_t is just an aggregate:

cccl/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh

Lines 65 to 78 in 025ed18

	template <typename PrecedingKeyItT, typename AccumT, typename GlobalOffsetT>
	struct streaming_context
	{
	bool first_partition;
	bool last_partition;
	PrecedingKeyItT preceding_key_it;
	// We use a double-buffer to track the aggregate of the last run of the previous partition
	AccumT* preceding_prefix;
	AccumT* prefix_out;
	// We use a double-buffer to track the number of runs of previous partition
	GlobalOffsetT* d_num_previous_uniques_in;
	GlobalOffsetT* d_num_accumulated_uniques_out;

The only reason it cannot be default constructed is if PrecedingKeyItT is not default constructible, which is:

thrust::THRUST_300200_SM_890_NS::transform_iterator<xgboost::linalg::cuda_impl::<unnamed>::IndOp, thrust::THRUST_300200_SM_890_NS::counting_iterator<unsigned long, thrust::THRUST_300200_SM_890_NS::use_default, thrust::THRUST_300200_SM_890_NS::use_default, thrust::THRUST_300200_SM_890_NS::use_default, thrust::THRUST_300200_SM_890_NS::compile_time_value<1>>, thrust::THRUST_300200_SM_890_NS::use_default, thrust::THRUST_300200_SM_890_NS::use_default>

I assume xgboost::linalg::cuda_impl::<unnamed>::IndOp is not default constructible.

[bernhardmgruber]

@trivialfis where can I find the matching source code of xgboost? The error message says:

            instantiation of "cudaError_t cub::CUB_300200_SM_890::DeviceRunLengthEncode::Encode ... at line 51 of /home/jiamingy/ws/xgboost_dev/xgboost/src/common/linalg_op.cu

But if I go to that file on master it has only 49 lines:
https://github.com/dmlc/xgboost/blob/master/src/common/linalg_op.cu
Also, the call to DeviceRunLengthEncode passes a transform_iterator with a lambda. I dont' see where IndOp is coming from.

[elstehle]

Thanks for the hint, @miscco and thank you for pinning down the issue, @bernhardmgruber. I will work on a fix.

[trivialfis]

Apologies, the IndOp is just the lambda function there extracted into a struct. I can repost the error message without the change if that's helpful. The same error happens with the original code.

[bernhardmgruber]

@trivialfis can you tell us whether IndOp or the lambda are default constructible?

[trivialfis]

No, it's not default constructible. It captures variables that are not default constructible.