Can we copy a "normal" GPU memory to a "unified" memory?

Question

We have two GPU memories, one is allocated with cuMalloc as normal device memory, the other is allocated with cuMallocManaged as unified memory. Is it possible to copy between them? and if we use driver API, what direction should I use?

float* normalMem, unifiedMem;
cuMalloc(&normalMem, 100);
cuMallocManaged(&unifiedMem, 100);
cuMemcpyD2D(unifiedMem, normalMem, 100); // ? D2D? or D2H? or else?

Answer 1

Yes you can. Look at the following code for instance.

• It declared a normal pointer a managed pointer and an host pointer all of them of 100 float.
• It then initializes the values in the host pointer and then copy the values using cudaMemCpy to the normal pointer.
• Values are now copied to the managed pointer
• The managed pointer is used in a kernel to shows that values have been copied from the two buffers.

I think that the code is pretty self-explanatory

__global__ 
void test(float* d_ptr){
    for(int i=0;i<100;i++)
        printf("%f \n",d_ptr[i]);
    printf("\n");
}

////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{

    size_t size = sizeof(float)*100;
    float* h_p =(float*) malloc(size);
    float* d_p, dm_p ; 
    cudaMalloc(&d_p,size);
    cudaMallocManaged(&dm_p,size);

    for(int i=0;i<100;i++)
        h_p[i]=2*(float)i;

    cudaMemcpy(d_p,h_p,size,cudaMemcpyHostToDevice);

    cudaDeviceSynchronize();

    cudaMemcpy(dm_p,d_p,size,cudaMemcpyDeviceToDevice);

    cudaDeviceSynchronize();

    test<<<1,1>>>(dm_p);

    cudaDeviceSynchronize();

    cudaFree(dm_p);
    cudaFree(d_p);
    free(h_p);
    return 0;
}

Remember to read the Unified Memory access rules.