Организация параллельных вычислений в блоке внимания

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

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

  • создание программы OpenCL,
  • внесение изменений в основную программу.

Код программы OpenCL мы добавим в ранее созданный файл opencl_program.cl. Именно в этом файле мы собрали все кернелы программы OpenCL, используемые в работе предыдущих классов. Для организации работы нашего класса внимания потребуется создание шести кернелов. В этих кернелах нам предстоит организовать поток информации между используемыми внутренними нейронными слоями в прямом и обратном направлении

Первым мы создадим кернел прямого прохода AttentionFeedForward. Кратко напомню последовательность операций при прямом проходе блока Self-Attention:

  1. Исходные данные подаются на вход трем внутренним сверточным нейронным слоям m_cQuery, m_cKeys, m_cValues.

  1. Тензоры результатов m_cQuery и m_cKeys перемножаются для получения матрицы зависимостей m_cScores.

  1. Значения матрицы m_cScores делятся на квадратный корень от размера вектора описания одного элемента последовательности m_cKeys и нормализуются функцией Softmax в разрезе строк (запросов m_cQuery).

  1. Нормализованная матрица m_cScores перемножается с тензором результатов нейронного слоя m_cValues для получения результатов Self-Attention.

  1. Результаты блока Self-Attention складываются с исходными данными и нормализуются.

  1. Полученный тензор служит исходными данными для блока из двух сверточных слоев m_cFF1 и m_cFF2.

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

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

__kernel void AttentionFeedForward(__global TYPE *querys,
                                   __global TYPE *keys,
                                   __global TYPE *scores,
                                   __global TYPE *values,
                                   __global TYPE *outputs,
                                   int window,
                                   int key_size)
  {

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

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

   const int q = get_global_id(0);
   const int units = get_global_size(0);
   int shift_query = key_size * q;
   int shift_scores = units * q;
   TYPE summ = 0;

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

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

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

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

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

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

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

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

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

   shift_query = window * q;
   for(int i = 0i < windowi++)
     {
      TYPE query = 0;
      for(int v = 0v < unitsv++)
         query += values[window * v + i] * scores[shift_scores + v];
      outputs[shift_query + i] = query;
     }
  }

На этом мы завершим работу над первым кернелом прямого прохода. Следующим этапом создадим кернел сложения двух тензоров. Подобную работу порой выгоднее сделать с использованием матричных операций на стороне основной программы. Операция несложная, и накладные расходы в виде передачи данных вряд ли будут оправданы. У нас сейчас обратная ситуация. Мы организовываем весь процесс на стороне контекста OpenCL. Вся информация уже находится в памяти контекста, и для реализации операции на стороне основной программы нам потребуется копирование данных. Для вычислений в контексте перенос данных не нужен Поэтому мы создали кернел Sum, в теле которого просто складываем элементы из двух буферов с одним индексом и сохраняем в элемент третьего буфера с тем же индексом.

__kernel void Sum(__global TYPE *inputs1,
                  __global TYPE *inputs2,
                  __global TYPE *outputs)
  {
   const int n = get_global_id(0);
//---
   outputs[n] = inputs1[n] + inputs2[n];
  }

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

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

Процесс нормализации организован в кернеле LayerNormalize. В параметрах кернел получает указатели на 3 буфера:

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

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

Кроме того, мы передадим в кернел два параметра: общее количество элементов в нормализуемом буфере и смещение в буфере среднеквадратических отклонений. Напомню, в рамках одного нейронного слоя внимания нормализацию данных мы осуществляем два раза. Нормализуем результаты блоков Self-Attention и FeedForward.

__kernel void LayerNormalize(__global TYPEinputs,
                             __global TYPEoutputs,
                             __global TYPEstds,
                             const int total,
                             const int std_shift)
  {

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

   uint i = (uint)get_global_id(0);
   uint l = (uint)get_local_id(0);
   uint ls = min((uint)get_local_size(0), (uint)LOCAL_SIZE);
   __local TYPE temp[LOCAL_SIZE];

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

   uint count = 0;
   do
     {
      uint shift = count * ls + l;
      temp[l] = (count > 0 ? temp[l] : 0) + (shift < total ? inputs[shift] : 0);
      count++;
     }
   while((count * ls + l) < total);
   temp[l] /= (TYPE)total;
   barrier(CLK_LOCAL_MEM_FENCE);

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

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

   count = ls;
   do
     {
      count = (count + 1) / 2;
      temp[l] += (l < count ? temp[l + count] : 0);
      barrier(CLK_LOCAL_MEM_FENCE);
     }
   while(count > 1);
//---
   TYPE mean = (TYPEtemp[0];

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

После определения среднего арифметического повторим циклы и вычислим среднеквадратическое отклонение.

   count = 0;
   do
     {
      uint shift = count * ls + l;
      temp[l] = (count > 0 ? temp[l] : 0) + (shift < total ? (TYPE)pow(inputs[shift] - mean2) : 0);
      count++;
     }
   while((count * ls + l) < total);
   temp[l] /= (TYPE)total;
   barrier(CLK_LOCAL_MEM_FENCE);

   count = ls;
   do
     {
      count = (count + 1) / 2;
      temp[l] += (l < count ? temp[l + count] : 0);
      barrier(CLK_LOCAL_MEM_FENCE);
     }
   while(count > 1);
//---
   TYPE std = (TYPE)sqrt(temp[0]);
   if(l == 0)
      stds[std_shift] = std;

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

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

   count = 0;
   while((count * ls + l) < total)
     {
      uint shift = count * ls + l;
      outputs[shift] = (inputs[shift] - mean) / (std + 1e-37f);
      count++;
     }
  }

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

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

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

   const int q = get_global_id(0);
   const int units = get_global_size(0);
   int shift_value = window * q;
   int shift_score = units * q;

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

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

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

//--- Распределение градиента на Values
   for(int i = 0i < windowi ++)
     {
      TYPE grad = 0;
      for(int g = 0g < unitsg++)
         grad += scores[units * g + q] * outputs_grad[window * g + i];
      values_grad[shift_value + i] = grad;
     }

После отработки системы циклов можно считать выполненной первую часть нашей задачи — передачу градиентов ошибки во внутренний нейронный слой m_cValues.

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

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

Описанный выше алгоритм распределения градиента ошибки в матричном виде можно представить следующим образом:

  1. Градиент ошибки на уровне матрицы Score.

  1. Скорректируем градиент ошибки на производную функции Softmax.

Обобщим весь расчет в одну формулу:

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

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

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

Матричное выражение процесса корректировки на производную функции Softmax приведено выше. А для практической реализации мы создадим систему из двух вложенных циклов. Оба цикла имеют одинаковое число итераций, которое равно размеру нормализуемого вектора. В данном случае он равен количеству элементов в последовательности. При проведении операций потребуется накапливать сумму градиентов ошибки с каждого элемента нормализованного вектора. Для этого в теле внешнего цикла мы создадим частную переменную grad. Кроме того, чтобы сократить количество обращений к глобальной памяти, сохраним повторяющийся элемент в частную переменную score. Напомню, что обращение к глобальной памяти более затратное по времени. Сокращая количество обращений к буферам глобальной памяти, мы сокращаем общие затраты времени на проведение операций. В теле вложенного цикла будем производить операции умножения элементов и складывать полученные произведения в предварительно созданную частную переменную grad.

Обратите внимание, что единичную матрицу мы заменили выражением (int)(i==k). Логическое выражение даст нам истинное значение только на диагонали матрицы. Перевод логического значение в целочисленное подставит 1 вместо истинных значений и 0 для ложных. Таким образом, столь краткая запись позволяет нам получить значения единичной матрицы прямо в потоке операций без необходимости предварительного ее формирования и сохранения.

//--- Корректируем на производную 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;
     }
  }

После завершения итераций системы циклов на выходе мы получим градиенты ошибки на уровне матрицы коэффициентов зависимости скорректированные на производную функции 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,
                                          __global TYPE *querys_grad,
                                          __global TYPE *keys,
                                          __global TYPE *keys_grad,
                                          __global TYPE *scores_grad,
                                          int key_size)
  {

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

   const int q = get_global_id(0);
   const int units = get_global_size(0);
   int shift_query = key_size * q;
   int shift_score = units * q;

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

Определим значение константы и сохраним в частную переменную.

//--- Распределение градиента на Querys и Keys
   const TYPE k = 1 / sqrt((TYPE)key_size);

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

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

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

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

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

   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 + i] * scores_grad[shift_score + s];
         grad_k += querys[key_size * s + i] * scores_grad[units * s + q];
        }
      querys_grad[shift_query + i] = grad_q * k;
      keys_grad[shift_query + i] = grad_k * k;
     }
  }

На выходе из системы циклов мы получаем градиенты ошибок для двух вложенных внутренних нейронных слоев m_cQuerys и m_cKeys. То есть задача данного кернела решена. А с учетом выше рассмотренного кернела AttentionCalcScoreGradient, мы распределили градиент ошибки на все внутренние нейронные слои и дальнейшее распределение градиента ошибки до предыдущего слоя будет осуществляться уже протестированными методами внутренних нейронных слоев, как это реализовано стандартными средствами MQL5.

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

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

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

В параметрах кернела мы передаем указатели на четыре буфера данных:

  • outputs  — буфер результатов нормализации прямого прохода;
  • out_gradient — буфер градиентов на выходе блока нормализации;
  • inp_gradient — буфер для записи скорректированных градиентов;
  • stds — буфер среднеквадратических отклонений, вычисленных при прямом проходе.

Также в параметрах укажем размер буферов и смещение в буфере среднеквадратических отклонений.

__kernel void LayerNormalizeGradient(__global TYPEoutputs,
                                     __global TYPEout_gradient,
                                     __global TYPEinp_gradient,
                                     __global TYPEstds,
                                     const int total,
                                     const int std_shift)
  {
   uint i = (uint)get_global_id(0);
   uint l = (uint)get_local_id(0);

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

   uint ls = min((uint)get_local_size(0), (uint)LOCAL_SIZE);
   __local TYPE dSTD[LOCAL_SIZE];
   __local TYPE dMean1[LOCAL_SIZE];
   __local TYPE dMean2[LOCAL_SIZE];

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

   uint count = 0;
   do
     {
      uint shift = count * ls + l;
      dSTD[l] = (count > 0 ? dSTD[l] : 0) - 
                (shift < total ? out_gradient[shift] * outputs[shift] /
                (2 * (pow(stds[std_shift], (TYPE)2) + 1e-37f)) : 0);
      dMean1[l] = (count > 0 ? dMean1[l] : 0) - 
                (shift < total ? out_gradient[shift] /
                (stds[std_shift] + 1e-37f) : 0);
      dMean2[l] = (count > 0 ? dMean2[l] : 0) -
                  (shift < total ? 2 * outputs[shift] * stds[std_shift] /
                  (TYPE)total : 0);
      count++;
     }
   while((count * ls + l) < total);
   barrier(CLK_LOCAL_MEM_FENCE);

В следующем цикле соберем сумму в первых элементах массива.

   count = ls;
   do
     {
      count = (count + 1) / 2;
      dSTD[l] += (l < count ? dSTD[l + count] : 0);
      dMean1[l] += (l < count ? dMean1[l + count] : 0);
      dMean2[l] += (l < count ? dMean2[l + count] : 0);
      barrier(CLK_LOCAL_MEM_FENCE);
     }
   while(count > 1);
//---
   TYPE dstd = dSTD[0];
   TYPE dmean = dMean1[0] + dstd * dMean2[0];

Полученные значения перенесем в частные переменные. При вычислении производной среднеарифметического отклонения значение правого слагаемого умножим на производную среднеквадратического отклонения и сложим с левым слагаемым.

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

//---
   count = 0;
   while((count * ls + l) < total)
     {
      uint shift = count * ls + l;
      inp_gradient[shift] = out_gradient[shift] / (stds[std_shift] + 1e-32f) + 
                (2 * dstd * outputs[shift] * stds[std_shift]  + dmean) / total;
      count++;
     }
  }

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

Для начала добавим константы для работы с кернелами в файл defines.mqh. Нам нужно добавить константы идентификации самих кернелов и их переменных. Для именования констант используем ранее оговоренные правила, которые применяются ко всем константам в рамках нашего проекта:

  • все константы начинаются с префикса def;
  • кернелы начинаются с префикса def_k;
  • константы параметров после префикса def содержат указатель на кернел.

#define def_k_AttentionFeedForward     28
#define def_k_AttentionScoreGradients  29
#define def_k_AttentionHiddenGradients 30
#define def_k_Sum                      31
#define def_k_LayerNormalize           32
#define def_k_LayerNormalizeGradient   33

//--- прямой проход блока внимания
#define def_attff_querys               0
#define def_attff_keys                 1
#define def_attff_scores               2
#define def_attff_values               3
#define def_attff_outputs              4
#define def_attff_window               5
#define def_attff_key_size             6

//--- определение градиента на матрице коэффициентов зависимости блока внимания
#define def_attscr_scores              0
#define def_attscr_scores_grad         1
#define def_attscr_values              2
#define def_attscr_values_grad         3
#define def_attscr_outputs_grad        4
#define def_attscr_scores_temp         5
#define def_attscr_window              6

//--- распределение градиента через блок внимания
#define def_atthgr_querys              0
#define def_atthgr_querys_grad         1
#define def_atthgr_keys                2
#define def_atthgr_keys_grad           3
#define def_atthgr_scores_grad         4
#define def_atthgr_key_size            5

//--- сумма векторов 
#define def_sum_inputs1                0
#define def_sum_inputs2                1
#define def_sum_outputs                2

//--- нормализация вектора
#define def_layernorm_inputs           0
#define def_layernorm_outputs          1
#define def_layernorm_std              2
#define def_layernorm_vector_size      3
#define def_layernorm_std_shift        4

//--- градиент нормализации вектора
#define def_layernormgr_outputs        0
#define def_layernormgr_out_grad       1
#define def_layernormgr_inp_grad       2
#define def_layernormgr_std            3
#define def_layernormgr_vector_size    4
#define def_layernormgr_std_shift      5

После этого нам предстоит добавить объявление новых кернелов в код основного класса-диспетчера нейронной сети. Как и всех ранее созданных кернелов, объявление новых кернелов мы добавим в метод CNet::InitOpenCL. В нем мы сначала изменим общее количество используемых в программе кернелов.

   if(!m_cOpenCL.SetKernelsCount(34))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

После этого объявим и сами кернелы.

   if(!m_cOpenCL.KernelCreate(def_k_AttentionFeedForward,
                                        "AttentionFeedForward"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

   if(!m_cOpenCL.KernelCreate(def_k_AttentionScoreGradients,
                                        "AttentionCalcScoreGradient"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

   if(!m_cOpenCL.KernelCreate(def_k_AttentionHiddenGradients,
                                        "AttentionCalcHiddenGradient"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

   if(!m_cOpenCL.KernelCreate(def_k_Sum"Sum"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

   if(!m_cOpenCL.KernelCreate(def_k_LayerNormalize"LayerNormalize"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

   if(!m_cOpenCL.KernelCreate(def_k_LayerNormalizeGradient
                                             "LayerNormalizeGradient"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

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

Первым дополним метод прямого прохода CNeuronAttention::FeedForward. В этом методе нам предстоит организовать процедуру вызова кернела прямого прохода AttentionFeedForward. Мы уже не раз создавали аналогичные процессы. Но все же напомню алгоритм:

  1. Проверка наличия буферов данных в контексте OpenCL;
  2. Передача параметров кернелу, в том числе указатели на буферы данных;
  3. Постановка кернела в очередь выполнения операций.

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

bool CNeuronAttention::FeedForward(CNeuronBase *prevLayer)
  {
//--- вычисление векторов Query, Key, Value
   .....
//--- Разветвление алгоритма по вычислительному устройству
   MATRIX out;
   if(!m_cOpenCL)
     {
   // Блок MQL5
   .....
     }
   else // Блок OpenCL
     {
      //--- проверка буферов данных
      if(m_cQuerys.GetOutputs().GetIndex() < 0)
         return false;
      if(m_cKeys.GetOutputs().GetIndex() < 0)
         return false;
      if(m_cValues.GetOutputs().GetIndex() < 0)
         return false;
      if(m_cScores.GetIndex() < 0)
         return false;
      if(m_cAttentionOut.GetOutputs().GetIndex() < 0)
         return false;

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

      //--- передача параметров кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForwarddef_attff_keys,
                                                   m_cKeys.GetOutputs().GetIndex()))
         return false;

      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForwarddef_attff_outputs,
                                            m_cAttentionOut.GetOutputs().GetIndex()))
         return false;

      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForwarddef_attff_querys,
                                                  m_cQuerys.GetOutputs().GetIndex()))
         return false;

      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForwarddef_attff_scores,
                                                               m_cScores.GetIndex()))
         return false;

      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionFeedForwarddef_attff_values,
                                                  m_cValues.GetOutputs().GetIndex()))
         return false;

      if(!m_cOpenCL.SetArgument(def_k_AttentionFeedForwarddef_attff_key_size,
                                                                        m_iKeysSize))
         return false;

      if(!m_cOpenCL.SetArgument(def_k_AttentionFeedForwarddef_attff_window,
                                                                          m_iWindow))
         return false;

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

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

На этом мы выполнили алгоритм запуска кернела блока Self-Attention. Но нам еще предстоит сложить содержимое двух буферов и нормализовать данные буфера результатов. Следуя алгоритму, сначала найдем сумму двух векторов (исходных данных и результатов Self-Attention). Эта операция довольно общая и может широко использоваться за пределами нашего класса нейронного слоя внимания CNeuronAttention. Поэтому было принято решение добавить ее отдельным методом в класс буфера данных CBufferType::SumArray.

В параметрах методу SumArray будем передавать указатель на добавляемый буфер. Сразу в теле метода проверим полученный указатель и размер полученного буфера. Для успешного выполнения операции размер текущего буфера, который будет выступать первым слагаемым, и полученного буфера (второго слагаемого) должны быть равны.

bool CBufferType::SumArray(CBufferType *src)
  {
//--- проверка массива исходных данных
   if(!src || src.Total() != Total())
      return false;

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

   if(!m_cOpenCL)
     {
      //--- изменение размера матрицы
      MATRIX temp = src.m_mMatrix;
      if(!temp.Reshape(Rows(), Cols()))
         return false;
      //--- сложения матриц
      m_mMatrix += temp;
     }

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

За блоком контролей следует передача параметров кернелу и постановка его в очередь выполнения.

   else
     {
      if(src.GetIndex() < 0 && !BufferCreate(m_cOpenCL))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_Sumdef_sum_inputs1m_myIndex))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_Sumdef_sum_inputs2src.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_Sumdef_sum_outputsm_myIndex))
         return false;
      uint off_set[] = {0};
      uint NDRange[] = {(uint)Total()};
      if(!m_cOpenCL.Execute(def_k_Sum1off_setNDRange))
         return false;
     }
//---
   return true;
  }

Процесс нормализации данных организован в методе CNeuronAttention::NormlizeBuffer. Но при соблюдении общих правил построения алгоритма в данном методе есть два исключения. Во-первых, мы исключили блок проверки наличия буферов в контексте OpenCL. В данном случае риск использование незагруженных буферов минимален. Дело в том, что до вызова данного метода используемые буферы были уже проверены не один раз, и повторная проверка будет лишней.

bool CNeuronAttention::NormlizeBuffer(CBufferType *buffer,
                                      CBufferType *std
                                      uint std_shift)
  {
   if(!m_cOpenCL)
     {
    // Блок MQL5
   .....
     }
   else
     {
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LayerNormalize,
                                     def_layernorm_inputsbuffer.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LayerNormalize,
                                    def_layernorm_outputsbuffer.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LayerNormalize,
                                           def_layernorm_stdstd.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_LayerNormalize,
                              def_layernorm_vector_size, (int)buffer.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_LayerNormalize,
                                          def_layernorm_std_shiftstd_shift))
         return false;

Второй момент связан с использованием локального массива данных и синхронизацией потоков. Дело в том, что синхронизация потоков доступна только в пределах work group. Нам нужно явно указать ее размер. Алгоритм кернела нормализации построен таким образом, что размер рабочей группы не может быть больше размера локального массива. Напомню, размер локального массива определяется константой LOCAL_SIZE. В то же время количество потоков не может быть больше размера нормализуемого буфера. Следовательно, в массиве указания размерности пространства задач мы укажем меньшее из двух величин. Так как мы нормализуем значения всего буфера одним пакетом, то и размерность глобального и локального пространства задач будет одна.

Определившись с размерностью задач, мы ставим кернел в очередь выполнения.

      int NDRange[] = {(int)MathMin(buffer.Total(), LOCAL_SIZE)};
      int off_set[] = {0};
      if(!m_cOpenCL.Execute(def_k_LayerNormalize1off_setNDRangeNDRange))
         return false;
     }
//---
   return true;
  }

На этом завершается блок использования технологии OpenCL в методе прямого прохода нашего класса механизма внимания, и мы завершаем работу над этим методом. Далее его код остается без изменений. Полностью код приведен в разделе описания построения метода стандартными средствами MQL5.

Мы переходим к работе над одним из методов обратного прохода — методом распределения градиента ошибки через скрытый слой CNeuronAttention::CalcHiddenGradient. Алгоритм наших действий остается прежний. Сделаем лишь поправку на использование двух кернелов последовательно.

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

class CNeuronAttention    :  public CNeuronBase
  {
protected:
   .....
   int               m_cScoreGrad;
   int               m_cScoreTemp;
   .....
  };

Только в данном случае мы не будем объявлять экземпляры объектов буфера в основной памяти. Мы не будем использовать указанные буферы для обмена данными между контекстом OpenCL и основной программой. Они нужны лишь для временного хранения данных их передачи между кернелами. А значит, нам достаточно их наличие в памяти контекста OpenCL. В основной программе мы лишь объявим переменные для хранения указателей на буферы.

Но вернемся к работе над методом CNeuronAttention::CalcHiddenGradient. Вначале проверяем наличие и при необходимости создаем новые буферы данных в контексте OpenCL, используемые в первом кернеле. Мы намеренно не создаем сразу буферы данных для второго кернела, чтобы использование памяти было более рационально. Это позволит нам использовать буферы данных большего размера при ограниченных ресурсах памяти контекста OpenCL.

bool CNeuronAttention::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
   .....
//--- разветвление алгоритма по вычислительному устройству
   if(!m_cOpenCL)
     {
   // Блок MQL5
   .....
     }
   else // блок OpenCL
     {
      //--- проверка буферов данных
      if(m_cValues.GetOutputs().GetIndex() < 0)
         return false;
      if(m_cValues.GetGradients().GetIndex() < 0)
         return false;
      if(m_cScores.GetIndex() < 0)
         return false;
      if(m_cAttentionOut.GetGradients().GetIndex() < 0)
         return false;
      if(m_cScoreGrad < 0)
         return false;
      if(m_cScoreTemp < 0)
         return false;

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

      //--- передача параметров кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionScoreGradients
              def_attscr_outputs_gradm_cAttentionOut.GetGradients().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionScoreGradients,
                                         def_attscr_scoresm_cScores.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionScoreGradients,
                                            def_attscr_scores_gradm_cScoreGrad))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionScoreGradients,
                                            def_attscr_scores_tempm_cScoreTemp))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionScoreGradients,
                            def_attscr_valuesm_cValues.GetOutputs().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionScoreGradients,
                     def_attscr_values_gradm_cValues.GetGradients().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_AttentionScoreGradients,
                                                    def_attscr_windowm_iWindow))
         return false;

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

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

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

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

      if(m_cQuerys.GetOutputs().GetIndex() < 0)
         return false;
      if(m_cQuerys.GetGradients().GetIndex() < 0)
         return false;
      if(m_cKeys.GetOutputs().GetIndex() < 0)
         return false;
      if(m_cKeys.GetGradients().GetIndex() < 0)
         return false;

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

      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionHiddenGradients
                                 def_atthgr_keysm_cKeys.GetOutputs().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionHiddenGradients,
                          def_atthgr_keys_gradm_cKeys.GetGradients().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionHiddenGradients,
                             def_atthgr_querysm_cQuerys.GetOutputs().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionHiddenGradients,
                      def_atthgr_querys_gradm_cQuerys.GetGradients().GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_AttentionHiddenGradients,
                                             def_atthgr_scores_gradm_cScoreGrad))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_AttentionHiddenGradients,
                                                 def_atthgr_key_sizem_iKeysSize))
         return false;

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

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

       if(!m_cOpenCL.Execute(def_k_AttentionHiddenGradients1off_setNDRange))
         return false;

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