how to use thrust with unified memory?

[ztdepztdep]

Does thrust support the unified memory?

[gevtushenko]

Hello, @ztdepztdep! There's thrust universal vector that represents unified (managed) memory:

#include <thrust/universal_vector.h>

__global__ void kernel(int *data) {
  printf("%d\n", data[threadIdx.x]);
}

int main() {
  thrust::universal_vector<int> in(4);

  in[0] = 0;
  in[1] = 1;
  in[2] = 2;
  in[3] = 3;

  kernel<<<1, 4>>>(thrust::raw_pointer_cast(in.data()));
  cudaDeviceSynchronize();
}

[ztdepztdep]

#include <thrust/universal_vector.h>

Thanks a lot. I use cuda11.3, I can't find the include fine in my "/usr/local/cuda-11.3/include/thrust/"

[gevtushenko]

Unfortunately, universal vector was introduces in Thrust 1.12 (CTK 11.4).

[alliepiper]

Thrust algorithms will also work with unified memory that's allocated manually. Just pass the thrust::device execution policy as the first argument if using raw pointers to unified memory. This will tell Thrust to execute on the GPU -- otherwise it will default to host-side execution for raw pointers.