- Описание архитектуры и принципов реализации
- Построение модели GPT средствами MQL5
- Организация параллельных вычислений в модели GPT
- Сравнительное тестирование реализаций
Организация параллельных вычислений в модели GPT
Мы продолжаем работу над нашим классом модели GPT CNeuronGPT. В прошлых разделах мы уже воссоздали алгоритм модели стандартными средствами MQL5. Теперь пришло время дополнить модель возможностью выполнения математических операций в многопоточном режиме с использованием вычислительных мощностей GPU. Именно такую возможность предоставляет нам технология OpenCL.
Для организации этого процесса нам предстоит выполнить две подзадачи:
- создание программы OpenCL,
- организация вызова программы OpenCL из основной программы.
Начнем работу с создания выполняемой программы на стороне OpenCL. В данной программе нужно реализовать ту часть алгоритма, которая не покрывается использованием методов внутренних объектов. У нас два таких блока: один — в части прямого прохода, а второй, зеркальный первому, — в методе распространения градиента ошибки при выполнении обратного прохода.
Для выполнения алгоритма прямого прохода мы создадим кернел GPTFeedForward. Отчасти алгоритм кернела напомнит аналогичный кернел для классов с использованием механизмов внимания. Это и не удивительно — все они используют механизм Self-Attention. Но в каждой реализации есть свои нюансы. Если в прошлый раз вместо создания нового кернела для организации многоголового внимания мы смогли довольно быстро модифицировать уже имеющийся кернел алгоритма Self-Attention, то сейчас создание нового кернела мне показалось менее затратным по сравнению с попыткой создания универсального кернела для всех задач.
В отличие от реализации механизма Multi-Heads Self-Attention, когда мы переводили кернел в двухмерное пространство задач, в данной реализации мы возвращаемся в одно мерное пространство. Это связано с отсутствием возможности деления задачи на параллельные потоки в разрезе элементов последовательности тензора запросов Query, потому что в реализации модели GPT предусмотрена обработка только одного запроса за итерацию. В этом случае у нас остается деление по потокам только в разрезе голов внимания.
В параметрах кернела прямого прохода GPTFeedForward мы, как и прежде, будем передавать указатели на пять буферов данных. Но при этом увеличивается количество переменных: если раньше мы получали размер последовательности из размерности пространства задач, то теперь нам приходится его явно указывать в параметрах кернела. Здесь же добавляется параметр для указания текущего элемента в последовательностях ключей и значений.
__kernel void GPTFeedForward(__global TYPE *querys,
|
Как уже было сказано выше, создаваемый кернел будет работать в одномерном пространстве задача в разрезе голов внимания. Поэтому первое, что мы делаем в теле кернела, это определяем активную голову внимания по идентификатору выполняемого потока и общее количество голов внимания по общему количеству запущенных потоков.
const int h = get_global_id(0);
|
Сразу же определяем смещение в тензорах запросов Query и матрицы коэффициентов зависимости Score.
Далее, согласно выстраиваемому алгоритму Self-Attention, определим коэффициенты зависимости между элементами последовательности. Для этого умножим вектор запроса Query на тензор ключей Key. Для реализации этих операций создадим систему из двух вложенных циклов. Внешний цикл будет перебирать элементы последовательности тензора ключей Key и, соответственно, элементы вектора коэффициентов зависимости Score. В теле цикла мы определим смещение в тензоре ключей Key и подготовим локальную переменную для подсчета промежуточных значений.
После этого организуем вложенный цикл с числом итераций равным размеру вектора описания одного элемента последовательности. В теле данного цикла мы и осуществим операцию умножения пары векторов. Полученное значение делим на квадратный корень из размерности вектора, берем экспоненту. Результат операции записываем в соответствующий элемент вектора коэффициентов зависимости и прибавляем к накопительной сумме всех элементов векторы коэффициентов зависимости для последующей нормализации данных.
Следует учесть момент конкатенированного буфера результатов слоя Query. Дело в том, что значения вектора ключей и запросов текущего элемента последовательности еще не перенесены в соответствующие буферы. Поэтому мы проверяем элемент, к которому обращаемся в тензоре ключей. Перед обращением к текущему элементу мы сначала копируем данные в буфер. Конечно, для текущих операций мы могли бы взять данные из буфера querys. Но эти данные нам потребуются и при последующих итерациях. Следовательно, перенос их в буфер неизбежен.
TYPE summ = 0;
|
В результате выполнения полного цикла итераций созданной выше системы из двух циклов мы получим вектор коэффициентов зависимости. Согласно алгоритму Self-Attention, перед дальнейшим использованием полученных коэффициентов их предстоит нормализовать функцией Softmax. При получении экспоненты из произведений векторов мы уже выполнили часть алгоритма указанной функции. Для полного завершения операции нормализации нам остается разделить сохраненные в векторе значения на их общую сумму, которую мы предусмотрительно собрали в локальной переменной summ. Поэтому организуем еще один цикл с числом итераций равным размеру вектора коэффициентов зависимости. В теле данного цикла мы разделим все значения вектора на значение локальной переменной summ.
for(int s = 0; s < units; s++)
|
Таким образом, после завершения итераций цикла в векторе Score мы получи нормализованные значения коэффициентов зависимости с общей суммой всех элементов раной единице. Фактически полученные коэффициенты дают нам представление о доле влияния каждого элемента последовательности из тензора значений Value на итоговое значение анализируемого элемента последовательности в тензоре результатов текущей головы внимания.
А значит, для получения итоговых значений нам предстоит умножить вектор нормализованных коэффициентов зависимости Score на тензор значений Value. Для выполнения этой операции нам понадобится еще одна система из двух вложенных циклов. Но перед ее запуском мы определим смещение в тензоре результатов до начала вектора анализируемого элемента последовательности.
Внешний цикл с числом итераций равным размеру вектора описания одного элемента последовательности укажет нам на порядковый номер собираемого элемента в векторе результатов. Вложенный цикл с числом итераций равным числу элементов в последовательности поможет сопоставить векторы тензора значений Value и коэффициенты зависимости из вектора Score. В теле вложенного цикла мы будем выполнять непосредственно операцию умножения вектора из тензора значений Value на соответствующий коэффициент зависимости элементов. Полученные значения произведений мы суммируем в локальную переменную, а после завершения итераций вложенного цикла сохраним полученное значение в буфере результатов работы блока Self-Attention.
shift_query = key_size * h;
|
В результате выполнения полного цикла итераций системы циклов мы получим вектор описания одного элемента последовательности в тензоре результатов одной головы внимания. Поставленная перед данным кернелом задача выполнена, и мы выходим из него.
На этом мы заканчиваем работу с кернелом прямого прохода и двигаемся дальше. Теперь нам предстоит организовать процесс обратного прохода. Реализацию данной задачи разобьем на два кернела. В кернеле GPTCalcScoreGradient мы доведем градиент ошибки до вектора коэффициентов зависимости. Во кернеле GPTCalcHiddenGradient продолжим распределение градиента ошибки и доведем его до уровня тензоров запросов Query и ключей Key.
Но обо всем по порядку. Кернел GPTCalcScoreGradient в параметрах получает указатели на шесть буферов данных и три параметра:
- scores — буфер вектора коэффициентов зависимости;
- scores_grad — буфер вектора градиентов ошибки на уровне коэффициентов зависимости;
- values — буфер тензора значений Value;
- values_grad — буфер тензора градиентов ошибки на уровне Value;
- outputs_grad — буфер тензора градиентов ошибки на уровне результатов блока Self-Attention;
- scores_temp — буфер для записи промежуточных значений;
- window — размер вектора описания одного элемента последовательности в тензоре значений Value;
- units — количество элементов в последовательности;
- current — порядковый номер текущего элемента в стеке значений Value.
__kernel void GPTCalcScoreGradient(__global TYPE *scores,
|
Как и при прямом проходе, кернел мы будем запускать в одномерном пространстве по числу используемых голов внимания. В теле кернела мы сразу определяем активную голову внимания по идентификатору потока и общее количество голов внимания по общему количеству запущенных потоков.
const int h = get_global_id(0);
|
И сразу определяем смещение в тензорах градиентов ошибки на уровне значений Value и в векторе коэффициентов зависимости. Обратите внимание, что смещение в тензоре градиентов ошибки значений Value и в самом тензоре значений Value в данном случае будет различным.
Напомню, что в данной реализации модели GPT мы использовали один внутренний нейронный слой для генерации конкатенированного тензора, содержащего сразу значение тензоров Query, Key и Value для всех голов внимания. Соответственно, и градиент ошибки мы собираем в аналогичный конкатенированный тензор градиентов ошибки указанного нейронного слоя. Но данный тензор содержит только текущий элемент последовательности. В то же время в тензоре стека Value содержится полная информация обо всей последовательности, но только тензора значений Value.
После проведения подготовительной работы мы распределим градиент ошибки на тензор значений Value. Как уже было сказано выше, мы распределяем градиент ошибки только для текущего элемента последовательности. Для этого мы организуем цикл с числом итераций равным размеру вектора описания одного элемента последовательности в тензоре значений Value. В теле цикла мы умножим вектор градиентов ошибки на уровне результатов блока Self-Attention на соответствующий коэффициент зависимости из вектора Score. Полученные значения сохраняем в буфер конкатенированного тензора градиентов ошибки.
//--- Распределение градиента на Values
|
После расчета градиента ошибки на тензоре значений Value определим значение градиента ошибки на уровне вектора коэффициентов зависимости. Для выполнения этой операции нам понадобится система из двух циклов: внешнего с числом итераций равным числу элементов последовательности и вложенным с числом итераций равным размеру вектора описания одного элемента последовательности в тензоре значений Value. В теле вложенного цикла мы будем осуществлять произведение 2-х векторов (Value на градиент ошибки). А полученное значение сохраняем в буфер временных данных.
//--- Распределение градиента на Score
|
В результате выполнения полного цикла итераций системы циклов во временном буфере мы получим полностью заполненный вектор градиентов ошибки для вектора коэффициентов зависимости. Но для распределения градиента ошибки дальше, нам предварительно надо его скорректировать на производную функции Softmax.
Организуем еще одну систему из двух вложенных циклов. Оба цикла будут содержать количество итераций равное числу элементов в последовательности. В теле вложенного цикла мы посчитаем производную функции по формуле.
//--- Корректируем на производную Softmax
|
Полученные значения сохраним в буфер градиентов ошибки на уровне вектора коэффициентов зависимости. На данном этапе мы завершаем работу первого кернела и переходим к работе над вторым.
Во втором кернеле организации процесса обратного прохода GPTCalcHiddenGradient нам предстоит распределить градиент ошибки дальше и довести его до уровня тензоров запросов Query и ключей Key.
В параметрах кернел GPTCalcHiddenGradient получает указатели на 4 буфера данных и 3 параметра.
__kernel void GPTCalcHiddenGradient(__global TYPE *querys,
|
Обратите внимание, что мы говорили о распределении градиента на тензоры Query и Key. А в параметрах кернела есть указатель только на буфер градиентов ошибки Query. Такая ситуация стала возможной благодаря использованию конкатенированного буфера, в который мы уже сохранили градиент ошибки на уровне тензора значений Value. И теперь в тот же буфер мы добавим градиенты ошибок на уровне тензоров Query и Key.
В теле кернела мы определяем порядковый номер анализируемой головы внимания по идентификатору потока и количество используемых голов внимания по общему количеству запущенного пула задач.
const int h = get_global_id(0);
|
Здесь же мы определяем смещения в буферах данных до начала анализируемых векторов.
Далее мы организуем систему из двух вложенных циклов, в теле которой определим градиенты ошибок на уровне искомых тензоров. Для этого умножим градиент ошибки на уровне вектора коэффициентов зависимости на противоположный тензор.
//--- Распределение градиента на Querys и Keys
|
Стоит отметить, что мы определяем градиент ошибки только для текущего элемента последовательности и записываем полученные значения в соответствующие элементы буфера градиентов ошибок.
В результате выполнения всех итераций нашей системы циклов мы получаем полностью заполненный конкатенированный тензор градиентов ошибок всех трех сущностей (Query, Key, Value). Завершаем работу по построению программы OpenCL и переходим к построению функционала на стороне основной программы.
Чтобы обслуживать построенные кернелы на стороне основной программы было удобнее, создадим именованные константы вызова кернелов и доступа к его элементам. Для этого открываем наш файл констант defines.mqh и создаем в нем константы обращения к кернелам.
#define def_k_GPTFeedForward 34
|
Добавляем константы обращения к параметрам кернелов.
//--- прямой проход GPT
|
//--- определение градиента на матрице коэффициентов зависимости GPT
|
//--- распределение градиента через GPT
|
После этого переходим к диспетчерскому классу обслуживания модели нейронной сети CNet и в методе инициализации контекста OpenCL InitOpenCL изменяем общее количество кернелов в нашей программе. Затем инициализируем создание новых кернелов в контексте OpenCL.
bool CNet::InitOpenCL(void)
|
На этом мы завершаем подготовительную работу и переходим непосредственно к методам нашего класса CNeuronGPT. В них нам предстоит выполнить три этапа работы для вызова каждого кернела:
- подготовка исходных данных и передача их в память контекста OpenCL;
- постановка кернела в очередь выполнения;
- загрузка результатов выполнения программы в память основной программы.
Первым мы модифицируем метод прямого прохода CNeuronGPT::FeedForward. В блоке организации многопоточных вычислений с использованием технологии OpenCL мы сначала проверяем наличие уже созданного буфера в памяти контекста OpenCL.
bool CNeuronGPT::FeedForward(CNeuronBase *prevLayer)
|
Когда все буферы созданы, а те из них, что необходимы для работы кернела, переданы в память контекста OpenCL, мы предаем в параметры кернела указатели на используемые буферы данных и необходимые константы.
//--- передача параметров кернелу
|
На этом этап подготовительной работы завершен. Переходим к этапу постановки кернела в очередь выполнения. Здесь мы сначала создаем два динамических массива, в которых укажем смещение и количество выполняемых потоков в каждом подпространстве задач. Затем вызовем метод постановки кернела в очередь выполнения m_cOpenCL.Execute.
//--- постановка кернела в очередь выполнения
|
На этом мы завершаем работу с методом прямого прохода CNeuronGPT::FeedForward. Но нам еще предстоит выполнить аналогичную работу в методе алгоритма обратного прохода CNeuronGPT::CalcHiddenGradient.
Напомню, что для реализации метода обратного прохода мы создали два кернела, которые будем вызывать последовательно друг за другом. Следовательно, и работу по обслуживанию кернела нужно повторить для каждого из них.
Сначала создадим буферы данных для первого кернела.
bool CNeuronGPT::CalcHiddenGradient(CNeuronBase *prevLayer)
|
Следуя нашему алгоритму работы с контекстом OpenCL, после создания буферов данных и передачи всей необходимой информации в память контекста мы передаем в параметры запускаемого кернела указатели на используемые буферы данных и константы для выполнения алгоритма программы.
//--- передача параметров кернелу
|
Обратите внимание, что вместо буфера градиентов ошибки тензора значений Value мы передаем указатель на буфер градиентов внутреннего нейронного слоя Querys. Это вызвано использованием конкатенированного буфера градиентов ошибки для всех трех тензоров. Чтобы исключить последующую операцию копирования данных, мы будем сразу записывать данные в конкатенированный буфер.
После этого мы выполняем операцию постановки кернела в очередь. Напомню, что запускаем мы кернел на выполнение в одномерном пространстве задач в разрезе голов внимания.
Укажем смещение в пространстве задач и количество запускаемых потоков в соответствующих динамических массивах. После этого вызовем метод постановки нашего кернела в очередь задач.
//--- постановка кернела в очередь выполнения
|
На этом работа над первым кернелом завершена, и мы переходим к построению аналогичного алгоритма для второго кернела обратного прохода.
Проверяем дополнительные буферы в памяти контекста OpenCL.
if(Querys.GetOutputs().GetIndex() < 0)
|
Передаем параметры кернелу.
if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTHiddenGradients,
|
После этого мы ставим кернел в очередь выполнения. Обратите внимание, что на это раз мы не создаем массивы смещения и размерности пространства задач. Мы просто используем без изменений массивы, созданные при выполнении предыдущего кернела.
if(!m_cOpenCL.Execute(def_k_GPTHiddenGradients, 1, off_set, NDRange))
|
На этом мы завершаем работу по построению класса модели GPT и можем перейти к оценке результатов проделанной работы.