0

I have a device_vector H. I want to create a shallow copy of H using selected indices. I call it J. I want to modify elements of J thereby modifying corresponding elements of H.

My attempt below fails to modify the elements of H when I change elements of J. It appears that thrust allocates new memory to J, instead of using the memory allocated to H.

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/execution_policy.h>

#include <iostream>

int main(void)
{
  // H has storage for 4 integers
  thrust::device_vector<int> H(10);
  thrust::sequence(thrust::device, H.begin(), H.end(),1);

  std::cout << "H="<< std::endl;
  thrust::copy(H.begin(), H.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  thrust::device_vector<int> J(H.begin()+3,H.begin()+9);

  std::cout << "Before modifying J="<< std::endl;
  thrust::copy(J.begin(), J.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  thrust::sequence(thrust::device, J.begin(), J.end(),10);

  std::cout << "after modifying J="<< std::endl;
  thrust::copy(J.begin(), J.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  std::cout << "After modifying H="<< std::endl;
  thrust::copy(H.begin(), H.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  return 0;
}
user27665
  • 673
  • 7
  • 27

3 Answers3

1

This:

thrust::device_vector<int> J(H.begin()+3,H.begin()+9);

is copy construction. There is no way to do what you want without resorting to pointers to the underlying storage, and even then you need to be careful that the source vector never falls out of scope

Two vectors cannot use the same underlying allocation. That is true for std::vector and it is true for thrust vectors.

You can do something similar to what you suggest with thrust::device_ptr:

$ cat t4.cu
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/execution_policy.h>
#include <thrust/device_ptr.h>

#include <iostream>

int main(void)
{
  // H has storage for 4 integers
  thrust::device_vector<int> H(10);
  thrust::sequence(thrust::device, H.begin(), H.end(),1);

  std::cout << "H="<< std::endl;
  thrust::copy(H.begin(), H.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  thrust::device_ptr<int> J(H.data()+3);
  thrust::device_ptr<int> J_end = J+6;

  std::cout << "Before modifying J="<< std::endl;
  thrust::copy(J, J_end, std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  thrust::sequence(thrust::device, J, J_end,10);

  std::cout << "after modifying J="<< std::endl;
  thrust::copy(J, J_end, std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  std::cout << "After modifying H="<< std::endl;
  thrust::copy(H.begin(), H.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  return 0;
}
$ nvcc -o t4 t4.cu
$ ./t4
H=
1,2,3,4,5,6,7,8,9,10,
Before modifying J=
4,5,6,7,8,9,
after modifying J=
10,11,12,13,14,15,
After modifying H=
1,2,3,10,11,12,13,14,15,10,
$
Robert Crovella
  • 143,785
  • 11
  • 213
  • 257
talonmies
  • 70,661
  • 34
  • 192
  • 269
  • Wouldn't it be better to use iterators instead of device_ptr. – user27665 Jan 22 '19 at 05:58
  • @user27665: Your question code uses iterators. That puts you in copy construction territory which you already know doesn't work – talonmies Jan 22 '19 at 06:11
  • my code below using iterators seems to work. please let me know if anything is wrong – user27665 Jan 22 '19 at 11:46
  • `thrust::device_vector J(H.begin()+3,H.begin()+9);` there is no such code, it can compile but can not run: – Nicholas Jela Dec 18 '19 at 03:22
  • @NicholasJela: I have no idea what your comment and downvote is about, but you are wrong: https://pastebin.com/2fsMwgiJ . That code is perfectly valid and it works. – talonmies Dec 19 '19 at 09:37
1

I tried with iterators. It seems to work. The results are posted after the code. The memory location seems to be overwritten as well.

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <thrust/execution_policy.h>
#include <thrust/device_ptr.h>

#include <iostream>

int main(void)
{
  // H has storage for 4 integers
  thrust::device_vector<int> H(10);
  thrust::sequence(thrust::device, H.begin(), H.end(),1);

  std::cout << "H="<< std::endl;
  thrust::copy(H.begin(), H.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  thrust::device_vector<int>::iterator J = H.begin()+3;
  thrust::device_vector<int>::iterator J_end = J+6;

  std::cout << "Before modifying J="<< std::endl;
  thrust::copy(J, J_end, std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  thrust::sequence(thrust::device, J, J_end,10);

  std::cout << "after modifying J="<< std::endl;
  thrust::copy(J, J_end, std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  std::cout << "After modifying H="<< std::endl;
  thrust::copy(H.begin(), H.end(), std::ostream_iterator<int>(std::cout, ","));
  std::cout<< std::endl;

  return 0;
}

Results:

./a.out

H=
1,2,3,4,5,6,7,8,9,10,
Before modifying J=
4,5,6,7,8,9,
after modifying J=
10,11,12,13,14,15,
After modifying H=
1,2,3,10,11,12,13,14,15,10,
user27665
  • 673
  • 7
  • 27
1

I want to create a shallow copy of H using selected indices.

No, you don't want to create a shallow copy.

I call it J [and] to modify elements of J thereby modifying corresponding elements of H.

What you actually want to do - and ended up doing - is modifying a subrange of a container's range of elements. In C++ we do this using iterators; in many cases, these iterator are essentially just pointers.

Another way to do it - when the elements are contiguous in memory - is with an std::span - but that's a C++20 construct (and you may have some trouble with it due to lack of explicit CUDA support, i.e. potential lack of __device__ attributes; same goes for gsl::span in some implementations).

einpoklum
  • 118,144
  • 57
  • 340
  • 684