如何區分指向共享和全局內存的指針？

Question

在CUDA中，給定一個指針的值或變量的地址，是否有一個內部或另一個API來反省指針指向哪個地址空間？

Answer 1

的CUDA頭文件sm_20_intrinsics.h定義函數

__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; 
} 

該函數返回1如果通用地址ptr是全球內存空間。 如果ptr位於共享，本地或常量內存空間中，它將返回0。

PTX指令isspacep做了繁重的工作。看來我們應該可以通過這種方式建立類似的功能：

__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; 
}