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

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.

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

__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 все мимо значения к каждому блоку либо совместно используемой памяти или постоянной памяти (в зависимости от архитектуры GPU) автоматически , Нет никакого преимущества в том, что вы делаете.