cuda - 共有メモリとグローバルメモリへのポインタを区別する方法は?

Question

CUDA では、ポインターの値または変数のアドレスが与えられた場合、ポインターが参照するアドレス空間をイントロスペクトする組み込み API または別の 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がグローバル メモリ空間にある場合に戻ります。が共有、ローカル、または定数メモリ空間にある0かどうかを返します。ptr

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