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

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


c++ gpu-shared-memory cuda
person einpoklum    schedule 28.02.2017    source источник

Ответы (2)


arrow_upward
5
arrow_downward

Вы можете использовать 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;
}

поэтому ваш пример становится

__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
     }
}
person tera    schedule 01.03.2017
comment
Отлично :-) Можете ли вы объяснить, зачем мне нужна эта дополнительная специфическая для Win64 вуду в начале? Кроме того, в руководстве по PTX говорится, что у нас есть общие адреса памяти и, предположительно, специфичные для пространства адреса. Как это соотносится с вашим ответом? - person einpoklum; 01.03.2017
comment
Я добавил это вуду на случай, если кто-то попытается скомпилировать в 32-битном режиме. Если вы знаете, что ваш код всегда 64-битный, вы можете просто использовать ограничение "l" для указателя напрямую. - person tera; 01.03.2017
comment
Кроме того, зачем вам селп, если p уже получает либо 0, либо 1? Нельзя ли регистр результата использовать непосредственно как регистр предиката? - person einpoklum; 01.03.2017
comment
Да, его можно было использовать непосредственно в PTX. Но я не знаю способа передать предикат обратно компилятору C, поэтому мы получаем дополнительную нагрузку на другую инструкцию преобразования и должны надеяться, что ptxas снова сможет ее оптимизировать. - person tera; 01.03.2017

arrow_upward
0
arrow_downward

Это обобщение ответа @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".

person einpoklum    schedule 01.03.2017
comment
@tera: Fankly - только скомпилировал, в эту секунду я не на машине с GPU... - person einpoklum; 01.03.2017
comment
@tera: Ты прав, это не сработает. Я сделаю что-нибудь еще. - person einpoklum; 01.03.2017