Ответ 1
Массивы, локальная память и регистры
Здесь есть неправильное представление об определении "локальной памяти". "Локальная память" в CUDA на самом деле является глобальной памятью (и на самом деле ее следует называть "глобальной локальной памятью потока") с чередованной адресацией (что делает итерацию по массиву параллельно немного быстрее, чем блокировка данных каждого потока). Если вы хотите, чтобы все было действительно быстро, вы хотите использовать либо разделяемую память, либо, что еще лучше, регистры (особенно на последних устройствах, где вы получаете до 255 регистров на поток). Объяснение всей иерархии памяти CUDA выходит за рамки этого поста. Вместо этого давайте сконцентрируемся на быстром вычислении небольших массивов.
Небольшие массивы, как и переменные, могут храниться целиком в регистрах. Однако на современном оборудовании NVIDIA размещение массивов в регистрах затруднительно. Зачем? Потому что регистры требуют очень бережного отношения. Если вы сделаете это не совсем правильно, ваши данные окажутся в локальной памяти (которая, опять же, действительно является глобальной памятью, которая является самой медленной из имеющихся у вас). Руководство по программированию CUDA, раздел 5.3.2, сообщает вам, когда используется локальная память:
Локальная память
Доступ к локальной памяти происходит только для некоторых автоматических переменных, как указано в Спецификаторах типов переменных. Автоматические переменные, которые компилятор может поместить в локальную память:
- Массивы, для которых он не может определить, что они проиндексированы с постоянными величинами,
- Большие структуры или массивы, которые будут занимать слишком много места в регистре,
- Любая переменная, если ядро использует больше регистров, чем доступно (это также называется проливом регистров).
Как работает распределение регистра?
Обратите внимание, что распределение регистров - чрезвычайно сложный процесс, поэтому вы не можете (и не должны) вмешиваться в него. Вместо этого компилятор преобразует код CUDA в код PTX (своего рода байт-код), который предполагает наличие машины с бесконечным числом регистров. Вы можете написать встроенный PTX, но это не сделает слишком много, чтобы зарегистрировать распределение. PTX-код - это независимый от устройства код, и это только первый этап. На втором этапе PTX будет скомпилирован в код сборки устройства, называемый SASS. Код SASS имеет фактическое распределение регистров. Компилятор SASS и его оптимизатор также будут в конечном итоге определять, будет ли переменная находиться в регистрах или в локальной памяти. Все, что вы можете сделать, это попытаться понять, что делает компилятор SASS в определенных случаях, и использовать это в своих интересах. В этом может помочь просмотр корреляции кода в Nsight (см. Ниже). Однако, поскольку компилятор и оптимизатор продолжают изменяться, нет никаких гарантий относительно того, что будет или не будет в регистрах.
Недостаточно регистров
Приложение G, раздел 1, сообщает, сколько регистров может иметь поток. Ищите "Максимальное количество 32-битных регистров на поток". Чтобы интерпретировать эту таблицу, вы должны знать свои вычислительные возможности (см. Ниже). Не забывайте, что регистры используются для всех видов вещей, и не просто соотносятся с отдельными переменными. Регистры на всех устройствах до CC 3.5 являются 32-битными каждый. Если компилятор достаточно умен (и компилятор CUDA продолжает изменяться), он может, например, упаковать несколько байтов в один и тот же регистр. Представление корреляции кода Nsight (см. "Анализ доступа к памяти" ниже) также показывает это.
Постоянная и динамическая индексация
Несмотря на то, что ограничение пространства является очевидным препятствием для массивов в реестре, легко отслеживается тот факт, что на текущем оборудовании (Compute Capability 3.x и ниже) компилятор помещает любой массив в локальную память, доступ к которому осуществляется с помощью динамическая индексация. Динамический индекс - это индекс, который компилятор не может понять. Массивы, доступ к которым осуществляется с помощью динамических индексов, не могут быть помещены в регистры, поскольку регистры должны определяться компилятором, и, следовательно, фактический используемый регистр не должен зависеть от значения, определенного во время выполнения. Например, для массива arr
arr[k]
является индексированием констант тогда и только тогда, когда k
является константой или зависит только от констант. Если k
каким-либо образом зависит от некоторого непостоянного значения, компилятор не может вычислить значение k
и вы получили динамическое индексирование. В циклах, где k
начинается и заканчивается с (маленькими) постоянными числами, компилятор (наиболее вероятно) может развернуть ваш цикл и все еще может достичь постоянной индексации.
пример
Например, сортировка небольшого массива может быть выполнена в регистрах, но вы должны использовать сортировку сетей или аналогичные "аппаратные" подходы, и не можете просто использовать стандартный алгоритм, потому что большинство алгоритмов используют динамическую индексацию.
С высокой вероятностью в следующем примере кода компилятор хранит весь массив aBytes
в регистрах, потому что он не слишком велик, и циклы можно полностью развернуть (поскольку цикл повторяется в постоянном диапазоне). Компилятор (очень вероятно) знает, к какому регистру обращаются на каждом этапе, и, таким образом, может полностью хранить его в регистрах. Имейте в виду, что нет никаких гарантий. Лучшее, что вы можете сделать, это проверить его в каждом конкретном случае с помощью инструментов разработчика CUDA, как описано ниже.
__global__
void
testSortingNetwork4(const char * aInput, char * aResult)
{
const int NBytes = 4;
char aBytes[NBytes];
// copy input to local array
for (int i = 0; i < NBytes; ++i)
{
aBytes[i] = aInput[i];
}
// sort using sorting network
CompareAndSwap(aBytes, 0, 2); CompareAndSwap(aBytes, 1, 3);
CompareAndSwap(aBytes, 0, 1); CompareAndSwap(aBytes, 2, 3);
CompareAndSwap(aBytes, 1, 2);
// copy back to result array
for (int i = 0; i < NBytes; ++i)
{
aResult[i] = aBytes[i];
}
}
Анализ доступа к памяти
Как только вы закончите, вы, как правило, хотите проверить, действительно ли данные хранятся в регистрах или они поступили в локальную память. Первое, что вы можете сделать, это указать вашему компилятору предоставить вам статистику памяти, используя --ptxas-options=-v
. Более подробный способ анализа обращений к памяти - использование Nsight.
Nsight имеет много интересных функций. Nsight для Visual Studio имеет встроенный профилировщик и представление корреляции кода CUDA <-> SASS. Функция объясняется здесь. Обратите внимание, что версии Nsight для разных IDE, вероятно, разрабатываются независимо, и, следовательно, их функции могут различаться в разных реализациях.
Если вы будете следовать инструкциям в приведенной выше ссылке (не забудьте добавить соответствующие флаги при компиляции!), Вы можете найти кнопку "Операции с памятью CUDA" в самом низу нижнего меню. В этом представлении вы хотите обнаружить, что нет транзакции памяти, исходящей из строк, которые работают только с соответствующим массивом (например, строки CompareAndSwap в моем примере кода). Потому что, если он не сообщает о доступе к памяти для этих строк, вы (очень вероятно) смогли бы сохранить все вычисления в регистрах и могли бы просто получить ускорение в тысячи, если не в десятки тысяч процентов (вы также можете захотеть проверьте фактическое увеличение скорости, вы выходите из этого!).
Вычисление вычислительных возможностей
Чтобы выяснить, сколько регистров у вас есть, вам нужно знать возможности вашего устройства для вычислений. Стандартный способ получения такой информации об устройстве - запуск образца запроса устройства. Для CUDA 5.5 в 64-разрядной версии Windows, которая по умолчанию находится в каталоге C:\ProgramData\NVIDIA Corporation\CUDA Samples\v5.5\Bin\win64\Release\deviceQuery.exe (В Windows окно консоли будет закрыто немедленно, вы можете захотеть сначала открыть cmd
и запустить его оттуда). Он имеет аналогичное расположение в Linux и MAC.
Если у вас есть Nsight для Visual Studio, просто перейдите в Nsight → Windows → Информация о системе.
Не оптимизировать рано
Я делюсь этим сегодня, потому что совсем недавно столкнулся с этой проблемой. Однако, как уже упоминалось в этой теме, принудительное включение данных в регистры - это далеко не первый шаг, который вы хотите предпринять. Сначала убедитесь, что вы действительно понимаете, что происходит, затем шаг за шагом подойдите к проблеме. Глядя на ассемблерный код, безусловно, хороший шаг, но обычно он не должен быть вашим первым. Если вы новичок в CUDA, руководство по рекомендациям CUDA поможет вам разобраться в некоторых из этих шагов.