Принуждение CUDA к использованию регистра для переменной

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

__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 resisters, 40 байт cmem [0]

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

 .loc 2 14 2 ret; 

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

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

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

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

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

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

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

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

Как уже отмечалось, регистры (и PTX «param space») не могут индексироваться динамически. Для этого компилятор должен будет испустить код, как для блока switch...case block, чтобы сразу включить динамический индекс. Я не уверен, что это когда-либо происходит автоматически. Вы можете помочь ему, используя структуру кортежей с фиксированным размером и switch...case . Метапрограммирование C / C ++, вероятно, будет лучшим выбором для того, чтобы держать код таким, как этот.

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

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

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

  • Упор внутри написанных пользователем ядер
  • Копирование данных в структуру данных «cufftComplex»?
  • OpenCV 2.4.3rc и CUDA 4.2: «Ошибка OpenCV: поддержка GPU не поддерживается»
  • Проблемы с запуском ядер CUDA из статического кода инициализации
  • Как начать новый проект CUDA в Visual Studio 2008?
  • Лучший подход для GPGPU / CUDA / OpenCL в Java?
  • Как измерить внутреннее время ядра в NVIDIA CUDA?
  • Изменение реестра для увеличения тайм-аута графического процессора, Windows 7
  • CUDA несовместим с моей версией gcc
  • CUDA определяет streamи на блок, блоки на каждую сетку
  • Сообщение об ошибке: не удается найти или открыть файл PDB
  • Давайте будем гением компьютера.