Заставить CUDA использовать регистр для переменной

У меня есть много неиспользуемых регистров в моем ядре. Я бы хотел сказать CUDA использовать несколько регистров для хранения некоторых данных, а не выполнять глобальное чтение данных каждый раз, когда мне это нужно. (Я не могу использовать общий мем.)

__global__ void simple(float *gData) {
float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

скомпилировать w/: nvcc -arch sm_20 --ptxas-options=-v simple.cu, и я получаю
Кадр стека 0 байт, хранилища разлива 0 байт, разливы 0 байт
Использовано 2 регистра, 40 байт cmem[0]

__global__ void simple(float *gData) {
register float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

объявлениерегистра ничего не делает.
Кадр стека 0 байт, хранилища разлива 0 байт, разливы 0 байт
Использовано 2 регистра, 40 байт cmem[0]

__global__ void simple(float *gData) {
volatile float rData[1024];
for(int i=0; i<1024; i++) {
  rData[i]=gData[i];
  }
// work on the data here
}

volatile объявление создает стековое хранилище:
Кадр стека 4096 байт, хранилища разлива 0 байт, разливы 0 байт
Использовано 21 резистор, 40 байт cmem[0]

1) Есть ли простой способ сказать компилятору использовать регистровое пространство для переменной?
2) Где находится стековый фрейм: регистр, глобальная память, локальная память,...? Что такое кадр стека? (С каких пор у GPU есть стек? Виртуальный стек?)
3) Файл simple.ptx в основном пустой: (nvcc -arch sm_20 -ptx simple.cu)

.loc 2 14 2
ret;

Любая идея, где я могу найти реальный компьютер / скомпилированный код?

Источник

2 ответа

Решение
  • Динамически индексированные массивы не могут быть сохранены в регистрах, поскольку регистровый файл графического процессора не является динамически адресуемым.
  • Скалярные переменные автоматически сохраняются в регистрах компилятором.
  • Статически индексируемый (т. Е. Где индекс может быть определен во время компиляции), небольшие массивы (скажем, менее 16 с плавающей запятой) могут храниться в регистрах компилятором.

Графические процессоры SM 2.0 (Fermi) поддерживают только до 63 регистров на поток. Если это будет превышено, значения регистров будут вылиты / заполнены из локальной (вне чиповой) памяти, поддерживаемой иерархией кэша. Графические процессоры SM 3.5 расширяют это до 255 регистров на поток.

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

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

Пример, который вы приводите, может быть примером совместной памяти, если:

  1. Многие потоки в блоке используют одни и те же данные, или
  2. Размер массива для каждого потока достаточно мал, чтобы выделить достаточно места для всех потоков в многопоточных блоках (1024 операций с плавающей точкой на поток - это очень много).

Как уже упоминалось в njuffa, причина, по которой ваше ядро ​​использует только 2 регистра, состоит в том, что вы ничего не делаете с данными в ядре, и мертвый код был полностью исключен компилятором.

Источник

Как уже отмечалось, регистры (и PTX "пространство параметров") не могут быть проиндексированы динамически. Чтобы сделать это, компилятор должен выдавать код как для switch...case блок, чтобы превратить динамический индекс в немедленный. Я не уверен, что это когда-либо происходит автоматически. Вы можете помочь этому, используя структуру кортежа фиксированного размера и switch...case, Скорее всего, метапрограммирование на C/C++ будет лучшим выбором для того, чтобы код был таким управляемым.

Также для CUDA 4.0 используйте переключатель командной строки -Xopencc=-O3 чтобы все, кроме простых скаляров (таких как структуры данных) отображалось в регистры (см. этот пост). Для CUDA > 4.0 вы должны отключить поддержку отладки (нет -G опция командной строки - оптимизация происходит только при отключенной отладке).

Уровень PTX позволяет гораздо больше виртуальных регистров, чем аппаратное обеспечение. Они отображаются на аппаратные регистры во время загрузки. Указанный предел регистра позволяет установить верхний предел для аппаратных ресурсов, используемых сгенерированным двоичным файлом. Он служит эвристикой для компилятора, который решает, когда следует пролить регистры (см. Ниже) при компиляции в PTX, так что определенные требования параллелизма могут быть удовлетворены (см. "Границы запуска", "заполнение" и "параллельное выполнение ядра" в документации CUDA). - Вам также может понравиться эта самая интересная презентация).

Для графических процессоров Fermi существует не более 64 аппаратных регистров. 64-й (или последний - при использовании меньше, чем аппаратный максимум) используется ABI в качестве указателя стека и, таким образом, для "проливания регистров" (это означает освобождение регистров путем временного сохранения их значений в стеке и происходит, когда больше регистров нужны, чем доступны) так что это неприкосновенно.

Источник
Другие вопросы по тегам