Jak rozróżnić wskaźniki na pamięć współużytkowaną i globalną?

Question

W CUDA, biorąc pod uwagę wartość wskaźnika lub adres zmiennej, czy istnieje wewnętrzny lub inny interfejs API, który będzie introspekować, do której przestrzeni adresowej odnosi się wskaźnik?

Answer 1

CUDA plik header sm_20_intrinsics.h określa funkcję

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

Funkcja ta zwraca 1 jeśli generic adres ptr jest w globalnej przestrzeni pamięci. Zwraca 0, jeśli ptr jest w udostępnionym, lokalnym lub stałym miejscu pamięci.

Instrukcja obsługi PTX isspacep wykonuje podnoszenie ciężkie. Wygląda na to, że powinniśmy być w stanie zbudować analogiczną funkcję w następujący sposób:

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