dynamically point to constant memory hangs constant memory


T* p = NULL;



p = c_Array; //constant memory


p = d_Array; //device memory



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_ ?


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];


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