3

I am currently attempting to use the thrust::remove function on a thrust::device_vector of structs in my main function as shown bellow:

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/execution_policy.h>
#include <thrust/remove.h>

struct config {
    unsigned int ld, ve, rd, solution;
    config() = default;
};

int main() {
    int n = 0;

    //initialize starting configuration
    int curr_size = 1, next_size;
    thrust::device_vector<config> curr(1);
    curr[1] = {0, 0, 0, 0};
    thrust::device_vector<config> next;
    thrust::device_vector<int> sums;

    for (int level = 0; level < ceil(n/2.0); ++level) {
         sums = thrust::device_vector<int>(curr_size);
         next = thrust::device_vector<config>(curr_size * n);
         thrust::fill(next.begin(), next.end(), config{0,0,0,0});



         next_size = thrust::reduce(thrust::device, sums.begin(), sums.end());
         //erroring code
         thrust::remove(thrust::device, curr.begin(), curr.end(), 0);
         next.resize(next_size);
         curr = thrust::device_vector<config>(next);
    }
}

however compiling the program leads to 9 compiler errors indicating invalid type conversions. I have attempted common fixes to thrust::remove errors such as running as a 64-bit project and even attempted it on a device_vector of ints but keep getting the same errors. Errors Shown Bellow(reduced due to character count limitations)

FAILED: CMakeFiles/Solver.dir/main.cu.o  
/usr/local/cuda-13.0/bin/nvcc -forward-unknown-to-host-compiler   -g "--generate-code=arch=compute_89,code=[compute_89,sm_89]" -MD -MT CMakeFiles/Solver.dir/main.cu.o -MF CMakeFiles/Solver.dir/main.cu.o.d -x cu -rdc=true -c /home/aowyn/CLionProjects/Solver/main.cu -o CMakeFiles/Solver.dir/main.cu.o
/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: function "cuda::std::__4::__not_fn_t<_Fn>::operator()(_Args &&...) & [with _Fn=thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>>, _Args=<config &>, __cccl_true_=true, <unnamed>=0]" (declared at line 62 of /usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cuda/std/__functional/not_fn.h) cannot be referenced -- it is a deleted function
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: invalid type conversion
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: function "cuda::std::__4::__not_fn_t<_Fn>::operator()(_Args &&...) & [with _Fn=thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>>, _Args=<config &>, __cccl_true_=true, <unnamed>=0]" (declared at line 62 of /usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cuda/std/__functional/not_fn.h) cannot be referenced -- it is a deleted function
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: invalid type conversion
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: function "cuda::std::__4::__not_fn_t<_Fn>::operator()(_Args &&...) & [with _Fn=thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>>, _Args=<config &>, __cccl_true_=true, <unnamed>=0]" (declared at line 62 of /usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cuda/std/__functional/not_fn.h) cannot be referenced -- it is a deleted function
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: invalid type conversion
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: function "cuda::std::__4::__not_fn_t<_Fn>::operator()(_Args &&...) & [with _Fn=thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>>, _Args=<config &>, __cccl_true_=true, <unnamed>=0]" (declared at line 62 of /usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cuda/std/__functional/not_fn.h) cannot be referenced -- it is a deleted function
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/cub/agent/agent_select_if.cuh(407): error: invalid type conversion
          selection_flags[ITEM] = static_cast<bool>(select_op(items[ITEM]));

/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/thrust/detail/function.h(44): error: no instance of function template "thrust::THRUST_300001_SM_890_NS::detail::functional::actor<Eval>::operator() [with Eval=thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>]" matches the argument list
            argument types are: (config)
            object type is: thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>>
      return static_cast<Result>(m_f(thrust::raw_reference_cast(::cuda::std::forward<Ts>(args))...));
                                 ^
/usr/local/cuda-13.0/targets/x86_64-linux/include/cccl/thrust/detail/functional/actor.h(59): note #3327-D: candidate function template "thrust::THRUST_300001_SM_890_NS::detail::functional::actor<Eval>::operator() [with Eval=thrust::THRUST_300001_SM_890_NS::detail::functional::composite<thrust::THRUST_300001_SM_890_NS::detail::functional::operator_adaptor<cuda::std::__4::equal_to<void>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::argument<0U>>, thrust::THRUST_300001_SM_890_NS::detail::functional::actor<thrust::THRUST_300001_SM_890_NS::detail::functional::value<int>>>]" failed deduction
                     auto operator()(Ts&&... ts) const -> decltype(Eval::eval(::cuda::std::forward<decltype(ts)>(ts)...))

9 errors detected in the compilation of "/home/aowyn/CLionProjects/Solver/main.cu".
ninja: build stopped: subcommand failed.

Any suggestions on how to fix this would be greatly appreciated, thank you.

2
  • 1
    curr[1] = {0, 0, 0, 0}; is an off-by-one error. Given that vector elements (i.e. curr[0]) are already default constructed, the whole access is unnecessary. Same goes for the fill. Commented Nov 11 at 17:22
  • 1
    While some critical information about the actual code was probably lost when creating the MRE, the usage of vector-constructors in the loop seems very inefficient to me. Allocation of device memory is not cheap. I suggest working more with either views (cuda::std::span is backported to C++17), swap()/move() or resize() (it is unclear if sizes are going up or down in later iterations). Commented Nov 11 at 17:32

1 Answer 1

5

Signature of remove is

template<typename DerivedPolicy, typename ForwardIterator, typename T>
ForwardIterator thrust::remove(
const thrust::detail::execution_policy_base<DerivedPolicy> &exec,
ForwardIterator first,
ForwardIterator last,
const T &value,
)

The last argument should be of type T (or something implicitly convertible to T), but you used 0. If you mean to remove the config element that is {0, 0, 0, 0}, the quick fix could be:

struct config {
    unsigned int ld, ve, rd, solution;
    // default possible since c++20, otherwise you need to implement it
    bool operator==(const config&) const = default;
};

and then

thrust::remove(thrust::device, curr.begin(), curr.end(), config{0, 0, 0, 0});

Note the defaulted comparison operator. You don't need to explicitly declare default constructor.

Sign up to request clarification or add additional context in comments.

2 Comments

Thank you very much! notably when i used the suggested line my I got an error saying that the extra = default is not supported in my language dialect. but once i removed it and just used bool operator==(const config&) const; it compiled fine.
A yes, forgot to mention that defaulted comparison operators are supported since C++20 en.cppreference.com/w/cpp/language/default_comparisons.html

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Start asking to get answers

Find the answer to your question by asking.

Ask question

Explore related questions

See similar questions with these tags.