Wie unterscheidet man Zeiger auf gemeinsamen und globalen Speicher?

Question

In CUDA, gibt es eine intrinsische oder eine andere API, die den Wert eines Zeigers oder die Adresse einer Variablen angibt, die den Adressraum prüft, auf den sich der Zeiger bezieht?

Answer 1

Die CUDA-Header-Datei sm_20_intrinsics.h definiert die Funktion

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

Diese Funktion gibt 1 wenn generische Adresse ptr in globaler Speicherplatz ist. Es gibt 0 zurück, wenn sich ptr in gemeinsam genutztem, lokalem oder konstantem Speicherplatz befindet.

Die PTX-Anweisung isspacep übernimmt das schwere Heben. Es scheint, als ob wir in der Lage sein sollten, die analoge Funktion auf diese Weise aufzubauen:

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