Come differenziare i puntatori alla memoria globale e condivisa?

Question

In CUDA, dato il valore di un puntatore, o l'indirizzo di una variabile, esiste una API intrinseca o un'altra che possa indagare su quale spazio di indirizzi si riferisce il puntatore?

Answer 1

Il CUDA file di intestazione sm_20_intrinsics.h definisce la funzione

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

Questa funzione restituisce 1 se l'indirizzo generico ptr è nello spazio di memoria globale. Restituisce 0 se ptr è nello spazio di memoria condiviso, locale o costante.

L'istruzione PTX isspacep esegue il sollevamento pesante. Sembra che dovremmo essere in grado di costruire la funzione analoga in questo modo:

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