Организация параллельных вычислений в модели 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,
                             __global TYPE *keys,
                             __global TYPE *scores,
                             __global TYPE *values,
                             __global TYPE *outputs,
                             int key_size,
                             int units,
                             int current)
  {

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

   const int h = get_global_id(0);
   const int heads = get_global_size(0);
   int shift_query = key_size * h;
   int shift_scores = units * h;

Сразу же определяем смещение в тензорах запросов Query и матрицы коэффициентов зависимости Score.

Далее, согласно выстраиваемому алгоритму Self-Attention, определим коэффициенты зависимости между элементами последовательности. Для этого умножим вектор запроса Query на тензор ключей Key. Для реализации этих операций создадим систему из двух вложенных циклов. Внешний цикл будет перебирать элементы последовательности тензора ключей Key и, соответственно, элементы вектора коэффициентов зависимости Score. В теле цикла мы определим смещение в тензоре ключей Key и подготовим локальную переменную для подсчета промежуточных значений.

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

Следует учесть момент конкатенированного буфера результатов слоя Query. Дело в том, что значения вектора ключей и запросов текущего элемента последовательности еще не перенесены в соответствующие буферы. Поэтому мы проверяем элемент, к которому обращаемся в тензоре ключей. Перед обращением к текущему элементу мы сначала копируем данные в буфер. Конечно, для текущих операций мы могли бы взять данные из буфера querys. Но эти данные нам потребуются и при последующих итерациях. Следовательно, перенос их в буфер неизбежен.

   TYPE summ = 0;
   for(int s = 0s < unitss++)
     {
      TYPE score = 0;
      int shift_key = key_size * (s * heads + h);
      for(int k = 0k < key_sizek ++)
        {
         if(s == current)
            keys[shift_key + k] = querys[shift_query + k + heads * key_size];
         score += querys[shift_query + k] * keys[shift_key + k];
        }
      score = exp(score / sqrt((TYPE)key_size));
      summ += score;
      scores[shift_scores + s] = score;
     }

В результате выполнения полного цикла итераций созданной выше системы из двух циклов мы получим вектор коэффициентов зависимости. Согласно алгоритму Self-Attention, перед дальнейшим использованием полученных коэффициентов их предстоит нормализовать функцией Softmax. При получении экспоненты из произведений векторов мы уже выполнили часть алгоритма указанной функции. Для полного завершения операции нормализации нам остается разделить сохраненные в векторе значения на их общую сумму, которую мы предусмотрительно собрали в локальной переменной summ. Поэтому организуем еще один цикл с числом итераций равным размеру вектора коэффициентов зависимости. В теле данного цикла мы разделим все значения вектора на значение локальной переменной summ.

   for(int s = 0s < unitss++)
      scores[shift_scores + s] /= summ;

Таким образом, после завершения итераций цикла в векторе Score мы получи нормализованные значения коэффициентов зависимости с общей суммой всех элементов раной единице. Фактически полученные коэффициенты дают нам представление о доле влияния каждого элемента последовательности из тензора значений Value на итоговое значение анализируемого элемента последовательности в тензоре результатов текущей головы внимания.

А значит, для получения итоговых значений нам предстоит умножить вектор нормализованных коэффициентов зависимости Score на тензор значений Value. Для выполнения этой операции нам понадобится еще одна система из двух вложенных циклов. Но перед ее запуском мы определим смещение в тензоре результатов до начала вектора анализируемого элемента последовательности.

Внешний цикл с числом итераций равным размеру вектора описания одного элемента последовательности укажет нам на порядковый номер собираемого элемента в векторе результатов. Вложенный цикл с числом итераций равным числу элементов в последовательности поможет сопоставить векторы тензора значений Value и коэффициенты зависимости из вектора Score. В теле вложенного цикла мы будем выполнять непосредственно операцию умножения вектора из тензора значений Value на соответствующий коэффициент зависимости элементов. Полученные значения произведений мы суммируем в локальную переменную, а после завершения итераций вложенного цикла сохраним полученное значение в буфере результатов работы блока Self-Attention.

   shift_query = key_size * h;
   for(int i = 0i < key_sizei++)
     {
      TYPE query = 0;
      for(int v = 0v < unitsv++)
        {
         if(v == current)
            values[key_size * (v * heads + h) + i] = 
                                      querys[(2 * heads + h) * key_size + i];
         query += values[key_size * (v * heads + h) + i] * 
                                                    scores[shift_scores + v];
        }
      outputs[shift_query + i] = query;
     }
  }

В результате выполнения полного цикла итераций системы циклов мы получим вектор описания одного элемента последовательности в тензоре результатов одной головы внимания. Поставленная перед данным кернелом задача выполнена, и мы выходим из него.

На этом мы заканчиваем работу с кернелом прямого прохода и двигаемся дальше. Теперь нам предстоит организовать процесс обратного прохода. Реализацию данной задачи разобьем на два кернела. В кернеле 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,
                                   __global TYPE *scores_grad,
                                   __global TYPE *values,
                                   __global TYPE *values_grad,
                                   __global TYPE *outputs_grad,
                                   __global TYPE *scores_temp,
                                   int window,
                                   int units,
                                   int current)
  {

Как и при прямом проходе, кернел мы будем запускать в одномерном пространстве по числу используемых голов внимания. В теле кернела мы сразу определяем активную голову внимания по идентификатору потока и общее количество голов внимания по общему количеству запущенных потоков.

   const int h = get_global_id(0);
   const int heads = get_global_size(0);
   int shift_value = window * (2 * heads + h);
   int shift_score = units * h;

И сразу определяем смещение в тензорах градиентов ошибки на уровне значений Value и в векторе коэффициентов зависимости. Обратите внимание, что смещение в тензоре градиентов ошибки значений Value и в самом тензоре значений Value в данном случае будет различным.

Напомню, что в данной реализации модели GPT мы использовали один внутренний нейронный слой для генерации конкатенированного тензора, содержащего сразу значение тензоров Query, Key и Value для всех голов внимания. Соответственно, и градиент ошибки мы собираем в аналогичный конкатенированный тензор градиентов ошибки указанного нейронного слоя. Но данный тензор содержит только текущий элемент последовательности. В то же время в тензоре стека Value содержится полная информация обо всей последовательности, но только тензора значений Value.

После проведения подготовительной работы мы распределим градиент ошибки на тензор значений Value. Как уже было сказано выше, мы распределяем градиент ошибки только для текущего элемента последовательности. Для этого мы организуем цикл с числом итераций равным размеру вектора описания одного элемента последовательности в тензоре значений Value. В теле цикла мы умножим вектор градиентов ошибки на уровне результатов блока Self-Attention на соответствующий коэффициент зависимости из вектора Score. Полученные значения сохраняем в буфер конкатенированного тензора градиентов ошибки.

//--- Распределение градиента на Values
   for(int i = 0i < windowi ++)
      values_grad[shift_value + i] = scores[units * h + current] * 
                                                   outputs_grad[window * h + i];

После расчета градиента ошибки на тензоре значений Value определим значение градиента ошибки на уровне вектора коэффициентов зависимости. Для выполнения этой операции нам понадобится система из двух циклов: внешнего с числом итераций равным числу элементов последовательности и вложенным с числом итераций равным размеру вектора описания одного элемента последовательности в тензоре значений Value. В теле вложенного цикла мы будем осуществлять произведение 2-х векторов (Value на градиент ошибки). А полученное значение сохраняем в буфер временных данных.

//--- Распределение градиента на Score
   for(int k = 0k < unitsk++)
     {
      TYPE grad = 0;
      for(int i = 0i < windowi++)
         grad += outputs_grad[shift_value + i] * 
                                        values[window * (k * heads + h) + i];
      scores_temp[shift_score + k] = grad;
     }

В результате выполнения полного цикла итераций системы циклов во временном буфере мы получим полностью заполненный вектор градиентов ошибки для вектора коэффициентов зависимости. Но для распределения градиента ошибки дальше, нам предварительно надо его скорректировать на производную функции Softmax.

Организуем еще одну систему из двух вложенных циклов. Оба цикла будут содержать количество итераций равное числу элементов в последовательности. В теле вложенного цикла мы посчитаем производную функции по формуле.

//--- Корректируем на производную Softmax
   for(int k = 0k < unitsk++)
     {
      TYPE grad = 0;
      TYPE score = scores[shift_score + k];
      for(int i = 0i < unitsi++)
         grad += scores[shift_score + i] * ((int)(i == k) - score) *
                                                scores_temp[shift_score + i];
      scores_grad[shift_score + k] = grad;
     }
  }

Полученные значения сохраним в буфер градиентов ошибки на уровне вектора коэффициентов зависимости. На данном этапе мы завершаем работу первого кернела и переходим к работе над вторым.

Во втором кернеле организации процесса обратного прохода GPTCalcHiddenGradient нам предстоит распределить градиент ошибки дальше и довести его до уровня тензоров запросов Query и ключей Key.

В параметрах кернел GPTCalcHiddenGradient получает указатели на 4 буфера данных и 3 параметра.

__kernel void GPTCalcHiddenGradient(__global TYPE *querys,
                                    __global TYPE *querys_grad,
                                    __global TYPE *keys,
                                    __global TYPE *scores_grad,
                                    int key_size,
                                    int units,
                                    int current)
  {

Обратите внимание, что мы говорили о распределении градиента на тензоры Query и Key. А в параметрах кернела есть указатель только на буфер градиентов ошибки Query. Такая ситуация стала возможной благодаря использованию конкатенированного буфера, в который мы уже сохранили градиент ошибки на уровне тензора значений Value. И теперь в тот же буфер мы добавим градиенты ошибок на уровне тензоров Query и Key.

В теле кернела мы определяем порядковый номер анализируемой головы внимания по идентификатору потока и количество используемых голов внимания по общему количеству запущенного пула задач.

   const int h = get_global_id(0);
   const int heads = get_global_size(0);
   int shift_query = key_size * h;
   int shift_key = key_size * (heads + h);
   int shift_score = units * h;

Здесь же мы определяем смещения в буферах данных до начала анализируемых векторов.

Далее мы организуем систему из двух вложенных циклов, в теле которой определим градиенты ошибок на уровне искомых тензоров. Для этого умножим градиент ошибки на уровне вектора коэффициентов зависимости на противоположный тензор.

//--- Распределение градиента на Querys и Keys
   const TYPE k = 1 / sqrt((TYPE)key_size);
//---
   for(int i = 0i < key_sizei++)
     {
      TYPE grad_q = 0;
      TYPE grad_k = 0;
      for(int s = 0s < unitss++)
        {
         grad_q += keys[key_size * (s * heads + h) + i] *
                                               scores_grad[shift_score + s];
         if(s == current)
            grad_k += querys[key_size * h + i] *
                                           scores_grad[units * h + current];
        }
      querys_grad[shift_query + i] = grad_q * k;
      querys_grad[shift_key + i] = grad_k * k;
     }
  }

Стоит отметить, что мы определяем градиент ошибки только для текущего элемента последовательности и записываем полученные значения в соответствующие элементы буфера градиентов ошибок.

В результате выполнения всех итераций нашей системы циклов мы получаем полностью заполненный конкатенированный тензор градиентов ошибок всех трех сущностей (Query, Key, Value). Завершаем работу по построению программы OpenCL и переходим к построению функционала на стороне основной программы.

Чтобы обслуживать построенные кернелы на стороне основной программы было удобнее, создадим именованные константы вызова кернелов и доступа к его элементам. Для этого открываем наш файл констант defines.mqh и создаем в нем константы обращения к кернелам.

#define def_k_GPTFeedForward           34
#define def_k_GPTScoreGradients        35
#define def_k_GPTHiddenGradients       36

Добавляем константы обращения к параметрам кернелов.

//--- прямой проход GPT
#define def_gptff_querys               0
#define def_gptff_keys                 1
#define def_gptff_scores               2
#define def_gptff_values               3
#define def_gptff_outputs              4
#define def_gptff_key_size             5
#define def_gptff_units                6
#define def_gptff_current              7

//--- определение градиента на матрице коэффициентов зависимости GPT
#define def_gptscr_scores              0
#define def_gptscr_scores_grad         1
#define def_gptscr_values              2
#define def_gptscr_values_grad         3
#define def_gptscr_outputs_grad        4
#define def_gptscr_scores_temp         5
#define def_gptscr_window              6
#define def_gptscr_units               7
#define def_gptscr_current             8

//--- распределение градиента через GPT
#define def_gpthgr_querys              0
#define def_gpthgr_querys_grad         1
#define def_gpthgr_keys                2
#define def_gpthgr_scores_grad         3
#define def_gpthgr_key_size            4
#define def_gpthgr_units               5
#define def_gpthgr_current             6

После этого переходим к диспетчерскому классу обслуживания модели нейронной сети CNet и в методе инициализации контекста OpenCL InitOpenCL изменяем общее количество кернелов в нашей программе. Затем инициализируем создание новых кернелов в контексте OpenCL.

bool CNet::InitOpenCL(void)
  {
   ......
   if(!m_cOpenCL.SetKernelsCount(37))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   ......
   if(!m_cOpenCL.KernelCreate(def_k_GPTFeedForward"GPTFeedForward"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   if(!m_cOpenCL.KernelCreate(def_k_GPTScoreGradients"GPTCalcScoreGradient"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   if(!m_cOpenCL.KernelCreate(def_k_GPTHiddenGradients"GPTCalcHiddenGradient"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
//---
   return true;
  }

На этом мы завершаем подготовительную работу и переходим непосредственно к методам нашего класса CNeuronGPT. В них нам предстоит выполнить три этапа работы для вызова каждого кернела:

  • подготовка исходных данных и передача их в память контекста OpenCL;
  • постановка кернела в очередь выполнения;
  • загрузка результатов выполнения программы в память основной программы.

Первым мы модифицируем метод прямого прохода CNeuronGPT::FeedForward. В блоке организации многопоточных вычислений с использованием технологии OpenCL мы сначала проверяем наличие уже созданного буфера в памяти контекста OpenCL.

bool CNeuronGPT::FeedForward(CNeuronBase *prevLayer)
  {
   ......
   for(int layer = 0layer < m_iLayerslayer++)
     {
   ......
      //--- разветвление алгоритма по вычислительному устройству
      if(!m_cOpenCL)
        {
         // Блок программы стандартными средствами MQL5
   ......
        }
      else // блок OpenCL
        {
         //--- проверка буферов данных
         if(Querys.GetOutputs().GetIndex() < 0)
            return false;
         if(Keys.GetOutputs().GetIndex() < 0)
            return false;
         if(Values.GetOutputs().GetIndex() < 0)
            return false;
         if(Scores.GetOutputs().GetIndex() < 0)
            return false;
         if(AttentionOut.GetOutputs().GetIndex() < 0)
            return false;

Когда все буферы созданы, а те из них, что необходимы для работы кернела, переданы в память контекста OpenCL, мы предаем в параметры кернела указатели на используемые буферы данных и необходимые константы.

         //--- передача параметров кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTFeedForward
                                def_gptff_keysKeys.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTFeedForward,
                     def_gptff_outputsAttentionOut.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTFeedForward
                            def_gptff_querysQuerys.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTFeedForward,
                            def_gptff_scoresScores.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTFeedForward
                            def_gptff_valuesValues.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTFeedForward,
                                             def_gptff_key_sizem_iKeysSize))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTFeedForward
                                                   def_gptff_unitsm_iUnits))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTFeedForward,
                                       def_gptff_currentm_iCurrentPosition))
            return false;

На этом этап подготовительной работы завершен. Переходим к этапу постановки кернела в очередь выполнения. Здесь мы сначала создаем два динамических массива, в которых укажем смещение и количество выполняемых потоков в каждом подпространстве задач. Затем вызовем метод постановки кернела в очередь выполнения m_cOpenCL.Execute.

         //--- постановка кернела в очередь выполнения
         int off_set[] = {0};
         int NDRange[] = {m_iHeads};
         if(!m_cOpenCL.Execute(def_k_GPTFeedForward1off_setNDRange))
            return false;
        }

На этом мы завершаем работу с методом прямого прохода CNeuronGPT::FeedForward. Но нам еще предстоит выполнить аналогичную работу в методе алгоритма обратного прохода CNeuronGPT::CalcHiddenGradient.

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

Сначала создадим буферы данных для первого кернела.

bool CNeuronGPT::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
   ......
   for(int layer = m_iLayers - 1layer >= 0layer--)
     {
   ......
      //--- разветвление алгоритма по вычислительному устройству
      attention_grad = AttentionOut.GetGradients();
      if(!m_cOpenCL)
        {
         // Блок программы стандартными средствами MQL5
   ......
        }
      else // блок OpenCL
        {
         //--- проверка буферов данных
         if(Values.GetOutputs().GetIndex() < 0)
            return false;
         if(Querys.GetGradients().GetIndex() < 0)
            return false;
         if(Scores.GetOutputs().GetIndex() < 0)
            return false;
         if(attention_grad.GetIndex() < 0)
            return false;
         if(Scores.GetGradients().GetIndex() < 0)
            return false;
         if(m_iScoreTemp < 0)
            return false;

Следуя нашему алгоритму работы с контекстом OpenCL, после создания буферов данных и передачи всей необходимой информации в память контекста мы передаем в параметры запускаемого кернела указатели на используемые буферы данных и константы для выполнения алгоритма программы.

         //--- передача параметров кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTScoreGradients
                           def_gptscr_outputs_gradattention_grad.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTScoreGradients,
                            def_gptscr_scoresScores.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTScoreGradients,
                     def_gptscr_scores_gradScores.GetGradients().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTScoreGradients,
                                         def_gptscr_scores_tempm_iScoreTemp))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTScoreGradients,
                            def_gptscr_valuesValues.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTScoreGradients,
                     def_gptscr_values_gradQuerys.GetGradients().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTScoreGradients,
                                               def_gptscr_windowm_iKeysSize))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTScoreGradients,
                                                    def_gptscr_unitsm_iUnits))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTScoreGradients,
                                         def_gptscr_currentm_iCurrentPosition))
            return false;

Обратите внимание, что вместо буфера градиентов ошибки тензора значений Value мы передаем указатель на буфер градиентов внутреннего нейронного слоя Querys. Это вызвано использованием конкатенированного буфера градиентов ошибки для всех трех тензоров. Чтобы исключить последующую операцию копирования данных, мы будем сразу записывать данные в конкатенированный буфер.

После этого мы выполняем операцию постановки кернела в очередь. Напомню, что запускаем мы кернел на выполнение в одномерном пространстве задач в разрезе голов внимания.

Укажем смещение в пространстве задач и количество запускаемых потоков в соответствующих динамических массивах. После этого вызовем метод постановки нашего кернела в очередь задач.

         //--- постановка кернела в очередь выполнения
         int off_set[] = {0};
         int NDRange[] = {m_iHeads};
         if(!m_cOpenCL.Execute(def_k_GPTScoreGradients1off_setNDRange))
            return false;

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

Проверяем дополнительные буферы в памяти контекста OpenCL.

         if(Querys.GetOutputs().GetIndex() < 0)
            return false;
         if(Keys.GetOutputs().GetIndex() < 0)
            return false;

Передаем параметры кернелу.

         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTHiddenGradients,
                                 def_gpthgr_keysKeys.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTHiddenGradients,
                             def_gpthgr_querysQuerys.GetOutputs().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTHiddenGradients,
                      def_gpthgr_querys_gradQuerys.GetGradients().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_GPTHiddenGradients,
                      def_gpthgr_scores_gradScores.GetGradients().GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTHiddenGradients,
                                              def_gpthgr_key_sizem_iKeysSize))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTHiddenGradients,
                                                    def_gpthgr_unitsm_iUnits))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_GPTHiddenGradients,
                                        def_gpthgr_currentm_iCurrentPosition))
            return false;

После этого мы ставим кернел в очередь выполнения. Обратите внимание, что на это раз мы не создаем массивы смещения и размерности пространства задач. Мы просто используем без изменений массивы, созданные при выполнении предыдущего кернела.

         if(!m_cOpenCL.Execute(def_k_GPTHiddenGradients1off_setNDRange))
            return false;
        }

На этом мы завершаем работу по построению класса модели GPT и можем перейти к оценке результатов проделанной работы.