CUDA Kernel内でのポインタの指しているメモリ判定

CUDAではポインタを引数にとってそれがどの記憶領域を指しているかを返す関数があります．
それがこの4つです．

__isConstant
__isGlobal
__isLocal
__isShared

使い方としてはこんな感じです．

#include <iostream>

__constant__ float c;

template <class T>
__device__ void print_memory(const T* const ptr, const char* name) {
  if (__isConstant(ptr)) {
    printf("%7s [%p] : constant\n", name, ptr);
  } else if (__isGlobal(ptr)) {
    printf("%7s [%p] : global\n", name, ptr);
  } else if (__isLocal(ptr)) {
    printf("%7s [%p] : local\n", name, ptr);
  } else if (__isShared(ptr)) {
    printf("%7s [%p] : shared\n", name, ptr);
  } else {
    printf("%7s [%p] : else\n", name, ptr);
  }
}

__global__ void kernel(const float* const g, const float* const h) {
  __shared__ float s;
  float r;

  print_memory(g, "g");
  print_memory(h, "h");
  print_memory(&s, "s");
  print_memory(&r, "r");
  print_memory(&c, "c");
  print_memory(reinterpret_cast<void*>(0), "null");
}

int main() {
  float *g, *h;
  cudaMalloc(&g, sizeof(float));
  cudaMallocHost(&h, sizeof(float));
  kernel<<<1, 1>>>(g, h);
  cudaDeviceSynchronize();
}

実行するとこんな感じの結果が得られます．

      g [0x7f9adb400000] : global
      h [0x7f9adb600000] : global
      s [0x7f9b04000000] : shared
      r [0x7f9b06fffce0] : local
      c [0x7f9ae0800000] : constant
   null [(nil)] : global

nullptrはGlobalメモリ扱いなんですね．
当然コンパイル時には判別しようがないことなので実行時に判定するのですが，PTXではどうなっているかと言うと，それぞれ

isspacep.const
isspacep.global
isspacep.local
isspacep.shared

という命令になっていました．
しかしアドレスを見れば空間の判定なんてできそうな気がしますよね?
そんなことを思いSASSを見てみたところ，適当にアドレスにマスクをかけて判定する系のコードに変わっていました．
流石にハードウェア実装することはないようです．

おまけ

Host側で判定したい場合はこんな感じ．

cudaPointerAttributes attributes;
CUDA_CHECK_ERROR(cudaPointerGetAttributes(&attributes, ptr));
if (attributes.type == cudaMemoryTypeDevice) {
  // ...
} else if (attributes.type == cudaMemoryTypeHost) {
  // ...
} else if (attributes.type == cudaMemoryTypeManaged) {
  // ...
}