how to change device id of thrust::device_vector

Question

is it possible to change the device id of thrust::device vector? I'm thinking to code like this

cudaSetDevice(0);                                //set to device:0
thrust::device_vector<int> a(10);                //define the device vector
a.clear(); thrust::device_vector<int>().swap(a); //deallocate the vector

cudaSetDevice(1);                                //set to device:1
a.resize(10);

is it possible? thanks

Answer 1

I don't know whether this is a correct answer. so please correct me if I'm wrong because i don't whether the test is proper enough. So I decided to do test on vector addition as follows

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

__global__ void
vectorAdd(const int *A, const int *B, int *C, int numElements) {
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < numElements) C[i] = A[i] + B[i];
};

int main(void)
{
    int numElements = 1024;
    int randacc = 30;

    cudaSetDevice(0);
    thrust::device_vector<int> a(numElements, 1);
    thrust::device_vector<int> b(numElements, 2);
    thrust::device_vector<int> c(numElements);

    int* a_d = thrust::raw_pointer_cast(&a[0]);
    int* b_d = thrust::raw_pointer_cast(&b[0]);
    int* c_d = thrust::raw_pointer_cast(&c[0]);

    int threadsPerBlock = 64;
    int blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(a_d, b_d, c_d, numElements);
    cudaError_t err = cudaGetLastError();

    if (err != cudaSuccess) std::cerr << cudaGetErrorString(err) << std::endl;
    std::cout << "random access on dev 0, c = " << c[randacc] << std::endl;

    a.clear(); thrust::device_vector<int>().swap(a); //deallocate the vector
    b.clear(); thrust::device_vector<int>().swap(b); //deallocate the vector
    c.clear(); thrust::device_vector<int>().swap(c); //deallocate the vector

    cudaSetDevice(1);                                //set to device:1
    a.resize(numElements, 1);
    b.resize(numElements, 2);
    c.resize(numElements);

    a_d = thrust::raw_pointer_cast(&a[0]);
    b_d = thrust::raw_pointer_cast(&b[0]);
    c_d = thrust::raw_pointer_cast(&c[0]);

    threadsPerBlock = 64;
    blocksPerGrid =(numElements + threadsPerBlock - 1) / threadsPerBlock;

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(a_d, b_d, c_d, numElements);
    err = cudaGetLastError();

    if (err != cudaSuccess) std::cerr << cudaGetErrorString(err) << std::endl;
    std::cout << "random access on dev 1, c = " << c[randacc] << std::endl;

    return 0;
}

and I get result:

random access on dev 0, c = 3

random access on dev 1, c = 3

Note: you need at least 2 GPUs on the same host to test it. i tested on my GTX690

Answer 2

I do not know if and how it works exactly with thrust and if you tried to change the device id of a device array not using thrust and if you are aware of peer-to-peer memory access.

It is an interesting question but I can do my own experiments. However, according to the CUDA Programming Guide (section 3.2.6.4) peer-to-peer memory access is possible (i.e. two devices can address each other's memory) if you use compute capability 2.x and above using Tesla cards.

This is an example from the Programming Guide:

cudaSetDevice(0);
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size);
MyKernel<<<1000, 128>>>(p0);
cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0, 0);  // <- this enables peer to peer access
MyKernel<<<1000,128>>>(p0);

Concerning memory copy from one device to another, the Programming Guide says that cudaMemcpyPeer() can do the job and provides an example. I can not find something in the thrust documentation that corresponds to your question, so I think the best way is to try it.