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

Мы продолжаем реализацию технологи Dropout. В предыдущих разделах мы уже полностью реализовали алгоритм работы данной технологии стандартными средствами MQL5. Теперь переходим к реализации алгоритма с использованием возможности многопоточных операций на GPU средствами OpenCL. В рамках этой книги мы уже не раз выполняли подобную операцию. Но все же напомню, что для ее реализации нам предстоит выполнить работу в двух направлениях. Вначале мы создаем программу OpenCL, а затем нам необходимо выполнить работу на стороне основной программы — организовать обмен данными между основной программой и контекстом OpenCL, в котором будет выполняться программа, а также осуществить вызов самой программы OpenCL.

Как всегда, данную работу начинаем с создания программы OpenCL. В данном случае нам не придется писать много кода на стороне OpenCL-программы. Более того, для реализации и прямого, и обратного прохода мы будем использовать один и тот же кернел. Как такое стало возможным? Давайте вспомним, какие операции нам надо реализовать.

При прямом проходе мы осуществляем маскирование данных. Непосредственно вектор-маску мы создаем средствами MQL5 на стороне основной программы. Здесь же нам предстоит замаскировать исходные данные. Для этого будем поэлементно умножать тензор исходных данных на вектор-маску.

Следовательно, для прямого прохода нужно создать кернел для поэлементного умножения двух тензоров одинакового размера.

В процессе обратного прохода необходимо провести градиент ошибки через операцию маскирования. Давайте внимательно посмотри на формулу операции маскирования. 1/q представляет собой константу, которая определяется на стадии инициализации класса и не меняется в течение всего процесса обучения и эксплуатации модели. xi — элемент вектора маскирования, который может принимать только два значения: 1 или 0. Следовательно, весь процесс маскирования можно представить как умножение некоего исходного значения на константу. Как вы знаете, производная от такой операции является константа, на которую осуществлялось умножением.

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

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

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

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

__kernel void MaskMult(__global TYPE *inputs,
                       __global TYPE *mask,
                       __global TYPE *outputs,
                       int outputs_total)
  {
   const int n = get_global_id(0) * 4;
//---

TYPE4 out = ToVect4(inputs, n, 1, outputs_total, 0) *
               ToVect4(mask, n, 1, outputs_total, 0);
   D4ToArray(outputs, out, n, 1, outputs_total, 0);
  }

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

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

#define def_k_MaskMult 40

//--- маскирование данных
#define def_mask_inputs                0
#define def_mask_mask                  1
#define def_mask_outputs               2
#define def_mask_total                 3

Затем мы перейдем к диспетчерскому классу модели. В методе инициализации контекста OpenCL изменим общее количество кернелов, а затем создадим кернел в контексте.

bool CNet::InitOpenCL(void)
  {
   ......
   if(!m_cOpenCL.SetKernelsCount(41))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
   ......

   if(!m_cOpenCL.KernelCreate(def_k_MaskMult, "MaskMult"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }
//---
   return true;
  }

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

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

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

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

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

Я нашел другое решение. Мы уже имеем кернел для выполнения линейной функции активации. Напомню ее математическое представление.

Если рассмотреть частный случай при a=1 и b=0, то получим простое копирование данных.

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

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

bool CNeuronDropout::FeedForward(CNeuronBase *prevLayer)
  {
   ......
//--- разветвление алгоритма в зависимости от устройства выполнения операций
   if(!m_cOpenCL)
     {
   ......
     }
   else  // Блок OpenCL
     {
      //--- проверка флага режима работы
      if(!m_bTrain)
        {
         //--- проверка буферов данных
         CBufferType *inputs = prevLayer.GetOutputs();
         if(inputs.GetIndex() < 0)
            return false;
         if(m_cOutputs.GetIndex() < 0)
            return false;
         //--- передача параметров кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LineActivation,
                                             def_activ_inputs, inputs.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LineActivation,
                                        def_activ_outputs, m_cOutputs.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_LineActivation,
                                                      def_activ_param_a, (TYPE)1))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_LineActivation,
                                                      def_activ_param_b, (TYPE)0))
            return false;
         uint offset[] = {0};
         uint NDRange[] = {(uint)m_cOutputs.Total()};
         if(!m_cOpenCL.Execute(def_k_LineActivation, 1, offset, NDRange))
            return false;
        }

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

      else
        {
         //--- проверка буферов данных
         CBufferType *inputs = prevLayer.GetOutputs();
         if(inputs.GetIndex() < 0)
            return false;
         if(!m_cDropOutMultiplier.BufferCreate(m_cOpenCL))
            return false;
         if(m_cOutputs.GetIndex() < 0)
            return false;
         //--- передача параметров кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_MaskMult,
                                             def_mask_inputs, inputs.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_MaskMult,
                                 def_mask_mask, m_cDropOutMultiplier.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_MaskMult,
                                        def_mask_outputs, m_cOutputs.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_MaskMult, def_mask_total, total))
            return false;
         //--- постановка в очередь выполнения
         int off_set[] = {0};
         int NDRange[] = { (int)(total + 3) / 4};
         if(!m_cOpenCL.Execute(def_k_MaskMult, 1, off_set, NDRange))
            return false;
        }
     }
//---
   return true;
  }

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

bool CNeuronDropout::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
   ......
//--- разветвление алгоритма в зависимости от устройства выполнения операций
   ulong total = m_cOutputs.Total();
   if(!m_cOpenCL)
     {
   ......
     }

   else  // блок OpenCL
     {
      //--- проверка флага режима работы
      if(!m_bTrain)
        {
         //--- проверка буферов данных
         CBufferType *grad = prevLayer.GetGradients();
         if(grad.GetIndex() < 0)
            return false;
         if(m_cGradients.GetIndex() < 0)
            return false;
         //--- передача параметров кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LineActivation,
                                def_activ_inputs, m_cGradients.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_LineActivation,
                                       def_activ_outputs, grad.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_LineActivation,
                                               def_activ_param_a, (TYPE)1))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_LineActivation,
                                               def_activ_param_b, (TYPE)0))
            return false;
         uint offset[] = {0};
         uint NDRange[] = {(uint)m_cOutputs.Total()};
         if(!m_cOpenCL.Execute(def_k_LineActivation, 1, offset, NDRange))
            return false;
        }

И режим работы в процессе обучения.

      else
        {
         //--- проверка буферов данных
         CBufferType* prev = prevLayer.GetGradients();
         if(prev.GetIndex() < 0)
            return false;
         if(m_cDropOutMultiplier.GetIndex() < 0)
            return false;
         if(m_cGradients.GetIndex() < 0)
            return false;
         //--- передача параметров кернелу
         if(!m_cOpenCL.SetArgumentBuffer(def_k_MaskMult,
                                         def_mask_inputs, m_cGradients.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_MaskMult,
                                   def_mask_mask, m_cDropOutMultiplier.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgumentBuffer(def_k_MaskMult,
                                                def_mask_outputs, prev.GetIndex()))
            return false;
         if(!m_cOpenCL.SetArgument(def_k_MaskMult, def_mask_total, total))
            return false;
         //--- постановка в очередь выполнения
         int off_set[] = {0};
         int NDRange[] = { (int)(total + 3) / 4 };
         if(!m_cOpenCL.Execute(def_k_MaskMult, 1, off_set, NDRange))
            return false;
        }
     }
//---
   return true;
  }

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

Поздравляю, мы завершили работу над методами класса реализации алгоритма Dropout. Мы проделали довольно большую работу и реализовали алгоритм Dropout и средствами MQL5, и в режиме моногопоточных операций с использованием технологии OpenCL. Теперь можем протестировать модели. Но прежде я предлагаю посмотреть на реализацию такого подхода на языке Python в библиотеке TensorFlow.

6.2.1.3 Методы работы с файлами

Реализация Dropout в Python