CODEWITHSUNDEEP
c++memorypointerscuda
Kristian D'Amato
Kristian D'Amato

Reputation: 4046

CUDA pointers to device constants

I have the following constant defined on the CUDA device:

__constant__ int 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?

Upvotes: 1

Views: 2071

Answers (2)

Ben Voigt
Ben Voigt

Reputation: 283901

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).

Upvotes: 4

fabmilo
fabmilo

Reputation: 48330

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");

Upvotes: 0

Related Questions