Говоря об оптимизации CUDA

Добавить Автора

Перевод: как если бы был свет

Предисловие:

Несколько месяцев назад я написал свой собственный автоэнкодер на основе статьи Симончелли 2016 года для исследовательских целей. Вначале я хотел использовать для своих экспериментов некоторые популярные фреймворки глубокого обучения, такие как TensorFlow, Caffe2 или MXNet. Однако, изучив все эти фреймворки в течение нескольких недель, я обнаружил очень неприятную проблему — масштабируемость. Я не говорю, что эти фреймворки плохо спроектированы, но что пользователям запрещено разрабатывать сторонние операторы, так же, как и при написании плагина, вы даете мне функцию без каких-либо параметров. Тогда единственный способ изменить поведение функции — изменить исходный код, что, несомненно, является огромным проектом из-за плохой организации документации. (Кажется, это обычная проблема с программным обеспечением с открытым исходным кодом.) Итак, поскольку необычный оператор GDN не включен во все эти фреймворки, разработка новой фреймворка кажется единственным решением.

Обратите внимание, сосредоточьтесь на технической сводке и обмене компьютерным зрением

GDN

Этот оператор является основной нелинейной функцией в теории и выражается следующим образом (формула не имеет значения, если вам не нравятся чертовы символы, вы можете просто пропустить этот раздел):

Верхние индексы (k) и (k+1) представляют количество слоев, w и u — вход и выход многоканального изображения, а нижний индекс i — количество каналов. β и γ — параметры, которые я хочу тренировать. Предположим, у нас есть N каналов, тогда γ — матрица размера N × N, а β — вектор размера N × 1. На первый взгляд, эта функциональность очень похожа на пакетную нормализацию (BN) или локальную нормализацию отклика (LRN), которые хорошо поддерживаются cudnn и всеми фреймворками глубокого обучения. Но поверьте мне, не позволяйте своим глазам обмануть вас. Это совсем другое. (Обратите внимание, что большое деление поэлементно.)

Вперед не потребляет много вычислительной мощности, а назад потребляет большую часть мощности моего графического процессора. Теперь давайте посмотрим на спину. Мне нужно вычислить 3 градиента, ∇β, ∇γ и ∇u.

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

Во-первых, легко заметить, что входные данные можно рассматривать как вектор длины m x n. Во-вторых, (blabla...)^(-3/2) появляется во всех этих градиентах. Это означает, что мы можем просто вычислить термин 1 раз и кэшировать его для последующего использования. Назовем это матрицей "(blabla...)^(-1/2)" D . Наконец, δ — это ошибка, распространяющаяся на предыдущий уровень.

Fig 1. Computation of

После некоторого упрощения становится понятнее, не так ли? Я знаю, что некоторые объяснения все еще необходимы. В правой части уравнения каждый прямоугольник представляет собой вектор, составленный из матриц, упомянутых выше. D — это член знаменателя в формуле контекстно-медийной сети. Помните «(blabla...)^(-1/2)», о котором мы только что упоминали?

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

1. Не только наивные алгоритмы.

Я называю этот метод «не просто наивным», потому что это первый метод, который я когда-либо использовал. Даже с изображением небольшого размера в качестве входных данных он израсходовал почти всю память моего графического процессора и показал самую низкую производительность. Не используя преимущества повторного использования памяти, я просто продублировал все эти маленькие прямоугольники по вертикали и горизонтали, чтобы получить большую матрицу, как показано на изображении ниже, и запустил множество одномерных организованных ядер. Затем добавьте их.

Fig 2. Less than naive Algo.

Единственным преимуществом этого алгоритма является то, что индекс не нужно вычислять в каждом потоке CUDA, поскольку идентификатор потока является лишь единственным соответствующим индексом памяти. Итак, все, что вам нужно сделать, это немного умножить, а затем использовать cublas, чтобы добавить скалярное произведение каждого маленького цветного прямоугольника к вектору 1 (вектор, заполненный всеми единицами). Но, как видите, размер прямоугольника не такой маленький, как то, что я нарисовал здесь, размер такой же, как на изображении. Для каждого вектора в этом изображении размер будет равен N x N x imageSize x batchSize. Очевидно, мы потратили впустую (N-1) x N x imageSize x batchSize x 4 байта, не говоря уже о времени, потраченном впустую на доступ ко всей этой избыточной глобальной памяти.

2. Наивный алгоритм.

Для первого алгоритма я смог обучить свою сеть менее чем на 4 изображениях размером 128 x 128 за итерацию почти за 2 секунды. (Мой графический процессор — GTX 1080.) Эта реальность заставила меня улучшить свой алгоритм, иначе мне пришлось бы ждать почти 2 месяца, чтобы получить свои результаты.

Поскольку количество ядер, которые мне нужно запустить, определенно намного больше, чем ядер CUDA в моем графическом процессоре, драйвер cuda будет сериализовать эти задачи независимо от того, какой метод я использую. Тогда я решил не копировать все эти воспоминания. Вместо этого я запущу N x 1D организованных ядер N x imageSize N раз (N — общее количество каналов).

Fig 3. Without memory replication

Как видите, улучшение налицо. Потому что нам больше не нужно много копировать данные. Доступ к глобальной памяти в графических процессорах очень дорог. Шаблон доступа к памяти также прост, потому что, когда вы получаете идентификатор потока, вы можете получить индекс памяти с помощью всего одной операции модификации (индекс памяти = идентификатор потока % imageSize). Однако при таком подходе, поскольку ядра по-прежнему организованы одномерно, и мы используем цикл for для их запуска, мы можем не получить выгоды от более умного алгоритма планирования графического процессора, хотя я попробовал кровь. Теперь, с этим небольшим изменением 2 месяца обучения можно сократить почти до 2 недель.

3. Умные алгоритмы организации.

До сих пор я не рассматривал возможности разделяемой памяти, потому что для меня проектирование хорошего режима ядра обычно скучно и вызывает головную боль. Очевидно, что режим ядра 1D — это самый простой код для написания. Однако более высокая производительность требует более тщательного проектирования. К моему удивлению, алгоритм в этом разделе в 3 раза быстрее второго алгоритма.

Возвращаясь к рисунку 1, можно увидеть, что первые строки δ0, w0 и D0 первых трех правых матриц одинаковы. Таким образом, мы можем вычислить строку γ в блоке, для каждого блока мы можем запустить потоки imageSize, и для каждого потока мы можем вычислить все каналы, используя цикл for.

Fig 5. Computation in one block

Таким образом, на рисунке 5 довольно интуитивно понятно поместить δ0, w0 и D0 в общую память, в то время как для потока i он считывает один пиксель в N каналах от 0 до N-1 с δ0, w0 и D0. Умножение и совместное использование памяти. Псевдокод выглядит следующим образом:

blockId = blockIdx.x; 
threadId = threadIdx.x;shareDelta <- delta[blockId];  
shareW <- W[blockId];
shareD <- D[blockId];
_synchronize();for(i = 0; i < N-1; i++)
{
   result[threadIdx i*imgSize] = shareDelta[threadId] *
                                 shareW[threadId] *
                                 shareD[threadId] * 
                                 W[threadId + i*imgSize];
}

Алгоритм 2 предпочитает вычисление по строкам вместо вычислений по столбцам, потому что при вычислении строки в сетке мы можем совместно использовать 3 вектора δ0, w0 и D0. Но если мы вычисляем столбец, как в Algo, мы можем использовать только 1 вектор w0. (Снова см. рис. 1).

В этом фрагменте кода нет блока if...else.... Это очень важно при параллельных вычислениях. Поскольку все потоки выполняются параллельно, в идеале все эти потоки выполняют свою работу одновременно. Но если есть блокировка if... else..., форк позволяет этим потокам выполнять разные задачи, чтобы они заканчивались в разное время. Тогда время вычислений будет определяться самым медленным потоком.

Безиндексные вычисления также являются преимуществом. При разработке одномерной схемы мы должны использовать идентификатор потока для вычисления индекса памяти, но здесь нам не нужно преобразовывать blockId и threadId в одномерный индекс памяти для доступа к данным.

Наконец, поскольку мои данные хранятся в основном столбце, это означает, что, как и в случае с вектором δ0, все элементы в этом векторе хранятся непрерывно. Таким образом, он выигрывает от механизма объединения глобальной памяти. Глобальная память также является важным понятием в cuda.

С аппаратной стороны 16 ядер cuda организованы в виде варпа. Когда один из потоков обращается к данным, таким как a1 на приведенном выше рисунке, шина данных не только передает a1, но также передает a1~a32 в кэш, чтобы ускорить доступ к данным других 15 ядер. Поэтому, когда я читаю глобальные данные для совместного использования памяти, я читаю только каждые 32 байта, все остальные байты считываются из кеша, что в сотни раз быстрее. Благодаря теории пространственно-временной локальности.

4. Еще немного улучшений

Сегодня я вдруг обнаружил, что мне не нужна разделяемая память, но можно использовать константную память. Потому что каждому потоку в блоке нужно получить доступ к векторам δ0, w0 и D0 только один раз. Таким образом, перед циклом for мы можем кэшировать элемент в константной памяти. Еще одна прелесть заключается в том, что, поскольку каждый поток обращается только к одному элементу, синхронизация потоков не требуется.

код показывает, как показано ниже:

blockId = blockIdx.x; 
threadId = threadIdx.x;const float constDelta = delta[blockId * imgSize + threadId];  
const float constW = W[blockId * imgSize + threadId];
const float constD = D[blockId * imgSize + threadId];for(i = 0; i < N-1; i++)
{
   result[threadIdx + i*imgSize] = constDelta * constW *
                                   constD * 
                                   W[threadId + i*imgSize];
}

Как видно из приведенного выше кода, constDelta, constW, constD можно N раз переиспользовать из локальной памяти, которая всегда хранится в локальных регистрах. Следовательно, пропускная способность больше, чем общая память.

Reduce Operation

Все алгоритмы, о которых я говорю, не завершены, потому что все, что я получаю из приведенных выше алгоритмов, на самом деле представляет собой необработанные гаммы, подобные этой:

Мне нужно накопить каждый вектор слева, чтобы получить один элемент. Первый вариант — API cublas, cublasSsbmv. Эта функция будет выполнять умножение матрицы на вектор. Таким образом, мы можем думать о векторе слева как о матрице и умножить его на вектор, состоящий из всех единиц, чтобы получить однострочный градиент γ. и повторить N раз, чтобы получить окончательный результат. Но я заметил, что есть и другие API cublasSgemmBatched. Эта функция может выполнять пакетное умножение матрицы на вектор. Затем я провел эксперимент, чтобы проверить, что быстрее:

Цикл for для N умножений матрицы на вектор против пакетного умножения матрицы на вектор.

Оказывается, цикл for намного быстрее. Но я не знаю, в чем причина, может быть потому, что у меня здесь N слишком мало (N=256).

Я не буду показывать, как вычислять ∇β и ∇u, потому что они аналогичны ∇γ. Я знаю, что должна быть дальнейшая оптимизация или лучший дизайн, чем у меня. Оптимизация CUDA часто сложна для людей без глубокого понимания организации графического процессора. Программисты, знакомые с процессорами, всегда выигрывают от современных операционных систем и мощных компиляторов. Тем не менее, графические процессоры довольно отличаются и сложны, чем центральные процессоры, с точки зрения написания достаточного количества кода, хотя это намного удобнее, чем раньше, использовать для вычислений графические шейдеры. На улучшение экологической обстановки уйдет несколько лет.

Оригинальная ссылка:

medium.com/@lawheadhunter032…

Эта статья взята из серии публикаций в Техническом руководстве CV для общедоступных аккаунтов.

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

Ответьте на ключевое слово «Техническое резюме» в официальном аккаунте, чтобы получить резюме следующих статей в формате pdf.