Продолжение статьи1 и статьи 2.

Ниже под катом, расскажу об опыте автора по использованию GPU для расчетов, в том числе в рамках создания бота для участия в AI mini cup. Но скорее это эссе на тему GPU.

-Имя у вас волшебное…
-Знаете что, Джоэл?.. Волшебство уходит...

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

А потом появилась физмат школа с компьютерами yamaha msx, программируемый калькулятор МК дома и стало не до химии. Интересы ребенка сместились к компьютерам. И что не хватало автору с первого знакомства с компьютером так это реакции горения, его программы тлели, не было этого ощущения природной мощи. Можно было увидеть процесс оптимизации вычислений в играх, но в то время заменять вычисление sin() на таблицу значений этой функции автор не умел, интернета не было...

Так вот чувство радости от вычислительной мощности, чистого горения, автор смог получить использую в расчетах GPU.

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

Начнем с простых форм. Вычисления на GPU поддерживают несколько фреймворков, но самые известные NVIDIA CUDA и OpenCL. Мы возьмем CUDA и сразу нам придется сузить свой набор языков программирования до C++. Существуют библиотеки подключения к CUDA других языков программирования, например ALEA GPU на С#, но это скорее тема отдельной обзорной статьи.

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

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

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

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

Пример записи основного цикла в CUDA (kernel):

    int  tid = blockIdx.x * blockDim.x + threadIdx.x;
    int  threadN = gridDim.x * blockDim.x;
    for (int pos = tid; pos < numElements; pos += threadN)
    {
        //  вычисления по параметру pos,  итерации цикла будут выполняться параллельно, другими словами цикл распадется на отдельные thread  каждый со своим параметром pos. 
Важное замечание: порядок выполнения отдельных  thread не зависит от вас, поэтому thread с номером pos=1146  может выполняться раньше чем thread c номером pos=956. Это надо помнить при работе с параллельными алгоритмами. Здесь как в зазеркалье много вещей непривычных для последовательно исполняемых программ.
     }

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

Думаю читателю уже становится понятно, что класс задач применительно к массовому параллельному программированию стремительно сужается. Если мы говорим про создание игр, 3d движки для рендеринга, нейронные сети, редактирование видео и другие подобные задачи, то полянка для самостоятельных действий читателя сильно истоптана, существуют для этих задач большие программы, маленькие программы, фреймворки, библиотеки известные и неизвестные. То есть остается область как раз из темы, создать свою маленькую вычислительную ракету, не SpaceX и Роскосмос, а что-то домашнее, но вполне себе злое до вычислений.

Вот картинка вполне себе пламя ракеты изображено.

К слову о задачах, которые параллельный цикл в ваших руках не сможет решить. И об этом уже подумали создатели CUDA в лице разработчиков NVIDIA.

Существует библиотека Thrust местами полезная до "без вариантов" сделать по другому. Кстати, не нашел полноценного ее обзора на Хабре.

Чтобы понять принцип ее работы, предварительно нужно сказать три предложения о принципах работы CUDA. Если нужно больше слов, то по ссылке можно прочитать.

Принципы работы CUDA:

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

Так вот, библиотека Thrust снимает рутину и берет на себя часть "сложных" для CUDA задач, вроде суммирования массивов или их сортировки. Вам больше не нужно писать отдельное ядро, загружать в память указатели и копировать данные по этим указателям в память GPU. Всë таинство произойдет у вас на глазах в основной программе и со скоростью немного уступающей CUDA. Библиотека Thrust написана на CUDA, так что это одного поля ягоды в части производительности.

Что нужно сделать в Thrust так это создать массив(thrust::vector) в рамках его библиотеки, который совместим с обычным массивов(std::vector). То есть конечно не так все просто, но смысл сказанного автором похож на правду. Действительно два массива, один на GPU(device), другой в основной программе (host).

Пример покажет простоту синтаксиса ( X, Y, Z массивы):

// initialize X to 0,1,2,3, ....
    thrust::sequence(X.begin(), X.end());

    // compute Y = -X
    thrust::transform(X.begin(), X.end(), Y.begin(), thrust::negate<int>());

    // fill Z with twos
    thrust::fill(Z.begin(), Z.end(), 2);

    // compute Y = X mod 2
    thrust::transform(X.begin(), X.end(), Z.begin(), Y.begin(), thrust::modulus<int>());

    // replace all the ones in Y with tens
    thrust::replace(Y.begin(), Y.end(), 1, 10);

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

У автора небольшой опыт использования CUDA и С++, месяца два. До этого около года C#. Это конечно немного противоречит началу статьи о его раннем знакомстве с компьютерами, физмат школе и прикладной математики в качестве образования. Скажу, так получилось. Но для чего пишу эту статью, нет не о том, что вот так все освоил, а о том, что С++ оказался комфортным языком(раньше я его побаивался, на фоне статей в хабре типа "лямбда функции → перегруз внутренних операторов, как все переопределить"), видно что годы его развития привели к вполне дружелюбным средам разработки (IDE). Сам язык в его последней версии, вроде как и мусор из памяти собирает, не знаю как было до этого. По крайней мере, программы написанные автором на самых простых алгоритмических конструкциях, сутками гоняли вычислительные алгоритмы для ботов и не было утечек памяти и других сбоев при высокой нагрузке. Это касается и CUDA, поначалу кажется сложным, но в основе лежат простые принципы и конечно местами трудоемко инициализировать массивы на GPU если их много, но ведь потом у вас будет своя маленькая ракета, с дымком с видеокарты.

Из классов объектов для тренировки с GPU автор рекомендует клеточные автоматы. Одно время был рост популярности и моды на них, но потом нейронные сети перехватили умы разработчиков.
Вплоть до:

"всякая величина в физике, включая время и пространство, является конечной и дискретной. "
чем не клеточный автомат.

А ведь красиво, когда три несложных формулы могут создать такое:

Если будет интересно почитать про клеточные автоматы на CUDA пишите в комментариях, там материала на небольшую статью наберется.
А это первоисточник по клеточным автоматам (под видео есть ссылки на исходники):

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

Создание бота для участия в AI mini cup. Опыт применения GPU