cuda - How to differentiate between pointers to shared and global memory? -

in cuda, given value of pointer, or address of variable, there intrinsic or api introspect address space pointer refers to?

the cuda header file sm_20_intrinsics.h defines function

__device__ unsigned int __isglobal(const void *ptr) {   unsigned int ret;   asm volatile ("{ \n\t"                 "    .reg .pred p; \n\t"                 "    isspacep.global p, %1; \n\t"                 "    selp.u32 %0, 1, 0, p;  \n\t" #if (defined(_msc_ver) && defined(_win64)) || defined(__lp64__)                 "} \n\t" : "=r"(ret) : "l"(ptr)); #else                 "} \n\t" : "=r"(ret) : "r"(ptr)); #endif    return ret; } 

this function returns 1 if generic address ptr in global memory space. returns 0 if ptr in shared, local or constant memory space.

the ptx instruction isspacep heavy lifting. seems should able build analogous function way:

__device__ unsigned int __isshared(const void *ptr) {   unsigned int ret;   asm volatile ("{ \n\t"                 "    .reg .pred p; \n\t"                 "    isspacep.shared p, %1; \n\t"                 "    selp.u32 %0, 1, 0, p;  \n\t" #if (defined(_msc_ver) && defined(_win64)) || defined(__lp64__)                 "} \n\t" : "=r"(ret) : "l"(ptr)); #else                 "} \n\t" : "=r"(ret) : "r"(ptr)); #endif    return ret; }