Copy from std::deque to thrust device_vector

Question

I'm trying to make a sample code where I copy from a std::deque to a thrust::device_vector but the compiler warns calling a __host__ function from a __host__ __device__ function is not allowed (I tried to copy-paste the entire error here but it is beyond the limit of characters in questions). I can post more details if needed.

The code compiles successfully, but I'm really annoyed by these errors as they do not happen with other stl containers as std::list and std::vector. I want to know why they are happening and how do I fix these warning.

Here is the nvcc --version result:

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Tue_Jan_10_13:22:03_CST_2017
Cuda compilation tools, release 8.0, V8.0.61

Here is my sample code

#include <iostream>
#include <algorithm>
#include <deque>
#include <thrust/device_vector.h>

const uint size = 100;

__global__
void hello (int *a) {
  a[threadIdx.x] += threadIdx.x;
}

int main (int argc, const char **argv) {
  std::deque<int> a(size);
  std::fill(a.begin(), a.end(), 1);

  thrust::device_vector<int> a_cuda(a.begin(), a.end());

  dim3 dimBlock(size, 1);
  dim3 dimGrid(1, 1);
  hello<<< dimGrid, dimBlock >>>(thrust::raw_pointer_cast(&a_cuda[0]));

  thrust::copy(a_cuda.begin(), a_cuda.end(), a.begin());

  for (const int i : a) {
    std::cout << i << ' ';
  }

  std::cout << std::endl;

  return EXIT_SUCCESS;
}

And here is the command I'm using to compile:

nvcc file.cu -std=c++11 -O3 -o hello

Thanks in advance

Answer 1

It's seven years later and the current toolchain (12.4) still reports the same warnings. Thrust is now dead as a separate project and has been integrated into a new CUDA C++ library called CCCL.

As I wrote in 2017:

It is obvious that there is something broken in the thrust copy constructor which is emitting code when it probably shouldn't (it also does it with thrust::host_vector copy construction with a std::deque as the source)

I saw no evidence on the Thrust Github tracker that this was ever raised as a bug, and clearly this was never fixed.

You can suppress these specific warnings by passing -diag-suppress 20014 -diag-suppress 20011 as an option to nvcc.

As a final comment, I don't believe that std::deque actually makes any guarantees about the memory it uses for storage to be contiguous, unlike std::vector. That means that using a deque for the source memory to construct a device_vector is technically invalid, because thrust will use a host to device memcpy call on the memory pointed to by begin() to copy rather than call the iterator successively from begin() to end() to traverse the memory. It might work in trivial cases, but it is undefined behaviour to do so.

[Answer assembled from comments and added as a community wiki entry to get the question off the unanswered list for the CUDA tag]