Могу ли я проверить, находится ли адрес в общей памяти?

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 часто может знать, используются ли они с общей памятью или нет.

1

c++cuda gpu-shared-memory

Решение

Другие решения

_{Это обобщение @ tera’s ответ.}

использование 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",

0

Источник

Accepted Answer

Вы можете использовать isspacep Инструкция PTX через немного встроенной «сборки»:

// 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
}
}

5