5
在CUDA中,给定一个指针的值或变量的地址,是否有一个内部或另一个API来反省指针指向哪个地址空间?如何区分指向共享和全局内存的指针?
在CUDA中,给定一个指针的值或变量的地址,是否有一个内部或另一个API来反省指针指向哪个地址空间?如何区分指向共享和全局内存的指针?
的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;
}
请注意,本地内存也有'isspacep.local'。 – BenC