如何区分指针到共享和全局内存？ [英] How to differentiate between pointers to shared and global memory?

问题描述

在CUDA中，给定一个指针的值或变量的地址，是否有一个内在的或另一个API，将会指示指针指向哪个地址空间？

In CUDA, given the value of a pointer, or the address of a variable, is there an intrinsic or another API which will introspect which address space the pointer refers to?

推荐答案

CUDA头文件 sm_20_intrinsics.h 定义函数

The CUDA header file sm_20_intrinsics.h defines the 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;
}

此函数返回 1 如果通用地址 ptr 在全局内存空间中。
如果 ptr 在共享，本地或常量内存空间中，则返回 0 。

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

PTX指令 isspacep 做得很繁重。看起来我们应该能够以这种方式构建类似的函数：

The PTX instruction isspacep does the heavy lifting. It seems like we should be able to build the analogous function this 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;
}

这篇关于如何区分指针到共享和全局内存？的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！