OpenCL как средство параллельных вычислений в MQL5

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

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

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

OpenCL (Open Computing Language) — это открытый бесплатный стандарт кроссплатформенного параллельного программирования различных ускорителей, используемых в суперкомпьютерах, облачных серверах, персональных компьютерах, мобильных устройствах и встроенных платформах.

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

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

В терминологии OpenCL микропроцессор компьютера (CPU) является Host. Он управляет всеми процессами выполняемой программы. Все микропроцессоры с поддержкой технологии OpenCL в составе СPU и GPU являются устройствами Device. Каждое устройство имеет свой уникальный номер внутри платформы.

Один Device может иметь несколько вычислительных единиц Computer Units. Их число определяется количеством физических и виртуальных ядер микропроцессора. Для видеокарт это будут SIMD-ядра. Каждое SIMD-ядро содержит несколько потоковых процессоров Stream Cores. Каждый потоковый процессор имеет несколько вычислительных элементов Processing Elements (или ALU).

Конкретное количество Computer Units, SIMD-ядер, Stream Cores и Processing Elements зависит от архитектуры конкретного устройства.

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

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

Host, как и классическое приложение, запускается и работает на CPU. Для организации многопоточности вычислений выделяется контекст — среда для выполнения специальных объектов программы OpenCL. Контекст объединяет набор OpenCL-устройств для запуска программы, сами программные объекты с их исходными кодами, а также набор объектов памяти, видимых хосту и OpenCL-устройствам. За создание контекста в MQL5 отвечает функция CLContextCreate, в параметрах которой указывается устройство для выполнения программы. Функция возвращает хендл контекста.

int CLContextCreate(
  int          device       // порядковый номер OpenCL устройства или макрос
  );

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

int CLProgramCreate(
  int          context,     // хендл на контекст OpenCL
  const string source       // исходный код
  );

OpenCL-программа делится на отдельные кернелы — исполняемые функции (ядра). Для объявления кернела предусмотрена функция CLKernelCreate. В параметрах функции указывается хендл ранее созданной программы и наименование кернела в ней. На выходе получаем хендл кернела.

int CLKernelCreate(
  int          program,     // хендл на объект OpenCL
  const string kernel_name  // имя кернела
  );

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

Каждая копия кернела выполняется для каждого индекса из этого пространства и называется «Work-Item» (рабочей единицей). Каждой рабочей единице предоставляется глобальный индекс ID. Кроме того, каждая такая единица выполняет один и тот же код, но данные для выполнения могут быть различными.

Рабочие единицы организовываются в рабочие группы Work-Groups. Группы предоставляют собой более крупное разбиение в пространстве индексов. Каждой группе присваивается групповой индекс ID. Размерность групп соответствует размерности для адресации отдельных элементов. Каждому элементу сопоставляется уникальный в рамках группы локальный индекс ID. Таким образом, рабочие единицы могут быть адресованы как по глобальному индексу ID, так и по комбинации группового и локального.

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

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

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

Глобальная память предоставляет доступ на чтение и запись элементам всех групп. Каждая рабочая единица Work-Item может писать и читать из любой части глобальной памяти.

Локальная память — локальная для группы область памяти, в ней можно создавать переменные, разделяемые всей группой. Она может быть реализована как отдельная память на OpenCL-устройстве или размечена как область в глобальной памяти.

Частная (private) память — область, видимая только рабочей единице Work-Item. Переменные, определенные в частной памяти одной рабочей единицы, не видны другим.

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

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

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

//+------------------------------------------------------------------+
//|  Функция умножения векторов на CPU                               |
//+------------------------------------------------------------------+
bool MultCPU(matrix<TYPE> &source1vector<TYPE> &source2vector<TYPE> &result)
  {
//---
   ulong rows = source1.Rows();
   ulong cols = source1.Cols();
   if(cols != source2.Size())
     {
      PrintFormat("Size of vectors not equal: %d != %d", colssource2.Size());
      return false;
     }
//---
   result = vector<TYPE>::Zeros(rows);
   for(ulong r = 0r < rowsr++)
     {
      result[r] = 0;
      for(ulong c = 0c < colsc++)
         result[r] += source1[rc] * source2[c];
     }
//---
   return true;
  }
//+------------------------------------------------------------------+

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

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

Решить это поможет использование OpenCL. Конечно, такая реализация будет немного сложнее. Вначале, напишем OpenCL-программу и сохраним в файле mult_vect_ocl.cl. Расширение *.cl общепринято для OpenCL-программ, но необязательно для реализации в среде MQL5 — в данном случае мы будем использовать файл только для хранения текста программы, а загружать программу будем в виде текста.

В коде программы включим поддержку типа double. Здесь стоит обратить внимание, что не все GPU поддерживают тип double. А даже если и поддерживают, то в большинстве случаев этот функционал отключен по умолчанию.

//--- by default some GPUs don't support doubles
//--- cl_khr_fp64 directive is used to enable work with doubles
#pragma OPENCL EXTENSION cl_khr_fp64 : enable

Еще один момент. MetaTrader 5 позволяет использовать для вычислений OpenCL-устройства как с поддержкой типа double, так и без. Поэтому при использовании типа double в своей OpenCL-программе необходимо проверить совместимость используемого устройства. В противном случае мы рискуем получить ошибку в процессе выполнения OpenCL-программы и прекращение ее работы.

В то же время MetaTrader 5 не ограничивает возможность использования всех доступных типов данных. Тут надо сказать, что язык OpenCL позволяет использовать различные скалярные типы данных:

  • логический: bool,
  • целочисленные: char, uchar, short, ushort, int, uint, long, ulong,
  • с плавающей запятой: float, double.

Поддержка аналогичных типов данных есть и в MQL5. Здесь надо помнить, что каждый тип данных имеет свои ограничения возможного диапазона значений, а вместе с тем и используемый объем памяти для хранения данных. Поэтому если в вашей программе не требуется особая точность или диапазон возможных значений не слишком велик, то рекомендуется использовать менее ресурсоемкие типы данных. Это позволит более эффективно использовать память девайса и снизить затраты на копирование данных между основной памятью и памятью контекста OpenCL. В частности, тип double можно заменить на float. Он, конечно, дает меньшую точность, но при этом занимает в 2 раза меньше памяти и поддерживается всеми современными OpenCL-устройствами. Это позволяет сократить затраты на передачу данных между устройствами и расширить область использования программы.

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

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

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

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

В теле кернела функция get_global_id укажет глобальный индекс ID запущенной единицы Work-Item. Индекс в данном случае является аналогом счетчика итераций внешнего цикла в классической реализации. Он указывает порядковый номер строи матрицы и элемент вектора результатов. Далее посчитаем сумму значений соответствующего потока аналогично расчету внутри вложенного цикла классической реализации. Но здесь есть нюанс. Для вычислений мы воспользуемся векторными операциями из четырех элементов. В свою очередь, для использования векторных операций нам необходимо подготовить данные. Из Host программы мы получаем массив из скалярных элементов, поэтому перенесем необходимые элементы в наши приватные векторные переменные с помощью функции ToVect (ее код рассмотрим ниже). Затем с помощью векторной операции dot получим значение умножения двух векторов из четырех элементов. То есть одной операцией мы получаем сумму произведений четырех пар значений. Полученное значение складываем в локальную переменную, в которой накапливается значение произведения строки матрицы на вектор.

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

//+------------------------------------------------------------------+
//| Mult of vectors                                                  |
//+------------------------------------------------------------------+
__kernel void MultVectors(__global TYPE *source1,
                          __global TYPE *source2,
                          __global TYPE *result,
                          int cols)
  {
   int shift = get_global_id(0) * cols;
   TYPE z = 0;
   for(int i = 0i < colsi+=4)
     {
      TYPE4 x = ToVect(source1icolsshift);
      TYPE4 y = ToVect(source2icols0);
      z += dot(x,y);
     }
   result[get_global_id(0)] = z;
  }

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

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

TYPE4 ToVect(__global TYPE *arrayint startint sizeint shift)
  {
   TYPE4 result = (TYPE4)0;
   if(start < size)
     {
      switch(size - start)
        {
         case  1:
            result = (TYPE4)(array[shift+start], 000);
            break;
         case  2:
            result = (TYPE4)(array[shift+start], array[shift+start + 1], 00);
            break;
         case  3:
            result = (TYPE4)(array[shift+start], array[shift+start + 1],
                             array[shift+start + 2], 0);
            break;
         default:
            result = (TYPE4)(array[shift+start], array[shift+start + 1],
                             array[shift+start + 2], array[shift+start + 3]);
            break;
        }
     }
   return result;
  }

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

На этом завершается работа с OpenCL-программой. Далее мы продолжим работу на стороне основной программы (Host). Для работы с OpenCL в MQL5 предлагается класс COpenCL в стандартной библиотеке OpenCL.mqh.

Вначале проведем подготовительную работу: подключим стандартную библиотеку, подгрузим ресурсом созданную ранее OpenCL-программу и объявим константы для индексов кернела, буферов и параметров программы. Также укажем используемый в программе тип данных. Я указал тип float, так как интегрированная GPU моего ноутбука не поддерживает тип double.

#include <OpenCL/OpenCL.mqh>
#resource "mult_vect_ocl.clas string OCLprogram
#define TYPE                        float
const string ExtType = StringFormat("#define TYPE %s\r\n"
                                    "#define TYPE4 %s4\r\n",
                                    typename(TYPE), typename(TYPE));
//+------------------------------------------------------------------+
//|  Defines                                                         |
//+------------------------------------------------------------------+
#define cl_program                  ExtType+OCLprogram
//---
#define k_kernel                    0
#define k_source1                   0
#define k_source2                   1
#define k_result                    2
#define k_cols                      3

Объявим экземпляр класса для работы с OpenCL и переменные для хранения хендлов буферов данных.

COpenCL*                  cOpenCL;
int                       buffer_Source1;
int                       buffer_Source2;
int                       buffer_Result;

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

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

bool OpenCL_Init(matrix<TYPE> &source1vector<TYPE> &source2)
  {
//--- создание OpenCL программы, кернела и буферов
   cOpenCL = new COpenCL();
   if(!cOpenCL.Initialize(cl_programtrue))
      return false;
   if(!cOpenCL.SetKernelsCount(1))
      return false;
   if(!cOpenCL.KernelCreate(k_kernel, "MultVectors"))
      return false;
   buffer_Source1 = CLBufferCreate(cOpenCL.GetContext(), 
                                      (uint)(sizeof(TYPE) * source1.Rows() * 
                                             source1.Cols()), CL_MEM_READ_ONLY);
   buffer_Source2 = CLBufferCreate(cOpenCL.GetContext(),
                                     (uint)(sizeof(TYPE) * source2.Size()),
                                            CL_MEM_READ_ONLY);

   buffer_Result = CLBufferCreate(cOpenCL.GetContext(),
                                     (uint)(sizeof(TYPE) * source1.Rows()),
                                            CL_MEM_WRITE_ONLY);
   if(buffer_Result <= 0 || buffer_Source1 <= 0 || buffer_Source2 <= 0)
      return false;
   if(!CLBufferWrite(buffer_Source1,0,source1) ||
      !CLBufferWrite(buffer_Source2,0,source2))
     return false;
//---
   return true;
  }

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

Вначале передадим в кернел указатели на буферы данных и параметры размеров буферов. Данные операции выполняются методами CLSetKernelArgMem и SetArgument. Пространство индексов зададим в массиве NDRange по числу строк в матрице исходных данных. Запустим кернел на выполнение методом Execute. После выполнения всего массива экземпляров кернела считываем результаты вычислений из памяти устройства с помощью метода CLBufferRead.

bool MultOCL(int rowsint colsvector<TYPE> &result)
  {
   result=vector<TYPE>::Zeros(rows);
//--- Set parameters
   if(!CLSetKernelArgMem(cOpenCL.GetKernel(k_kernel), k_source1buffer_Source1))
      return false;
   if(!CLSetKernelArgMem(cOpenCL.GetKernel(k_kernel), k_source2buffer_Source2))
      return false;
   if(!CLSetKernelArgMem(cOpenCL.GetKernel(k_kernel), k_resultbuffer_Result))
      return false;
   if(!cOpenCL.SetArgument(k_kernelk_cols, cols))
      return false;
//--- Run kernel
   int off_set[] = {0};
   int NDRange[] = {rows};
   if(!cOpenCL.Execute(k_kernel1off_setNDRange))
      return false;
//--- Get result
   uint data_read = CLBufferRead(buffer_Result0result);
   if(data_read <= 0)
      return false;
//---
   return true;
  }

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

void OpenCL_Deinit()
  {
   if(!cOpenCL)
      return;
//---
   cOpenCL.Shutdown();
   delete cOpenCL;
  }

Очевидно, что при использовании OpenCL объем работы программиста увеличивается. Что же мы получаем взамен?

Для оценки результативности создадим небольшой скрипт opencl_test.mq5. Во внешних параметрах скрипта укажем размер матрицы исходных данных.

//+------------------------------------------------------------------+
//| Внешние параметры                                                |
//+------------------------------------------------------------------+
sinput int Rows = 100000;   // Строк в матрице
sinput int Colms = 100;     // Столбцов в матрице

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

//+------------------------------------------------------------------+
//| Программа скрипта                                                |
//+------------------------------------------------------------------+
void OnStart()
  {
   matrix<TYPEX = matrix<TYPE>::Zeros(RowsColms);
   vector<TYPEY = vector<TYPE>::Zeros(Colms);
   vector<TYPEZ;
   for(int i = 0i < Colmsi++)
     {
      for(int r = 0r < Rowsr++)
         X[ri] = MathRand() / (TYPE)32767;
      Y[i] = MathRand() / (TYPE)32767;
     }

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

   if(!OpenCL_Init(XY))
      return;

Теперь мы можем замерить скорость выполнения операций в OpenCL-контексте. С помощью функции GetTickCount получим количество миллисекунд от старта системы до и после вычислений. Вычисления осуществим в ранее рассмотренной функции MultOCL.

   uint start = GetTickCount();
   if(!MultOCL(RowsColmsZ))
      Print("Error OCL function");
   uint end = GetTickCount();
   PrintFormat("%.1e OCL duration %0 000d msec, result %.5e",
                           Rows * Colmsend - startZ.Sum());
   OpenCL_Deinit();

После выполнения операций очищаем контекст OpenCL.

Аналогичным способом замерим время выполнения операций классическим способом на CPU.

   start = GetTickCount();
   if(!MultCPU(XYZ))
      Print("Error CPU function");
   end = GetTickCount();
   PrintFormat("%.1e CPU duration %0 000d msec, result %.5e",
                            Rows * Colmsend - startZ.Sum());

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

   start = GetTickCount();
   Z = X.MatMul(Y);
   end = GetTickCount();
   PrintFormat("%.1e matrix operation duration %0 000d msec, result %.5e", 
                                        Rows * Colmsend - startZ.Sum());
  }

Тестирование описанного скрипта проводилось на ноутбуке с CPU Intel Core i7-1165G7 и интегрированном графическом процессоре GPU Intel(R) Iris(R) Xe. По результатам замера времени выполнения вычислений победила технология OpenCL. Наиболее медленной оказалась классическая реализация с использованием системы вложенных циклов. При этом результат вычислений во всех трех вариантах был идентичен.

Результаты сравнительного тестирования выполнения вычислений с использованием OpenCL и без

Результаты сравнительного тестирования выполнения вычислений с использованием OpenCL и без

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