Анализ производительности нейросетей на GPU с помощью NVIDIA Visual Profiler

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

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

Оптимизация кода на GPU

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

 __device__ inline void interpolate(
        const RGBApixel* pixels, RGBApixel& output,
        int width, float x, float y)
{
        int px = (int)x; // floor of x
        int py = (int)y; // floor of y
        const int stride = width;
        const RGBApixel* p0 = &pixels[0] + px + py * stride;
 
        // load the four neighboring pixels
        const RGBApixel& p1 = p0[0 + 0 * stride];
        const RGBApixel& p2 = p0[1 + 0 * stride];
        const RGBApixel& p3 = p0[0 + 1 * stride];
        const RGBApixel& p4 = p0[1 + 1 * stride];
 
        // Calculate the weights for each pixel
        float fx = x - px;
        float fy = y - py;
        float fx1 = 1.0f - fx;
        float fy1 = 1.0f - fy;
 
        int w1 = fx1 * fy1 * 256.0f + 0.5f;
        int w2 = fx  * fy1 * 256.0f + 0.5f;
        int w3 = fx1 * fy  * 256.0f + 0.5f;
        int w4 = fx  * fy  * 256.0f + 0.5f;
 
        // Calculate the weighted sum of pixels (for each color channel)
        int outr = p1.Red * w1 + p2.Red * w2 + p3.Red * w3 + p4.Red * w4;
        int outg = p1.Green * w1 + p2.Green * w2 + p3.Green * w3 + p4.Green * w4;
        int outb = p1.Blue * w1 + p2.Blue * w2 + p3.Blue * w3 + p4.Blue * w4;
        int outa = p1.Alpha * w1 + p2.Alpha * w2 + p3.Alpha * w3 + p4.Alpha * w4;
 
        output.Red = (outr + 128) >> 8;
        output.Green = (outg + 128) >> 8;
        output.Blue = (outb + 128) >> 8;
        output.Alpha = (outa + 128) >> 8;
}
 
__global__ void bilinear(const int width, const int height,
        RGBApixel* input, RGBApixel* output)
{
        int j = blockDim.y * blockIdx.y + threadIdx.y;
        int i = blockDim.x * blockIdx.x + threadIdx.x;
 
        if (j >= 2 * height) return;
        if (i >= 2 * width) return;
 
        float x = width * (i - 0.5f) / (float)(2 * width);
        float y = height * (j - 0.5f) / (float)(2 * height);
 
        interpolate(input, output[i + j * 2 * width], width, x, y);
}

__device__ inline void interpolate(

const RGBApixel* pixels, RGBApixel& output,

int width, float x, float y)

{

int px = (int)x; // floor of x

int py = (int)y; // floor of y

const int stride = width;

const RGBApixel* p0 = &pixels[0] + px + py * stride;

// load the four neighboring pixels

const RGBApixel& p1 = p0[0 + 0 * stride];

const RGBApixel& p2 = p0[1 + 0 * stride];

const RGBApixel& p3 = p0[0 + 1 * stride];

const RGBApixel& p4 = p0[1 + 1 * stride];

// Calculate the weights for each pixel

float fx = x - px;

float fy = y - py;

float fx1 = 1.0f - fx;

float fy1 = 1.0f - fy;

int w1 = fx1 * fy1 * 256.0f + 0.5f;

int w2 = fx * fy1 * 256.0f + 0.5f;

int w3 = fx1 * fy * 256.0f + 0.5f;

int w4 = fx * fy * 256.0f + 0.5f;

// Calculate the weighted sum of pixels (for each color channel)

int outr = p1.Red * w1 + p2.Red * w2 + p3.Red * w3 + p4.Red * w4;

int outg = p1.Green * w1 + p2.Green * w2 + p3.Green * w3 + p4.Green * w4;

int outb = p1.Blue * w1 + p2.Blue * w2 + p3.Blue * w3 + p4.Blue * w4;

int outa = p1.Alpha * w1 + p2.Alpha * w2 + p3.Alpha * w3 + p4.Alpha * w4;

output.Red = (outr + 128) >> 8;

output.Green = (outg + 128) >> 8;

output.Blue = (outb + 128) >> 8;

output.Alpha = (outa + 128) >> 8;

}

__global__ void bilinear(const int width, const int height,

RGBApixel* input, RGBApixel* output)

{

int j = blockDim.y * blockIdx.y + threadIdx.y;

int i = blockDim.x * blockIdx.x + threadIdx.x;

if (j >= 2 * height) return;

if (i >= 2 * width) return;

float x = width * (i - 0.5f) / (float)(2 * width);

float y = height * (j - 0.5f) / (float)(2 * height);

interpolate(input, output[i + j * 2 * width], width, x, y);

}

Такое ядро скомпилируется и будет правильно работать. Однако, если бы мы умели читать ассемблерные листинги, то увидели бы кое-что интересное:

$ cuobjdump -sass bilinear_gpu | grep LD.*U8
        /*0290*/                   LDG.E.U8.SYS R9, [R4+0x6] ;
        /*02a0*/                   LDG.E.U8.SYS R6, [R4+0x5] ;
        /*02b0*/                   LDG.E.U8.SYS R0, [R4+0x4] ;
        /*02c0*/                   LDG.E.U8.SYS R8, [R4+0x7] ;

$ cuobjdump -sass bilinear_gpu | grep LD.*U8

/*0290*/ LDG.E.U8.SYS R9, [R4+0x6] ;

/*02a0*/ LDG.E.U8.SYS R6, [R4+0x5] ;

/*02b0*/ LDG.E.U8.SYS R0, [R4+0x4] ;

/*02c0*/ LDG.E.U8.SYS R8, [R4+0x7] ;

Оказывается, компилятор не распознал процесс загрузки пикселей и сгенерировал ассемблерный код для однобайтовых значений вместо целого RGBA-вектора. Негативный эффект будет частично самортизирован L1-кэшем, с учётом того, в какой мере он занят поддержкой других операций. Но подобное решение всё равно может существенно замедлить приложение. На большом изображении (6496px × 6618px × 24bpp) для GPU Tesla V100 разница в скорости составляет примерно 40%:

biinear/unoptimized$ ./bilinear_gpu ../PIA00004.bmp 
GPU kernel time = 0.003350 sec
bilinear/unoptimized$ ./bilinear_gpu ../PIA00004.bmp 
GPU kernel time = 0.003363 sec
bilinear/coalescing$ ./bilinear_gpu ../PIA00004.bmp 
GPU kernel time = 0.002358 sec
bilinear/coalescing$ ./bilinear_gpu ../PIA00004.bmp 
GPU kernel time = 0.002353 sec

biinear/unoptimized$ ./bilinear_gpu ../PIA00004.bmp

GPU kernel time = 0.003350 sec

bilinear/unoptimized$ ./bilinear_gpu ../PIA00004.bmp

GPU kernel time = 0.003363 sec

bilinear/coalescing$ ./bilinear_gpu ../PIA00004.bmp

GPU kernel time = 0.002358 sec

bilinear/coalescing$ ./bilinear_gpu ../PIA00004.bmp

GPU kernel time = 0.002353 sec

В этом случае чтение пикселей оптимизировано вручную с помощью union:

union RGBApixel_
{
        RGBApixel p;
        int i;
};
 
const RGBApixel* p0 = &pixels[0] + px + py * stride; // pointer to first pixel
 
// Load the four neighboring pixels
RGBApixel_ p1_; p1_.i = *(int*)&p0[0 + 0 * stride];
RGBApixel_ p2_; p2_.i = *(int*)&p0[1 + 0 * stride];
RGBApixel_ p3_; p3_.i = *(int*)&p0[0 + 1 * stride];
RGBApixel_ p4_; p4_.i = *(int*)&p0[1 + 1 * stride];
 
const RGBApixel& p1 = p1_.p;
const RGBApixel& p2 = p2_.p;
const RGBApixel& p3 = p3_.p;
const RGBApixel& p4 = p4_.p;

union RGBApixel_

{

RGBApixel p;

int i;

};

const RGBApixel* p0 = &pixels[0] + px + py * stride; // pointer to first pixel

// Load the four neighboring pixels

RGBApixel_ p1_; p1_.i = *(int*)&p0[0 + 0 * stride];

RGBApixel_ p2_; p2_.i = *(int*)&p0[1 + 0 * stride];

RGBApixel_ p3_; p3_.i = *(int*)&p0[0 + 1 * stride];

RGBApixel_ p4_; p4_.i = *(int*)&p0[1 + 1 * stride];

const RGBApixel& p1 = p1_.p;

const RGBApixel& p2 = p2_.p;

const RGBApixel& p3 = p3_.p;

const RGBApixel& p4 = p4_.p;

В новой версии приложения однобайтовые загрузки превратились в одну четырёхбайтовую:

 $ cuobjdump -sass bilinear_gpu | grep LD.*
        /*0290*/                   LDG.E.SYS R9, [R12+0x4] ;

1 2	$ cuobjdump -sass bilinear_gpu \| grep LD.* /0290/ LDG.E.SYS R9, [R12+0x4] ;

Но главный вопрос — как разработчику заметить такой случай? Опыт и знания стоят времени и денег. Регулярное чтение ассемблерного кода могло бы помочь, но не очень реалистично на практике. Поэтому существуют более универсальные средства анализа, такие как:

Аппаратные счётчики в GPU для сбора статистики о работе программы;
Метрики, которые обрабатывают данные счётчиков и объединяют их в качественные характеристики приложения.

Грубо говоря, такая важная метрика, как GPU Occupancy, — это что-то похожее на рейтинг вашего смартфона в программе AnTuTu. Сбор и анализ метрик называется профилированием. Сведения о параметрах работы приложения, полученные в результате профилирования, гораздо легче анализировать и использовать для повышения эффективности. Так, программа NVIDIA Visual Profiler превращает счётчики и метрики в графики и диаграммы, и даже даёт советы по оптимизации. Рассмотрим её основные возможности.

*Пример анализа использования памяти в NVIDIA Visual Profiler*

Обзор NVIDIA Visual Profiler

NVIDIA Visual Profiler — это графический инструмент профилирования, который отображает хронологию загрузки CPU и GPU во время работы вашего приложения. Программа автоматически анализирует GPU-ядра и помогает определить возможности для оптимизации.

Подготовка приложения к профилированию

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

Основная рекомендация при работе с Visual Profiler — использовать небольшие участки для профилирования. По умолчанию данные собираются в течение всего времени запуска, но лучше анализировать только критические места. Так вы сможете сосредоточить внимание на коде, оптимизация которого приведёт к существенному увеличению производительности. Использование Visual Profiler с большими и сложными приложениями может вызывать зависание движка JVM, поэтому лучше запускать его не более чем на 30 секунд.

Также стоит присвоить пользовательские названия ресурсам CPU и CUDA, поскольку на временной шкале Visual Profiler имена по умолчанию не слишком информативные. Если использовать более понятные и говорящие наименования, то можно улучшить понимание поведения приложения, особенно когда в нём присутствует много устройств, контекстов или потоков.

Создание сеанса профилирования

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

После применения настроек Visual Profiler немедленно запустит ваше приложение и начнёт собирать данные, необходимые для первого этапа управляемого анализа (если только при создании сеанса не была выбрана опция Don't run guided analysis). Систему управляемого анализа можно использовать для получения рекомендаций по улучшению производительности приложения.

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

Разбор результатов

— Временная шкала

В Visual Profiler можно одновременно открыть несколько временных шкал на разных вкладках. На следующем рисунке показана шкала для приложения CUDA:

В верхней части находятся горизонтальные отметки времени, прошедшего с начала профилирования приложения. В левой части отображены единицы исполнения: процесс (process), потоки (thread), GPU (device), контексты (context), ядра (kernel), стримы (stream) и т.д. (с полным списком можно ознакомиться в документации). В центре показаны строки, отражающие активность отдельных элементов. Каждая строка отображает интервалы времени между началом и окончанием каких-либо процессов. Например, строки напротив ядер показывают время начала и окончания выполнения этого ядра.

— Анализ

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

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

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

— Дизассемблер

Source-Disassembly View используется для отображения результатов анализа на уровне ассемблера. В исходном коде отмечается интенсивность и эффективность выполнения отдельных команд. Соответствующие маркеры окрашиваются в разные цвета в зависимости от уровня критичности — низкого, среднего или высокого. Эта информация — основа для базового приёма оптимизации: выделение разработчиком наиболее критичных («тяжёлых») на общем фоне команд (hotpoints) и их переработка — изменение программной логики, понижение точности и так далее.

— Просмотр сведений о GPU

GPU Details View показывает таблицу с информацией о каждом копировании памяти (memcpy) и запуске ядра (kernel) в профилируемом приложении. Для ядер в столбцах показаны соответствующие метрики и события.

Подробный справочник по метрикам можно найти в документации.

Анализ производительности ML-приложения на GPU

Рассмотрим пример профилирования ML-приложения, которое использует мощности GPU. Предположим, перед нами стоит задача обучить LSTM-нейросеть для генерации связного текста. Для этого был написан скрипт на Python, основанный на Keras. Обратите внимание, что сам python-код явно GPU не использует. Но «выход» на GPU может происходить уже внутри функций библиотеки Keras без непосредственного участия пользователя. Поэтому смело запускайте NVIDIA Profiler, указывая в качестве приложения сам Python-скрипт, и профилировщик доберётся в нём до GPU-кода. Тем не менее, следует ограничить длительность профилирования небольшим интервалом, например, 30 секунд. Поскольку алгоритм обучения нейросети однообразен и периодичен, мы можем считать, что этот интервал характеризует поведение приложения в целом, упрощая при этом работу профилировщика по сбору данных.

Для эффективной загрузки GPU должны выполняться два условия:

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

На рисунке ниже показан процесс профилирования нашего ML-скрипта:

Слева направо идёт время выполнения, маленькие жёлтые «кирпичики» в центре — API-вызовы в потоке CPU, который запускает расчёты, а зелёные отрезки снизу — загрузка GPU. 100% использования GPU, как нетрудно догадаться, отображалось бы не маленькими квадратиками, а одним непрерывным блоком. Но на практике этого добиться невозможно, и лучшее, на что можно рассчитывать — чтобы между «кирпичиками» почти не было зазоров. Таким образом, условие (1) в этом примере не выполняется.

Рассматриваемое приложение использует ядро умножения матриц с одинарной точностью (Single precision floating General Matrix Multiply, SGEMM). В окне свойств (справа) указано, что ядро запущено в конфигурации «16 блоков по 128 потоков». Это дало бы хороший результат на GPU начального уровня, но в мощных графических процессорах с архитектурой VOLTA доступно до тысячи таких блоков! Из этого следует, что ядро использует менее 1/10 доступных ресурсов.

⌘⌘⌘

Теперь вы знаете, как написать неэффективное приложение для GPU (лучше, конечно же, этого не делать). Ваш заказчик может возразить, что код всё-таки работает на 10% быстрее с видеокартой NVIDIA RTX. Такой аргумент звучит как «Приора оказалась чуть быстрее Феррари в гонке по болоту». Другими словами, если приложение не раскрывает преимуществ графического процессора, то производительность может быть почти любой, а в абсолютных значениях — слабой во всех вариантах. Чтобы получить существенное ускорение на мощных GPU, необходимо решать по-настоящему высоконагруженные задачи.

Серверы с GPU для обучения нейросетей

Редакция

18 октября 2019