Блог

Главная - Вопросы по программированию - Ядро CUDA нужно ли указывать const во всех параметрах передачи по значению?

Ядро CUDA нужно ли указывать const во всех параметрах передачи по значению?

#c #cuda #constants #compiler-optimization #gpu-constant-memory

#c #cuda #константы #оптимизация компилятора #gpu-постоянная память

Вопрос:

 __global__
void diff2D(int top, int left, (const int windowSize or int windowSize), float* a1, float* a2)
{
    (const int mid or int mid) = windowsize / 2;
    ... // templateSize, windowSize are all const. 
}

Мой вопрос в том, нужно ли мне явно использовать ключевое слово «const int windowSize», чтобы компилятор помещался mid в постоянную память / кэш? Я не хочу ставить «const» перед примитивными типами передачи по значению из-за стандартов кодирования.

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

Ответ №1:

Все параметры ядра любого типа передаются через постоянную память. Смотрите здесь . Тип или оформление параметров не имеет значения для этой функции.

Никакие другие константы любого типа не помещаются в постоянную память автоматически, если вы не используете соответствующее __constant__ оформление / использование.

Вне параметров ядра, указывающих:

 const int mid = windowsize/2;

не приведет к тому, что что-либо будет помещено в постоянную память.

Один из способов думать о постоянной памяти заключается в том, что она находится в глобальном масштабе. Это имеет смысл как для использования параметров ядра (все потоки могут получить к ним доступ), так __constant__ и для объявлений (все потоки могут получить к ним доступ).

Когда вы объявляете обычную локальную переменную:

 const int mid = windowsize/2;

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

Логическим продолжением этого является то, что вы не можете сделать что-то вроде:

 __constant__ int mid = windowsize/2;

посреди вашего кода ядра внезапно объявить переменную в __constant__ памяти. Эти объявления должны быть сделаны в глобальной области / file .

Комментарии:

1. Итак, если я хочу, чтобы «mid» был помещен в постоянную память вместо register , я должен передать его как параметр, правильно ли это?

2. @DukeLe вы пытаетесь уменьшить использование регистра?

3. ДА. Я не знаю, стоит ли помещать часто используемые переменные (например, количество циклов, высоту / ширину фильтра и т. Д.) В постоянную память (вместо регистра) Или нет.

4. Неразумно говорить «вместо register». По большей части инструкции GPU SASS извлекают операнды из регистров . Данные в памяти перемещаются в регистр, прежде чем вы сможете сделать с ним что-нибудь полезное. На самом деле есть несколько исключений из этого, в частности, для операндов, извлекаемых непосредственно из постоянной памяти, или для постоянных операндов, извлекаемых непосредственно из потока команд, но IMO не имеет смысла пытаться беспокоиться об этом. По сути, вы не можете контролировать, будет ли операнд, который начал жизнь в постоянной памяти, в конечном итоге использовать регистр или нет. Очень часто так и будет.

Метки: