[BUG]: Thrust code used to compile on CUDA 12.3, no longer compiles on 12.5

[shreyasashok]

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

Thrust

Describe the bug

#include <iostream>
#include <thrust/iterator/transform_output_iterator.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/tuple.h>

class InfoStruct {
};

struct Transformer : public thrust::unary_function<double, double> {
    Transformer(InfoStruct in) 
    : test(in) {
    }  

    __device__
    int operator() (const int& input) {
        return input;
    }
    InfoStruct test;
};

struct ExamplePredicate {
    ExamplePredicate() {}

    __device__
    bool operator() (const int& input) {
        return true;
    }
};

int main() {
    printf("Running MWE test.\n");
    using T = double;
    using TVec = thrust::device_vector<T>;

    TVec input;
    TVec output;

    thrust::copy_if(input.begin(), input.end(), 
                    thrust::make_transform_output_iterator(output.begin(), 
                        Transformer(InfoStruct())),
                    ExamplePredicate());
                    
    printf("Done!\n");
    return 0;
}

The above MWE fails to compile on CUDA toolkit 12.5, but compiles and runs fine on CUDA toolkit 12.3. There appears to be some problem with the make_transform_output_iterator() in combination with copy_if.

The error message is as follows:

/usr/local/cuda-12.5/bin/nvcc example.cu -o example
/usr/local/cuda-12.5/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/copy_if.h(199): error: the default constructor of "thrust::THRUST_200400_520_NS::transform_output_iterator<Transformer, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>>" cannot be referenced -- it is a deleted function
    OutputIt output_end{};
                       ^
          detected during:
            instantiation of "OutputIt thrust::THRUST_200400_520_NS::cuda_cub::detail::copy_if(thrust::THRUST_200400_520_NS::cuda_cub::execution_policy<Derived> &, InputIt, InputIt, StencilIt, OutputIt, Predicate) [with Derived=thrust::THRUST_200400_520_NS::cuda_cub::tag, InputIt=thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>, StencilIt=cub::CUB_200400_520_NS::NullType *, OutputIt=thrust::THRUST_200400_520_NS::transform_output_iterator<Transformer, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>>, Predicate=ExamplePredicate]" at line 244
            instantiation of "OutputIterator thrust::THRUST_200400_520_NS::cuda_cub::copy_if(thrust::THRUST_200400_520_NS::cuda_cub::execution_policy<Derived> &, InputIterator, InputIterator, OutputIterator, Predicate) [with Derived=thrust::THRUST_200400_520_NS::cuda_cub::tag, InputIterator=thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>, OutputIterator=thrust::THRUST_200400_520_NS::transform_output_iterator<Transformer, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>>, Predicate=ExamplePredicate]" at line 49 of /usr/local/cuda-12.5/bin/../targets/x86_64-linux/include/thrust/detail/copy_if.inl
            instantiation of "OutputIterator thrust::THRUST_200400_520_NS::copy_if(const thrust::THRUST_200400_520_NS::detail::execution_policy_base<DerivedPolicy> &, InputIterator, InputIterator, OutputIterator, Predicate) [with DerivedPolicy=thrust::THRUST_200400_520_NS::cuda_cub::tag, InputIterator=thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>, OutputIterator=thrust::THRUST_200400_520_NS::transform_output_iterator<Transformer, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>>, Predicate=ExamplePredicate]" at line 88 of /usr/local/cuda-12.5/bin/../targets/x86_64-linux/include/thrust/detail/copy_if.inl
            instantiation of "OutputIterator thrust::THRUST_200400_520_NS::copy_if(InputIterator, InputIterator, OutputIterator, Predicate) [with InputIterator=thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>, OutputIterator=thrust::THRUST_200400_520_NS::transform_output_iterator<Transformer, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<double>>>, Predicate=ExamplePredicate]" at line 39 of example.cu

1 error detected in the compilation of "example.cu".

How to Reproduce

Compile the MWE with nvcc example.cu -o example using nvcc version 12.5.

Expected behavior

The MWE should compile successfully on CUDA 12.5.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2024 NVIDIA Corporation
Built on Wed_Apr_17_19:19:55_PDT_2024
Cuda compilation tools, release 12.5, V12.5.40
Build cuda_12.5.r12.5/compiler.34177558_0

[shreyasashok]

Further info: the code will compile on NVCC version 12.5 if you add a default constructor to Transformer. This wasn't necessary with NVCC 12.3.

#include <iostream>
#include <thrust/iterator/transform_output_iterator.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/tuple.h>

class InfoStruct {
};

struct Transformer : public thrust::unary_function<double, double> {
    Transformer() = default;
    
    Transformer(InfoStruct in) 
    : test(in) {
    }  

    __device__
    int operator() (const int& input) {
        return input;
    }
    InfoStruct test;
};

struct ExamplePredicate {
    ExamplePredicate() {}

    __device__
    bool operator() (const int& input) {
        return true;
    }
};

int main() {
    printf("Running MWE test.\n");
    using T = double;
    using TVec = thrust::device_vector<T>;

    TVec input;
    TVec output;

    thrust::copy_if(input.begin(), input.end(), 
                    thrust::make_transform_output_iterator(output.begin(), 
                        Transformer(InfoStruct())),
                    ExamplePredicate());
                    
    printf("Done!\n");
    return 0;
}

[bernhardmgruber]

Hi! Thanks for reporting this issue! I can find the offending line in the CUDA Toolkit 12.5, but it no longer exists on our develop branch. Also, your exmaple compiles on Godbolt using CCCL trunk: https://godbolt.org/z/b5sPxbEd5. The issue should thus be fixed in newer versions of CCCL. I don't know which CUDA Toolkit will ship the first fixed version, but you can always use CCCL from GitHub directly.

Also note that many thrust/CUB algorithms require random access iterators, which are required to be default-constructible. Since a transform_iterator stores its transformation function, a non-default constructible function will also make its wrapping transform_iterator not default-constructible and thus no longer a random access iterator.

As a sidenote, thrust::unary_function is being deprecated and will eventually be removed. Your example already works without it in Cuda Toolkit 12.5. Please don't use it anymore.