How to access members of thrust::device_vector

Question

CUDA has some documentation found here: https://docs.nvidia.com/cuda/thrust/index.html#vectors which allows the use of vector in device memory/code. I am trying to create a vector of a struct type to use for general processing. Here is the sample code:

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

struct Data
{
  double first, second, total;
};

__global__
void add(thrust::device_vector<Data> *d_matrix)
{
  &d_matrix[1].total = &d_matrix[1].first + &d_matrix[1].second;
}

int main()
{
  thrust::host_vector<Data> matrix;
  thrust::device_vector<Data> *d_matrix;
  int size = sizeof(thrust::host_vector<Data>);

  matrix[1].first = 2100;
  matrix[1].second = 100;

  cudaMalloc(&d_matrix, size);

  cudaMemcpy(d_matrix, &matrix, size, cudaMemcpyHostToDevice);

  add<<<1,1>>>(d_matrix);

  cudaMemcpy(&matrix, d_matrix, size, cudaMemcpyDeviceToHost);

  cudaFree(d_matrix);

  std::cout << "The sum is: " << matrix[1].total;

  return 0;
}

I get the following error:

gpuAnalysis.cu(13): error: class "thrust::device_vector>" has no member "total"

gpuAnalysis.cu(13): error: class "thrust::device_vector>" has no member "first"

gpuAnalysis.cu(13): error: class "thrust::device_vector>" has no member "second"

3 errors detected in the compilation of "/tmp/tmpxft_000013c9_00000000-8_gpuAnalysis.cpp1.ii".

According to the documentation provided on the nvidia site, these vectors are able to store all data types as std::vector. Is there a way to fix this error to access the members of the struct with each vector element?

There is a lot wrong here. Even after you fix the allocation and operator precedence issues in the code, it won't work because you can't use thrust device vectors in device code, — talonmies, Jun 15 '18 at 18:04
Is there any way for me to use a vector in device code? I assumed that the documentation from Nvidia was saying that thrust device vectors allowed the use of vectors in device code. — Lumbeezl, Jun 15 '18 at 18:11
Nowhere does the documentation state or imply that. You can access the memory backing a device vector in device code, but you cannot use the class itself. — talonmies, Jun 15 '18 at 18:14
Ok. I'm new to CUDA and not the most experienced programmer. Would you mind providing an answer with an examples of that? — Lumbeezl, Jun 15 '18 at 18:17
Thank you my friend. The duplicate is what I was looking for. I searched for a while but could not find a reference to my issue. — Lumbeezl, Jun 15 '18 at 18:20

Daniel Langr · Answer 1 · 2018-06-15T18:03:52.093

void add(thrust::device_vector<Data> *d_matrix) {
   &d_matrix[1].total = &d_matrix[1].first + &d_matrix[1].second;
}

In this code, d_matrix parameter is actually a pointer to an object of type thrust::device_vector<Data>. The expression &d_matrix[1].total is due to C++ operator precedence evaluated such that d_matrix[1] is considered to be the 2nd element of some non-existing array of elements of type thrust::device_vector<Data>, since a pointer can be treated as an array automatically. This (non-existing) 2nd element is then subject of .total member access, which does not exist.

Try (*d_matrix)[1].total = ... instead.

Also I am not sure that your code is correct. For example, you don't specify the size (number of elements, not size of an object) of neither your host_vector nor device_vector. You also cudaMemcpy vector object themselves; does it copy their content as well? Is it even allowed? I have no experience with Thrust, but according this page, there are much simpler way how to create device_vector.

I tried that but the error persists. I'm just not sure how to access the vector memory in device code using CUDA. — Lumbeezl, Jun 15 '18 at 18:17
Even after all this is done (and it is all perfectly correct), the whole idea in the original code is completely wrong and can never be made to work. — talonmies, Jun 16 '18 at 13:02

How to access members of thrust::device_vector

1 Answers1