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

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

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

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

  • написание дополнительных кернелов в ранее созданной программе OpenCL (opencl_program.cl);
  • организация процесса взаимодействия с контекстом OpenCL на стороне основной программы.

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

Подвыборочный слой

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

  • inputs — вектор исходных данных;
  • outputs — вектор для записи результатов.

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

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

Ну и конечно, в параметрах кернелу укажем размер окна для анализа исходных данных (window), шаг перемещения окна (step), количество фильтров (window_out) и функцию активации (activation).

__kernel void ProofFeedForward(__global double *inputs,
                               __global double *outputs,
                               int inputs_total,
                               int input_neurons,
                               int window,
                               int step,
                               int activation)

Данный кернел мы будем запускать в двухмерном пространстве задач. А значит, в каждом кернеле мы будем обрабатывать один элемент массива результатов в одном фильтре. Номер обрабатываемого элемента будем определять по идентификатору потока в измерении с индексом 0. Следовательно, общее количество потоков покажет нам количество элементов на выходе одного фильтра (neurons). По этим данным определим смещения до начала окна анализируемых данных в разрезе фильтра массива исходных данных (shift).

  {
   const int n = get_global_id(0);
   const int w = get_global_id(1);
   const int neurons = get_global_size(0);
   const int window_out = get_global_size(1);
   int shift = n * step;

Второе измерение с индексом 1 укажет нам на индекс анализируемого фильтра. Соответственно, определим смещение в массивах исходных данных (shift_inp) и результатов (out) до начала обрабатываемого фильтра. Не забываем проверить на предмет выхода за пределы диапазона массива результатов.

Подготовим переменную для хранения промежуточных значений текущего элемента вектора результатов (s).

   int out = w * neurons + n;
   int shift_inp = w * input_neurons;
   TYPE s = 0;
   TYPE k = (TYPE)1 / (TYPE)window;
   TYPE4 k4 = (TYPE4)(k);

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

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

  • Average pooling — среднее арифметической элементов окна исходных данных;
  • Max pooling — максимальный элемент окна исходных данных.

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

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

   for(int i = 0i < windowi += 4)
      switch(activation)
        {
         case 0:
            s += dot(ToVect4(inputsi1min(shift_inp+input_neurons,inputs_total),
                             shift_inp + shift), k4);
            break;
         case 1:
            s = Max4(ToVect4(inputsi1min(shift_inp+input_neurons,inputs_total), 
                             shift_inp + shift), s);
            break;
         default:
            break;
        }
   outputs[out] = s;
  }

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

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

С внешней программой новый кернел будет «общаться» через четыре буфера данных:

  • inputs — буфер результатов предшествующего слоя;
  • gradient_inputs — буфер градиентов предшествующего слоя (в данном случае служит для записи результатов работы кернела);
  • outputs — буфер результатов прямого прохода текущего слоя;
  • gradients — буфер градиентов на уровне результатов текущего слоя.

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

Не будем забывать, что в отличие от полносвязного слоя, нейроны подвыборочного слоя имеют ограниченные связи с нейронами предыдущего слоя. Определять зоны связей будем с помощью параметров window, step. Можно заметить, что одноименные параметры были объявлены в кернеле прямого прохода. Мы сохранили и их функциональное значение.

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

__kernel void ProofCalcHiddenGradient(__global TYPE *inputs,
                                      __global TYPE *gradient_inputs,
                                      __global TYPE *outputs,
                                      __global TYPE *gradients,
                                      int inputs_total,
                                      int outputs_total,
                                      int window,
                                      int step,
                                      int neurons,
                                      int activation)

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

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

В начале тела кернела определим порядковый номер искомого элемента и фильтра по идентификаторам потока. Общее количество потоков даст нам количество элементов одного фильтра в буфере исходных данных (input_neurons) и количество фильтров (window_out). По этим данным определим первый (start) и последний (stop) элементы вектора результатов, на которые влияет обрабатываемый элемент. При определении зоны влияния помним об ограничениях размерности буфера данных каждого фильтра. Следовательно, первый элемент не может быть меньше 0, а последний не может быть больше числа элементов в одном фильтре neurons.

  {
   const int n = get_global_id(0);
   const int w = get_global_id(1);
   const int input_neurons = get_global_size(0);
   const int window_out = get_global_size(1);
//---
   int start = n - window + step;
   start = max((start - start % step) / step0);
   int stop = min((n - n % step) / step + 1neurons);

Далее определим смещение анализируемого элемента в общем буфере исходных данных. При этом не забываем проверить на предмет выхода за пределы массива исходных данных.

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

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

   TYPE grad = 0;
   int shift_inp = w * input_neurons + n;
   if(shift_inp >= inputs_total)
      return;
   TYPE inp = inputs[shift_inp];

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

Для Average pooling мы просто разделим значение градиента ошибки на размер окна исходных данных и полученное значение прибавим к накопленному градиенту ошибки анализируемого элемента исходных данных. Обратите внимание, что делить градиент ошибки будем на размер окна исходных данных, а не зоны влияния. Ведь ошибка получена в результате прямого прохода, и на ее значение повлияли все элементы исходных данных, оказывающие влияние на конкретное значение.

В случае Max pooling мы сначала сравним значение соответствующих элементов на выходе и входе нейронного слоя. Только в случае их соответствия передадим градиент ошибки в полном объеме.

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

   for(int o = starto < stopo ++)
     {
      int shift_g = w * neurons + o;
      if(shift_g >= outputs_total)
         break;
      switch(activation)
        {
         case 0:
            grad += gradients[shift_g] / (TYPE)window;
            break;
         case 1:
            grad += (outputs[shift_g] == inp ? gradients[shift_g] : 0);
            break;
         default:
            break;
        }
     }
   gradient_inputs[shift_inp] = grad;
  }

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

Сверточный слой

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

  • inputs — буфер исходных данных;
  • weights — матрица весов;
  • sums — вектор взвешенных сумм исходных данных перед функцией активации;
  • outputs — вектор результатов.

Помимо буферов, для нормального функционирования новому кернелу потребуются параметры:

  • inputs_total — размер массива исходных данных;
  • window — размер анализируемого окна исходных данных;
  • step — шаг окна исходных данных;
  • window_out — количество фильтров в слое.

__kernel void ConvolutionFeedForward(__global TYPE *inputs,
                                     __global TYPE *weights,
                                     __global TYPE *outputs,
                                     int inputs_total,
                                     int window,
                                     int step,
                                     int window_out)

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

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

  {
   const int n = get_global_id(0);
   const int neurons = get_global_size(0);
   const int weights_total = (window + 1) * window_out;
   int shift = n * step;

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

   for(int w = 0w < window_outw++)
     {
      int out = (transposed_out == 1 ? w + n * window_out : w * neurons + n);
      int shift_weights = w * (window + 1) ;
      if((shift_weights + window) >= weights_total)
         break;
      TYPE s = weights[shift_weights + window];

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

      for(int i = 0i < windowi += 4)
         s += dot(ToVect4(inputsi1inputs_totalshift),
                  ToVect4(weightsi1shift_weights + windowshift_weights));
      outputs[out] = s;
     }
  }

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

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

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

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

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

Помимо буферов данных, для правильной работы кернела потребуются ряд параметров:

  • outputs_total — общее количество элементов в буфере результатов (градиентов на выходе текущего нейронного слоя);
  • window — размер окна исходных данных (количество элементов исходных данных анализируемых одним нейроном текущего слоя);
  • step — шаг перемещения окна по массиву исходных данных;
  • window_out — количество фильтров в текущем сверточном слое;
  • neurons — количество элементов на выходе одного фильтра.

__kernel void ConvolutionCalcHiddenGradient(__global TYPE *gradient_inputs,
                                            __global TYPE *weights,
                                            __global TYPE *gradients,
                                            int window,
                                            int step,
                                            int window_out,
                                            int neurons)

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

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

  {
   const int n = get_global_id(0);
   const int inputs_total = get_global_size(0);
   int weights_total = (window + 1) * window_out;

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

   TYPE grad = 0;
   int w_start = n % step;
   int r_start = max((n - window + step) / step0);
   int total = (window - w_start + step - 1) / step;
   total = min((n + step) / steptotal);

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

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

   for(int i = 0i < totali ++)
     {
      int row = r_start + i;
      if(row >= neurons)
         break;

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

      for(int wo = 0wo < window_outwo++)
        {
         int shift_g = (transposed_out == 1 ? row * window_out + wo :
                                                        row + wo * neurons);
         int shift_w = w_start + (total - i - 1) * step + wo * (window + 1);
         grad += gradients[shift_g] * weights[shift_w];
        }
     }
   gradient_inputs[n] = grad;
  }

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

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

Для корректной работы кернела потребуется использование 3-х буферов данных:

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

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

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

  • inputs_total — общее число элементов в буфере результатов и, соответственно, буфере градиентов ошибки;
  • step — шаг перемещения окна анализируемых данных по массиву исходных данных;
  • neurons — количество элементов на выходе одного фильтра.

__kernel void ConvolutionCalcDeltaWeights(__global TYPE *inputs,
                                          __global TYPE *delta_weights,
                                          __global TYPE *gradients,
                                          int inputs_total,
                                          int step,
                                          int neurons)

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

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

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

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

  {
   const int inp_w = get_global_id(0);
   const int w = get_global_id(1);
   const int window = get_global_size(0) - 1;
   const int window_out = get_global_size(1);
//---
   int shift_delt = w * (window + 1) + inp_w;
   TYPE value = 0;

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

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

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

   if(inp_w == window)
     {
      for(int n = 0n < neuronsn ++)
         value += gradients[w * neurons + n];
     }
   else
      for(int n = 0n < neuronsn ++)
        {
         int shift_inp = n * step + inp_w;
         if(shift_inp >= inputs_total)
            break;
         value += inputs[shift_inp] * gradients[w * neurons + n];
        }
   delta_weights[shift_delt] += value;
  }

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

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

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

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

#define def_k_ProofFeedForward            21
#define def_k_ProofHiddenGradients        22
#define def_k_ConvolutionFeedForward      23
#define def_k_ConvolutionHiddenGradients  24
#define def_k_ConvolutionDeltaWeights     25

Также объявим константы параметров каждого кернела. Константы параметров должны строго соответствовать порядковому номеру параметра в кернеле программы OpenCL. Нумерация параметров начинается с нуля.

//--- прямой проход подвыборочного слоя
#define def_prff_inputs                   0
#define def_prff_outputs                  1
#define def_prff_inputs_total             2
#define def_prff_input_neurons            3
#define def_prff_window                   4
#define def_prff_step                     5
#define def_prff_activation               6

//--- распределение градиента через подвыборочный слой
#define def_prhgr_inputs                  0
#define def_prhgr_gradient_inputs         1
#define def_prhgr_outputs                 2
#define def_prhgr_gradients               3
#define def_prhgr_inputs_total            4
#define def_prhgr_outputs_total           5
#define def_prhgr_window                  6
#define def_prhgr_step                    7
#define def_prhgr_neurons                 8
#define def_prff_activation               9

//--- прямой проход сверточного слоя
#define def_cff_inputs                    0
#define def_cff_weights                   1
#define def_cff_outputs                   2
#define def_cff_inputs_total              3
#define def_cff_window                    4
#define def_cff_step                      5
#define def_cff_window_out                6

//--- распределение градиента через сверточный слой
#define def_convhgr_gradient_inputs       0
#define def_convhgr_weights               1
#define def_convhgr_gradients             2
#define def_convhgr_window                3
#define def_convhgr_step                  4
#define def_convhgr_window_out            5
#define def_convhgr_neurons               6

//--- распределение градиента на матрицу весов сверточного слоя
#define def_convdelt_inputs               0
#define def_convdelt_delta_weights        1
#define def_convdelt_gradients            2
#define def_convdelt_inputs_total         3
#define def_convdelt_step                 4
#define def_convdelt_neurons              5

После объявления констант нам нужно обновить список используемых кернелов из программы OpenCL. Напомню, что эта работа осуществляется в методе CNet::InitOpenCL. Здесь нам нужно изменить количество используемых кернелов на 26.

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

Создадим точки входа для новых кернелов.

   if(!m_cOpenCL.KernelCreate(def_k_ProofFeedForward"ProofFeedForward"))
     {
      m_cOpenCL.Shutdown();
      delete m_cOpenCL;
      return false;
     }

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

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

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

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

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

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

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

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

      //--- Передача параметров в кернел
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ProofFeedForwarddef_prff_inputs,
                                                         input_data.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ProofFeedForwarddef_prff_outputs,
                                                         m_cOutputs.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofFeedForwarddef_prff_inputs_total,
                                                            input_data.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofFeedForwarddef_prff_window,
                                                                     m_iWindow))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofFeedForwarddef_prff_stepm_iStep))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofFeedForwarddef_prff_activation,
                                                            (int)m_eActivation))
         return false;
      uint input_neurons = (input_data.Total()+m_iWindowOut-1) / m_iWindowOut;
      if(!m_cOpenCL.SetArgument(def_k_ProofFeedForwarddef_prff_input_neurons,
                                                                 input_neurons))
         return false;

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

       //--- Постановка кернела в очередь на выполнение
      uint off_set[] = {00};
      uint NDRange[] = {m_iNeuronsm_iWindowOut};
      if(!m_cOpenCL.Execute(def_k_ProofFeedForward2off_setNDRange))
         return false;
     }
//---
   return true;
  }

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

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

Первые два буфера используются для определение нужных элементов при использовании Max pooling.

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

bool CNeuronProof::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
//--- Блок контролей
   if(!prevLayer || !m_cOutputs ||
      !m_cGradients || !prevLayer.GetOutputs() ||
      !prevLayer.GetGradients())
      return false;
   CBufferType *input_data = prevLayer.GetOutputs();
   CBufferType *input_gradient = prevLayer.GetGradients();
   if(!input_gradient.BufferInit(input_data.Rows(), input_data.Cols(), 0))
      return false;
//---  Разветвление алгоритма в зависимости от устройства выполнения операций
   if(!m_cOpenCL)
     {
     // Здесь пропущен Блок MQL5
     }
   else    // Блок работы с OpenCL
     {
      //--- проверяем наличие буферов в контекст OpenCL
      if(input_data.GetIndex() < 0)
         return false;
      if(m_cOutputs.GetIndex() < 0)
         return false;
      if(input_gradient.GetIndex() < 0)
         return false;
      if(m_cGradients.GetIndex() < 0)
         return false;

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

      //--- Передача параметров в кернел
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ProofHiddenGradients,
                                         def_prhgr_inputsinput_data.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ProofHiddenGradients,
                                        def_prhgr_outputsm_cOutputs.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ProofHiddenGradients,
                                    def_prhgr_gradientsm_cGradients.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ProofHiddenGradients
                            def_prhgr_gradient_inputsinput_gradient.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofHiddenGradients,
                                      def_prhgr_inputs_totalinput_data.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofHiddenGradients,
                                                     def_prhgr_windowm_iWindow))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofHiddenGradients,
                                                         def_prhgr_stepm_iStep))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofHiddenGradients,
                                        def_prhgr_activation, (int)m_eActivation))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofHiddenGradients
                                                   def_prhgr_neuronsm_iNeurons))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ProofHiddenGradients,
                                     def_prhgr_outputs_totalm_cOutputs.Total()))
         return false;

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

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

      //--- Постановка кернела в очередь на выполнение
      uint input_neurons = (input_data.Total() + m_iWindowOut - 1) / m_iWindowOut;
      uint off_set[] = {00};
      uint NDRange[] = {input_neuronsm_iWindowOut};
      if(!m_cOpenCL.Execute(def_k_ProofHiddenGradients2off_setNDRange))
         return false;
     }
//---
   return true;
  }

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

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

  • исходные данные,
  • матрица весов,
  • дополнительный буфер функции активации (используется для функции активации Swish),
  • буфер результатов.

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

bool CNeuronConv::FeedForward(CNeuronBase *prevLayer)
  {
//--- блок контролей
   if(!prevLayer || !m_cOutputs || !m_cWeights || !prevLayer.GetOutputs())
      return false;
   CBufferType *input_data = prevLayer.GetOutputs();
   ulong total = input_data.Total();
//--- разветвление алгоритма в зависимости от устройства выполнения операций
   if(!m_cOpenCL)
     {
     // Здесь пропущен Блок MQL5
     }
   else
     {
      //--- проверка буферов данных
      if(input_data.GetIndex() < 0)
         return false;
      if(m_cWeights.GetIndex() < 0)
         return false;
      if(m_cOutputs.GetIndex() < 0)
         return false;

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

      //--- передача аргументов кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionFeedForward,
                                          def_cff_inputsinput_data.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionFeedForward,
                                         def_cff_weightsm_cWeights.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionFeedForward,
                                         def_cff_outputsm_cOutputs.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionFeedForward,
                                       def_cff_inputs_totalinput_data.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionFeedForward,
                                                      def_cff_windowm_iWindow))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionFeedForward,
                                                           def_cff_stepm_iStep))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionFeedForward,
                                                 def_cff_window_outm_iWindowOut))
         return false;

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

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

В завершение вызываем функцию активации и завершаем работу метода.

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

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

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

За распределение градиента ошибки через сверточный слой отвечает метод CNeuronConv::CalcHiddenGradient. Корректное выполнение алгоритма данного метода требует наличия трех буферов данных:

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

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

bool CNeuronConv::CalcHiddenGradient(CNeuronBase *prevLayer)
  {
//--- блок контролей
   if(!prevLayer || !prevLayer.GetOutputs() || !prevLayer.GetGradients() ||
      !m_cGradients || !m_cWeights)
      return false;
//--- корректировка градиентов ошибки на производную функции активации
   if(m_cActivation)
     {
      if(!m_cActivation.Derivative(m_cGradients))
         return false;
     }
//--- разветвление алгоритма в зависимости от устройства выполнения операций
   CBufferTypeinput_gradient = prevLayer.GetGradients();
   if(!m_cOpenCL)
     {
     // Здесь пропущен Блок MQL5
     }

   else // Блок работы с OpenCL
     {
      //--- проверка буферов данных
      if(m_cWeights.GetIndex() < 0)
         return false;
      if(input_gradient.GetIndex() < 0)
         return false;
      if(m_cGradients.GetIndex() < 0)
         return false;

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

      //--- передача аргументов кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionHiddenGradients,
                        def_convhgr_gradient_inputsinput_gradient.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionHiddenGradients,
                                    def_convhgr_weightsm_cWeights.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionHiddenGradients,
                                def_convhgr_gradientsm_cGradients.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionHiddenGradients,
                                               def_convhgr_neuronsm_iNeurons))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionHiddenGradients,
                                                 def_convhgr_windowm_iWindow))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionHiddenGradients,
                                                     def_convhgr_stepm_iStep))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionHiddenGradients,
                                          def_convhgr_window_outm_iWindowOut))
         return false;

Далее укажем количество потоков равное количеству элементов в буфере исходных данных и отправим кернел в очередь на выполнение.

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

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

Алгоритм работы метода распределения градиента ошибки до матрицы весов требует наличие трех буферов:

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

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

bool CNeuronConv::CalcDeltaWeights(CNeuronBase *prevLayer)
  {
//--- блок контролей
   if(!prevLayer || !prevLayer.GetOutputs() || !m_cGradients || !m_cDeltaWeights)
      return false;
//--- разветвление алгоритма в зависимости от устройства выполнения операций
   CBufferType *input_data = prevLayer.GetOutputs();
   if(!m_cOpenCL)
     {
     // Здесь пропущен Блок MQL5
     }
   else // Блок работы с OpenCL
     {
      //--- проверка буферов данных
      if(m_cGradients.GetIndex() < 0)
         return false;
      if(m_cDeltaWeights.GetIndex() < 0)
         return false;
      if(input_data.GetIndex() < 0)
         return false;

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

      //--- передача аргументов кернелу
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionDeltaWeights,
                        def_convdelt_delta_weightsm_cDeltaWeights.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionDeltaWeights,
                                    def_convdelt_inputsinput_data.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgumentBuffer(def_k_ConvolutionDeltaWeights,
                               def_convdelt_gradientsm_cGradients.GetIndex()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionDeltaWeights,
                                 def_convdelt_inputs_totalinput_data.Total()))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionDeltaWeights,
                                              def_convdelt_neuronsm_iNeurons))
         return false;
      if(!m_cOpenCL.SetArgument(def_k_ConvolutionDeltaWeights,
                                                    def_convdelt_stepm_iStep))
         return false;

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

  • по количеству фильтров,
  • по количеству весовых коэффициентов в одном фильтре.

Для этого мы указываем два параметра в массиве NDRange. Каждый параметр задает размерность соответствующей области задач. Отправляем кернел в очередь выполнения.

      //--- постановка кернела в очередь выполнения
      uint NDRange[] = {m_iWindow + 1m_iWindowOut};
      uint off_set[] = {00};
      if(!m_cOpenCL.Execute(def_k_ConvolutionDeltaWeights2off_setNDRange))
         return false;
     }
//---
   return true;
  }

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