CODEWITHSUNDEEP
cudanvidiathrust
Rk1
Rk1

Reputation: 178

accessing pointer from __Global__ via thrust

How can we accessing pointer from Global via thrust.

thrust::device_vector<int> testvec;

int *send_data = thrust::raw_pointer_cast(&testvec[0]);

<< <1, 1 >> > (send_data, raw_ptr );

I am able to use send_data in global

work with. i have trouble retrieving it 

thrust::device_ptr<int> dev_ptr = thrust::device_pointer_cast(raw_ptr);

printf("addr: %p\n", dev_ptr);   <<< i am able to retrieve pointer address but not data

Upvotes: 0

Views: 103

Answers (1)

Robert Crovella
Robert Crovella

Reputation: 152003

There are many ways to access the device data from the host. You could:

  1. Copy it back implicitly (demonstrated in the code below)
  2. Copy it back using thrust::copy
  3. Copy it back cudaMemcpy
  4. You can also grab the device_vector elements directly in host code - thrust will do the necessary copying for you under the hood, and then print them out.

The following code demonstrates these four methods - I'm sure there are other approaches or variants as well:

$ cat t897.cu
#include <thrust/device_vector.h>
#include <thrust/copy.h>

#define DSIZE 10

__global__ void kernel(int *send_data, int sz){

  int idx = threadIdx.x+blockDim.x*blockIdx.x;

  if (idx < sz) send_data[idx]++;
}

int main(){

  thrust::device_vector<int> testvec(DSIZE);
  int *send_data = thrust::raw_pointer_cast(&(testvec[0]));

  kernel<<<1,DSIZE>>>(send_data, DSIZE);
// method 1
  thrust::host_vector<int> hvec1 = testvec;
  for (int i = 0; i < DSIZE; i++) printf("%d: %d\n", i, hvec1[i]);
// method 2
  thrust::host_vector<int> hvec2(DSIZE);
  thrust::copy(testvec.begin(), testvec.end(), hvec2.begin());
  for (int i = 0; i < DSIZE; i++) printf("%d: %d\n", i, hvec2[i]);
// method 3
  int *hvec3 = (int *)malloc(DSIZE*sizeof(int));
  cudaMemcpy(hvec3, send_data, DSIZE*sizeof(int), cudaMemcpyDeviceToHost);
  for (int i = 0; i < DSIZE; i++) printf("%d: %d\n", i, hvec3[i]);
// method 4
  for (int i = 0; i < DSIZE; i++) { int temp = testvec[i]; printf("%d: %d\n", i, temp);}
}
$ nvcc -o t897 t897.cu
$ ./t897
0: 1
1: 1
2: 1
3: 1
4: 1
5: 1
6: 1
7: 1
8: 1
9: 1
0: 1
1: 1
2: 1
3: 1
4: 1
5: 1
6: 1
7: 1
8: 1
9: 1
0: 1
1: 1
2: 1
3: 1
4: 1
5: 1
6: 1
7: 1
8: 1
9: 1
0: 1
1: 1
2: 1
3: 1
4: 1
5: 1
6: 1
7: 1
8: 1
9: 1
$

Upvotes: 2

Related Questions