Làm cách nào để phân biệt các con trỏ với bộ nhớ chung và toàn cục?

Question

Trong CUDA, được đưa ra giá trị của một con trỏ, hoặc địa chỉ của một biến, có một nội tại hoặc một API khác mà sẽ xem xét không gian địa chỉ mà con trỏ đề cập đến?

Answer 1

Các CUDA phần đầu tập tin sm_20_intrinsics.h xác định chức năng

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

Hàm này trả 1 nếu địa chỉ chung chung ptr là trong không gian bộ nhớ toàn cầu. Nó trả về 0 nếu ptr nằm trong không gian bộ nhớ được chia sẻ, cục bộ hoặc không đổi.

Hướng dẫn PTX isspacep thực hiện việc nâng hạng nặng. Có vẻ như chúng ta có thể xây dựng hàm tương tự theo cách này:

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