1) Когда ядро начинает проливать регистры в локальную память?
Это полностью контролируется компилятором. Это не выполняется во время выполнения, и нет никаких динамических c решений во время выполнения об этом. Когда ваш код достигает точки разлива, это означает, что компилятор вставил команду вроде:
STL [R0], R1
В этом случае R1 сохраняется в локальной памяти, адрес локальной памяти задается в R0. Это был бы магазин разливов. (После этой инструкции R1 может использоваться для / загружаться с чем-то еще.) Конечно, компилятор знает, когда он это сделал, и поэтому может сообщать о количестве загрузок и хранилищах, которые он выбрал для использования / делать. Вы можете получить эту информацию (наряду с использованием регистров и другую информацию), используя переключатель компилятора -Xptxas=-v.
Компилятор (если вы не ограничиваете его, см. Ниже) принимает решения об использовании регистров, в основном ориентированные на производительность, обращая меньше внимания на то, сколько регистров фактически используется. Первым приоритетом является производительность.
2) Когда регистров недостаточно, как среда выполнения CUDA решает не запускать ядро и выдает слишком много ресурсов по запросу об ошибке? Сколько регистров достаточно для запуска ядра?
Во время компиляции, когда код вашего ядра компилируется, компилятор не знает, как он будет запущен. Он не имеет представления о том, на что будет похожа ваша конфигурация запуска (количество блоков, количество потоков в блоке, объем динамически выделяемой общей памяти и т. Д. c). Фактически процесс компиляции в основном происходит так, как если бы компилируемая вещь была единой thread.
Во время компиляции компилятор принимает ряд статических c решений о назначении регистров (как и где регистры будут использоваться). CUDA имеет бинарные утилиты , которые могут помочь в понимании этого. Назначения регистров не изменяются во время выполнения, никоим образом не являются динамическими c и, следовательно, полностью определяются во время компиляции. Следовательно, по завершении компиляции для данной функции кода устройства, как правило, можно определить, сколько регистров необходимо. Компилятор включает эту информацию в двоичный скомпилированный объект.
Во время выполнения, в момент запуска ядра, среда выполнения CUDA теперь знает:
- Сколько регистров (на поток) необходимо для данного ядра
- На каком устройстве мы работаем и, следовательно, какие совокупные ограничения
- Что такое конфигурация запуска (блоки, потоки)
Сборка этих 3-х частей информации означает, что среда выполнения может сразу узнать, имеется или будет достаточно «регистрационного пространства» для запуска. Грубо говоря, арифметика прохождения / неудачи c заключается в том, что запуск удовлетворял бы этому неравенству:
registers_per_thread*threads_per_block <= max_registers_per_multiprocessor
В этом уравнении также необходимо учитывать гранулярность. Регистры часто распределяются группами по 2 или 4 во время выполнения, т. Е. Количество registers_per_thread может потребоваться округлить до следующего целого числа, кратного чему-то вроде 2 или 4, до применения теста на неравенство. Количество registers_per_thread устанавливается компилятором, как уже описано. Количество threads_per_block зависит от конфигурации запуска вашего ядра. Количество max_registers_per_multiprocessor является машиночитаемым (т.е. это функция графического процессора, на котором вы работаете). Вы можете узнать, как получить это количество самостоятельно, если вы будете sh, изучив пример кода deviceQuery CUDA.
3) Поскольку существует механизм различий регистров, не все ядра CUDA должны быть установлены запускаться, даже если регистров недостаточно?
Я повторяю, что присвоение регистров (и решения о разливе регистров) является / является полностью устаревшим процессом компиляции c. Не принимается никаких решений или изменений во время выполнения. Присвоение регистра полностью проверяется из скомпилированного кода. Поэтому, поскольку никакие корректировки не могут быть сделаны во время выполнения, никакие изменения не могут быть сделаны, чтобы позволить произвольный запуск. Любое такое изменение потребует перекомпиляции кода. Хотя это теоретически возможно, в настоящее время это не реализовано в CUDA. Кроме того, он может приводить как к переменным, так и, возможно, к непредсказуемому поведению (по производительности), поэтому могут быть причины этого не делать.
Возможно сделать все ядра "запускаемыми" (с учетом ограничений регистра). ) соответствующим образом ограничивая выбор компилятора в отношении назначения регистров. __launch_bounds__ и переключатель компилятора -maxrregcount - это два способа добиться этого. CUDA предоставляет как калькулятор занятости , так и API занятости , чтобы помочь с этим процессом.