CUDA pointers to device constants

I have the following constant defined on the CUDA device:

__constant__ i开发者_开发百科nt deviceTempVariable = 1;

Now I tried getting the address of deviceTempVariable using two methods, and I'm getting different results. The first is a direct memory access from a CUDA kernel as follows:

__global__ void cudaPointers(pointerStruct* devicePointer)
{
    devicePointer->itsPointer = &deviceTempVariable;
}

The other is through host code as follows:

cudaGetSymbolAddress((void**) &pointerCuda, "deviceTempVariable");

I was curious and checked the address values; the first gives something like 00000008, and the second is 00110008. The offset seems to be the same in all cases (the number 8), but the rest is different. What's going on here, and which address would I have to use?

A pointer created in the kernel is clearly usable on the device. It's probably a physical address, although some GPUs might be starting to add virtualization (MMU and TLB).

It looks like cudaGetSymbolAddress is giving you an address usable from the host processor. It's different because device memory has been mapped into the host's virtual address space with an offset.

Host code should use the address returned by cudaGetSymbolAddress, kernel code should use the address-of operator &.

Pointers embedded in shared-data structures need to use based addressing (basically the same as array indexing, you store the offset from a known location that both host and kernel can find).

This is interesting. How do you print the two address ? are you sure you are not printing the address of the pointer ?

You need to pass a

void * address;

to

cudaGetSymbolAddress( &address, "deviceTempVariable");