Графические акселераторы для высокопроизводительных вычислений. Часть 2

Эта статья подготовлена на основе доклада Андрея Чередарчука и Александра Судакова на Root Linux Conference 2017 — ежегодной конференции embedded- и Linux-разработчиков.

Андрей Чередарчук — ИТ-инструктор и администратор. Сертифицированный инструктор учебных программ HP, IBM, VMware, ранее также и Cisco. Разрабатывает авторские учебные курсы. Занимается поддержкой HPC-инфраструктуры в НАН Украины.

Александр Судаков — глава лаборатории параллельных вычислений, доцент КНУ им. Тараса Шевченко. Одно из основных направлений его научной деятельности — высокопродуктивные вычислительные компьютерные системы. Александр является разработчиком и руководителем вычислительного кластера информационно-вычислительного центра КНУ им. Тараса Шевченко. Принимал участие в создании первых в Украине сайтов Grid-систем.

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

Современные тренды в HPC

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

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

На чем разрабатывать

Какие сегодня есть фреймворки для разработки под GPU и под процессоры в общей памяти? Во-первых, OpenCL. Эта технология рассчитана на использование как GPU, так и процессоров хоста. Из всех технологий, которые сегодня есть, OpenCL поддерживает наибольшее количество видов аппаратных ресурсов. Но есть мнение (иногда спорное), что программы, написанные под OpenCL, не всегда работают на NVidia GPU так же быстро, как программы, написанные с использованием CUDA.

Название технологии CUDA расшифровывается как Compute Unified Device Architecture. Ее большой недостаток — то, что она поддерживает только устройства Nvidia. С другими устройствами, скажем AMD, она не работает. Но зато большинство вычислительных приложений под NVidia GPU разрабатывается именно с использованием CUDA.

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

OpenACC — это еще один фреймворк, который похож на OpenMP и поддерживается некоторыми коммерческими компиляторами. В OpenACC есть набор готовых библиотек для работы с GPU. OpenACC не распространяется свободно.

Если сравнивать CUDA и OpenCl, то у них очень похожая идеология разработки. Однако API у них отличается. Если разобраться с одним, то понять другой потом будет несложно. До появления OpenCL 2.0 главным принципиальным отличием CUDA и OpenCl была поддержка Managed Memory в технологии CUDA. Технологии быстро развиваются и сегодня OpenCl имеет аналог, который называется Shared Virtual Memory. В дальнейшем остановимся более подробно на технологии CUDA. Далее все изображения и термины взяты с сайта Nvidia, где документация для разработчиков под CUDA есть в свободном доступе.

Модель вычислений

Основная терминология CUDA такая: хост — это вычислительный узел, на котором запускается ваша программа. В этот вычислительный узел, у которого есть один или несколько процессоров, вставлен один или несколько графических акселераторов, которые называются девайсами или устройствами. Программа, которая работает на хосте и устройстве, состоит из двух частей: это код хоста, который выполняется на процессорах хоста, и так называемый кернел (Device Kernel), или ядро устройства. Ядро выполняется на графическом акселераторе. Любая программа для GPU имеет и ядро, и код хоста.

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

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

В чем особенность выполнения кернелов на устройстве? В кернелах используется так называемая SIMD-модель (Single Instruction Multiple Data). Это модель массового параллелизма, в которой все потоки (thread) выполняют один и тот же код, но для разных данных. Если у вас есть какой-то массив, то каждый поток обрабатывает один элемент этого массива. При этом набор из нескольких потоков, которые называются варпом (Warp), физически выполняют одну машинную инструкцию.

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

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

Блоки объединяются в грид (Grid). Грид — это фактически набор различных блоков, которые могут выполняться на устройстве. Каждый поток имеет свой номер в блоке, каждый блок имеет свой номер в гриде. Эти номера могут быть представлены в виде одно-, двух- или трехмерной решетки. Именно по номеру потоков и блоков обычно осуществляется адресация элементов массивов, с которыми работает поток. Всего в гриде может быть очень большое количество потоков, больше миллиарда. Понятно, что выполняться будут только те потоки, для которых хватит ресурсов физических устройств (мультипроцессоров). Остальные потоки будут планироваться на выполнение, но выполняться позже.

То есть в варпе все выполняется одновременно. Блок — почти одновременно. Грид — параллельно или псевдопараллельно, как «захочет» планировщик.

Код кернела

Что собой представляет самый простой код кернела? Простой код кернела может представлять собой вот такую функцию на языке С:

О том, что это кернел говорит слово «global». В эту функцию передаются указатели на массивы данных a, b, c в памяти устройства. Задача кернела — сложить массивы a + b и записать результат в массив c. Все потоки кернела выполняют одинаковые операции сложения и присваивания, но для разных элементов массивов. Для каждого элемента массива в каждом потоке вычисляется свой индекс на основании номера блока, номера потока и размера блока. Разные потоки одновременно складывают разные элементы двух массивов. Это типичная идеология для разработки под CUDA.

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

Код хоста для запуска кернела вызывает вот такую вот функцию: mykernel<<< N_Blocks, N_Threads_Per_Block >>>(a, b, c,); Сюда передается количество блоков (размер грида), количество потоков в блоке (размер блока) и аргументы кернела. Этот синтаксис поддерживается компиляторами CUDA.

Распределенность

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

Если программа хоста выполняется на виртуальной машине, то появляются еще промежуточные этапы передачи данных, поэтому скорость обычно не превышает 4-8 ГБ в секунду. Это, конечно, быстрее, чем по локальной сети, но по сравнению со скоростью передачи внутри устройства, это очень медленно. Внутри устройства скорость передачи может быть больше 100 ГБайт/сек. Поэтому если вам удастся написать программу так, что время передачи данных на устройство будет очень маленьким по сравнению со временем обработки, то можно получить очень высокую производительность. Чтобы избежать задержек при передаче данных, необходимо максимально кэшировать данные на хосте и устройстве.

Память

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

Глобальная и локальная память — это относительно медленная память (обычно DDR4 или DDR5). В этой памяти выделяются области константной памяти, текстурной памяти и в зависимости от computing capability — L2-кэш. L1-кэш есть у всех устройств.

Вторая особенность — на GPU объем памяти значительно меньше, чем на хосте. Это один из недостатков, с которыми приходится сталкиваться. Ниже показаны примеры, как можно при написании своего кода указать, какую память вам нужно использовать. Ключевое слово «global» относится к кернелам, т. е. у кернела указывается ключевое слово «global» при написании кода. Если указано ключевое слово «device», то это значит, что эти данные должны находиться в глобальной памяти. Т. е. вы можете явно указать, где хранить ваши данные на GPU. «Constant» означает, что данные будут храниться в константной памяти. Это read-only память, которая кэшируется в L1 кэше. Изменять эти данные нельзя, они задаются на этапе компиляции или на этапе запуска ядра. «Shared» означает, что данные должны храниться в быстрой общей памяти блока. Ниже указаны типичные размеры областей памяти, с которыми приходится иметь дело.

__global__ void mykernel(int)
    – Код кернела
__device__ float data;
    – Данные в голобальной памяти
__constant__ float data [<=64 Kbytes];
    - Константа в кешируемой глобальной памяти
__shared__ float data[<=48 Kbytes];
    – Общая память для блока – очень быстрая!
Текстурная память
    – Кэшируется в L1 , редко используется для GPGPU
__restrict__ float *pdata
    – Только чтение! Кэширование в L2
Включение/выключение кэширования при компиляции
    - Кэшировать глобальную память в L1 –Xptxas –dlcm=cg
       может быть SEGFAULT!
    - Размер L1, общей памяти, блока регистров –maxrregcount=N
    - cudaFuncSetCacheConfig Размер кэша/общей памяти
Локальная память – стек в глобальной памяти GPU, медленно
    - Маленький объем, кэшируется в L1

Текстурная память для расчетов сейчас уже не так актуальна. В старых версиях GPU текстурную память можно было использовать как быструю кэшируемую память. Сейчас она больше актуальна для обработки 3D-графики. Ключевое слово «restrict» означает, что данные хранятся в L2-кэше в read-only режиме. Если у вас какие-то структуры данных часто используются как read-only, вы можете указать компилятору, чтобы он обращал на это внимание и сгенерировал более эффективный код. Также можно включать и отключать кэширование глобальной памяти в L1-кэше, но это поддерживают не все GPU.

Дело в том, что если у вас GPU, скажем, Tesla с computing capability 3,5 и выше, то там такое кэширование поддерживается. Если же у вас GeForce с той же computing capability 3,5 — то там такого кэширования нет. С L1 кэшированием глобальной памяти нужно быть осторожным. Так как L1-кэш свой для каждого вычислительного процессора GPU, иногда возникают проблемы с общими структурами данных в глобальной памяти при кэшировании в L1. L2-кэширование глобальной памяти работает по-умолчанию, если доступно. На этапе компиляции можно изменять количество регистров, и даже на этапе выполнения можно изменять соотношения между общей памятью, глобальной памятью и L1-кэшем.

Следующий этап — память нужно выделить. Графический акселератор — это отдельное устройство. Если вы пишете программу, то нужно четко понимать, где выделяется память — в оперативной памяти хоста или где-то на устройстве — и как копировать данные с хоста на графический акселератор и обратно. В чем тут может быть проблема? Если программа разрабатывается «с нуля», все достаточно просто. Если же вы портируете существующую программу, в которой уже есть сложные структуры данных, то скопировать их с одного устройства на другое или, хуже того, с устройства с одной архитектурой на устройство с другой архитектурой — проблема. Если этим приходится заниматься, то нужно использовать все возможности работы с памятью.

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

Например, у вас есть готовый массив, вы его на хосте заполнили, потом отобразили на устройство и там используете. Можно отобразить память одного устройства на память другого. Кроме всего прочего, можно также выполнять выделение памяти в самом кернеле. Там возможностей значительно меньше, но все равно есть функция malloc и оператор new. Эти все вещи в зависимости от computing capability и версии CUDA будут немного разными.

Для работы с памятью существует такая подсистема, как Unified Virtual Address Space (UVA). Это программно-аппаратная технология, которая позволяет очень сильно облегчить жизнь программисту. Фактически память хоста и память устройства видится как одна общая память. Что UVA позволяет делать? Она позволяет прозрачно для программиста отображать память с хоста на устройство и с одного устройства на другое. Таким образом, программисту нет необходимости выполнять сложное копирование данных с хоста на устройство и обратно. Можно просто заполнить массив структур, связанный список, или хэш-таблицу на хосте, отобразить ее на устройство и там сразу же использовать. Но есть ряд нюансов.

Первый вариант использования UVA это — zero-copy pinned memory. Это очень быстрый и удобный способ. Единственная проблема: адрес отображенной памяти на устройстве не совпадает с адресом памяти на хосте. Поэтому если в памяти массив данных типа float — то проблем нет. Если в памяти структура данных с указателями — проблемы есть. Как это можно использовать? Объявляем массив на хосте, отображаем этот массив на память устройства, получаем адрес на устройстве и запускаем кернел с этим адресом в качестве параметра.

Еще один инструмент, который есть в CUDA (и которого совсем недавно не было, но уже есть в OpenCL) — это managed memory. Это еще более удобная штука. При использовании managed memory даже указатели сохраняются. То есть адрес памяти на хосте и адрес памяти на устройстве при отображении будет один и тот же. В сложной хэш-таблице с указателями все указатели на хосте и на устройстве будут одинаковыми и будут одинаково работать. Это очень удобно, но гарантированно будет работать, если кернел не обращается к отображаемой памяти одновременно с кодом хоста.

Опять же, жизнь не стоит на месте, и начиная с Nvidia Pascal (computing capability 6) и CUDA-8, поддерживается параллельное использование managed memory кодом хоста и устройства, однако для более старых устройств все остается по старому! Т. е. вы на хосте заполнили массив, запустили кернел, хост подождал, кернел посчитал — хост забрал результаты.

Если вы отобразили массив как managed memory, то при попытке хоста записать одну его часть, а устройства — другую, возникнет ошибка. Конечно, есть возможность «переотобразить» части данных для использования хостом или устройством, но это неудобно и медленно. Поэтому если программа пишется под десктоп, где один слабый процессор или программа работает в клауде под виртуализацией на слабых процессорах, но есть GPU, то managed memory — это спасение для тех, кто хочет портировать свои программы на CUDA и получить высокую производительность.

Как Nvidia рекомендует использовать managed memory в структурах данных? Определяется класс, в котором есть оператор new и оператор delete, выделяющие и освобождающие память с использованием managed memory.

// class for transparent allocations
struct cuda_mapped {
    void *operator new(size_t len) {
      void* ptr;
      cudaMallocManaged(&ptr, len);
      cudaDeviceSynchronize();
      return ptr;
    }
    void operator delete(void *ptr) {
      cudaDeviceSynchronize();
      cudaFree(p);
    }
    …
};
// allocated at device and host
struct complicated: public cuda_mapped {
    complicated* next;
    …
    void func(){
        next = new complicated;
    }
};

Оператор new вызывает API-функцию cudaMallocManaged, оператор delete вызывает API-функцию cudaFree. После этого все необходимые классы объявляются производными от этого класса, и дальше оператор new и оператор delete будут прекрасно работать и прозрачно использовать managed memory (если, конечно, структур данных не слишком много, так как объем managed memory ограничен). Можно с помощью препроцессора (#ifdef __CUDACC__) включать и выключать наследование ваших классов от cuda_mapped в зависимости от того, поддерживает ваш компилятор CUDA или нет. Больше ничего в структурах данных менять не нужно — структуры данных портируются очень просто.

Самый сложный и неприятный способ работы с памятью — это низкоуровневые API: выделение, копирование с устройства на хост, с хоста на устройство, с устройства на устройства. Однако если необходимо получить действительно высокую производительность, то приходится использовать именно этот вариант.

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

template <typename T>
int copy_type_to_cuda(T* cuda_to, T* host_from ){
cudaMemcpy(cuda_to,host_from,sizeof(T),cudaMemcpyHostToDevice) ;
}
-----------------------------------------------------------------------
template <typename T>
int copy_vector_to_cuda(T* to, T* from){
    typename T::value_type **p = (typename T::value_type**)to;
    if(from->size()){
    cudaMalloc(to,from->size()*sizeof(typename T::value_type) );
    p[1] = p[0]+from->size(); p[2] = p[1];
    cudaMemcpy(p[0],from->data(),from->size()*sizeof(typename
T::value_type),cudaMemcpyHostToDevice) ;
    } else { p[0]=p[1]=p[2]=NULL;}
}
--------------------------------------------
template <typename T>
int copy_pointers_vector_to_cuda(T* to, T* from){…}

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

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

Стримы и планирование

Еще одна интересная вещь, которую предоставляет CUDA, — асинхронное планирование с использование CUDA streams. Кто запускал задачи на кластере в batch режиме хорошо знает, что такое очередь, планирование и синхронизация между очередями заданий. Фактически CUDA streams — это набор очередей для выполнения заданий, таких как копирование данных, запуск кернелов и др. Есть дефолтный стрим 0 — это синхронный стрим, куда все действия ставятся в очередь по умолчанию и выполняются последовательно в порядке записи в синхронном режиме.

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

Использование CUDA streams — достаточно удобная ведь для того, чтобы получить максимальную производительность вашей программы как на хосте, так и на устройстве. Если у вас была MPI-программа, то вместо MPI send или MPI receive можно использовать CUDA streams. Т. е. отправка данных на устройства, запуск кернела и так далее. Синхронизация между стримами выполняется с помощью ивентов (событий). При синхронном выполнении вы сначала копируете данные на устройство, потом запускаете кернел, а потом копируете данные обратно.

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

Ниже показан пример планировщика, который использовался при портировании программы, ранее написанной в стандарте OpenMP для одновременной работы на GPU и CPU.

#pragma omp parallel for
for(int i=-gpus_num; i<host_data_size; i++){
    if(i<0)
        start_gpu_job(i);
    else
        start_host_job(i)
}
if(gpus_finish_job())
    increase_gpu_job();
else
    increase_host_job();
wait for gpus();

Этот код содержит цикл по всем данным хоста, которые мы хотим обработать. Строка #pragma omp parallel for — это директива компилятора для автораспараллеливания цикла на несколько потоков хоста. В оригинале цикл начинался с нуля и заканчивался общим количеством данных. После портирования цикл начинается с отрицательного числа, которое по модулю равно общему количеству акселераторов в системе и заканчивается количеством данных, которое обрабатывает хост. Для отрицательных значений параметра цикла запускается обработка данных на GPU, для положительных — вычислительный поток хоста. Операции выполняются параллельно. После завершения цикла проверяется, какая работа закончилась раньше — устройств или хоста, и соответственно корректируется объем данных, с которыми работают хост и каждое из устройств.

Несколько акселераторов

Нет ничего сложного в использовании нескольких GPU. Нужно определить, сколько у вас устройств, какие их характеристики и годятся ли они для вашей задачи. Это можно сделать с помощью функций cudaGetDeviceCount и cudaGetDevice. Перед запуском операций на какое-либо устройство вызывается функция cudaSetDevice, в которую передается номер устройства. Главное, это нужно не забывать делать перед созданием стримов, перед выделением памяти, перед запуском кернела и т. д. Компиляторы CUDA дают возможность сгенерировать код для устройств с различными значениями computing capability и скомпоновать этот код в одной программе, которая будет работать с разными устройствами. Главное, не забыть указать соответствующие опции компилятора.

Компиляторы и библиотеки

На чем и как писать код? В фреймворке CUDA поддерживаются языки C и C++. Некоторые коммерческие компиляторы, как от PGI или от IBM, поддерживают язык Fortran. При отсутствии компилятора, который поддерживает написание кернелов на языке Fortran, можно написать кернел и его вызов на C или C++ и скомпоновать с программой на языке Fortran. Код, который компилируется для устройства и хоста, обозначается ключевыми словами __host__ и __device__ соответственно. Можно с помощью препроцессора определять разный код, который будет компилироваться или только для хоста, или только для устройства.

Далеко не все функции, которые написаны для хоста нужно портировать на устройство. Например, если в классе есть функция «прочитать конфигурационный файл», то она нужна на хосте, а в кернеле совершенно не нужна. Главная проблема портирования кода состоит в том, что для устройства нет всех тех библиотек, которые есть для хоста. Если программа использует, например STL, то для того, чтобы ее портировать для выполнения на устройстве, нужно или переписать код без STL, или реализовать те функции STL, которые вам нужны самостоятельно. В CUDA нет реализации STL для использования в кернелах. Еще больше проблем возникает при необходимости копирования полиморфных объектов, которые создаются на хосте, но вызывать виртуальные методы для них нужно на устройстве. Хотя эту задачу тоже можно решить.

Эффективность

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

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

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

Живой пример

В качестве примера кода, который был успешно портирован с процессоров общего назначения на GPU, можно привести программу авторов для расчета динамики сложных биологических нейросетей от сотен тысяч до сотен миллионов нейронов. С помощью GPU и этой программы впервые удалось промоделировать динамику трехмерных систем в модели Курамото-Сакагучи из примерно 100 000 000 нейронов и открыть новые типы состояний больших связанных систем [Link 1, Link 2].

Графический акселератор GeFroce GT 640, достаточно недорогое устройство, для этой задачи показал производительность примерно равную производительности 12-ядерного вычислительного узла с процессорами Intel Xeon 2620. Примерно такую же производительность показали и более новые энергоэффективные акселераторы Quadro K620. Производительность Tesla K40 (не новый, но дорогой и производительный акселератор) соответствует примерно 50 таким процессорным ядрам, то есть этот акселератор заменяет кластер из 4-х таких узлов [Link 1].

Для портированного кода удалось достичь примерно 60% пиковой производительности Tesla (порядка 600 гигафлопсов). Для GPU это очень неплохо. На CPU программа использовала 75-80% пиковой производительности вычислительного узла.

Выводы

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

Главные недостатки — не очень простое портирование кода, отсутствие библиотек и маленький объем памяти.

LinkedIn

31 комментарий

Подписаться на комментарииОтписаться от комментариев Комментарии могут оставлять только пользователи с подтвержденными аккаунтами.

Как же притяно иногда открыть DOU и читать такого рода статьи. А еще приятнее когда преподаватель твоего вуза , дал нам основы CUDA и открыв статью не так сильно плывешь в терминах и концепциях, ведь уже понимаешь разницу между __shared__, __global__и __constant__ и тд. , понимаешь что это за <<< >>> старанный C++ такой , и как выполнить одну операцию на огромном количестве данных с помощью нитей и блоков. Я понимаю, что это только вершина айсберга, но все же очень полезно студенту познать общие концеции и реализовать несколько алгоритмов которые можно распаралелить. Реализовывали в kernel функциях различные сортировки (bitonic, обменную сортировку со слиянием (параллельную сортировку Бэтчера), поразрядную обменную сортировку) и тд. Конечно это все простейшие алгоримы для профессионала в это области, но очень полезно для новичков.

Большая блгодарность преподаватеям и просто неравнодушим людям за то,что доносите такие знания до масс.

Глобальная и локальная память — это относительно медленная память (обычно DDR4 или DDR5). В этой памяти выделяются области константной памяти, текстурной памяти и в зависимости от computing capability — L2-кэш. L1-кэш есть у всех устройств.

Ещё одно замечание, по всей статье идут какие-то сравнение памятей на скорость и т.п. Вот только со стороны это выглядит как, эта машина жёлтого цвета, поэтому она быстрая. Основной критерий скорости работы кернела с памятью — это количество времени, которое Execution Unit проводит в Stall состоянии ожидая более медленные устройства. У разной памяти разное latency доступа. Константная память обычно идёт прицепом к кернелу, поэтому доступ к ней из кернела максимально быстрый, он не тормозит весь ворп, и он быстрее shared или любой другой. Из вашей статьи всё следует наоборот. Мало того, при должной сноровке можно добиться того, что все константы будут вкомпилены в код напрямую, например коэффициенты FIR или IIR фильтров.

по всей статье идут какие-то сравнение памятей на скорость и т.п.

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

Константная память обычно идёт прицепом к кернелу, поэтому доступ к ней из кернела максимально быстрый

Никаких прицепов, на сколько я знаю, на GPU нет. Инструкции кернела физически хранятся в глобальной памяти, а перед выполнении загружаются в L1-кеш инструкций. L1-кеш, регистры, shared память одну и ту же область. То что доступ к костантной памяти может выпоняться быстрее, чем к shared сваязано с тем, что єта память read-only, это известно этапе компиляции и генерируется более эффективный код доступа к памяти, что обеспесивает меньше

времени, которое Execution Unit проводит в Stall состоянии

.

Рекомендую ознакомиться с документом Demystifying GPU Microarchitecture through Microbenchmarking: www.stuffedcow.net/research/cudabmk — там и ссылка на статью и примеры.

Спасибо большое! Очень интересная информация

В кернелах используется так называемая SIMD-модель (Single Instruction Multiple Data). Это модель массового параллелизма, в которой все потоки (thread) выполняют один и тот же код, но для разных данных

И что, никто за всё время не указал на очевидное в корне неверное высказывание? nVidia использует SIMT (Single-Instruction-Multiple-Thread). SIMD/VLIW используют Intel и AMD.

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

Как выполнится код, если один поток обрабатывает элемент массива со значением 2, а другой со значением 8:

if (a[thread] > 5) {
// делаем одно очень долго
} else {
// делаем другое очень долго
}

?

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

Это как? Что есть результат работы потока?

Майк угомонись, это популярная статья в качестве легкого введения, что такое есть.
Ну и на всю Украину хорошо, ели 5 спецов в GPGPU найдешь.
Так что для просвещения масс статья вполне неплоха.

А нюансы и детали уже изучат те, кто свалит в развитую страну писать на GPGPU.

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

Майк угомонись, это популярная статья в качестве легкого введения, что такое есть.

А кто против?

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

Вот эта?

Для портированного кода удалось достичь примерно 60% пиковой производительности Tesla (порядка 600 гигафлопсов). Для GPU это очень неплохо

С какий пор 40% в GPU stall — это норма?

А нюансы и детали уже изучат те, кто свалит в развитую страну писать на GPGPU.

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

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

Да, всё верно, но разве в статье речь шла про посчитать один раз?

У тебя этой почве развился бред уже.

Может ты и прав. Я уже стараюсь преувеличивать в своем восприятии хоть что-то позитивное в постсовке. А тут вон ребята даже GPU воспользовались сами и не для майнинга битков.

С какий пор 40% в GPU stall — это норма?

Нам бы только иностранное ругать ©. Может 40% в GPU stall и плохо, но речь в статье шла про 60% пиковой производительности GPU, а єто не одно и тоже и для расчетов динамики очень даже нормально, учитавая, что большая часть данных в память GPU не вмещается, идет периодичесакий обмен GPU/host, программа не писалась для GPU специально а портровалась из OpenMP версии с минимальными правками и многое другое. Вкомпилировать «константы прицепом в кернел», извините проблематично, так как данные получаются и меняются на этапе выполнения, объекты разных нейронов создаются динамически в процессе выполнения могут и обрабатываться как хостом так и GPU в зависимости от размера системы и типов моделей.
На всякий случай напомню, что производительность считается как количество операций/на время выполнения. А пиковая произволдительность — максимальное количество операций за такт* тактовую частоту и на практике достичь пиковой производительности нельзя по определению

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

А в чём смысл тогда? Портирование — это не просто скомпилить, абы работало.

Вкомпилировать «константы прицепом в кернел», извините проблематично, так как данные получаются и меняются на этапе выполнения

Тогда какие это константы?

На всякий случай напомню, что производительность считается как количество операций/на время выполнения. А пиковая произволдительность — максимальное количество операций за такт* тактовую частоту и на практике достичь пиковой производительности нельзя по определению

Как вы определяете количество произведённых операций в суперскалярных архитектурах?

И что, никто за всё время не указал на очевидное в корне неверное высказывание? nVidia использует SIMT (Single-Instruction-Multiple-Thread). SIMD/VLIW используют Intel и AMD

А чего тут нужно было указывать? Знатоки проприетарной терминологии не знают общеизвестные вещи? Согласно классификации Флинна, которой все пользуются с 1966 года, en.wikipedia.org/wiki/Flynn’s_taxonomy и SIMT и SPMD и векторные процессорные инструкции (SSE, AVX) и другие подобные штуки относятнся к SIMD классу. Придумыванием новых классов занимаются маркетологи.

Как выполнится код, если один поток обрабатывает элемент массива со значением 2, а другой со значением 8:

if (a[thread] > 5) {
// делаем одно очень долго
} else {
// делаем другое очень долго
}
?

все потоки выполнят if (a[thread] > 5) . Тот где 8 (и другие для которых условие выполняется) выполнят (// делаем одно очень долго) . Тот где 2 (и другие где условие не выполняется) деактивируется. После завершения всех потоков, которые выполняли (// делаем одно очень долго) запустится поток где 2 (и другие где if (a[thread] <= 5)) и будут выполнять (// делаем другое очень долго), а первые деактивируются, пока не закончится (// делаем другое очень долго).
Время выполнения в худшем случае равно время(// делаем одно очень долго)+ время (// делаем другое очень долго)+ оверхед. Если условие выполнялось/не выполнялось для все потоков, то дивергенции кода нет и все выполняется сразу всеми потоками варпа, что быстрее, но не всегда реалтизуется.

Это как? Что есть результат работы потока?

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

Статья посвящалась особенностям портирования существующего кода с CPU на GPU и к єтой части у критика вроде претензий не было.

Знатоки проприетарной терминологии не знают общеизвестные вещи? Согласно классификации Флинна, которой все пользуются с 1966 года, en.wikipedia.org/wiki/Flynn’s_taxonomy и SIMT и SPMD и векторные процессорные инструкции (SSE, AVX) и другие подобные штуки относятнся к SIMD классу. Придумыванием новых классов занимаются маркетологи.

Там же чёрным по английскому написано:
Single instruction, multiple threads (SIMT) is an execution model used in parallel computing where single instruction, multiple data (SIMD) is combined with multithreading. This is not originally part of Flynn’s taxonomy but a proposed addition

Чтобы понять, почему nVidia не называет это SIMD, достаточно взглянуть на машинный код.

все потоки выполнят if (a[thread] > 5) . Тот где 8 (и другие для которых условие выполняется) выполнят (// делаем одно очень долго) . Тот где 2 (и другие где условие не выполняется) деактивируется. После завершения всех потоков, которые выполняли (// делаем одно очень долго) запустится поток где 2 (и другие где if (a[thread] <= 5)) и будут выполнять (// делаем другое очень долго), а первые деактивируются, пока не закончится (// делаем другое очень долго).

А если бранчей в коде более 1, скажем 20? Будем запускать 2^20 потоков для всех кейсов?

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

Результатом процесса производства масла является масло %) Кто принимает решение «но результаты одних потоков будут использоваться, а других нет»?

Статья посвящалась особенностям портирования существующего кода с CPU на GPU и к єтой части у критика вроде претензий не было.

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

Thanks guys.
Very interesting

Стримы и планирование

Interesting if there is an authomatic planning to utilize full gpu concarently.
And also traid off between multiple low cost gpus or one high cost with stream.
And do you have comparison of using streams with different gpus.
If so please share.
Hope i’m not very annoying).

I do not know if any universal solutions for automatic host/gpus scheduling with I/O and execution concurrency exit. IMHO no. It is for the application developer
In our case we did not write the code for GPU specially but we ported the existing CPU code to GPU. So this code it not the optimal for GPU. Streams helped us significantly because we sending a lot of data to/from GPU and GPU kernels are relatively slow. The performance increase was about two fold
We tried streams with Tesla K40, Tesla K20, GeForce G640, Qaudro K620. In most cases a single GPU was used. About 20 streams was an optimal for us. This number was about the same for different GPUs. Only the number of neurons scheduled for GPU changed significantly
IMHO you need streams if your kernels receives and sends a lot of data to/from host and you can partition this data into chunks to not interfere with your algorithm. In this case both low and high cost GPUs will provide better performance.

На счет —

В CUDA нет реализации STL для использования в кернелах

Это таки да, но юзал я одну библиотечку, (docs.nvidia.com/cuda/thrust/index.html) которая позволяет юзать что то похожее на STL для CUDA.
Намного удобне, чем долбаться с С-стайл кодом или писать велосипеды.

Спасибо! Єто ценная информация. В свое время сходу подходящую библиотеку не нашел, но посмотрел, что дописівать не много и дописал велосипед

Народ рекомендует PyTorch или TensorFlow для исследовательской работы. Что Вы о них скажете?

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

Для своих задач великолепны, для других плохи.
Но выбрать что подходит под твою задачу не просто. Надо понимать, что именно делают эти инструменты и отсюда вытекает их применимость.
Нейронные сети в одних задачах хороши в других плохи.
Как пример в речи, только 4-5 лет начали пробовать DNN и уже начали получать неплохие результаты. Но в речи эти DNN по сути развитие HMM. Т.е. они отличаются от тех, что используют в распознавании других образов.
Но есть куча задач, где HMM просто не подходят.
Посему, прежде чем пользоваться инструментом нужно его изучить.
Ну и CUDA используются в той же kaldi при обучении. Но использовать этот движок без его понимания практически невозможно. Уверен, что всё то же самое с необходимостью изучения и с упомянутыми тобой инструментами.
А многие задачи эффективнее решаются совсем другими методами. Всё сильно зависит от задачи.

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

Например у меня до них пока руки не доходят. Голова не резиновая и все в мире из области ML в одну голову не впихаешь.

As for me customize TensorFlow Op in c++ is easy more flexible and have better performance.This is important if you will define special kernels.
Also i use MXnet or TensorFlow if i need highly customizable layers.But if i need to debug i will use Theano or PyTorch it is easy to debug as python.
sorry have no cyrillic keyboard.

Эту было уже приятно читать.
Только вы не указали важного момента. Либ и инструментов, что умеют CUDA сильно больше, чем тех, что умеют OpenCL (Спасибо Интелу и АМД за это).
Так, что сейчас, если стоит выбор между NVIDIA и остальными лучше выбирать NVIDIA.

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

И вы не указали, что С в GPGPU не стандартный и различный в CUDA и OpenCL. И собственно кернелы хоть и сильно похожие, но различные для этих двух решений.

Спасибо! Согласен с комментарием.

Если бы две части были вместе было бы еще лучше. Разбить статью на куски было политическим решением редакции в связи с какими-то техническими особенностями

Когда первый раз посмотрел на CUDA и OpenCL, то как-то сразу захотелось использовать CUDA по многим причинам. Кроме того альтернативы Nvidia GPU у нас не было. В принципе есть возможность получить доступ на один Xeon Phi, но как-то не дошли руки поробовать

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

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

В GPGPU мне пока не нравится, что нет удобного инструмента для компиляции и отладки кернелов. Типичный путь. Написали, потом в проге скомпилировали и запустили, ну или куча трюков с их сохранением и т.п.
Нельзя просто к make прилепить, как другие исходники и вызвать компилятор для них, что подготовит какой-то контейнер, что я уже в проге заюзаю. Точнее что-то есть, но там танец с бубном — меня не напугал.

Да при компиляции в проге с отадкой плохо. У нас компиляция в проге не используется. Есть один злой и универсальный кернел. Кернел копилируется один раз под разное железо и содержит все необходимое. Перекомпилировать нужно только под новое железо или при добавлении новых моделей. Зато отладили и забыли.
Считаем в гриде, поэтому приходится компилировать много разных версий программы под разное железо и софт в том числе и статически линкованные. Кстати с CUDA полностью статически слинковать не вышло. Даже статическая либа CUDA хочет dl_open, а для этого прога должна подерживать динамический линкер. Компилируем єтот зоопарк через automake — помогает, но Makefile.am не тревиальный и громозкий.

И это большой недостаток GPGPU. Долго возиться и шаманить с кернелом, потом его юзать.
Если же у тебя несколько кернелов предполагается, то такое шаманство с каждым по отдельности и отладка на разных GPU (а там архитектура меняется).
В итоге юзание GPGPU становиться дорогим за счет затрат на разработку и к нему прибегают, когда в самом деле это экономически выгодно.
Т.е. нет такого, как CPU, где стоимость разработки сильно дешевле.
Но сейчас ни АМД ни NVIDIA нет смысла снижать стоимость разработки, потому как биткомайнеры их обеспечивают прибылью полностью.
Так что, как мне кажется, эта задница с разработкой под GPGPU продолжиться, пока битки не лопнут.

Согласен. Пока напишешь, пока отладишь, потом оказівается, что на другом железе вілетает. Опять правишь...
По опіту знаю, что на кластерах, где пользуют gpu ,подавляющее большинство приложений — єто уже готовіе программі, которіе где-то сделали за деньги и раздают бесплатно. Мало кто пишет свое для себя — напряжно, много работі, а навара особого нет. На процессорах дольше, но проще. Нашу программу мі начали портировать потому, что для больших систем приходилось одну задачу по две недели считать на одной ноде. MPI мало помогал из-за слабой сети , а кроме того дешевіе gpu таки значительно дешевле сервера, а скорость та же.
На счет майнеров и стоимости разработки может оно и так. Может что-то и бістее поменяется. В последнее время массово ударились в deep learning и что-то подобное, даже железо стали под єто специально делать. Може в связи с єтим и средтва разработки, библиотеки и д.р. появятся под gpu для более массовой аудитории, чтобі от жизни не отстать

Подписаться на комментарии