Стоит ли передавать параметры ядра через разделяемую память?

Question

Предположим, что у нас есть массив int * dataкаждый поток получит доступ к одному элементу этого массива. Поскольку этот массив будет общим для всех потоков, он будет сохранен в глобальной памяти.

Давайте создадим тестовое ядро:

 __global__ void test(int *data, int a, int b, int c){ ... }

Я точно знаю, что data массив будет в глобальной памяти, потому что я выделил память для этого массива, используя cudaMalloc, Теперь, что касается других переменных, я видел несколько примеров, которые передают целое число без выделения памяти непосредственно в функцию ядра. В моем случае такие переменные ab а также c,

Если я не ошибаюсь, хотя мы не звоним напрямую cudaMalloc Чтобы выделить 4 байта для каждых трех целых чисел, CUDA автоматически сделает это за нас, поэтому в конце переменные ab а также c будет выделяться в глобальной памяти.

Теперь эти переменные являются только вспомогательными, потоки только читают их и ничего больше.

Мой вопрос: не лучше ли перенести эти переменные в общую память?

Я полагаю, что если бы мы имели, например, 10 блоки с 1024 нам нужны темы 10*3 = 30 читает о 4 байты для хранения номеров в общей памяти каждого блока.

Без разделяемой памяти и если каждый поток должен прочитать все эти три переменные один раз, общий объем чтения глобальной памяти будет 1024*10*3 = 30720 что очень неэффективно.

Теперь вот проблема, я немного новичок в CUDA, и я не уверен, возможно ли передать память для переменных ab а также c в общую память каждого блока, без того, чтобы каждый поток считывал эти переменные из глобальной памяти и загружал их в общую память, так что в итоге общий объем чтения глобальной памяти был бы 1024*10*3 = 30720 и не 10*3 = 30,

На следующем сайте есть этот пример:

 __global__ void staticReverse(int *d, int n)
 {
    __shared__ int s[64];
    int t = threadIdx.x;
    int tr = n-t-1;
    s[t] = d[t];
    __syncthreads();
   d[t] = s[tr];
 }

Здесь каждый поток загружает разные данные внутри общей переменной s, Таким образом, каждый поток, в соответствии с его индексом, загружает указанные данные в общую память.

В моем случае я хочу загружать только переменные ab а также c в общую память. Эти переменные всегда одинаковы, они не меняются, поэтому они не имеют никакого отношения к самим потокам, они являются вспомогательными и используются каждым потоком для запуска какого-либо алгоритма.

Как мне подойти к этой проблеме? Можно ли добиться этого, только делая total_amount_of_blocks*3 глобальная память читает?

Answer 1

Среда выполнения GPU уже делает это оптимально без необходимости что-либо делать (и ваше предположение о том, как передача аргументов работает в CUDA, неверно). Вот что сейчас происходит:

• В устройствах с возможностями вычислений 1.0 / 1.1 / 1.2 / 1.3 аргументы ядра передаются средой выполнения в разделяемой памяти.
• В вычислительных возможностях 2.x / 3.x / 4.x / 5.x / 6.x аргументы ядра передаются средой выполнения в зарезервированный банк постоянной памяти (который имеет выделенный кэш с широковещательной передачей).

Так в вашем гипотетическом ядре

__global__ void test(int *data, int a, int b, int c){ ... }

data, a, b, а также c все они передаются по значению каждому блоку в общей памяти или в постоянной памяти (в зависимости от архитектуры графического процессора) автоматически. Нет никакого преимущества в том, что вы предлагаете.