我想编写以下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通常可以知道它们是否与共享内存一起使用或不。

最佳答案

您可以通过一些内联“汇编”使用 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;
}

所以你的例子变成
__device__ 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
     }
}

关于c++ - 我可以检查地址是否在共享内存中?,我们在Stack Overflow上找到一个类似的问题:https://stackoverflow.com/questions/42519766/

10-12 12:21
查看更多