М. Е. Жуковский, Р. В. Усков о применении графических процессоров видеоускорителей в прикладных задачах
|
Скачать 181.72 Kb.
|
|
РОССИЙСКАЯ АКАДЕМИЯ НАУК ИНСТИТУТ ПРИКЛАДНОЙ МАТЕМАТИКИ ИМ. М.В.КЕЛДЫША РАН М.Е.Жуковский, Р.В.Усков О ПРИМЕНЕНИИ ГРАФИЧЕСКИХ ПРОЦЕССОРОВ ВИДЕОУСКОРИТЕЛЕЙ В ПРИКЛАДНЫХ ЗАДАЧАХ (Часть I. Обзор технологии.) Москва, 2010 М.Е.Жуковский, Р.В.Усков О ПРИМЕНЕНИИ ГРАФИЧЕСКИХ ПРОЦЕССОРОВ ВИДЕОУСКОРИТЕЛЕЙ В ПРИКЛАДНЫХ ЗАДАЧАХ АННОТАЦИЯ В работе рассмотрены основы применения технологии nVidia© CUDA для распараллеливания вычислений с использованием графических процессоров. Обсуждаются особенности программной реализации технологии, в частности, использование потоковых процессоров и памяти видеоадаптера. Эффективность распараллеливания вычислений показана на примере. M.E.Zhukovskiy, R.V.Uskov USE OF GRAPHICAL PROCESSORS OF VIDEO ACCELERATORS IN APPLICATIONS. ABSTRACT Basics of application of nVidia© CUDA technology are considered for paralleling the calculations with use of graphical processors. Particularities of code developing by use of the technology are discussed. In particular, the use of thread processors and video memory of video adapter is taken up. Effectiveness of using the technology is shown by example. Введение По мере развития игровой индустрии и плавного приближения картинки в современных играх к фотографическому качеству, объем вычислений необходимых для прорисовки картинки резко увеличивался. На определённом этапе этих вычислений стало так много, что для дальнейшей эволюции графики, мощностей тогдашних центральных процессоров стало не хватать. В силу того, что вычисления необходимые для прорисовки игровой картинки, были однотипны для всех игр и носили весьма специфический характер, решением проблемы стало создание узкоспециализированных устройств – видеоускорителей. За счет узкой специализации удалось добиться создания устройств с внушительными вычислительными способностями и относительно небольшой стоимостью. В настоящее время вычислительная мощность современных видеоускорителей в десятки и даже сотни раз превышает аналогичные показатели центральных процессоров при примерно равной стоимости. Долгое время эти мощности были доступны лишь для разработки графических приложений и не могли использоваться для решения прикладных программ. Однако потенциал в данной области был очевиден, поэтому предпринимались многочисленные попытки приспособить видеоускорители для решения прикладных задач. К настоящему времени в корпорациях - лидерах отрасли по производству видеоускорителей (nVidia и ATI) - разработаны технологии, позволяющие использовать мощности видеоускорителей для проведения расчетов в прикладных задачах. Этими технологиями являются nVidia© CUDA и ATI© Stream. Существуют и другие технологии, семейства GPGPU1, то есть позволяющие использовать ресурсы видеоускорителей в прикладных программах: Microsoft© Direct3D 11 - Compute Shader, а также полностью открытая технология OpenCL, разработку которой поддержали такие компании как Apple©, nVidia©, AMD©, IBM©, Intel©, ARM©, Motorola© и другие. Общим для этих технологий является то, что, затратив определенные усилия на адаптацию прикладных программ к использованию в качестве вычислителей графические процессоры (GPU), можно значительно понизить стоимость вычислений, получив при этом, схожую производительность. Ниже в настоящей работе речь пойдёт о технологии от компании nVidia – CUDA (Common Unified Development Architecture). 1. Обзор технологии. С программной точки зрения, технология CUDA представляет собой расширение языка Си, которое можно представить в виде двух взаимно зависимых составляющих [1]:
Первая из указанных составляющих (т.н. «управляющая часть») позволяет производить такие операции как выделение памяти на GPU и копирование данных между GPU и хостом, то есть служит для организации запусков вычислений на GPU [3] - "вводить" начальные данные и получать результаты расчетов (рис.1). На рисунке 1 RAM – это память центрального процессора, VRAM – память видеоускорителя. Согласно изображенной схеме: под управлением центрального процессора осуществляется передача данных из оперативной памяти CPU в память GPU, затем по «команде» из CPU запускаются вычисления на GPU (запуск «ядра», см. ниже). Результаты вычислений «считываются» из памяти видеоадаптера. Кроме того, при необходимости центральный процессор используется, конечно, и в качестве вычислителя. Вторая составляющая рассматриваемого расширения Си («реализующая часть») дает возможность проводить вычисления непосредственно на GPU [1].
Для описания этой части технологии CUDA нам понадобятся следующие определения [2]. Потоком исполнения или просто потоком (thread) называется совокупность логических и арифметических операций (исполняемый код), выполняемых в определённой последовательности на вычислительном устройстве. При наличии множества потоков будем предполагать, что они выполняются параллельно, то есть без предписанного порядка во времени. Ядро - это совокупность потоков, запускаемая на выполнение с хоста и выполняемая на GPU. Потоки в ядре группируются в так называемые блоки (block), совокупность блоков представляет собой решетку (grid). Потоки в блоке организованы в трёхмерный массив. Блоки в решетке – в двумерный. Все потоки имеют идентичный исполняемый код, поэтому, чтобы отличать потоки друг от друга, каждому присвоен номер его блока в решетке (двумерный индекс) и его номер в блоке (трёхмерный индекс). На рисунке 2 изображена структура решетки, состоящей из набора блоков, и схема блока, состоящего из набора потоков. Для наглядности блок представлен двумерным массивом, то есть размерность блока по третьему индексу равна единице. Warp - это совокупность потоков блока, исполняемых одновременно в текущий момент. Число потоков в warp'е определяется системой и обычно равно 32.
Аппаратная структура видеоадаптера. Основные элементы видеоадаптера: видео память (VRAM) и графический процессор (GPU). Графический процессор состоит из набора мультипроцессоров (рис.3). В отличие от обычных процессоров, имеющих одно устройство контроля исполнения, кэш и несколько арифметико-логических устройств, графический процессор состоит из набора мультипроцессоров (рис.4), каждый из которых имеет своё устройство контроля исполнения (Control), область разделяемой (Cache) памяти, области регистров и значительно расширенный по сравнению с обычным процессором набор арифметико-логических устройств (ALU).
Такая структура позволяет добиться не только параллельного выполнения алгоритма на разных мультипроцессорах, но и параллельной обработки данных в рамках одного мультипроцессора. При запуске ядра блоки распределяются между мультипроцессорами (при нехватке мультипроцессоров становятся в очередь) и выполняются независимо. С каждым мультипроцессором может быть ассоциирован один и более блоков (в зависимости от ресурсов, используемых блоком). Один блок не может выполняться более чем на одном мультипроцессоре. На рис.5 показано соответствие программных средств и аппаратных устройств при проведении расчетов с использованием графических ускорителей:
Важной особенностью графического адаптера является наличие трех видов памяти: локальной (local), связанной с конкретным потоковым процессором, разделяемой (shared) памяти, доступной в рамках мультипроцессора, и глобальной (VRAM) – определенной для видеоадаптера. На рис. 6 показаны доступность и «время жизни» данных, находящихся в различных видах памяти. Так, локальная память доступна только потоку, который выполняется на данном вычислителе. Время жизни данных в такой памяти равно времени жизни потока. Разделяемая память доступна всем потокам блока и её время жизни определяется временем жизни блока. Глобальная память доступна всем потокам во всех блоках решетки. Время жизни областей глобальной памяти определяется программой: области глобальной памяти могут быть выделены или освобождены только с хоста, то есть только между вызовами ядер. Таким образом, время жизни области в глобальной памяти больше, чем время жизни одного или нескольких ядер (Grid0 и Grid1 на рисунке 6 — решетки различных ядер) и меньше, чем время жизни программы. Глобальная память, при всём её удобстве, не кэшируется и является относительно медленной по сравнению с разделяемой памятью. Вообще говоря, использование в процессе решения задачи разделяемой памяти вместо глобальной позволяет добиться значительного увеличения эффективности использования графических процессоров. Особенности использования указанных видов памяти рассмотрены ниже. 2. Особенности разработки приложений для CUDA Из описанного выше следует, что CUDA-приложения фактически являются многопоточными приложениями. Однако, в силу специфики потоков, а именно, большого их количества и крайней легковесности (то есть, крайне малых затрат на создание, уничтожение потоков и переключение между ними и, как следствие, ограниченных возможностей управления потоками), программирование CUDA-приложений значительно отличается от программирования обычных многопоточных приложений. Отметим основные из этих отличий.
Первое. В силу того, что задачам, под которые изначально создавались современные видеоадаптеры, характерна практически абсолютная декомпозиция данных (то есть для параллельно обрабатываемых данных результаты обработки одной их части не зависят от результатов обработки другой), в CUDA-приложениях практически отсутствуют методы контроля конкурентного доступа к данным. Существуют минимальные возможности такого контроля для потоков в рамках одного блока, однако их зачастую бывает недостаточно. Такие ограничения требуют введения определённой степени декомпозиции данных в прикладных приложениях. То есть, для параллельно обрабатываемых данных требуется по возможности максимально уменьшить, а то и вовсе исключить, зависимость результатов обработки этих данных друг от друга. Другая особенность также обусловлена архитектурой аппаратных средств, нацеленных на задачи обработки графики, в которых алгоритмы линейны и не содержат условных переходов. То есть, условные переходы не поддерживаются аппаратной частью вовсе. При этом условные переходы и циклы хоть и поддерживаются программными средствами CUDA, тем не менее, являются самыми «тяжелыми» операциями для GPU, поэтому при написании приложений следует старательно избегать таких операций. Работа с различными видами памяти. При разработке CUDA-приложений для управления и оптимизации доступны два вида памяти графического адаптера: глобальная и разделяемая (существует также локальная память, но её распределением занимается компилятор, и программист, вообще говоря, не может управлять этим видом памяти). Глобальная память (VRAM). Скорость доступа к этой памяти, вообще говоря, является достаточно низкой, однако при определённом подходе (см. ниже) к организации доступа, может быть значительно повышена. Такое повышение достигается при помощи «связывания» нескольких логических запросов к памяти (запросов от разных потоков) в один аппаратный (физический) запрос. «Связанным» (“coalesced” [2]) запросом к ресурсам глобальной памяти (VRAM) называется такая совокупность логических запросов, при которой половина потоков warp'а осуществляет чтение/запись за один физический запрос. Для того чтобы запросы стали связанными необходимо выполнение следующих условий (для случая доступа к четырёхбайтным элементам)2:
В случае, если хотя бы одно из условий нарушается, логические запросы от разных потоков перестают быть связанными, и для каждого из них формируется отдельный физический запрос3. Таким образом, максимальное число запросов, которые можно связать, равно половине размера warp'а, то есть максимальная скорость чтения из глобальной памяти может быть в 16 раз выше минимальной скорости. Рассмотрим примеры. На рис.7 все потоки читают из ячеек расположенных последовательно. Номера ячеек соответствуют номеру потока. В обоих изображенных на рисунке 7 случаях, для получения данных всеми потоками будет сформирован один единственный физический запрос к памяти. На рис.8 изображен пример «несвязанного» доступа к памяти. Слева - непоследовательное обращение к ячейкам, расположенным в области, допускающей «связанные» запросы. Справа - последовательное обращение к ячейкам, выходящее за область, допускающую «связанные» запросы. В обоих изображенных на рисунке 8 случаях, для получения данных всеми потоками будут сформированы 16 запросов к памяти4.
Современные устройства могут частично «исправлять» такие ситуации: обращение слева (рис.8) «связывается» и приводится к единственному запросу к памяти; обращение справа приводится к двум запросам к памяти.
Разделяемая (быстрая) память. Организация этого вида памяти отличается от организации глобальной памяти. Разделяемая память разбита на так называемые «банки». При обращении нескольких потоков к ячейке памяти внутри одного и того же банка, происходит так называемый «банк-конфликт», в результате которого одиночный запрос на чтение превращается в N запросов, где N – максимальное количество потоков, обращающихся одному банку. Для уменьшения количества банк-конфликтов при осуществлении запроса на чтение, один из банков объявляется системой «широковещательным» (обычно это тот, к которому обращается максимальное количество потоков). Запрос на чтение из «широковещательного» банка всегда происходит как одиночный запрос, вне зависимости от количества обращающихся потоков. Тем самым разрешается один банк-конфликт максимального порядка.
На рис.9 в вариантах 1 и 2 банк-конфликтов не возникнет, так как все потоки обращаются к различным банкам. В варианте 3 банк-конфликтов также удается избежать, так как Bank 3 будет выбран системой «широковещательным». В четвёртом варианте банк конфликтов не возникнет лишь в том случае, если Bank 5 будет выбран «широковещательным». На рис.10 пример с банк-конфликтами: слева - второго порядка, справа – восьмого порядка. Банк-конфликт N-го порядка приводит к N физическим запросам к памяти, в то время как бесконфликтные обращение всегда проходят за единственный физический запрос.
3. Пример использования технологии В качестве примера применения технологии CUDA рассмотрим решение задачи о вычислении корней большого числа квадратных уравнений. Пусть необходимо решить  квадратных уравнений  , для каждого из них заданы  . Требуется найти действительные  и  , а также мнимые  и  части корней уравнений.Рассмотрим сначала однопроцессорную реализацию расчетного алгоритма: void computeGold(eqdata* data) { int i; float d; for(i = 0; i < BLOCK_SIZE*GRID_SIZE; i++) { d = -data[i].b*data[i].b + 4*data[i].a*data[i].c; data[i].re_x1 = -data[i].b; data[i].re_x2 = -data[i].b; if(d > 0) { data[i].re_x1 += sqrtf(d); data[i].re_x2 -= sqrtf(d); } else { data[i].im_x1 = sqrtf(-d); data[i].im_x2 = -sqrtf(-d); } data[i].re_x1 /= 2*data[i].a; data[i].re_x2 /= 2*data[i].a; data[i].im_x1 /= 2*data[i].a; data[i].im_x2 /= 2*data[i].a; } } Где data – массив структур eqdata, состоящий из  элементов.struct eqdata { //Коэффициенты квадратного уравнения float a; float b; float c; //Действительные части корней уравнения float re_x1; float re_x2; //Мнимые части корней уравнения float im_x1; float im_x2; }; Время выполнения этого расчета на AMD Athlon 64 X2 3800+ (2.0Ггц) составляет 1038мс. Рассмотрим теперь реализацию расчетного алгоритма, использующую технологию CUDA: __global__ void compute(eqdata* data) { int tx = threadIdx.x; int bx = blockIdx.x; int i = tx+bx*BLOCK_SIZE; float d; d = -data[tx].b*data[tx].b + 4*data[tx].a*data[tx].c; data[i].re_x1 = -data[tx].b; data[i].re_x2 = -data[tx].b; data[i].im_x1 = 0; data[i].im_x2 = 0; if(d > 0) { data[i].re_x1 += sqrtf(d); data[i].re_x2 -= sqrtf(d); } else { data[i].im_x1 += sqrtf(-d); data[i].im_x2 -= sqrtf(-d); } data[i].re_x1 /= (2*data[tx].a); data[i].re_x2 /= (2*data[tx].a); data[i].im_x1 /= (2*data[tx].a); data[i].im_x2 /= (2*data[tx].a); } В данной реализации, каждый поток решает отдельное квадратное уравнение, всего потоков, они разбиты примерно на 16000 блоков, по 512 потоков в каждом.Время выполнения расчета с использованием видеоускорителя NVidia GeForce 9600GT составляет 118мс, то есть простой подход с использованием в расчете видеокарты, что называется «в лоб», дает ускорение около 10 раз. Однако для такого простого примера, где данные для разных уравнений никак не связаны между собой эффективность простого применения рассматриваемой технологии оказалась весьма низкой. В данном случае дело заключается в крайне неэффективном использовании возможностей GPU. Основной проблемой в данном примере является работа с памятью. Нетрудно заметить, что в приведенной программе используется только глобальная (медленная) память GPU (__global__ void compute(eqdata* data)), в то время как использование быстрой разделяемой памяти могло бы значительно увеличить производительность. Перепишем реализацию кода так, чтобы минимизировать работу с глобальной памятью: __global__ void compute(eqdata* data) { __shared__ float a[BLOCK_SIZE]; __shared__ float b[BLOCK_SIZE]; __shared__ float c[BLOCK_SIZE]; __shared__ float reX1[BLOCK_SIZE]; __shared__ float reX2[BLOCK_SIZE]; __shared__ float imX1[BLOCK_SIZE]; __shared__ float imX2[BLOCK_SIZE]; int tx = threadIdx.x; int bx = blockIdx.x; int i = tx+bx*BLOCK_SIZE; float d; a[tx] = data[i].a; b[tx] = data[i].b; c[tx] = data[i].c; __syncthreads(); d = -b[tx]*b[tx] + 4*a[tx]*c[tx]; reX1[tx] = -b[tx]; reX2[tx] = -b[tx]; imX1[tx] = 0; imX2[tx] = 0; if(d > 0) { reX1[tx] += sqrtf(d); reX2[tx] -= sqrtf(d); } else { imX1[tx] += sqrtf(-d); imX2[tx] -= sqrtf(-d); } data[i].re_x1 = reX1[tx]/(2*a[tx]); data[i].re_x2 = reX2[tx]/(2*a[tx]); data[i].im_x1 = imX1[tx]/(2*a[tx]); data[i].im_x2 = imX2[tx]/(2*a[tx]); } Как видно из приведенного выше фрагмента, в данной версии работа с глобальной памятью происходит лишь в момент загрузки коэффициентов уравнения в быструю память, а также сохранения результатов. Все остальные операции происходят с использованием быстрой разделяемой памяти. Во втором варианте кода время выполнения расчета на видеокарте NVidia GeForce 9600GT составляет 37мс, что в 3 раза быстрее предыдущей версии. Общее ускорение по сравнению с однопроцессорным вариантом составило при этом около 30 раз. Но и это, как мы увидим ниже, не предел. Еще одним недостатком рассматриваемой реализации является то, что организация хранения данных, описанная выше (см. описание структуры eqdata), не позволяет осуществлять «связанные» запросы к памяти (см. выше), так как данные, читаемые различными потоками, расположены в памяти неупорядоченно. Для оптимизации обращения к памяти изменим организацию хранения данных следующим образом: struct eqdata { float a[N]; float b[N]; float c[N]; float re_x1[N]; float re_x2[N]; float im_x1[N]; float im_x2[N]; }; То есть вместо использования массива структур, мы используем структуру, хранящую в себе массивы коэффициентов. В этом случае «соседние» потоки будут запрашивать данные, располагающиеся в глобальной памяти подряд, что сделает все запросы «связанными», а это, в свою очередь, повысит скорость считывания. Ниже приведена реализация алгоритма, в которой сделаны соответствующие изменения в организации хранения данных: __global__ void compute(eqdata* data) { __shared__ float a[BLOCK_SIZE]; __shared__ float b[BLOCK_SIZE]; __shared__ float c[BLOCK_SIZE]; __shared__ float reX1[BLOCK_SIZE]; __shared__ float reX2[BLOCK_SIZE]; __shared__ float imX1[BLOCK_SIZE]; __shared__ float imX2[BLOCK_SIZE]; int tx = threadIdx.x; int bx = blockIdx.x; int i = tx+bx*BLOCK_SIZE; float d; a[tx] = data->a[i]; b[tx] = data->b[i]; c[tx] = data->c[i]; __syncthreads(); d = -b[tx]*b[tx] + 4*a[tx]*c[tx]; reX1[tx] = -b[tx]; reX2[tx] = -b[tx]; imX1[tx] = 0; imX2[tx] = 0; if(d > 0) { reX1[tx] += sqrtf(d); reX2[tx] -= sqrtf(d); } else { imX1[tx] += sqrtf(-d); imX2[tx] -= sqrtf(-d); } data->re_x1[i] = reX1[tx]/(2*a[tx]); data->re_x2[i] = reX2[tx]/(2*a[tx]); data->im_x1[i] = imX1[tx]/(2*a[tx]); data->im_x2[i] = imX2[tx]/(2*a[tx]); } В таком варианте время выполнения расчета составляет около 9мс, что в 4 раза быстрее предыдущей версии и в 115 раз быстрее однопроцессорного варианта. Таким образом, «правильная» организация работы с различными видами памяти позволяет более чем на порядок повысить эффективность применения технологии CUDA даже в таком примитивном примере. При получении коэффициентов ускорения, мы сравнивали различные CUDA-реализации с первоначальной реализацией на CPU. Стоит, однако, отметить, что правильная организация памяти с точки зрения вычислений GPU, не является таковой с точки зрения классических вычислений на CPU и, будучи применена к изначальной CPU-реализации, увеличивает время выполнения на CPU более чем в два раза. Это отличие ещё раз подчеркивает разницу подходов в программировании на GPU и на CPU. Аналогичный пример был рассмотрен с использованием более мощного видеоускорителя nVidia GeForce GTX 275. Результаты исследования эффективности применения обсуждаемой технологии приведены на рисунке 11.
Заключение. Рассмотренные вопросы применения технологии nVidia© CUDA показали следующее.
Важно отметить, что, несмотря на усложнение алгоритмов распараллеливания по сравнению с традиционными подходами, разработанными для решения задач на совокупности центральных процессоров, привлекательность использования графических процессоров обусловлена в первую очередь их относительной дешевизной и доступностью. В следующей части работы предполагается детально рассмотреть подходы к применению CUDA для решения практической задачи о поглощении энергии в оболочке многокомпонентного объекта, облучаемого потоком гамма-квантов. Использованные материалы 1. CUDA 2.1 Quick start Guide – http://www.nvidia.com/object/cuda_develop.html 2. CUDA 2.1 Programming Guide – http://www.nvidia.com/object/cuda_develop.html 3. CUDA 2.1 Reference Manual – http://www.nvidia.com/object/cuda_develop.html 1 General-Purpose computing on Graphics Processing Units – вычисления общего назначения с использованием графических процессоров. 2 данные условия актуальны для устройств с Compute Capability 1.1 и ниже, так как для более поздних устройств выполнение условия 3 не является обязательным. 3 актуально лишь для устройств с Compute Capability 1.1 и ниже, так как для более поздних устройств возможен «сбор» логических запросов в группы, внутри которых выполняются условия 1 и 2. 4 Действительно только для устройств с Compute Capability 1.1 и ниже |
Похожие:
 |
Информация о социально-экономическом развитии муниципального района «Жуковский район» «Жуковский район» расположен на северо-востоке Калужской области. Территория муниципального района «Жуковский район» (далее мр) составляет... |
 |
Положение о Фестивале компьютерного моделирования, графики и дизайна Фестиваль компьютерного моделирования, графики и дизайна (фестиваль) – это смотр знаний и умений студентов, школьников, работающей... |
|
Положение об олимпиаде по компьютерному моделированиЮ, графикЕ и дизайнУ Олимпиада по компьютерному моделированию, графике и дизайну (олимпиада) – это смотр знаний и умений студентов, школьников, работающей... |
|
Батыревский район Кроме основных графических материалов, в пояснительной записке дополнительно приложены схемы графических материалов, а также выданы... |
|
Реферат Визуализация результатов расчетов в Scilab. Построение двух-... Пакет прикладных математических программ, предоставляющий открытое окружение для инженерных (технических) и научных расчётов. Это... |
|
Главный государственный санитарный врач по Раменскому муниципальному... ... |
|
Инструкция по эксплуатации и техническому уходу ... |
|
Широкова Татьяна я пью свое бессмертие вода для молодости и долголетия Защиту интеллектуальной собственности и прав Издательской Компании «Невский проспект» осуществляет юридическая компания «Усков и... |
|
Задачах на 2014 год Итоги работы Отдела социальной защиты населения г. Рошаль Отдела социальной защиты населения г. Рошаль Министерства социальной защиты населения Московской области и подведомственных учреждений... |
|
Отчёт о результатах деятельности отдела загс администрации Лужского... Во исполнение поручения главы администрации Лужского муниципального района Малащенко Олега Михайловича предоставляем отчет о результатах... |
|
Надежные обмуровочные материалы на основе сухих смесей ... |
|
Аэрофобия ... |
|
Ежедневный мониторинг сми 04 июля 2017 Авиакомпания Pegas Fly откроет рейсы из подмосковного аэропорта «Жуковский» в Баку, говорится на сайте аэропорта |
|
Дипломная работа содержит таблиц, литературных источников 11 и графических приложений 4 Дипломная работа содержит таблиц, литературных источников – 11 и графических приложений – 4 |
|
Руководство по эксплуатации жуковский При покупке аппарата требуйте заполнения торгующей организацией талона на гарантийный ремонт, проверьте комплектность и товарный... |
|
Руководство программиста Пакет средств разработки С/С++ программ для процессоров NeuroMatrix для операционных платформ Linux и Win32 |