cuda - 如何区分共享内存和全局内存的指针？

Question

在 CUDA 中，给定指针的值或变量的地址，是否有内部函数或其他 API 可以内省(introspection)指针引用的地址空间？

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

如果通用地址ptr位于全局内存空间中，则该函数返回1。如果 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;
}

更新:

__isGlobal() 和其他地址空间谓词函数在 CUDA C++ Programming Guide 中进行了描述。 .