Detection if a pointer is pointing in the device or host in CUDA

Question

In CUDA are there any way of knowing if the pointer is pointing in memory on the device or the host.

A made op example of this could be:

int *dev_c, *host_c;
cudaMalloc( (void**)&dev_c, sizeof(int) );
host_c = (int*) malloc(sizeof(int));

I can of course look at the name, but are there any way of looking ad the pointer dev_c and host_c and saying, host_c points ad the host and dev_c points ad the device.

Answer 1

Starting (I think) CUDA 4 and Fermi GPUs. Nvidia supports UVA (Unified Virtual Address Space). Function cudaPointerGetAttributes seems to do exactly what you're asking for. Note that I believe it only works for host pointers allocated with cudaHostAlloc (and not c malloc).

Answer 2

This is a small example showing how Unified Virtual Addressing can be used to detect if a pointer is pointing to host or device memory space. As @PrzemyslawZych has pointed out, it works only for host pointers allocated with cudaMallocHost.

#include<stdio.h>

#include<cuda.h>   
#include<cuda_runtime.h>

#include<assert.h>
#include<conio.h>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
    if (code != cudaSuccess) 
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        getch();
        if (abort) { exit(code); getch(); }
    }
}

int main() {

    int* d_data;
    int* data; // = (int*)malloc(16*sizeof(int));
    cudaMallocHost((void **)&data,16*sizeof(int));

    gpuErrchk(cudaMalloc((void**)&d_data,16*sizeof(int)));

    cudaDeviceProp prop;
    gpuErrchk(cudaGetDeviceProperties(&prop,0));

    printf("Unified Virtual Addressing %i\n",prop.unifiedAddressing);

    cudaPointerAttributes attributes;
    gpuErrchk(cudaPointerGetAttributes (&attributes,d_data));
    printf("Memory type for d_data %i\n",attributes.memoryType);
    gpuErrchk(cudaPointerGetAttributes (&attributes,data));
    printf("Memory type for data %i\n",attributes.memoryType);

    getch();

    return 0;
}

Answer 3

Not directly. One approach would be to write an encapsulating class for device pointers so that is it absolutely clear that device and host pointers are different in your code. You can see a model of this idea in the Thrust template library, which has a type called device_ptr to clearly delineate device and host pointer types.