我可以检查地址是否在共享内存中？

Question

我想写以下CUDA功能：

void foo(int* a, size_t n) 
{ 
    if (/* MAGIC 1 */) { 
     // a is known to be in shared memory, 
     // so use it directly 
    } 
    else { 
     // make a copy of a in shared memory 
     // and use the copy 
    } 
} 

在主机方面，我们有一个轻微相关设施中的cudaPointerGetAttributes形式，它可以告诉我们一个指针是否为设备内存或主机内存;也许还有一些方法可以区分设备代码中的指针，也可以从全局指针中辨别共享。或者，甚至可能更好 - 也许有编译时机制来做到这一点，因为毕竟，设备函数只被编译到内核中，并不是独立的，所以nvcc通常可以知道它们是用于共享内存还是不。

Answer 1

您可以通过内嵌位的 “集结号” 使用isspacep PTX instruction：

// First, a pointer-size-related definition, in case 
// this code is being compiled in 32-bit rather than 
// 64-bit mode; if you know the code is always 64-bit 
// you can just use the "l" 

#if defined(_WIN64) || defined(__LP64__) 
# define PTR_CONSTRAINT "l" 
#else 
# define PTR_CONSTRAINT "r" 
#endif 

__device__ int isShared(void *ptr) 
{ 
    int res; 
    asm("{" 
     ".reg .pred p;\n\t" 
     "isspacep.shared p, %1;\n\t" 
     "selp.b32 %0, 1, 0, p;\n\t" 
     "}" : 
     "=r"(res): PTR_CONSTRAINT(ptr)); 
    return res; 
} 

所以你的榜样变得

void foo(int* a, size_t n) 
{ 
    if (isShared(a)) { 
     // a is known to be in shared memory, 
     // so use it directly 
    } else { 
     // make a copy of a in shared memory 
     // and use the copy 
    } 
} 

Answer 2

_{这是@万亿的answer的推广。}

使用is_in_shared_memory()从下面的代码，这对于所有可能的设备上的存储空间定义了类似的功能：

#ifndef STRINGIFY 
#define STRINGIFY(_q) #_q 
#endif 

#define IS_IN_MEMORY_SPACE(_which_space) \ 
__forceinline__ __device__ int is_in_ ## _which_space ## _memory (const void *ptr) \ 
{ \ 
    int result; \ 
    asm ("{" \ 
     ".reg .pred p;\n\t" \ 
     "isspacep." STRINGIFY(_which_space) " p, %1;\n\t" \ 
     "selp.b32 %0, 1, 0, p;\n\t" \ 
     "}" \ 
     : "=r"(result) : "l"(ptr)); \ 
    return result; \ 
} 

IS_IN_MEMORY_SPACE(const) 
IS_IN_MEMORY_SPACE(global) 
IS_IN_MEMORY_SPACE(local) 
IS_IN_MEMORY_SPACE(shared) 

#undef IS_IN_MEMORY_SPACE 

如果你正在构建32位代码，更换"l"约束（一64位地址）与"r"。