CUDA константа памяти

Question

У меня есть массив в постоянной памяти (это глобальная переменная) и получил ссылку на него с помощью вызова функции cudaGetSymbolAddress. Мое ядро ​​работает медленно, когда я использую эту ссылку для извлечения постоянных данных, а не с помощью глобальной переменной. Что является причиной этого?

__constant__ int g[2] = {1,2};
// __device__ int g[2] = {1,2};

// kernel: use by reference
__global__ void add_1( int *a, int *b, int *c, int *f )
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = f[0] * a[tid] + f[1] * b[tid];
}

// kernel: use global variable
__global__ void add_2( int *a, int *b, int *c, int *f )
{
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    c[tid] = g[0] * a[tid] + f[1] * b[tid];
}

int main()
{
    ......
    // a,b,c are large arrays in device memory of size 40960.

    int *f;
    cudaGetSymbolAddress( (void **)&f, (char *)&g);

    add_1 <<< 160, 256 >>> ( a, b, c, f );

    ......
}

Это пример кода, и все потоки в варпе загружаются в одно и то же место одновременно. Код с комментариями осуществляется путем прямого доступа к постоянной памяти

Объяснение того, почему постоянная кэш-память не используется (talonmies)

Причина в отсутствии постоянного кеша. Кэшированный доступ происходит только тогда, когда компилятор выдает определенную инструкцию PTX (ld.const) для переменной, явно помеченной как находящаяся в пространстве постоянных состояний. И компилятор знает, как это сделать, когда объявлена ​​переменная __constant__ - это статический атрибут времени компиляции, который влияет на генерацию кода. Тот же процесс не может произойти во время выполнения.

Если вы передадите указатель в глобальную память, а компилятор не сможет определить, что указатель находится в пространстве постоянных состояний, он не сгенерирует правильный PTX для доступа к этой памяти через постоянный кэш. В результате доступ будет медленнее.

Неотвеченный вопрос

Почему даже когда массив g объявлен как __device__ переменная, код медленнее, когда ссылка на него используется. Видя PTX код для загрузки глобальной памяти в регистры:

• 2 инструкции ld.global.s32 используются, который загружает 4 байта в регистр. (в коде с использованием ссылки)
• 1 инструкция ld.global.v2.s32 используется, который загружает 8 байтов в 2 регистра (в коде с использованием глобальной переменной)

В чем разница, и любая ссылка на документацию будет оценена?

Answer 1

В отличие от глобальной памяти, доступ к постоянной памяти будет сериализован (разделен на несколько транзакций), если они будут неоднородными (все потоки деформации (наполовину для вычислительной возможности 1.x) обращаются к одному и тому же адресу).

Поэтому используйте постоянную память только в том случае, если доступы будут одинаковыми.