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

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

  • создание программы OpenCL;
  • внесение изменений в методы основной программы для организации обмена данными с контекстом и вызова программы OpenCL.

Начинаем работу с создания программы OpenCL. Сначала реализуем кернел прямого прохода BatchNormFeedForward. В параметрах кернелу будем передавать указатели на четыре буфера и две константы:

  • inputs — буфер исходных данных (результатов предыдущего слоя);
  • options — буфер параметров нормализации;
  • weights — буфер матрицы обучаемых параметров (назван по аналогии с буфером класса);
  • output — буфер результатов;
  • batch — размер пакета нормализации;
  • total — размер буфера результатов.

__kernel void BatchNormFeedForward(__global TYPE *inputs,
                                   __global TYPE *options,
                                   __global TYPE *weights,
                                   __global TYPE *output,
                                   int batch,
                                   int total)
  {

Последний параметр необходим, потому что для оптимизации процесса вычислений мы используем векторные переменные типа TYPE4. Подобный подход позволяет распараллелить вычисления не на программном уровне, а на уровне микропроцессора. Использование вектора из четырех элементов типа double позволяет полностью заполнить 256-битный регистр микропроцессора и за один такт произвести вычисления над всем вектором. Таким образом, за один такт микропроцессора мы осуществляем операции над четырьмя элементами нашего массива данных. OpenCL поддерживает векторные переменные из 2, 3, 4, 8 и 16 элементов. Перед выбором размерности вектора ознакомьтесь с техническими характеристиками вашего оборудования.

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

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

   int n = get_global_id(0);
   if(batch <= 1)
     {
      D4ToArray(outputToVect4(inputsn * 41total0), n * 41total0);
      return;
     }

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

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

   int shift = n * 4;
   int shift_options = n * 3 * 4;
   int shift_weights = n * 2 * 4;

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

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

   TYPE4 inp = ToVect4(inputsshift1total0);
   TYPE4 mean = ToVect4(optionsshift3total * 30) * ((TYPE)batch - 1) + inp ;
   if(options[shift_options ] != 0 && options[shift_options + 1] > 0)
      mean /= (TYPE4)batch;
   TYPE4 delt = inp - mean;
   TYPE4 variance = ToVect4(optionsshift3total * 31) * ((TYPE)batch - 1) + pow(delt2);
   if(options[shift_options + 1] > 0)
      variance /= (TYPE4)batch;

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

   TYPE4 nx = delt / sqrt(variance + 1e-37f);

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

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

   if(weights[shift_weights] == 0)
      D4ToArray(weights, (TYPE4)1shift2total * 20);

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

После такой нехитрой операции мы можем смело осуществить масштабирование и сдвиг.

   TYPE4 res = ToVect4(weightsshift2total * 20) * nx + 
               ToVect4(weightsshift2total * 21);

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

   D4ToArray(optionsmeanshift3total * 30);
   D4ToArray(optionsvarianceshift3total * 31);
   D4ToArray(optionsnxshift3total * 32);
   D4ToArray(outputresshift1total0);
  }

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

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

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

  • inputs — буфер исходных данных (результатов предыдущего слоя);
  • options — буфер параметров нормализации;
  • weights — буфер матрицы обучаемых параметров (назван по аналогии с буфером класса);
  • gradient — буфер градиентов ошибки на уровне результатов текущего слоя;
  • gradient_inputs — буфер градиентов ошибки на уровне результатов предыдущего слоя (в данном случае результат работы кернела);
  • batch — размер пакета нормализации;
  • total — размер буфера результатов.

__kernel void BatchNormCalcHiddenGradient(__global TYPE *options,
                                          __global TYPE *gradient,
                                          __global TYPE *inputs,
                                          __global TYPE *gradient_inputs,
                                          __global TYPE *weights,
                                          int batch,
                                          int total)
  {

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

   int n = get_global_id(0);
   int shift = n * 4;
   if(batch <= 1)
     {
      D4ToArray(gradient_inputsToVect4(gradientshift1total0),
                                                   shift1total0);
      return;
     }

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

   TYPE4 inp = ToVect4(inputsshift1total0);
   TYPE4 gnx = ToVect4(gradientshift1total0) * 
               ToVect4(weightsshift2total * 20);
   TYPE4 temp = 1 / sqrt(ToVect4(optionsshift3total * 31) + 1e-37f);
   TYPE4 delt = inp - ToVect4(optionsshift3total * 30);
   TYPE4 gvar = delt / (-2 * pow(ToVect4(optionsshift3total * 31) +
                                              1.0e-37f3.0f / 2.0f)) * gnx;
   TYPE4 gmu = (-temp) * gnx - gvar * 2 * delt / (TYPE4)batch;
   TYPE4 gx = temp * gnx + gmu/(TYPE4)batch + gvar * 2 * delt/(TYPE4)batch;

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

   D4ToArray(gradient_inputsgxshift1total0);
  }

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

В параметрах данному кернелу будем передавать три буфера данных:

  • options — буфер параметров нормализации;
  • delta_weights — буфер градиентов ошибки на уровне матрицы обучаемых параметров (в данном случае результат работы кернела);
  • gradient — буфер градиентов ошибки на уровне результатов текущего слоя.

__kernel void BatchNormCalcDeltaWeights(__global TYPE *options,
                                        __global TYPE *delta_weights,
                                        __global TYPE *gradients)
  {

В данном кернеле нужно реализовать только две математические формулы:

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

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

   const int n = get_global_id(0);
   int shift_options = n * 3;
   int shift_weights = n * 2;

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

   TYPE grad = gradients[n];
   delta_weights[shift_weights] += grad * options[shift_options + 2];
   delta_weights[shift_weights + 1] += grad;
  }

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

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

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

#define def_k_BatchNormFeedForward        37
#define def_k_BatchNormCalcHiddenGradient 38
#define def_k_BatchNormCalcDeltaWeights   39

А затем добавляем идентификаторы параметров кернелов.

//--- прямой проход пакетной нормализации
#define def_bnff_inputs                0
#define def_bnff_options               1
#define def_bnff_weights               2
#define def_bnff_outputs               3
#define def_bnff_batch                 4
#define def_bnff_total                 5

//--- распределение градиента через слой пакетной нормализации
#define def_bnhgr_options              0
#define def_bnhgr_gradient             1
#define def_bnhgr_inputs               2
#define def_bnhgr_gradient_inputs      3
#define def_bnhgr_weights              4
#define def_bnhgr_batch                5
#define def_bnhgr_total                6

//--- распределение градиента до оптимизируемых параметров пакетной нормализации
#define def_bndelt_options             0
#define def_bndelt_delta_weights       1
#define def_bndelt_gradient            2

На следующем этапе надо инициализировать новые кернелы в программе. Для этого переходим в метод инициализации программы OpenCL основного класса-диспетчера нашей модели CNet::InitOpenCL. Сначала изменяем общее количество используемых кернелов.

   if(!m_cOpenCL.SetKernelsCount(40))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   if(!m_cOpenCL.KernelCreate(def_k_BatchNormFeedForward,
                                   "BatchNormFeedForward"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   if(!m_cOpenCL.KernelCreate(def_k_BatchNormCalcHiddenGradient,
                                   "BatchNormCalcHiddenGradient"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   if(!m_cOpenCL.KernelCreate(def_k_BatchNormCalcDeltaWeights,
                                   "BatchNormCalcDeltaWeights"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

Теперь, когда кернелы созданы, а мы можем к ним обращаться, переходим к работе с методами нашего класса пакетной нормализации.

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

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

bool CNeuronBatchNorm::FeedForward(CNeuronBase *prevLayer)
  {
   ......
//--- разветвление алгоритма по вычислительному устройству
   if(!m_cOpenCL)
     {
//--- Реализация средствами MQL5
   ......
     }
   else  // Блок OpenCL
     {
      //--- проверка буферов данных
      CBufferType *inputs = prevLayer.GetOutputs();
      if(inputs.GetIndex() < 0)
         return false;
      if(m_cBatchOptions.GetIndex() < 0)
         return false;
      if(m_cWeights.GetIndex() < 0)
         return false;
      if(m_cOutputs.GetIndex() < 0)
         return false;

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

      //--- передача параметров кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormFeedForward,
                                         def_bnff_inputsinputs.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormFeedForward,
                                    def_bnff_weightsm_cWeights.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormFeedForward,
                               def_bnff_optionsm_cBatchOptions.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormFeedForward
                                    def_bnff_outputsm_cOutputs.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_BatchNormFeedForward,
                                    def_bnff_total, (int)m_cOutputs.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_BatchNormFeedForward,
                                               def_bnff_batchm_iBatchSize))
         return false;

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

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

     //--- постановка в очередь выполнения
      uint off_set[] = {0};
      uint NDRange[] = { (int)(m_cOutputs.Total() + 3) / 4 };
      if(!m_cOpenCL.Execute(def_k_BatchNormFeedForward1off_setNDRange))
         return false;
     }
//---
   if(!m_cActivation.Activation(m_cOutputs))
      return false;
//---
   return true;
  }

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

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

bool CNeuronBatchNorm::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
   ......
//--- разветвление алгоритма по вычислительному устройству
   if(!m_cOpenCL)
     {
//--- Реализация средствами MQL5
   ......
     }
   else  // блок OpenCL
     {
      //--- проверка буферов данных
      CBufferTypeinputs = prevLayer.GetOutputs();
      CBufferTypeinputs_grad = prevLayer.GetGradients();
      if(inputs.GetIndex() < 0)
         return false;
      if(m_cBatchOptions.GetIndex() < 0)
         return false;
      if(m_cWeights.GetIndex() < 0)
         return false;
      if(m_cOutputs.GetIndex() < 0)
         return false;
      if(m_cGradients.GetIndex() < 0)
         return false;
      if(inputs_grad.GetIndex() < 0)
         return false;

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

      //--- передача параметров кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormCalcHiddenGradient
                                             def_bnhgr_inputsinputs.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormCalcHiddenGradient,
                                        def_bnhgr_weightsm_cWeights.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormCalcHiddenGradient,
                                   def_bnhgr_optionsm_cBatchOptions.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormCalcHiddenGradient,
                                     def_bnhgr_gradientm_cGradients.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_BatchNormCalcHiddenGradient,
                               def_bnhgr_gradient_inputsinputs_grad.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_BatchNormCalcHiddenGradient,
                                        def_bnhgr_total, (int)m_cOutputs.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_BatchNormCalcHiddenGradient,
                                                   def_bnhgr_batchm_iBatchSize))
         return false;

После передачи всех параметров мы готовим кернел к постановке в очередь выполнения. Напомню, что при создании кернела мы определили использование векторных операций с типом TYPE4. Соответственно, мы сокращаем в четыре раза количество выполняемых потоков. Вызываем метод постановки кернела в очередь.

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

На этом мы завершаем работу с методом распределения градиента ошибки через скрытый слой CNeuronBatchNorm::CalcHiddenGradient. Нам остается повторить операции для второго метода обратного прохода CNeuronBatchNorm::CalcDeltaWeights.

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

bool CNeuronBatchNorm::CalcDeltaWeights(CNeuronBase *prevLayerbool read)
  {
   ......
//--- разветвление алгоритма по вычислительному устройству
   if(!m_cOpenCL)
     {
//--- Реализация средствами MQL5
   ......
     }
   else
     {
      //--- проверка буферов данных
      if(m_cBatchOptions.GetIndex() < 0)
         return false;
      if(m_cGradients.GetIndex() < 0)
         return false;
      if(m_cDeltaWeights.GetIndex() < 0)
         return false;

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

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

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

      //--- постановка в очередь выполнения
      int off_set[] = {0};
      int NDRange[] = {(int)m_cOutputs.Total()};
      if(!m_cOpenCL.Execute(def_k_BatchNormCalcDeltaWeights1off_setNDRange))
         return false;
      if(read && !m_cDeltaWeights.BufferRead())
         return false;
     }
//---
   return true;
  }

На этом мы заканчиваем работу с нашим классом пакетной нормализации данных CNeuronBatchNorm. Он готов к использованию — сейчас в нем полностью реализован алгоритм нормализации данных. При этом мы повторили алгоритм в двух вариантах: стандартными средствами MQL5 и с использованию технологии многопоточных вычислений OpenCL. Тем самым мы предоставляем пользователю право выбора используемой технологии в соответствии с его требованиями.

Сейчас я предлагаю посмотреть на реализацию метода пакетной нормализации на языке Python.