我可以检查一个地址是否在共享内存中吗?
Can I check whether an address is in shared memory?
我想编写以下 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
通常可以知道它们是否被使用是否共享内存。
您可以通过一些内联 "assembly":
使用 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
}
}
这是@tera 的概括。
使用以下代码中的 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"
。
我想编写以下 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
通常可以知道它们是否被使用是否共享内存。
您可以通过一些内联 "assembly":
使用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
}
}
这是@tera
使用以下代码中的 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"
。