- Описание архитектуры и принципов реализации
- Построение Self-Attention средствами MQL5
- Организация параллельных вычислений в блоке внимания
- Тестирование механизма внимания
Организация параллельных вычислений в блоке внимания
В предыдущих разделах мы построили рабочий алгоритм блока внимания стандартными средствами языка MQL5. Вы уже можете добавить блок внимания в сою модель и испытать качество работы механизма Self-Attention. Но посмотрите на структуру блока. В его работе мы использовали пять внутренних слоев, создали алгоритм передачи данных между ними в прямом и обратном направлении. И немаловажен тот факт, что каждый элемент последовательности, описанный вектором значений, обрабатывается с использованием единых матриц весовых коэффициентов, но независимо друг от друга. Это позволяет легко распределить операции по параллельным потокам, что даст нам возможность совершать полный набор операций в более короткие временные интервалы. И да, мы изначально определились, что будем создавать библиотеку с возможностью использования двух технологий. Тем самым мы даем пользователю возможность самостоятельно протестировать и выбрать наиболее подходящую технологию для применения в каждом конкретном случае.
Как и ранее, блок параллельных вычислений организуем с использованием технологии OpenCL. Для использования данной технологии нам потребуется выполнить два этапа работы:
- создание программы OpenCL,
- внесение изменений в основную программу.
Код программы OpenCL мы добавим в ранее созданный файл opencl_program.cl. Именно в этом файле мы собрали все кернелы программы OpenCL, используемые в работе предыдущих классов. Для организации работы нашего класса внимания потребуется создание шести кернелов. В этих кернелах нам предстоит организовать поток информации между используемыми внутренними нейронными слоями в прямом и обратном направлении
Первым мы создадим кернел прямого прохода AttentionFeedForward. Кратко напомню последовательность операций при прямом проходе блока Self-Attention:
- Исходные данные подаются на вход трем внутренним сверточным нейронным слоям m_cQuery, m_cKeys, m_cValues.
- Тензоры результатов m_cQuery и m_cKeys перемножаются для получения матрицы зависимостей m_cScores.
- Значения матрицы m_cScores делятся на квадратный корень от размера вектора описания одного элемента последовательности m_cKeys и нормализуются функцией Softmax в разрезе строк (запросов m_cQuery).
- Нормализованная матрица m_cScores перемножается с тензором результатов нейронного слоя m_cValues для получения результатов Self-Attention.
- Результаты блока Self-Attention складываются с исходными данными и нормализуются.
- Полученный тензор служит исходными данными для блока из двух сверточных слоев m_cFF1 и m_cFF2.
Если пункты 1 и 6 покрываются использованием ранее рассмотренного класса сверточных слоев, в котором уже реализован блок многопоточных вычислений, то остальные пункты нам предстоит реализовать в новом кернеле.
Для организации указанных операций нам потребуется передать в кернел шесть буферов данных и два параметра. Чтобы код программы был более читабелен, наименования буферов и переменных будут созвучны с наименованием соответствующих матриц в описании алгоритма.
__kernel void AttentionFeedForward(__global TYPE *querys,
|
Как можно было заметить из описания алгоритма Self-Attention, основной аналитической единицей в данном методе является элемент последовательности, описываемый вектором значений. Для языковых моделей это чаще всего слово. В случае анализа финансовых рынков это бар. Именно между этими элементами последовательности определяются коэффициенты взаимных зависимостей. С учетом этих коэффициентов корректируются значения векторов описания элементов. Поэтому вполне логично разделить операции на потоки именно по элементам последовательности.
Следовательно, в теле кернела первым делом определим анализируемый элемент последовательности по идентификатору нашего потока. В то же время общее количество запущенных потоков подскажет нам количество элементов в последовательности. Здесь же мы сразу определим смещение в тензоре запросов и матрице коэффициентов зависимостей до первого анализируемого значения.
const int q = get_global_id(0);
|
Напомню, что для нормализации данных функцией Softmax нам потребуется сумма экспонент всех нормализуемых значений. Для ее подсчета добавим переменную с начальным нулевым значением.
После проведения подготовительной работы определим значения одного вектора из матрицы коэффициентов зависимостей, относящегося к расчетам зависимостей анализируемого элемента последовательности. Для этого организуем цикл с числом итераций равным количеству элементов в последовательности. В теле цикла мы будем поочередно перемножать вектор Query анализируемого элемента последовательности со всеми векторами тензора Key. Для каждого результата умножения векторов возьмем экспоненциальное значение и запишем его в соответствующий элемент матрицы Score. И конечно прибавим к нашей накопительной сумме всех значений вектора для последующей нормализации значений.
for(int s = 0; s < units; s++)
|
После завершения работы цикла в нашей переменной summ будет накоплена сумма всех элементов нашего вектора из тензора коэффициентов зависимостей. Для завершения нормализации значений данного вектора нам остается лишь разделить значение каждого его элемента на общую сумму всех значений вектора.
for(int s = 0; s < units; s++)
|
Таким образом, в анализируемом векторе мы получили коэффициенты зависимостей анализируемого элемента последовательности от остальных ее элементов. Сумма всех коэффициентов будет равна единице.
Далее, согласно алгоритму, нам необходимо каждый вектор тензора значений Value умножить на соответствующий элемент полученного вектора коэффициентов зависимостей. Полученные векторы нужно сложить. Итоговый вектор значений и будет результатом работы блока Self-Attention.
Перед передачей данных дальше нужно сложить полученные данные с тензором исходных данных и нормализовать их. В теле кернела я предлагаю остановиться на определении результатов работы блока Self-Attention. Сложение же матриц и нормализацию данных в пределах всего нейронного слоя будет эффективнее организовать отдельно.
Посмотрим на реализацию такого решения. Чтобы не пересчитывать на каждой итерации, сначала определим смещение в тензорах исходных данных и результатов. Тензоры имеют одинаковую размерность, поэтому и смещение будет одно для обоих случаев. Затем организуем систему из двух вложенных циклов: во внешнем будем перебирать элементы вектора анализируемого элемента последовательности, а во внутреннем будем осуществлять непосредственно вычисление значений каждого элемента вектора результатов. Для этого количество итераций внутреннего цикла будет равняться количеству элементов последовательности. В теле данного цикла будем умножать значения элементов тензора Value на соответствующие коэффициенты зависимостей из матрицы Score. Полученные произведения будем накапливать в локальной переменной query. После завершения итераций вложенного цикла результат запишем в соответствующий элемент тензора результатов.
shift_query = window * q;
|
На этом мы завершим работу над первым кернелом прямого прохода. Следующим этапом создадим кернел сложения двух тензоров. Подобную работу порой выгоднее сделать с использованием матричных операций на стороне основной программы. Операция несложная, и накладные расходы в виде передачи данных вряд ли будут оправданы. У нас сейчас обратная ситуация. Мы организовываем весь процесс на стороне контекста OpenCL. Вся информация уже находится в памяти контекста, и для реализации операции на стороне основной программы нам потребуется копирование данных. Для вычислений в контексте перенос данных не нужен Поэтому мы создали кернел Sum, в теле которого просто складываем элементы из двух буферов с одним индексом и сохраняем в элемент третьего буфера с тем же индексом.
__kernel void Sum(__global TYPE *inputs1,
|
Процесс нормализации данных имеет более сложную архитектуру. Как вы знаете, ее процесс выражается следующими математическими формулами:
Как можно заметить, для вычисления нормализованного значения каждого элемента последовательности необходимо среднее арифметическое и среднеквадратическое отклонение всей последовательности. Для их вычисления нам потребуется организация передачи данных между отдельными потоками. Данную задачу будем решать аналогично многопоточной реализации функции активации Softmax — через массив в локальной памяти. Только нам потребуется организовать два блока суммирования значений по всему вектору, ведь до определения средней арифметической мы не можем вычислить дисперсию. Кроме того, до определения дисперсии нельзя вычислить нормализованное значение.
Процесс нормализации организован в кернеле LayerNormalize. В параметрах кернел получает указатели на 3 буфера:
- буфер исходных данных,
- буфер результатов,
- буфер для записи параметров среднеквадратического отклонения.
Последний буфер среднеквадратических отклонений нам понадобился для сохранения и передачи данных в кернел обратного прохода.
Кроме того, мы передадим в кернел два параметра: общее количество элементов в нормализуемом буфере и смещение в буфере среднеквадратических отклонений. Напомню, в рамках одного нейронного слоя внимания нормализацию данных мы осуществляем два раза. Нормализуем результаты блоков Self-Attention и FeedForward.
__kernel void LayerNormalize(__global TYPE* inputs,
|
В теле кернела определяем идентификаторы потоков и инициализируем локальный массив данных.
uint i = (uint)get_global_id(0);
|
Первым мы определим среднее арифметическое значение элементов буфера. Для этого организуем цикл, в котором каждый поток просуммирует свои значения и результат сохранит в своем элементе локального массива. Так как мы вычисляем среднее арифметическое всего буфера, то и полученное значение разделим на количество элементов в буфере.
uint count = 0;
|
Синхронизировать работу потоков будем с помощью функции barrier. Так как вычисления потоков не пересекаются, нам достаточно одного барьера в конце блока.
Далее нам необходимо собрать части общей суммы в единое целое. Мы организуем еще один цикл, в теле которого соберем среднее арифметическое буфера в одном элементе локального массива с индексом 0. Результат сохраним в локальную переменную.
count = ls;
|
Хочу еще раз обратить внимание на расстановку барьеров. Здесь нужно уделить особое внимание работе алгоритма, ведь до каждого барьера должны дойти все потоки. Причем последовательность их посещений также должна быть соблюдена.
После определения среднего арифметического повторим циклы и вычислим среднеквадратическое отклонение.
count = 0;
|
count = ls;
|
Полученное среднеквадратическое отклонение сохраним в буфер. Чтобы исключить одновременную запись значения всеми потоками, сохранять будем только в одном потоке. Для этого сделаем проверку индекса потока перед операцией записи значения в буфер.
Теперь, когда мы вычислили значения средних, можно нормализовать исходные данные. При этом надо учесть, что ограничение размера work group может не позволить организовать отдельный поток для каждого элемента буфера исходных данных. Поэтому нормализацию данных тоже будем осуществлять в цикле.
count = 0;
|
На этом мы завершаем работу с кернелами прямого прохода. Продолжая работу над внесением дополнений в программу OpenCL, мы переходим к построению обратного прохода. Его алгоритм полностью повторяет проделанный выше путь, но в обратном направлении. В нем нам предстоит распределить градиент ошибки от выхода блока Self-Attention до внутренних нейронных слоев m_cQuery, m_cKeys, m_cValues.
Наиболее простым представляется расчет градиента ошибки для внутреннего нейронного слоя m_cValues. Напомню, для получения результата блока Self-Attention мы перемножали матрицу коэффициентов зависимостей m_cScores на тензор результатов нейронного слоя m_cValues. Следовательно, для получения градиента ошибки на уровне выхода указанного нейронного слоя нам необходимо умножить полученный от предыдущих операций градиент ошибки на производную последней операции. В данном случае нам предстоит умножить матрицу коэффициентов зависимостей на тензор градиентов ошибки от предыдущих операций.
После определения градиента ошибки на внутреннем нейронном слое m_cValues нужно распределить градиент ошибки еще на два внутренних нейронных слоя, m_cQuerys и m_cKeys. Но для того, чтобы довести градиент ошибки до уровня указанных нейронных слоев, нужно провести его через матрицу коэффициентов зависимостей.
Но если при реализации средствами MQL5 мы не создаем дополнительный буфер для градиентов ошибки на уровне матрицы коэффициентов зависимостей, то в OpenCL существует сложность с выделением динамического массива для записи промежуточных данных о значениях градиента ошибки на уровне матрицы коэффициентов зависимостей. Поэтому здесь создадим два временных буфера данных: один для градиента ошибки нормализованных данных, а второй для градиентов ошибки, скорректированных на производную функции Softmax.
Обратите внимание: когда мы пересчитываем градиент ошибки до уровня нейронных слоев m_cQuerys и m_cKeys, одни и те же элементы матрицы градиентов ошибки коэффициентов зависимостей используются в разных потоках операций. Поэтому весь алгоритм обратного распределения ошибки внутри слоя внимания мы разделим на два блока. В первом мы проведем градиент ошибки до уровня внутреннего нейронного слоя значений m_cValues и матрицы коэффициентов m_cScores. Во втором блоке доведем градиент ошибки до двух других нейронных слоев, m_cQuerys и m_cKeys.
Первый блок операций реализуем в кернеле AttentionCalcScoreGradient. В параметрах данного кернела передадим указатели на пять буферов данных и один параметр:
- scores — буфер матрицы коэффициентов зависимости;
- scores_temp — буфер градиентов ошибки на уровне нормализованной матрицы коэффициентов зависимости;
- scores_grad — буфер градиентов ошибки на уровне матрицы коэффициентов зависимости, скорректированный на производную функции нормализации;
- values — буфер тензора значений Values (буфер результатов нейронного слоя m_cValues);
- values_grad — буфер тензора градиентов ошибки на уровне результатов нейронного слоя m_cValues;
- outputs_grad — буфер градиентов ошибки на уровне выхода блока Self-Attention;
- window — размер вектора описания одного элемента последовательности в тензоре значений Values.
Обратите внимание, что у буферов scores_temp и scores_grad нет аналогов на стороне основной программы. Дело в том, что градиенты ошибки на уровне матрицы коэффициентов зависимости нам потребуется только для выполнения операций текущего обратного прохода. Но в OpenCL нет возможности создавать динамические массивы. Вместо них мы создали указанные буферы.
__kernel void AttentionCalcScoreGradient(__global TYPE *scores,
|
Алгоритм прямого прохода предусматривает нормализацию матрицы коэффициентов зависимости Score функцией Softmax в разрезе запросов Query. Поэтому после определения градиентов ошибки на уровне матрицы коэффициентов нужно скорректировать полученные значения на производную операции нормализации данных. Следовательно, вполне логично будет разделить операции на потоки в таком же ключе. К тому же такое распределение операций на потоки будет вполне уместно и для распределения градиента ошибки на уровень значений внутреннего нейронного слоя.
В начале кернела мы, как всегда, проведем небольшую подготовительную работу. Определим порядковый номер анализируемых вектора значений и строки матрицы коэффициентов зависимостей по идентификационному номеру потока. Общее количество запущенных потоков подскажет нам размерность тензоров. Сразу же определим смещение в буферах данных до первого элемента анализируемых векторов значений.
const int q = get_global_id(0);
|
Далее мы распределим градиент ошибки до уровня внутреннего нейронного слоя m_cValues. Как уже было сказано выше, для определения градиента ошибки нам необходимо умножить транспонированную матрицу коэффициентов зависимостей на тензор градиентов на выходе блока Self-Attention.
В рамках кернела мы определим градиент ошибки лишь для одного вектора описания элемента. Как вы знаете, при прямом проходе каждый элемент последовательности в тензоре значений Value оставляет свой след в формировании всех элементов последовательности результатов блока Self-Attention. Следовательно, и каждый элемент тензора Value должен получить свою долю градиента ошибки со всех элементов тензора результатов блока Self-Attention. Мерой влияния будет соответствующий коэффициент зависимости из матрицы Score. Таким образом, каждому элементу последовательности тензора значений Value соответствует один столбец в матрице коэффициентов зависимостей Score. Этим и объясняется использование транспонированной матрицы Score в формуле, приведенной выше.
Для организации этого процесса создадим систему из двух вложенных циклов. Число итераций в первом цикле равно размеру вектора описаний одного элемента последовательности в тензоре Value. Надо отметить, что тензор градиентов ошибки на выходе блока Self-Attention имеет такие же размеры. Во вложенном цикле с числом итераций равным количеству элементов в последовательности мы будем перебирать значения соответствующего столбца матрицы коэффициентов зависимости Score и вектора градиентов ошибки на уровне результатов блока Self-Attention. При этом будем перемножать соответствующие элементы, а полученные произведения суммировать в частную переменную. После завершения итераций внутреннего цикла скопируем накопленную сумму произведений в буфер градиентов ошибки внутреннего сверточного слоя m_cValues.
//--- Распределение градиента на Values
|
После отработки системы циклов можно считать выполненной первую часть нашей задачи — передачу градиентов ошибки во внутренний нейронный слой m_cValues.
Вторая часть нашего кернела посвящена определению градиента ошибки на уровне матрицы коэффициентов зависимости.
При прямом проходе каждый элемент последовательности запросов Query перемножается со всеми элементами последовательности ключей Key для формирования одного вектора матрицы коэффициентов зависимостей Score. Каждый такой вектор нормализуется функцией Softmax. После этого умножаем его на тензор значений Value. В результате этих операций получаем скорректированный вектор описания одного элемента последовательности в тензоре результатов блока Self-Attention. Таким образом, один элемент последовательности запросов Query взаимодействует со всеми элементами тензоров Key и Value для формирования вектора описания одного элемента последовательности результатов. Следовательно, для распределения градиента ошибки до отдельно взятого вектора из тензора запросов Query нам необходимо взять один соответствующий вектор градиентов ошибки одного элемента последовательности на уровне результатов блока Self-Attention и сначала умножить его на транспонированный тензор значений Value. Тем самым мы получим вектор ошибки на уровне матрицы коэффициентов зависимости Score. Далее нам необходимо скорректировать полученный вектор на производную функции Softmax. Именно эту часть распределения градиента ошибки мы реализуем в данном кернеле. Для дальнейшего распределения градиента ошибки до уровня внутренних нейронных слоев m_cQuerys и m_cKeys немного позже мы создадим еще один кернел.
Описанный выше алгоритм распределения градиента ошибки в матричном виде можно представить следующим образом:
- Градиент ошибки на уровне матрицы Score.
- Скорректируем градиент ошибки на производную функции Softmax.
Обобщим весь расчет в одну формулу:
Вначале перенесем градиент ошибки на уровень матрицы коэффициентов зависимости Score. Так как благодаря разделению операций на параллельные потоки в рамках кернела мы будем определять градиент ошибки только одной строки, то и для определения этого вектора градиентов ошибки нам необходимо взять вектор градиентов ошибки для одного элемента последовательности на уровне результатов бока Self-Attention и умножить его на транспонированный тензор результатов внутреннего слоя m_cValues. Практически же мы будем использовать алгоритм, описанный выше при вычислении градиентов ошибки на слой m_cValues. Мы создадим систему из двух вложенных циклов. Но на этот раз количество итераций внешнего цикла будет равно количеству элементов в последовательности. Вложенный цикл повторит свои операции по числу элементов в векторе описания одного элемента последовательности. Такая разница объясняется величиной вектора результатов и подтверждается логикой производимых операций. Вспомните, ведь при прямом проходе каждый элемент в строке матрицы коэффициентов зависимостей соответствует одному вектору описания элемента последовательности в тензоре значений Values.
//--- Распределение градиента на Score
|
После переноса градиента ошибки на уровень матрицы коэффициентов зависимостей нам остается скорректировать полученные значения на производную функцию нормализации Softmax. Как и при прямом проходе для получения одного нормализованного значения требовалась обработка всего вектора нормализуемых значений, так и для вычисления одного скорректированного значения нам необходимо использование всех элементов обоих векторов (градиентов ошибок на уровне матрицы коэффициентов зависимости и самого нормализованного вектора коэффициентов).
Матричное выражение процесса корректировки на производную функции Softmax приведено выше. А для практической реализации мы создадим систему из двух вложенных циклов. Оба цикла имеют одинаковое число итераций, которое равно размеру нормализуемого вектора. В данном случае он равен количеству элементов в последовательности. При проведении операций потребуется накапливать сумму градиентов ошибки с каждого элемента нормализованного вектора. Для этого в теле внешнего цикла мы создадим частную переменную grad. Кроме того, чтобы сократить количество обращений к глобальной памяти, сохраним повторяющийся элемент в частную переменную score. Напомню, что обращение к глобальной памяти более затратное по времени. Сокращая количество обращений к буферам глобальной памяти, мы сокращаем общие затраты времени на проведение операций. В теле вложенного цикла будем производить операции умножения элементов и складывать полученные произведения в предварительно созданную частную переменную grad.
Обратите внимание, что единичную матрицу мы заменили выражением (int)(i==k). Логическое выражение даст нам истинное значение только на диагонали матрицы. Перевод логического значение в целочисленное подставит 1 вместо истинных значений и 0 для ложных. Таким образом, столь краткая запись позволяет нам получить значения единичной матрицы прямо в потоке операций без необходимости предварительного ее формирования и сохранения.
//--- Корректируем на производную Softmax
|
После завершения итераций системы циклов на выходе мы получим градиенты ошибки на уровне матрицы коэффициентов зависимости скорректированные на производную функции Softmax.
На этом мы завершаем первый кернел обратного прохода и переходим к созданию второго кернела AttentionCalcHiddenGradient, в котором доведем градиент ошибки до внутренних нейронных слоев m_cQuerys и m_cKeys. Для этого нам потребуется передать в параметрах кернела указатели на пять буферов данных и одну константу:
- querys — буфер результатов внутреннего нейронного слоя m_cQuerys;
- querys_grad — буфер градиентов ошибки внутреннего нейронного слоя m_cQuerys;
- keys — буфер результатов внутреннего нейронного слоя m_cKeys;
- keys_grad — буфер градиентов ошибки внутреннего нейронного слоя m_cKeys;
- scores_grad — буфер градиентов ошибки матрицы коэффициентов зависимостей m_cScores;
- key_size — размер вектора ключа одного элемента.
__kernel void AttentionCalcHiddenGradient(__global TYPE *querys,
|
По аналогии со всеми выше рассмотренными кернелами, распределение на протоки операций будем осуществлять в разрезе одного элемента последовательности. В начале кернела проведем подготовительную работу и определим смещение в буферах данных до первого элемента вектора анализируемого элемента.
const int q = get_global_id(0);
|
В рассмотренном выше кернеле AttentionCalcScoreGradient мы уже скорректировали градиент ошибки матрицы коэффициентов зависимости на производную функции нормализации Softmax. Но во время прямого прохода перед нормализацией матрицы мы разделили все ее элементы на квадратный корень из размерности вектора ключа. Сейчас нам предстоит корректировать градиент ошибки на производную указанной операции. По аналогии с операцией умножения нужно будет разделить все значения буфера градиентов ошибки матрицы коэффициентов зависимости на туже константу.
Определим значение константы и сохраним в частную переменную.
//--- Распределение градиента на Querys и Keys
|
На этом подготовительная работа завершена. Можно приступить непосредственно к пересчету градиентов ошибки. Для получения коэффициентов зависимостей мы перемножали два тензора (запросов Query и ключей Key). С производными операций умножения мы уже не раз сталкивались. Для получения градиентов ошибки на одном из тензоров нам необходимо умножить тензор градиентов ошибки на уровне матрицы коэффициентов зависимостей на второй тензор. А благодаря тому, что тензоры Query и Key имеют одинаковые размерности, мы можем посчитать градиенты ошибки для обоих тензоров в одной системе циклов.
Создадим систему из двух вложенных циклов. Внешний цикл имеет количество итераций равное размеру вектора ключа одного элемента последовательности. А во вложенном цикле мы будем перебирать векторы противоположного тензора и соответствующие градиенты ошибки матрицы коэффициентов зависимости. Следовательно, число его итераций будет равняться количеству элементов в анализируемой последовательности.
Как вы понимаете, в теле вложенного цикла мы будем умножать соответствующие элементы. Результаты этих произведений нужно будет суммировать. Для накопления этой суммы мы создадим две частные переменные grad_q и grad_k перед объявлением вложенного цикла.
И еще один момент. Для сокращения количества операций вычисления мы не будем добавлять в произведения вложенного цикла наш ранее посчитанный коэффициент для корректировки градиента ошибки. Воспользуемся математическими свойствами функций и вынесем постоянный множитель за скобки.
Таким образом, нет необходимости умножать значение каждый раз на поправочный коэффициент в теле вложенного цикла. Вместо этого достаточно один раз умножить на поправочный коэффициент итоговую сумму перед записью в буфер данных.
for(int i = 0; i < key_size; i++)
|
На выходе из системы циклов мы получаем градиенты ошибок для двух вложенных внутренних нейронных слоев m_cQuerys и m_cKeys. То есть задача данного кернела решена. А с учетом выше рассмотренного кернела AttentionCalcScoreGradient, мы распределили градиент ошибки на все внутренние нейронные слои и дальнейшее распределение градиента ошибки до предыдущего слоя будет осуществляться уже протестированными методами внутренних нейронных слоев, как это реализовано стандартными средствами MQL5.
Рассмотренные выше кернелы обратного прохода обошли процессы сложения буферов результатов и нормализации данных, которые мы осуществляли при прямом проходе. И если производная двух функций будет равна сумме производных этих функций, то для операции сложения градиентов мы можем воспользоваться аналогичным кернелом прямого прохода. Надо лишь указать правильные буферы данных.
В случае же корректировки градиента ошибки на функцию нормализации данных нам предстоит создать дополнительный кернел. Напомню формулы корректировки градиента ошибки.
Как можно заметить, в приведенных выше формулах при вычислении производных средних используется сумма значений по всему буферу значений. Но в отличие от прямого прохода, у нас есть возможность все три суммы вычислять параллельно.
В параметрах кернела мы передаем указатели на четыре буфера данных:
- outputs — буфер результатов нормализации прямого прохода;
- out_gradient — буфер градиентов на выходе блока нормализации;
- inp_gradient — буфер для записи скорректированных градиентов;
- stds — буфер среднеквадратических отклонений, вычисленных при прямом проходе.
Также в параметрах укажем размер буферов и смещение в буфере среднеквадратических отклонений.
__kernel void LayerNormalizeGradient(__global TYPE* outputs,
|
В теле кернела определяем идентификаторы потока и сразу объявляем локальные массивы данных. Их будет три. В одном будем собирать производную среднеквадратического отклонения, два других предназначены для слагаемых формулы производной среднего арифметического.
uint ls = min((uint)get_local_size(0), (uint)LOCAL_SIZE);
|
Как и при прямом проходе, сначала каждый поток соберет свою часть от общей суммы.
uint count = 0;
|
В следующем цикле соберем сумму в первых элементах массива.
count = ls;
|
Полученные значения перенесем в частные переменные. При вычислении производной среднеарифметического отклонения значение правого слагаемого умножим на производную среднеквадратического отклонения и сложим с левым слагаемым.
На данном этапе у нас достаточно данных для корректировки градиента ошибки по каждому элементу буфера. Организуем еще один цикл, в теле которого и будет выполнена эта работа.
//---
|
На этом заканчиваем работу с программой OpenCL. Теперь нам предстоит сделать вторую часть и организовать подготовительную работу для запуска многопоточных вычислений на стороне основной программы.
Для начала добавим константы для работы с кернелами в файл defines.mqh. Нам нужно добавить константы идентификации самих кернелов и их переменных. Для именования констант используем ранее оговоренные правила, которые применяются ко всем константам в рамках нашего проекта:
- все константы начинаются с префикса def;
- кернелы начинаются с префикса def_k;
- константы параметров после префикса def содержат указатель на кернел.
#define def_k_AttentionFeedForward 28
|
//--- прямой проход блока внимания
|
//--- определение градиента на матрице коэффициентов зависимости блока внимания
|
//--- распределение градиента через блок внимания
|
//--- сумма векторов
|
//--- нормализация вектора
|
//--- градиент нормализации вектора
|
После этого нам предстоит добавить объявление новых кернелов в код основного класса-диспетчера нейронной сети. Как и всех ранее созданных кернелов, объявление новых кернелов мы добавим в метод CNet::InitOpenCL. В нем мы сначала изменим общее количество используемых в программе кернелов.
if(!m_cOpenCL.SetKernelsCount(34))
|
После этого объявим и сами кернелы.
if(!m_cOpenCL.KernelCreate(def_k_AttentionFeedForward,
|
if(!m_cOpenCL.KernelCreate(def_k_AttentionScoreGradients,
|
if(!m_cOpenCL.KernelCreate(def_k_AttentionHiddenGradients,
|
if(!m_cOpenCL.KernelCreate(def_k_Sum, "Sum"))
|
if(!m_cOpenCL.KernelCreate(def_k_LayerNormalize, "LayerNormalize"))
|
if(!m_cOpenCL.KernelCreate(def_k_LayerNormalizeGradient
|
Затем перейдем к классу механизма внимания CNeuronAttention и внесем изменения в его методы в части работы с технологией OpenCL.
Первым дополним метод прямого прохода CNeuronAttention::FeedForward. В этом методе нам предстоит организовать процедуру вызова кернела прямого прохода AttentionFeedForward. Мы уже не раз создавали аналогичные процессы. Но все же напомню алгоритм:
- Проверка наличия буферов данных в контексте OpenCL;
- Передача параметров кернелу, в том числе указатели на буферы данных;
- Постановка кернела в очередь выполнения операций.
При этом не забываем контролировать выполнение операций, чтобы исключить возможные критические ошибки в процессе дальнейшего выполнения программы.
bool CNeuronAttention::FeedForward(CNeuronBase *prevLayer)
|
При наличии всех необходимых буферов в контексте OpenCL мы организуем передачу указателей на них в параметры кернела.
//--- передача параметров кернелу
|
if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForward, def_attff_outputs,
|
if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForward, def_attff_querys,
|
if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForward, def_attff_scores,
|
if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForward, def_attff_values,
|
if(!m_cOpenCL.SetArgument(def_k_AttentionFeedForward, def_attff_key_size,
|
if(!m_cOpenCL.SetArgument(def_k_AttentionFeedForward, def_attff_window,
|
Далее идет процедура постановки кернела в очередь выполнения. Для начала укажем количество необходимых потоков для запуска и смещение. Только потом вызовем функцию запуска кернела с передачей в нее информации о количестве запускаемых копий.
//--- постановка кернела в очередь выполнения
|
На этом мы выполнили алгоритм запуска кернела блока Self-Attention. Но нам еще предстоит сложить содержимое двух буферов и нормализовать данные буфера результатов. Следуя алгоритму, сначала найдем сумму двух векторов (исходных данных и результатов Self-Attention). Эта операция довольно общая и может широко использоваться за пределами нашего класса нейронного слоя внимания CNeuronAttention. Поэтому было принято решение добавить ее отдельным методом в класс буфера данных CBufferType::SumArray.
В параметрах методу SumArray будем передавать указатель на добавляемый буфер. Сразу в теле метода проверим полученный указатель и размер полученного буфера. Для успешного выполнения операции размер текущего буфера, который будет выступать первым слагаемым, и полученного буфера (второго слагаемого) должны быть равны.
bool CBufferType::SumArray(CBufferType *src)
|
Как и все ранее рассмотренные методы, алгоритм данного метода разделяется на два потока в зависимости от устройства выполнения операций. В блоке выполнения операций средствами MQL5 мы сначала приведем в соответствие форматы матриц обоих буферов. Затем выполним операцию матричного сложения. Результат операции сохраним в матрице текущего буфера.
if(!m_cOpenCL)
|
Алгоритм блока многопоточных операций аналогичен рассмотренному выше. Мы сначала проверяем наличие данных в контексте OpenCL и при необходимости загружаем данные полученного буфера. Обратите внимание, что мы проверяем только полученный буфер. Выше, при разделение алгоритма в зависимости от вычислительного устройства, мы уже проверили указатель на контекст OpenCL текущего буфера. Следовательно, считаем данные текущего буфера уже переданными в контекст OpenCL.
За блоком контролей следует передача параметров кернелу и постановка его в очередь выполнения.
else
|
Процесс нормализации данных организован в методе CNeuronAttention::NormlizeBuffer. Но при соблюдении общих правил построения алгоритма в данном методе есть два исключения. Во-первых, мы исключили блок проверки наличия буферов в контексте OpenCL. В данном случае риск использование незагруженных буферов минимален. Дело в том, что до вызова данного метода используемые буферы были уже проверены не один раз, и повторная проверка будет лишней.
bool CNeuronAttention::NormlizeBuffer(CBufferType *buffer,
|
Второй момент связан с использованием локального массива данных и синхронизацией потоков. Дело в том, что синхронизация потоков доступна только в пределах work group. Нам нужно явно указать ее размер. Алгоритм кернела нормализации построен таким образом, что размер рабочей группы не может быть больше размера локального массива. Напомню, размер локального массива определяется константой LOCAL_SIZE. В то же время количество потоков не может быть больше размера нормализуемого буфера. Следовательно, в массиве указания размерности пространства задач мы укажем меньшее из двух величин. Так как мы нормализуем значения всего буфера одним пакетом, то и размерность глобального и локального пространства задач будет одна.
Определившись с размерностью задач, мы ставим кернел в очередь выполнения.
int NDRange[] = {(int)MathMin(buffer.Total(), LOCAL_SIZE)};
|
На этом завершается блок использования технологии OpenCL в методе прямого прохода нашего класса механизма внимания, и мы завершаем работу над этим методом. Далее его код остается без изменений. Полностью код приведен в разделе описания построения метода стандартными средствами MQL5.
Мы переходим к работе над одним из методов обратного прохода — методом распределения градиента ошибки через скрытый слой CNeuronAttention::CalcHiddenGradient. Алгоритм наших действий остается прежний. Сделаем лишь поправку на использование двух кернелов последовательно.
Напомню, что при создании кернелов обратного прохода мы определили необходимость использования двух дополнительных буферов для записи промежуточных значений градиентов ошибки матрицы коэффициентов зависимости. Поэтому сделаем шаг назад и объявим дополнительные буферы: m_cScoreGrad и m_cScoreTemp.
class CNeuronAttention : public CNeuronBase
|
Только в данном случае мы не будем объявлять экземпляры объектов буфера в основной памяти. Мы не будем использовать указанные буферы для обмена данными между контекстом OpenCL и основной программой. Они нужны лишь для временного хранения данных их передачи между кернелами. А значит, нам достаточно их наличие в памяти контекста OpenCL. В основной программе мы лишь объявим переменные для хранения указателей на буферы.
Но вернемся к работе над методом CNeuronAttention::CalcHiddenGradient. Вначале проверяем наличие и при необходимости создаем новые буферы данных в контексте OpenCL, используемые в первом кернеле. Мы намеренно не создаем сразу буферы данных для второго кернела, чтобы использование памяти было более рационально. Это позволит нам использовать буферы данных большего размера при ограниченных ресурсах памяти контекста OpenCL.
bool CNeuronAttention::CalcHiddenGradient(CNeuronBase *prevLayer)
|
После проверки всех необходимых буферов передадим указатели на них в кернел.
//--- передача параметров кернелу
|
В дополнение к указателям на буферы данных передадим в кернел размер вектора описания одного элемента последовательности.
После передачи всех параметров, укажем количество необходимых параллельных потоков и вызовем функцию постановки кернела в очередь.
//--- Постановка кернела в очередь выполнения
|
Переходим к работе над следующим кернелом. Проверим наличие буферов, необходимых для нового кернела.
if(m_cQuerys.GetOutputs().GetIndex() < 0)
|
После проверки всех необходимых буферов данных передадим указатели на них в кернел.
if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionHiddenGradients,
|
В дополнение к указателям на буферы данных в параметры кернела передадим размер вектора ключа одного элемента последовательности.
После завершения передачи всех необходимых данных в кернел мы инициализируем постановку его в очередь исполнения. Массивы с указанием смещения и количества необходимых копий кернела для запуска уже готовы после запуска предыдущего кернела и нам нет необходимости их задавать повторно. Поэтому мы просто вызываем функцию постановки кернела в очередь.
if(!m_cOpenCL.Execute(def_k_AttentionHiddenGradients, 1, off_set, NDRange))
|
На этом мы заканчиваем работу над построением методов нашего класса внимания и можем перейти к тестированию его работы.