Учебник по нейросетямРазделы
Нейросети в алготрейдинге на MQL5Базовые типы нейронных слоёвРекуррентные нейронные сетиОрганизация параллельных вычислений средствами OpenCL 
Скачать одним файлом:
Книга по нейросетям в PDF
Книга по нейросетям в CHM

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

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

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

  • поток консолидации и обработки данных от внутренних нейронных слоев в рамках прямого прохода;
  • распределение градиента ошибки от выхода LSTM блока до внутренних нейронных слоев в рамках обратного прохода.

Это дает нам понимание задания.

Мы уже имеем реализацию процесса средствами MQL5. Это дает понимание процесса и алгоритма выполнения операций.

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

  1. Программа выполняемых операций.
  2. Исходные данные для выполнения операций.
  3. Команды управления процессом (момент запуска программы, количество создаваемых потоков и т.д.)

Давайте рассмотрим выполнение этих пунктов.

 

4.Внесение дополнений в программу OpenCL

Первым пунктом у нас указана программы выполняемых операций. Это значит, что нам необходимо дополнить нашу программу OpenCL новыми кернелами для выполнения требуемых нам дополнительных операций. Весь код программы OpenCL мы собрали в файле opencl_program.cl. Открываем данный файл и дополняем его двумя новыми кернелами: LSTMFeedForward и LSTMCalcHiddenGradient. Названия кернелов созвучны с названиями методов наших классов. Поэтому несложно догадаться, что первый будет дополнять метод прямого прохода, а второй — метод распределения градиента ошибки.

Схеиа рекуррентного LSTM блока

Схема рекуррентного LSTM блока

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

  • forgetgate — указатель на буфер врат забвения (исходные данные);
  • inputgate — указатель на буфер входных врат  (исходные данные);
  • outputgate — указатель на буфер врат результатов  (исходные данные);
  • newcontent — указатель на буфер нового контента  (исходные данные);
  • memory — указатель на поток памяти (буфер результатов);
  • hiddenstate — указатель на поток скрытого состояния (буфер результатов);
  • outputs_total — количество элементов в потоке данных (константа).

__kernel void LSTMFeedForward(__global TYPE *forgetgate,
                              __global TYPE *inputgate,
                              __global TYPE *outputgate,
                              __global TYPE *newcontent,
                              __global TYPE *memory,
                              __global TYPE *hiddenstate,
                              int outputs_total)

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

  {
   const int n = get_global_id(0);
   const int shift = n * 4;

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

   TYPE4 fg = ToVect4(forgetgate, shift, 1, outputs_total, 0);
   TYPE4 ig = ToVect4(inputgate, shift, 1, outputs_total, 0);
   TYPE4 og = ToVect4(outputgate, shift, 1, outputs_total, 0);
   TYPE4 nc = ToVect4(newcontent, shift, 1, outputs_total, 0);
   TYPE4 mem = ToVect4(memory, shift, 1, outputs_total, 0);

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

   TYPE4 temp = mem * fg;
   temp += ig * nc;
   D4ToArray(memory, temp, shift, 1, outputs_total, 0);

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

   temp = tanh(temp) * og;
   D4ToArray(hiddenstate, temp, shift, 1, outputs_total, 0);
  }

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

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

  • outputs — указатель на буфер вектора результатов (исходные данные);
  • gradients — указатель на буфер вектора градиентов текущего слоя (исходные данные);
  • inputgate — указатель на буфер входных врат (исходные данные);
  • outputgate — указатель на буфер врат результатов (исходные данные);
  • newcontent — указатель на буфер нового контента (исходные данные);
  • memory — указатель на поток памяти (исходные данные);
  • fg_gradients — указатель на буфер градиентов врат забвения (буфер результатов);
  • ig_gradients — указатель на буфер градиентов входных врат (буфер результатов);
  • og_gradients — указатель на буфер градиентов врат результатов (буфер результатов);
  • nc_gradients — указатель на буфер градиентов нового контента (буфер результатов);
  • outputs_total — количество элементов в потоке данных (константа).

__kernel void LSTMCalcHiddenGradient(__global TYPE *outputs,
                                     __global TYPE *gradients,
                                     __global TYPE *inputgate,
                                     __global TYPE *outputgate,
                                     __global TYPE *newcontent,
                                     __global TYPE *memory,
                                     __global TYPE *fg_gradients,
                                     __global TYPE *ig_gradients,
                                     __global TYPE *og_gradients,
                                     __global TYPE *nc_gradients,
                                     int outputs_total)

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

  {
   const int n = get_global_id(0);
   int shift = n * 4;

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

   TYPE4 out = ToVect4(outputs, shift, 1, outputs_total, 0);
   TYPE4 grad = ToVect4(gradients, shift, 1, outputs_total, 0);
   TYPE4 ig = ToVect4(inputgate, shift, 1, outputs_total, 0);
   TYPE4 og = ToVect4(outputgate, shift, 1, outputs_total, 0);
   TYPE4 nc = ToVect4(newcontent, shift, 1, outputs_total, 0);
   TYPE4 mem = ToVect4(memory, shift, 1, outputs_total, 0);

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

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

   TYPE4 m = out / (og + 1.0e-37f);

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

//--- OutputGate градиент
   TYPE4 temp = grad * m;
   D4ToArray(og_gradients, temp, shift, 1, outputs_total, 0);

Далее мы должны аналогичным путем определить градиент ошибки на другом множителе — нормализованное состояние памяти. То есть умножаем градиент ошибки на выходе нашего рекуррентного блока на состояние врат результатов.

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

//--- Градиент памяти корректируем на производную TANH
   grad = grad * og * (1 - pow(m, 2));

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

//--- InputGate градиент
   temp = grad * nc;
   D4ToArray(ig_gradients, temp, shift, 1, outputs_total, 0);
//--- NewContent градиент
   temp = grad * ig;
   D4ToArray(nc_gradients, temp, shift, 1, outputs_total, 0);
//--- ForgetGates градиент
   temp = grad * mem;
   D4ToArray(fg_gradients, temp, shift, 1, outputs_total, 0);
  }

После выполнения операций выходим из кернела.

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

 

4.Реализация функционала на стороне основной программы

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

#define def_k_LSTMFeedForward          26
#define def_k_LSTMHiddenGradients      27

//--- LSTM Feed Forward
#define def_lstmff_forgetgate          0
#define def_lstmff_inputgate           1
#define def_lstmff_outputgate          2
#define def_lstmff_newcontent          3
#define def_lstmff_memory              4
#define def_lstmff_hiddenstate         5
#define def_lstmff_outputs_total       6

При добавлении констант соблюдаем определенные ранее правила именования. Все константы кернелов начинаются с приставки def_k_, а константы параметров содержат аббревиатуру кернела: def_lstmff_ для параметров кернела прямого прохода и def_lstmhgr_ у параметров кернела распределения градиента.

//--- LSTM Hidden Gradients
#define def_lstmhgr_outputs            0
#define def_lstmhgr_gradients          1
#define def_lstmhgr_inputgate          2
#define def_lstmhgr_outputgate         3
#define def_lstmhgr_newcontent         4
#define def_lstmhgr_memory             5
#define def_lstmhgr_fg_gradients       6
#define def_lstmhgr_ig_gradients       7
#define def_lstmhgr_og_gradients       8
#define def_lstmhgr_nc_gradients       9
#define def_lstmhgr_outputs_total      10

Затем мы переходим в файл neuronnet.mqh, который содержит код нашего класса нейронной сети. Находим метод CNet::InitOpenCL — в нем нам надо изменить количество используемых кернелов и одновременно открытых буферов.

   if(!m_cOpenCL.SetKernelsCount(28))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   if(!m_cOpenCL.SetBuffersCount(10))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

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

Далее мы объявим кернелы для использования в рамках нашей программы. При этом не забываем контролировать процесс выполнения операций.

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

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

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

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

bool CNeuronLSTM::FeedForward(CNeuronBase *prevLayer)
  {
   ....   
//--- Разветвление алгоритма по вычислительному устройству
   CBufferType *fg = m_cForgetGate.GetOutputs();
   CBufferType *ig = m_cInputGate.GetOutputs();
   CBufferType *og = m_cOutputGate.GetOutputs();
   CBufferType *nc = m_cNewContent.GetOutputs();
   if(!m_cOpenCL)
     {
     // Здесь пропущен Блок MQL5
     }
   else // Блок работы с OpenCL
     {
      //--- проверяем буферы
      if(fg.GetIndex() < 0 || ig.GetIndex() < 0 || og.GetIndex() < 0 ||
         nc.GetIndex() < 0 || memory.GetIndex() < 0 || hidden.GetIndex() < 0)
         return false;

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

      //--- передаем параметры кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMFeedForward, def_lstmff_forgetgate,
                                                                      fg.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMFeedForward, def_lstmff_inputgate,
                                                                      ig.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMFeedForward, def_lstmff_newcontent,
                                                                      nc.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMFeedForward, def_lstmff_outputgate,
                                                                      og.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMFeedForward, def_lstmff_memory,
                                                                  memory.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMFeedForward, def_lstmff_hiddenstate,
                                                                  hidden.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_LSTMFeedForward, def_lstmff_outputs_total,
                                                                 m_cOutputs.Total()))
         return false;

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

Рассчитанное количество потоков запишем в массив NDRange, а нулевое смещение в буферах данных укажем в массиве off_set. Поставим наш кернел в очередь выполнения. В случае возникновения ошибки постановки кернела в очередь выполнения функция m_cOpenCL.Execute вернет результат false, который мы должны проверить и обработать.

      //--- запуск кернела
      int NDRange[] = {(int)(m_cOutputs.Total() + 3) / 4};
      int off_set[] = {0};
      if(!m_cOpenCL.Execute(def_k_LSTMFeedForward, 1, off_set, NDRange))
         return false;
     }

На этом работа над методом прямого прохода LSTM-блока завершена. Переходим к внесению дополнений в метод распределения градиента ошибки.

Как и в случае прямого прохода, работу в методе распределения градиента ошибки CNeuronLSTM::CalcHiddenGradient мы начнем с проверки наличия исходных данных в памяти контекста OpenCL.

bool CNeuronLSTM::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
   ....   
      //--- Разветвление алгоритма по вычислительному устройству
      if(!m_cOpenCL)
        {
         // Здесь пропущен Блок MQL5
        }

      else // Блок работы с OpenCL
        {
         //--- проверяем буферы
         if(hidden.GetIndex() < 0)
            return false;
         if(m_cGradients.GetIndex() < 0)
            return false;
         if(ig.GetIndex() < 0)
            return false;
         if(og.GetIndex() < 0)
            return false;
         if(nc.GetIndex() < 0)
            return false;
         if(memory.GetIndex() < 0)
            return false;
         if(fg_grad.GetIndex() < 0)
            return false;
         if(ig_grad.GetIndex() < 0)
            return false;
         if(og_grad.GetIndex() < 0)
            return false;
         if(nc_grad.GetIndex() < 0)
            return false;

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

         //--- передаем параметры кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                    def_lstmhgr_fg_gradients, fg_grad.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                  def_lstmhgr_gradients, m_cGradients.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                    def_lstmhgr_ig_gradients, ig_grad.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                            def_lstmhgr_inputgate, ig.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                           def_lstmhgr_memory, memory.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients, 
                                    def_lstmhgr_nc_gradients, nc_grad.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                           def_lstmhgr_newcontent, nc.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                    def_lstmhgr_og_gradients, og_grad.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                           def_lstmhgr_outputgate, og.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LSTMHiddenGradients,
                                          def_lstmhgr_outputs, hidden.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_LSTMHiddenGradients,
                                   def_lstmhgr_outputs_total, m_cOutputs.Total()))
            return false;

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

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

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

         //--- запуск кернела
         int NDRange[] = { (int)(m_cOutputs.Total() + 3) / 4 };
         int off_set[] = {0};
         if(!m_cOpenCL.Execute(def_k_LSTMHiddenGradients, 1, off_set, NDRange))
            return false;
        }

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

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

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

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