dynamically point to constant memory hangs constant memory

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

T* p = NULL;

..

if(..)

p = c_Array; //constant memory

else

p = d_Array; //device memory

hangs.

But

T* p = d_Array;

...

p = c_Array; //change type without condition

doesn’t hang.

Why? How to dynamically switch a pointer between d_ and the faster c_ ?

Thanks!

2

I believe it’s not possible because the GPU doesn’t dynamically figure out whether it’s a pointer to device or constant memory.

You can see this by looking at the PTX output code. Pointers to constant memory are statically dereferenced with a ld.const.type instruction where loads from the device memory have the ld.global.type instruction.

You can, however, make a conditional memory load without letting a pointer conditionally point to different memory types:

// conditionally dereference either the constant or the device pointer:

T p = (condition) ? c_Array[i] : d_Array[i];

/Pyry

Edit: Changed pointer indirection to array subscripts in code example.