@ Карта сайта News Автора!

Bog BOS: Использование модели массового параллелизма CUDA для разработки программ

Последние изменения:
2015.11.18: hard: обновлена статья про ИБП и их мониторинг

Последнее изменение файла: 2010.09.01
Скопировано с www.bog.pp.ru: 2017.03.25

Bog BOS: Использование модели массового параллелизма CUDA для разработки программ

CUDA 2.2 (Compute Unified Device Architecture) - ориентированная на массовый параллелизм модель разработки программ, а также набор необходимых инструментов и документации для GPU фирмы NVIDIA. Ранее фирма NVIDIA продвигала для этих целей модель GPGPU. Для разработки программ в рамках модели CUDA необходимо иметь:

Модель CUDA

GPU предназначены для выполнения интенсивных расчётов. Задачи с интенсивным обращением к памяти или сложной логикой будут выполняться неэффективно, т.к. GPU обладает слабыми средствами кеширования обращения к памяти (к тому же их требуется настраивать вручную) и "не переносит" ветвлений в программе, особенно если потоки одного блока расходятся по разным веткам.

Модель CUDA предполагает, что программист в начале разбивает задачу на независимые части (блоки), которые могут выполняться параллельно. Затем каждый блок разбивается на множество параллельно выполняющихся потоков (thread), которые могут зависеть друг от друга. CUDA обеспечивает средства расширения языка C для параллельного запуска множества потоков, выполняющих одну и ту же функцию (ядро, kernel). Максимальный размер ядра - 2 миллиона инструкций PTX (если сумеете скомпилировать - мне не удалось превысить 50 тысяч инструкций (блоб в 400KB) из-за падения nvopencc). Потоки объединяются в блоки (до 512 потоков), блоки объединяются в сетки (решётки, grid). Потоки внутри блока запускаются на одном мультипроцессоре (MP), имеют общую разделяемую память и могут (должны) синхронизовать ход выполнения задачи. Каждый поток имеет уникальный идентификатор внутри блока, выражаемый с помощью одномерного, двумерного или трёхмерного индекса (встроенная структурная переменная threadIdx типа dim3). Размерность блока доступна через встроенную переменную blockDim. Максимальные размерности: 512, 512, 64. Решётки могут быть одномерными или двумерными, максимальное значение элемента индекса: 65535. Индекс блока в решётке доступен через встроенную переменную blockIdx. Компоненты индексов нумеруются с нуля.

Порядок выполнения блоков не определён, блоки должны быть независимыми друг от друга. При запуске ядра блоки решётки нумеруются и распределяются по MP, имеющим достаточную свободную ёмкость региcтров, разделяемой памяти и ресурсов планировщика команд. MP состоит из 8 простых процессоров, 2 процессора для сложных операций (например, умножения ;), пула регистров, разделяемой памяти и планировщика команд. Планировщик команд последовательно разбивает потоки активного блока на порции (warp), по 4 на каждый простой процессор и выполняет по одной простой команде одновременно для всех потоков порции за 4 цикла. Для исполнения одной команды порции потоков MP должен загрузить операнды для всех потоков порции, выполнить команду (одновременно), записать результат. Если доступ к памяти вызывает задержку (до 600 циклов), то планировщик может перейти к следующей порции. Все потоки порции начинают исполнение программы с одного и того же адреса, но каждый простой процессор имеет собственный счётчик команд (?) и регистр состояния, что позволяет осуществить условное исполнение и ветвление. Однако, каждая ветка условия выполняется всеми потоками порции по очереди. Те потоки, для которых условие ветки не выполняется, "пропускают ход". По завершению расхождения все потоки порции опять одновременно выполняют полезную работу. Таким образом, ветвление при исполнении внутри порции сильно замедляет работу ядра. Потоки из разных порций могут выполнять различные ветки совершенно безнаказанно. По завершению всех потоков блока ресурсы MP освобождаются и на него может быть распределён следующий блок.

Планировщик MP имеет ограничения по максимальному числу одновременно запущенных блоков (до 8), максимальному числу порций (до 24 активных порций в версиях аппаратуры 1.0 и 1.1; до 32 - в версиях аппаратуры 1.2 и 1.3), максимальному числу потоков (до 768 активных потоков в версиях аппаратуры 1.0 и 1.1; до 1024 - в версиях аппаратуры 1.2 и 1.3).

Количество потоков в блоке и количество блоков в решётке выбирается программистом исходя из максимизации загрузки ресурсов MP и с учётом аппаратных ограничений (количество регистров, разделяемой памяти и т.д.). Блоков д.б. не менее числа MP, лучше с запасом, чтобы MP не простаивал во время чтения из памяти. Количество потоков в блоке должно быть кратно размеру порции (32).

Времена выполнения команд (напоминаю, что MP выполняет одновременно 4x8 потоков):

Типы памяти (слово - 32 бита):

Ядро выполняется на GPU, в отличие от остальной части программы, выполняемой на CPU хостовой системы. GPU и CPU может быть несколько. GPU должны быть одинаковыми (требование снято?), причём не в режиме SLI (требование снято?). Один хостовый поток (процесс) может использовать только один GPU - тот, который был задействован первым. У меня были проблемы с дочерними процессами, порождаемыми с помощью fork - достаточно опросить количество GPU в головной программе, чтобы дочерние процессы имели проблемы. Хостовая система имеет свою ОП, обмен с глобальной памятью GPU и запуск ядра осуществляется через специальный интерфейс. Запуск ядра является асинхронным, т.е. управление немедленно возвращается хостовой программе. Также асинхронным может являться обмен мехду хостом и GPU (при этом хостовая память д.б. залочена) и пересылки внутри GPU. Приложение может использовать либо низкоуровневый интерфейс драйвера CUDA (поставляется вместе с драйвером видеокарты), либо высокоуровневый API приложений CUDA (runtime, использует интерфейс драйвера CUDA). Поверх API приложений реализованы специализированные библиотеки CUDAPP (сортировка и псевдослучайный генератор), FFT (CUFFT, преобразование Фурье) и BLAS (CUBLAS, линейная алгебра).

Операции над вещественными числами обладают некоторыми отклонениями от IEEE-754:

Оборудование, подходящее для работы с CUDA

Новые поколения устройств, поддерживающих модель CUDA, обладают дополнительными возможностями (compute capability), в скобках указано количество MP, приведены только настольные версии GeForce, модели отличаются также частотой GPU, объёмом, типом и частотой ОП:

Драйвер видеокарты

В состав графического драйвера NVIDIA версии 177.13/117.67 (Linux, i386, Beta) входит и CUDA 2.0 runtime драйвер (/usr/lib/libcuda.so.177.67, /usr/include/cuda/cuda.h, /usr/include/cuda/cudaGL.h). Соответственно, для CUDA 2.1 - драйвер версии ?; для CUDA 2.2 - драйвер версии 185.18.08 (185.18.14). Запуск программ CUDA возможен в обычном режиме, без прерывания работы сервера X Window, но если ядро занимает GPU на большой отрезок времени, то становятся заметны задержки в реакции на события. Если запуск ядра на GPU, разделяемом с X сервером, не завершается в течении 5 секунд добровольно, то выполнение прерывается аварийно. Запуск X может потребоваться для инициализации устройства и загрузки модуля, хотя это можно сделать и вручную

modprobe nvidia
mknod -m 666 /dev/nvidia0 c 195 0
...
mknod -m 666 /dev/nvidiactl c 195 255

Среда разработки (CUDA Toolkit)

В составе CUDA Toolkit поставляется компилятор nvcc, документация (в т.ч. ptx_isa_1.4.pdf, nvcc_2.2.pdf), библиотека API приложений CUDA (runtime, libcudart.so), библиотека низкоуровневого интерфейса драйвера CUDA (/usr/include/cuda/cuda.h и /usr/lib/libcuda.so из комплекта драйвера графики), библиотеки CUBLAS и CUFFT, CUDA Profiler 2.2.

Основной программой среды разработки является компилятор nvcc, который управляет всей цепочкой преобразований. Исходная программа может быть на языке C (.c), C++ (.cc, .cxx, .cpp), расширении языка C (.cu; .cup для препроцессированных файлов). В начале, стандартный cpp осуществляет макроподстановки, затем утилита cudafe в 2 прохода разделяет общий код программы, написанный на расширении языка C (.cu; .cup для препроцессированных файлов), на части, предназначенные к исполнению на хосте (.c, компилируется обычным gcc в режиме C++ в .o), и части, предназначенные к исполнению на GPU (.gpu, C). Простенький развёрнутый crypt-des (25 итераций по 16 раундов по 8 S-блоков по 10 переменных) обрабатывался 15 минут. Последние преобразуются nvopencc в код виртуальной машины (.ptx). nvopencc имеет ограничение по умолчанию - не более 26000 регистров каждого типа (а их препроцессор генерирует бесжалостно). После снятия ограничения на оптимизацию потребил (в фазе be) 6GB памяти (реально используются!) и 90 минут CPU, после чего всё равно упал. Код виртуальной машины компилируется оптимизирующим ассемблером ptxas (распределение виртуальных регистров в реальные, имеет свои ограничения) в двоичный блоб (.cubin), который утилита fatbin в сочетании с кодовым именем устройства либо размещает во внешний репозитарий (.fatbin), либо добавляет к коду хостовой части (.cu.c), из которого стандартый компилятор gcc (версии от 3.4 до 4.2; кстати, в Fedora gcc 4.3!) компилирует и собирает программу.

К хостовой программе в дополнение к остальным библиотекам (.a, .so) присоединяется библиотека функций (высокого уровня cudart и/или низкого уровня cuda; cudart использует cuda; приложение может использовать только одну из них), обеспечивающих управление и доступ к GPU со стороны хоста. К программе, выполняемой на GPU, присоединяется библиотека функций, специфических для устройства. Имеется также общая часть - подмножество библиотеки C, реализованное как на хосте, так и на GPU. nvcc генерирует код, предназначенный для работы с cudart, при этом инициализация, управление контекстами и модулями скрыты от программиста.

С помощью макро __CUDACC__ можно определить, компилируется ли программа с помощью nvcc.

Несколько хостовых потоков могут запускать ядра на одном GPU (ресурсы одного потока недоступны другому). Один хостовый поток не может использовать несколько GPU.

Поставляется в виде привычного для NVIDIA исполняемого bash-скрипта, содержащего самораспаковываемый архив (--help, --info, --list, --check). По умолчанию, всё ставится в каталог /usr/local/cuda, но можно задать свой путь вместо /usr/local (я задал /usr/local/cuda2.2, чтобы можно было иметь несколько версий одновременно).

Я ставил NVIDIA CUDA Toolkit 2.0 (начиная с beta2) для RHEL5.1 на FC6 (32 бит) [NVIDIA CUDA Toolkit 2.2 для F10 (64 бит)]. Для задания пути к разделяемым библиотекам CUDA необходимо указать их в /etc/ld.so.conf и запустить ldconfig (и "chcon -t texrel_shlib_t" для SELinux). Перед использованием необходимо задать PATH (/usr/local/cuda2.2/cuda/bin) и LD_LIBRARY_PATH (/usr/local/cuda2.2/cuda/lib), например, с помощью системы модулей.

Для автономной компиляции хостовой подпрограммы с помощью gcc необходимо указывать ключ "-malign-double" (nvcc делает это автоматически).

Хостовые подпрограммы компилируются в режиме C++, чтобы компилировать в режиме C требуется указать ключ "--host-compilation=c".

64-битная версия среды разработки генерирует 64-битный код как для хостовой части программы, так и для GPU. Пришлось оттрасировать все вызовы nvcc и переиграть их, заменив -m64 на -m32 в предназначенных для генерации кода для GPU командах. У меня это дало выигрыш в 9% и меньшее число регистров, что вообще бесценно.

Нетерпеливые могут сразу попробовать примеры из SDK.

Расширения языка C

Расширения языка C:

Высокоуровневая библиотека управления GPU

Высокоуровневая библиотека управления GPU (CUDA runtime API, cudart) реализована с использованием низкоуровневой библиотеки (libcuda). Используется интерфейс C++ (именование функций, возможность использовать умолчания). Инициализация происходит неявно при вызове первой функции. Это необходимо учитывать при измерении производительности и разборе ошибок, а также при многозадачно и многопоточной обработке.

Большинство функций API имеет тип cudaError_t, что позволяет обрабатывать ошибки (cudaSuccess - отсутствие ошибок). С помощью функции cudaGetErrorString (cudaError_t) можно получить текстовое сообщение об ошибке. При обработке ошибок необходимо учитывать асинхронность многих функций API. Последнее значение кода возврата можно получить с помощью функции cudaGetLastError.

Управление GPU:

Запуск ядра:

Управление памятью и пересылки (размер в байтах):

В связи с асинхронностью запуска ядра введены дополнительные средства управления задачами (stream). Задача - это последовательность операций, выполняемых строго по очереди. Относительный порядок операций, принадлежащих различным задачам, не определён. По умолчанию, операция относится к задаче 0. При необходимости организовать выполнение нескольких паралелльных задач необходимо определить задачи и указывать их номера в командах запуска ядра и копирования из памяти в память. Установка переменной окружения CUDA_LAUNCH_BLOCKING в 1 блокирует асинхронное выполнение. Функции синхронизации:

Функции работы с событиями (event) позволяют измерять время выполнения операций:

Функции работы с текстурами.

Функции взаимодействия с OpenGL.

Низкоуровневая библиотека управления GPU

Низкоуровневая библиотека управления GPU (CUDA driver API, cuda, libcda). В отличие от cudart используется интерфейс C (вместо C++) и не требуется nvcc.

Функции, реализованные и для хоста и для GPU

Встроенные векторные типы данных (поля: x, y, z, w; конструкторы: make_имя):

Функции над вещественными числами одинарной точности:

Аналогичные функции реализованы для вещественных чисел двойной точности, но при использовании на GPU версии 1.2 и ниже выполняются как функции одинарной точности.

Целочисленный функции min и max реализованы в виде одной инструкции.

clock - при выполнении на GPU возвращает номер тика.

Функции для работы с текстурами.

Функции, реализованные для GPU

В основном, реализованы более быстрые, но менее точные аналоги стандартных функций. К имени стандартной функции при этом спереди добавляется строка "__". Если к имени в качестве суффикса добавлена строка "_rn", то результат округляется к ближайшему чётному. Если к имени в качестве суффикса добавлена строка "_rz", то результат округляется к нулю. Если к имени в качестве суффикса добавлена строка "_ru", то результат округляется вверх. Если к имени в качестве суффикса добавлена строка "_rd", то результат округляется вниз. Фунции __fadd_rn и __fmul_rn гарантированно не будут слиты в одну команду FMAD.

Функции над вещественными числами одинарной точности:

Функции над целыми числами:

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

Функции для работы с текстурами.

Атомарные (неделимые) функции обеспечивают неделимость цикла обработки числа чтение-изменение-запись (возвращают старое значение):

Функции голосования: __all и __any.

nvcc

Ключи nvcc делятся на:

Ключи могут иметь краткую ("-v", "-cuda") и длинную ("--verbose", "--cuda") форму.

Каталог, содержащий nvcc, должен иметь файл nvcc.profile, задающий значения переменных окружения (PATH, LD_LIBRARY_PATH, INCLUDES, LIBRARIES, CUDAFE_FLAGS, OPENCC_FLAGS, PTXAS_FLAGS), необходимые для сборки и запуска собранных с помощью nvcc программ. Альтернативным способом задания пути к разделяемым библиотекам CUDA является указание их в /etc/ld.so.conf и выполнение ldconfig (и chcon при использовании SELinux). Данный файл создаётся при установке и не предназначен для модификации пользователем.

Процесс компиляции и сборки программы с помощью nvcc делится на описанные выше фазы, которые выбираются с помощью суффикса, определяющего формат входного файла, и опций, задающих формат выходного файла (по умолчанию сборка исполняемого файла, "--link"):

Ключи задания файлов и путей:

Ключи для отладчика и оптимизации:

Ключи для передачи параметров для отдельных фаз обработки:

Ключи управления nvcc:

Управление генерацией хостового кода:

Управление генерацией кода GPU:

nvopencc - генератор PTX

nvopencc является урезанной и адаптированной версией Open64 (ранее Pro64, ещё ранее MIPSPro от SGI) от NVIDIA (исходные тексты nvopencc; документация open64).

nvopencc состоит из

Параметры передаются из nvcc с помощью ключа "--opencc-options" (компилируется имя.cpp3.i) или переменной окружения OPENCC_FLAGS (описания не нашёл, но есть --help, выдающий огромный список опций, а какие из них осмысленно для CUDE непонятно):

Оптимизирующий ассемблер ptxas

Оптимизирующий ассемблер ptxas. Ключи:

Стандартные библиотеки CUFFT и CUBLAS

Поставляемые в комплекте со средой разработки библиотеки CUFFT (cufft.h, libcufft.so, libcufftemu.so) и CUBLAS (cublas.h, libcublas.so, libcublasemu.so) реализуют набор функций быстрого преобразования Фурье (FFT) и операции линейной алгебры (BLAS).

Эмуляция

Ключ компилятора "-deviceemu" позволяет компилировать и запускать программу на CPU, что позволяет отлаживать её с помощью обычного отладчика. Макрос __DEVICE_EMULATION__ позволяет осуществлять условную компиляцию отладочной печати и т.п.. При эмуляции на каждый предполагаемый поток GPU создаётся поток CPU, а каждый поток требует 256KB стека!

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

SDK

CUDA Developer SDK содержит набор примеров реализации программ в модели CUDA и набор скриптов и утилит для быстрого "вхождения в предмет". Устанавливается запуском cudasdk_2.21_linux.run и указанием каталога установки (/usr/local/cuda2.2/SDK). Для тестирования установки драйвера, toolkit и SDK надо зайти в каталог установки SDK и выполнить команду make, затем углубиться в bin/linux/release и попробовать выполнить собранные там программы (начать с deviceQuery и bandwidthTest - у меня запускается только под gdb (иначе Segmentation fault где-то в StopWatch::create) и не на всех GPU (на 8600 GT: cudaSafeCall - out of memory), приходится всё задавать вручную: "--device=1 --mode=range --start=1048576 --end=8388608 --increment=1048576").

С помощью SDK можно буквально за день собрать свою первую пррограмму, работающую на GPU: скопировать projects/template в projects/свой-проект, заменить в Makefile слово template на имя своего проекта, добавить свой код (или взять из примеров) в свой-проект.cu и запустить make. При отсутствии ошибок компиляции можно запускать ../../bin/linux/release/свой-проект на выполнение.

Команда "make emu=1" позволяет откомпилировать и собрать программы под эмулятор GPU (складываются в bin/linux/emurelease).

Профилировщик для оптимизации производительности

Профилирование состоит из двух этапов:

Управление сбором информации производится с помощью переменных окружения:

Конфигурационный файл определяет какие счётчики будут задействованы (не более 4 одновременно, комментарии определяются символом '#'):

Собранные данные содержат (по умолчанию, используется формат имя=значение; времена в микросекундах):

Визуализация собранных данных возможна с помощью CUDA Visual Profiler. Поставляется в формате tar.gz (требуется libstdc++.so.6, например, compat-gcc-34-3.4.6-4 и compat-gcc-34-c++-3.4.6-4). В архиве, кроме самого визуализатора cudaprof поставляются разделяемые библиотеки libQtCore.so.4 и libQtGui.so.4, которые необходимо сделать доступными (например, LD_LIBRARY_PATH=.../bin). Можно импортировать журналы в формате CSV. Если профилируемая программа ожидает ввода, то запускать с галочкой "Run in separate window" (xterm).

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

Имеется возможность импорта журнала в формате CSV ограниченного объёма (1000 строк?).

Настройки сохраняются в файле $HOME/.config/NVIDIA/cudaprof.conf.

BtR-MySQL как пример использования CUDA

Постановка задачи: поиск пароля MySQL 3.xx по известному хешу методом полного перебора.

Исходное решение: John the Ripper в режиме mysql-fast (да, я знаю, что JtR не предназначен для полного перебора) - 32Mpps (миллионов паролей в секунду) на Q6600 разогнанном до 3.3GHz.

Специально написанная программа с использованием SSSE3 позволяет достичь на этом же процессоре - 1451Mpps на одном ядре, с использованием всех 4 ядер - 5714Mpps. Условия измерения: набор символов - 94 символа ASCII, длина предполагаемого пароля - 8 символов. Кстати, использование SSSE3 дало очень немного (около 50%), т.к. 6 итераций по улучшению SIMD набора команд фирме Intel не хватило для появления полноценной команды умножения (появилась только в SSE4).

BtR-MySQL (Brute force the Ripper - MySQL 3 password cracker) реализован как модельная задача для исследования возможностей и режимов эксплуатации GPU фирмы NVIDIA с использованием CUDA Toolkit. Программа считывает со стандартного ввода хеш искомого пароля, определяет число доступных CPU и GPU, запускает для каждого устройства отдельный исполнительный процесс. После этого основной процесс раздаёт пакеты заданий, а дочерние процессы выполняют их на CPU (требуется SSE3) или GPU (требуется GeForce 8xxx и CUDA).

Файл BtR-MySQL.h содержит следующие параметры компиляции:

Файл BtR-MySQL.h содержит следующие параметры, определяющие значения по умолчанию:

Ключи запуска:

Достигнутый результат:


	GeForce 8600 GT (безвентиляторный): 4 MOP at 1404 MHz 

        Charset length: 94
        Charset interval
        minimal password length: 8
        maximal password length: 16
        tasks queue length per thread: 6
        CPU numbers: 0
        GPU numbers: 1
        CPUs threads priority: 19
        GPUs threads priority: 1
        CUDA sync pause: -1
        CPUs threads affinity: 0
        GPU threads affinity to CPU: -1
        threads per GPU: 1
        tasks per GPU thread: 1
		6868 Mpps

	GeForce 8600 GT (безвентиляторный): 4 MOP at 1404 MHz
	GeForce GTX 260: 24 MOP at 1242 MHz
        Charset length: 94
        Charset interval
        minimal password length: 8
        maximal password length: 16
        tasks queue length per thread: 6
        CPU numbers: 0
        GPU numbers: 2
        CPUs threads priority: 19
        GPUs threads priority: 1
        CUDA sync pause: -1
        CPUs threads affinity: 0
        GPU threads affinity to CPU: -1
        threads per GPU: 1
        tasks per GPU thread: 1
		42149 Mpps

	Q6600 @ 3.3 GHz + GeForce 8600 GT + GTX 260
        Charset length: 94
        Charset interval
        minimal password length: 8
        maximal password length: 16
        tasks queue length per thread: 6
        CPU numbers: 4
        GPU numbers: 2
        CPUs threads priority: 19
        GPUs threads priority: 1
        CUDA sync pause: 0
        CPUs threads affinity: 0
        GPU threads affinity to CPU: 3
        threads per GPU: 1
		46145 Mpps

Подводные камни:

Итого: 46145 Mpps по сравнению с исходными 32 Mpps.

Изменения

Ссылки

@ Карта сайта News Автора!

Bog BOS: Использование модели массового параллелизма CUDA для разработки программ

Последние изменения:
2015.11.18: hard: обновлена статья про ИБП и их мониторинг

TopList

Copyright © 1996-2017 Sergey E. Bogomolov; www.bog.pp.ru (КГБ знает все, даже то что у Вас на диске ;)