calling a __host__ function from a __host__ __device__ when using gather with zip_iterator for vector of vectors
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
I would like to gather data from a struct of arrays, which contains some vector of vectors. `thrust::zip_iterator` is used to fuse the calls of individual arrays. This causes a series of warnings.
```
#include
#include
#include
int main(){
thrust::host_vector map(42);
thrust::host_vector> vecvec(42);
thrust::host_vector> outvecvec(42);
auto in = thrust::make_zip_iterator(thrust::make_tuple(
vecvec.begin()
));
auto out = thrust::make_zip_iterator(thrust::make_tuple(
outvecvec.begin()
));
#if 1
// warning calling a __host__ function from a __host__ __device__
thrust::gather(
map.begin(),
map.end(),
in,
out
);
#else
// no warning
thrust::gather(
map.begin(),
map.end(),
vecvec.begin(),
outvecvec.begin()
);
#endif
}
```
https://cuda.godbolt.org/z/4Ebo6Yn57
```
/opt/compiler-explorer/libs/thrustcub/trunk/thrust/detail/tuple.inl(458): warning: calling a __host__ function from a __host__ __device__ function is not allowed
detected during:
instantiation of "thrust::detail::cons::cons(const thrust::detail::cons &) [with HT=thrust::host_vector>, HT2=thrust::host_vector> &]"
/opt/compiler-explorer/libs/thrustcub/trunk/thrust/tuple.h(346): here
/opt/compiler-explorer/libs/thrustcub/trunk/thrust/detail/function.h(118): warning: calling a __host__ function("thrust::host_vector > ::~host_vector()") from a __host__ __device__ function("thrust::detail::cons< ::thrust::host_vector > , ::thrust::null_type> ::~cons") is not allowed
/opt/compiler-explorer/libs/thrustcub/trunk/thrust/detail/tuple.inl(458): warning: calling a __host__ function("thrust::host_vector > ::host_vector(const thrust::host_vector > &)") from a __host__ __device__ function("thrust::detail::cons< ::thrust::host_vector > , ::thrust::null_type> ::cons< ::thrust::host_vector > &> ") is not allowed
```
Contributor guide
Assessment
This issue has not been assessed yet.