Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[BUG]: thrust::remove_if on permutation iterator with counting iterator fails to compile #1974

Closed
1 task done
goodcucumber opened this issue Jul 11, 2024 · 1 comment
Closed
1 task done

[BUG]: thrust::remove_if on permutation iterator with counting iterator fails to compile #1974

goodcucumber opened this issue Jul 11, 2024 · 1 comment
Assignees
@elstehle
Labels
bug Something isn't working right.

Comments

@goodcucumber
Copy link

Is this a duplicate?

Type of Bug

Compile-time Error

Component

Thrust

Describe the bug

The following code compiles and runs fine on cuda 12.4 with thrust 2.3.2, but fails on cuda 12.5 with thrust 2.4.0.

// t.cu
#include <iostream>
#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/remove.h>
__global__ void f(){
	printf("%d\n",threadIdx.x);
}

int main(){
	thrust::device_vector<int> v1{10,11,12,13,20,21,22,23,30,31,32,33};
	thrust::device_vector<bool> v2{true,false,false};
	auto t1 = thrust::make_permutation_iterator(v1.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0),[]__host__ __device__(const int& i){return i*4+0;}));
	auto t2 = thrust::make_permutation_iterator(v1.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0),[]__host__ __device__(const int& i){return i*4+1;}));
	auto t3 = thrust::make_permutation_iterator(v1.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0),[]__host__ __device__(const int& i){return i*4+2;}));
	auto t4 = thrust::make_permutation_iterator(v1.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0),[]__host__ __device__(const int& i){return i*4+3;}));

	auto begin = thrust::make_zip_iterator(thrust::make_tuple(t1,t2,t3,t4,v2.begin()));
	auto end = thrust::remove_if(begin, begin+3,[]__device__(const thrust::tuple<int,int,int,int,bool>&v){return thrust::get<4>(v);});
	std::cout<<(end-begin)<<std::endl;
	v1.resize((end-begin)*4);
	v2.resize(end-begin);
	for(int i = 0; i < v1.size(); i++){
		std::cout<<static_cast<int>(v1[i])<<std::endl;
	}
}

I get the folloing very complicated error message:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\bin/../include\thrust/system/cuda/detail/copy_if.h(199): error: the default constructor of "thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>" 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 &, InputIt, InputIt, StencilIt, OutputIt, Predicate) [with Derived=thrust::THRUST_200400_520_NS::cuda_cub::tag, InputIt=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, StencilIt=cub::CUB_200400_520_NS::NullType *, OutputIt=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, Predicate=thrust::THRUST_200400_520_NS::detail::unary_negate<lambda [](const cuda::std::__4::tuple<int, int, int, int, __nv_bool> &)->__nv_bool>]" at line 244
instantiation of "OutputIterator thrust::THRUST_200400_520_NS::cuda_cub::copy_if(thrust::THRUST_200400_520_NS::cuda_cub::execution_policy &, InputIterator, InputIterator, OutputIterator, Predicate) [with Derived=thrust::THRUST_200400_520_NS::cuda_cub::tag, InputIterator=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, OutputIterator=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, Predicate=thrust::THRUST_200400_520_NS::detail::unary_negate<lambda [](const cuda::std::__4::tuple<int, int, int, int, __nv_bool> &)->__nv_bool>]" at line 72 of C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\bin/../include\thrust/system/cuda/detail/remove.h
instantiation of "InputIt thrust::THRUST_200400_520_NS::cuda_cub::remove_if(thrust::THRUST_200400_520_NS::cuda_cub::execution_policy &, InputIt, InputIt, Predicate) [with Derived=thrust::THRUST_200400_520_NS::cuda_cub::tag, InputIt=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, Predicate=lambda [](const cuda::std::__4::tuple<int, int, int, int, __nv_bool> &)->__nv_bool]" at line 80 of C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\bin/../include\thrust/detail/remove.inl
instantiation of "ForwardIterator thrust::THRUST_200400_520_NS::remove_if(const thrust::THRUST_200400_520_NS::detail::execution_policy_base &, ForwardIterator, ForwardIterator, Predicate) [with DerivedPolicy=thrust::THRUST_200400_520_NS::cuda_cub::tag, ForwardIterator=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, Predicate=lambda [](const cuda::std::__4::tuple<int, int, int, int, __nv_bool> &)->__nv_bool]" at line 185 of C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\bin/../include\thrust/detail/remove.inl
instantiation of "ForwardIterator thrust::THRUST_200400_520_NS::remove_if(ForwardIterator, ForwardIterator, Predicate) [with ForwardIterator=thrust::THRUST_200400_520_NS::zip_iterator<cuda::std::__4::tuple<thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::permutation_iterator<thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr>, thrust::THRUST_200400_520_NS::transform_iterator<lambda [](const int &)->int, thrust::THRUST_200400_520_NS::counting_iterator<int, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>, thrust::THRUST_200400_520_NS::use_default, thrust::THRUST_200400_520_NS::use_default>>, thrust::THRUST_200400_520_NS::detail::normal_iterator<thrust::THRUST_200400_520_NS::device_ptr<__nv_bool>>>>, Predicate=lambda [](const cuda::std::__4::tuple<int, int, int, int, __nv_bool> &)->__nv_bool]" at line 18 of t.cu

How to Reproduce

  1. make some permutation_iterators with a device_vector and a transformed counting iterator (thrust::make_permutation_iterator(v1.begin(), thrust::make_transform_iterator(thrust::make_counting_iterator(0),[]__host__ __device__(const int& i){return i*4+0;}));)
  2. do remove_if on this iterator
  3. compilation fails on the code with a very complicated error message.

Expected behavior

The code compiles without error

Reproduction link

No response

Operating System

windows 10 and Ubuntu 20.04

nvidia-smi output

windows 10:
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 555.85 Driver Version: 555.85 CUDA Version: 12.5 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Driver-Model | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 Tesla P4 TCC | 00000000:01:00.0 Off | 0 |
| N/A 36C P8 7W / 75W | 9MiB / 7680MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
+-----------------------------------------------------------------------------------------+

Ubuntu 20.04:
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 555.42.02 Driver Version: 555.42.02 CUDA Version: 12.5 |
|-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA GeForce RTX 3080 Ti Off | 00000000:02:00.0 Off | N/A |
| 30% 36C P8 20W / 350W | 30MiB / 12288MiB | 0% Default |
| | | N/A |
+-----------------------------------------+------------------------+----------------------+

+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| 0 N/A N/A 1261 G /usr/lib/xorg/Xorg 9MiB |
| 0 N/A N/A 1567 G /usr/bin/gnome-shell 4MiB |
+-----------------------------------------------------------------------------------------+

NVCC version

Windows 10:
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

Ununtu 20.04:
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
@goodcucumber goodcucumber added the bug Something isn't working right. label Jul 11, 2024
@github-project-automation github-project-automation bot moved this to Todo in CCCL Jul 11, 2024
@elstehle elstehle self-assigned this Jul 11, 2024
@elstehle
Copy link
Collaborator

Thank you for reporting the issue.

The issue has already been fixed by #1643. It should be fixed with the latest CCCL version, if you have a chance to use CCCL from GitHub. Otherwise, the fix will be included in the next CTK release.

I'm closing the issue as fixed. Please feel free to reopen, if you are still having trouble.

@github-project-automation github-project-automation bot moved this from Todo to Done in CCCL Jul 11, 2024
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Assignees

@elstehle elstehle

Labels
bug Something isn't working right.
Projects
CCCL
Archived in project
Milestone
No milestone
Development

No branches or pull requests
2 participants
@goodcucumber @elstehle