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

Последние изменения:
2024.11.11: sysadmin: Linux: пространства имён
2024.11.06: sysadmin: настройка TCP/IP в Linux: виртуальный интерфейс и виртуальный мост
2024.10.25: sysadmin: Linux VFS, атрибуты, расширенные атрибуты, ACL
2024.10.22: sysadmin: Монтирование файловых систем: bind, shared и OverlayFS

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

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

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

понимание модели CUDA
оборудование
драйвер видеокарты
CUDA драйвер API (libcuda, включён в состав драйвера видеокарты)
среду разработки (CUDA Toolkit)
стандартные библиотеки (поставляются CUDPP, FFT и BLAS)
SDK (набор скриптов и примеров)
профилировщик для оптимизации производительности
отладчик

эмулятор
пример использования

Модель 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 потоков):

работа с памятью - 4 цикла (задержки доступа к памяти учитываются отдельно в зависимости от типа памяти)
сложение целых (32 бита) - 4 цикла
умножение целых чисел - 16 циклов
умножение 24-битных целых (__[u]mul24) - 4 цикла; есть угроза, что в будущих версиях аппаратуры целые числа будут умножаться быстрее 24-битных
деление и получение остатка целых чисел - очень медленно (видимо, реализовано подпрограммой)
битовые операции - 4 цикла
сравнение - 4 цикла
min и max - 4 цикла
преобразование типов - 4 цикла (операнды типа char или short необходио преобразовать в int перед выполнением)
сложение, умножение и умножение-и-сложение (FMAD) вещественных чисел одинарной точности - 4 цикла
деление вещественных чисел одинарной точности - 36 циклов
быстрое деление вещественных чисел одинарной точности (__fdividef) - 20 циклов
обратная величина, обратная величина к квадратному корню, быстрая версия логарифма - 16 циклов
квадратный корень - 32 цикла
быстрые версии синуса, косинуса и экспоненты - 32 цикла
остальные арифметические команды реализованы как подпрограммы
команда синхронизации - 4 цикла (если некого ждать)

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

локальная память потока (reg, register, регистры); расположены в MP; распределяются по потокам из общего пула MP: 8192 регистра в версиях аппаратуры 1.0 и 1.1; 16384 регистра в версиях аппаратуры 1.2 и 1.3; не кешируется (нет смысла); время доступа равно нулю, если нет задержки "read-after-write" (эту проблему можно игнорировать, если есть хотя бы 192 активных потока на MP) или конфликта банков регистров (оптимизатор рассчитывает, что число потоков в блоке кратно 64, но гарантировать ничего не может); не более 128 регистров на поток (ptxas не даёт более 124 регистров, остальные переводятся в "локальную" память)
разделяемая память блока потоков (smem); находится в MP; освобождается по завершению блока; 16 банков по 1KB на каждый MP; доступ к банкам осуществляется параллельно; время доступа равно нулю, если нет конфликта банков между потоками; соседние 32-битные слова располагаются в соседних банках; пропускная способность банка - 32-бита за 2 цикла; учитывая, что обработка одной команды порции потоков занимает 4 цикла и один запрос к разделяемой памяти выдаётся первой полупорции, а второй запрос - второй, можно достичь полностью бесконфликтного обращения к банкам разделяемой памяти; имеются широковещательные запросы на чтения - если несколько потоков одной полупорции запрашивают данные из одного 32-битного слова, то все запросы будут обслужены одновременно
глобальная память (находится в видеокарте, освобождается по завершению приложения); не кешируется; часть памяти используется на другие нужды (видеобуфер и пр.); задержка доступа от 400 до 600 циклов; загрузка в регистры может происходить словами по 32, 64 или 128 бит (для этого тип переменной д.б. соответстующего размера или больше, а переменная в глобальной памяти должна д.б. выровнена на границу слова, встроенные типы выравниваются автоматически, другие можно выровнять спецификатором __align__); запросы к памяти потоков одной полупорции могут быть объединены в одну транзакцию размером 32 байта (только в GPU версии аппаратуры 1.2 и выше), 64 байта или 128 байт (требуется соответствующее выравнивание сегмента в глобальной памяти; потоки должны обращаться к 32/64/128 битным словам; все 16 слов должны лежать в одном сегменте; потоки полупорции должны обращаться к словам в сегменте последовательно, т.е. n-ый поток полупорции должен обращаться к n-ому слову сегмента; для 32-битных слов выигрыш достигает одного порядка, для 64-битных - 4 раза; для 128-битных - 2 раза); для GPU версии аппаратуры 1.2 и выше могут сливаться обращения к словам размера 8 и 16 бит, требование на последовательность обращений снято; для работы с двумерными массивами рекомендуется использовать функцию cudaMallocPitch, которая выравнивает строки массива правильным образом;
локальная память потока (не регистр, lmem, .local, ld.local); не кешируется; задержка доступа от 400 до 600 циклов; запросы к памяти потоков одной полупорции всегда объединены в одну транзакцию; судя по задержкам доступа и отсутствию ограничений на размер локальной памяти в MP, локальная память эмулируется с помощью глобальной
константы (cmem): часть глобальной памяти; из GPU доступна только на чтение; MP кеширует запросы; размер кеша - 8KB на MP; задержка доступа при отсутствии в кеше - от 400 до 600 циклов; задержка доступа при наличии в кеше - как при работе с регистрами, если все потоки полупорции читают одни и те же данные; адресное пространство - 64KB
текстуры: часть глобальной памяти; из GPU доступна только на чтение; MP кеширует запросы; размер кеша на MP - 6KB или 8KB; оптимизирован для работы с двумерными массивами; задержка доступа при отсутствии в кеше - от 400 до 600 циклов; задержка доступа при наличии в кеше - непонятна; доступ к текстурам осуществляется через текстурные процессоры, которые могут реализовать нестандартные схемы адресации (нормализация координат, клипирование координат, свёртка координат) или простую обработку (преобразование типа, интерполяция)

Ядро выполняется на 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:

исключения всегда замаскированы
оптимизатор комбинирует умножение и последующее сложение в одну инструкция (FMAD), в результате промежуточный результат округляется
деление реализовано через получение обратной величины и умножение
только статическое задание методов округления: до ближайшего чётного и к нулю
не поддерживаются ненормализованные числа

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

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

версия 1.0: 8800 Ultra (16), 8800 GTX (16), 8800 GTS (12)
версия 1.1 (добавлена поддержка неделимых (атомарных) операций над 32-битными словами глобальной памяти): GTS 250 (16), GTS 150 (16), 9800 GTX (16), 8800 GT (14), GT 130 (12), 9600 GSO (12), 8800 GS (12), 8800M GTX (12), 9600 GT (8), 8800M GTS (8), GT 120 (4), 9500 GT (4), 8600 GTS (4), 8600 GT (4), 8700M GT (4), 8600M GT (4), 8600M GS (4), G100 (2), 8500 GT (2), 8400 GS (2), 8400M GT (2), 8400M GS (2), 8400M G (1)
версия 1.2 (добавлена поддержка неделимых операций над 64-битным словами глобальной памяти и 32-битным словами разделяемой; функция голосования (warp vote)): не осталось
версия 1.3 (добавлена поддержка вещественных чисел с двойной точностью, на целый MP - один блок двойной точности): GTX 295 (2x30), GTX 285 (30), GTX 280 (30), GTX 260 (24)

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

В состав графического драйвера 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:

квалификатор типа функции, указывающий где должна выполняться функция и откуда её можно вызывать
- __device__ (функция выполняется на GPU, вызывается из GPU; рекурсия запрещена; статические переменные запрещены; фиксированное число параметров; указатель на функцию недопустим)
- __global__ (это ядро и должно вызываться специальной директивой запуска ядра; вызов является асинхронным; параметры передаются через разделяемую память (до 256 байт); функция выполняется на GPU, вызывается из хоста; рекурсия запрещена; статические переменные запрещены; фиксированное число параметров; тип функции только void)
- __host__ (функция выполняется на хосте, вызывается из хоста, по умолчанию)
- __device__ __host__ (код функции компилируется и для хоста, и для GPU)
квалификатор типа переменной, указывающий в памяти какого типа её размещать (можно использовать только в функциях, исполняемых на GPU):
- __device__ (переменная размещается в глобальной памяти GPU; существует пока существует приложение; доступна из любого потока на GPU и из хоста через библиотеку)
- __constant__ (переменная размещается в памяти констант GPU; существует пока существует приложение; доступна из любого потока на GPU и из хоста через библиотеку)
- __shared__ (переменная размещается в разделяемой памяти блока, существует пока существует блок; доступна из любого потока блока; запись становится гарантированно видима остальным потокам блока только после вызова __syncthreads() )
- внутренние переменные исполняемой на GPU функции без квалификатора обычно хранятся в регистре, но если оптимизатор из ptxas посчитает, что их слишком много (более 124), то в локальной нерегистровой памяти MP (по времени доступа неотличима от глобальной); узнать, что он туда перенёс, невозможно (в .ptx всё в регистрах, а в блоб не заглянешь); массивы с достпом по неконстантному индексу тоже попадают в локальную память
директива запуска ядра: вызов из хоста ядра (функции с квалификатором __global__) оформляется специальным образом - между именем функции и списком параметров в тройных угловых скобках ( "<<<...>>>" ) записываются через запятую размерность и размер решётки (тип dim3), размерность и размер блока (тип dim3), количество дополнительной динамически выделяемой разделяемой памяти на блок (опционально), номер задачи (stream, опционально); конфигурация ядра вычисляется до вычисления параметров и передаётся через разделяемую память
встроенные переменные, доступные из функции, работающих на GPU (только чтение, указывать на встроенные переменные недопустимо):
- gridDim (тип dim3) - размерность и размер решётки
- blockIdx (тип uit3) - индекс блока в решётке
- blockDim (тип dim3) - размерность и размер блока
- threadIdx (тип uit3) - индекс потока в блоке
- warpSize (тип int) - размер порции
__noinline__ (по умолчанию, все функции, работающие на устройстве, компилируются, как встроенные (inline); nvcc может проигнорировать это указание, если список параметров слишком велик или при передаче ссылки)
компилятор по умолчанию развёртывает маленькие циклы с известным числом шагов; непосредственно перед циклом можно явно указать, что делать с этим циклом директивой "#pragma unroll [число-шагов]"

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

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

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

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

cudaGetDeviceCount (int *) - узнать число устройств; при отсутствии GPU возвращает число 1 (эмулятор GPU)
cudaSetDevice (int) - выбрать устройство для использования; необходимо сделать до первого запуска или обращения к устройству; нумерация с 0
cudaGetDevice (int *) - узнать номер используемого GPU
cudaGetDeviceProperties (cudaDeviceProp*, dev) - узнать характеристики указанного устройства: имя, объём глобальной памяти, объём разделяемой паяти на MP, число регистров, размер порции, максимальное число потоков на блок, максимальная размерность блока, максимальная размерность сетки, объём памяти констант, версия GPU, частота GPU (непонятно какого функционального блока), количество MP, возможность асинхронной пересылки
cudaChooseDevice (int*, const struct cudaDeviceProp*) - выбор наиболее подходящего устройства

Запуск ядра:

cudaConfigureCall (dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, int tokens = 0) - задать размерность сетки и блока процессов для будущего запуска; помещаются в стек для дальнейшего использования в cudaLaunch
cudaSetupArgument (void* arg, size_t count, size_t offset) - задать аргументы запуска
cudaLaunch (T entry) - запустить ядро entry, которое задаётся либо функцией __global__, либо текстовой строкой с именем такой функции; размерности и параметры извлекаются из стека, создаваемого cudaConfigureCall

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

cudaMalloc (void**, size_t) - выделить линейный массив из глобальной памяти
cudaMallocPitch (void**, size_t* pitch, size_t widthInBytes, size_t height) - выделить двумерный массив из глобальной памяти (заботится о выравниваниях и прочей оптимизации)
cudaFree (void *)
cudaMallocHost (void**, size_t) - выделить из ОП хоста в режиме блокировки от подкачки; это сильно ускоряет обмен с GPU; асинхронные операции копирования допустимы только с блокированной памятью
cudaFreeHost (void *)
cudaMemset (void*, int, size_t) - заполнить память GPU константой
cudaMemcpy (void* dst, const void* src, size_t count, enum cudaMemcpyKind kind) - копирование, где cudaMemcpyKind определяет направление пересылки (перекрытие областей запрещено):
- cudaMemcpyHostToHost
- cudaMemcpyHostToDevice
- cudaMemcpyDeviceToHost
- cudaMemcpyDeviceToDevice
cudaMemcpyAsync (void* dst, const void* src, size_t count, enum cudaMemcpyKind kind, cudaStream_t) - асинхронное копирование, только для блокированной памяти
cudaMemcpyToSymbol (const T& symbol, const void* src, size_t count, size_t offset, enum cudaMemcpyKind kind) - копирование count байт с адреса src на адрес, определяемый переменной symbol, со смещением offset; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной; cudaMemcpyKind определяет направление пересылки (перекрытие областей запрещено):
- cudaMemcpyHostToDevice
- cudaMemcpyDeviceToDevice
cudaMemcpyToSymbol (void* dst, const T& symbol, size_t count, size_t offset, enum cudaMemcpyKind kind) - копирование count байт на адрес dst с адреса, определяемого переменной symbol, со смещением offset; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной; cudaMemcpyKind определяет направление пересылки (перекрытие областей запрещено):
- cudaMemcpyDeviceToHost
- cudaMemcpyDeviceToDevice
cudaGetSymbolAddress (void** devPtr, const T& symbol) - получить указатель на переменную в глобальной памяти GPU; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной
cudaGetSymbolSize (size_t* size, const T& symbol) - получить размер переменной в глобальной памяти GPU; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной
cudaMallocArray - выделение памяти под CUDA-массив
cudaFreeArray
cudaMemset2D
cudaMemcpy2D, cudaMemcpy2DAsync, cudaMemcpyToArray, cudaMemcpyToArrayAsync, cudaMemcpy2DToArray, cudaMemcpy2DToArrayAsync, cudaMemcpyFromArray, cudaMemcpy2DFromArray, cudaMemcpyArrayToArray, cudaMemcpy2DArrayToArray - копирование матриц
cudaMalloc3D, cudaMalloc3DArray
cudaMemset3D
cudaMemcpy3D

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

cudaStreamCreate (cudaStream_t*)
cudaStreamQuery (cudaStream_t) - проверить, завершились ли все операции задачи: cudaErrorNotReady (единственный способ синхронизации без 100% загрузки CPU, однако уменьшает эффективность работы)
cudaStreamSynchronize (cudaStream_t) - ждать завершения всех операций задачи (реализовано в виде непрерывного опроса устройства, что полностью загружает CPU)
cudaStreamDestroy (cudaStream_t) - ждать завершения всех операций задачи и завершить задачу
cudaThreadSynchronize (void) - ждать завершения всех операций всех задач (реализовано в виде непрерывного опроса устройства, что полностью загружает CPU)
cudaThreadExit (void) - ждать завершения всех операций всех задач, освободить все ресурсы; последующее обращение к одной из функций API повторно инициализирует API (глючит, не советую)

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

cudaEventCreate (cudaEvent_t*) - создать описание события
cudaEventRecord (cudaEvent_t, CUstream) - зафиксировать момент наступления события после выполнения всех операций указанной задачи; сама функция является асинхронной, так что для извлечения точного значения необходимо использовать cudaEventQuery или cudaEventSynchronize; извлечь информацию для не нулевой задачи нечем
cudaEventQuery (cudaEvent_t) - завершена ли запись наступления события (cudaErrorNotReady)
cudaEventSynchronize (cudaEvent_t) - ждать окончания записи наступления события
cudaEventElapsedTime (float*, cudaEvent_t, cudaEvent_t) - вычислить интервал времени между событиями (в мс), точность - 0.5 мкс; не работает для событий из не основной (не нулевой) задачи
cudaEventDestroy (cudaEvent_t) - удалить описание события

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

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

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

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

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

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

[u]{char|short|int|long}{1234}
float{1234}
double2
dim3 (uint3; неуказанные компоненты автоматически получают значения 1)

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

x+y, x*y, fmaf, x/y, 1/x, ldexpf, scalbnf, scalblnf, fmodf, remainderf, remquof, fdimf
fminf, fmaxf, fabsf
корни: rsqrtf, sqrtf, cbrtf
степени: expf, exp2f, exp10f, expm1f, logf, log2f, log10f, logp1f, powf
тригонометрические функции
erff, erfcf, lgammaf, tgammaf
округления: roundf - на GPU занимает 8 инструкций; rintf, truncf, ceilf, floorf - по одной
округления: nearbyintf, lrintf, lroundf, llrintf, llroundf, nextafterf
разборки: frexpf, logbf, ilogbf, modff, signbit, isinf, isnan, isfinite, copysignf, nanf

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

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

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

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

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

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

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

__fadd_[rn|rz], __fmul_[rn|rz], __fdividef
__expf, __exp10f, __logf, __log10f, __powf
__sinf, __cosf, __sincosf, __tanf
__saturate

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

__[u]mul24 (умножение 24 бит), __[u]mulhi, __[u]mul64hi
__[u]sad (сложение с модулем разности)
__clz, __clzll (число старших нулей, вариант с дополнительными "ll" - для 64 бит)
__ffs, __ffsll (номер позиции младшей единицы)
__popc, _popcll (число единиц)
__brev, __brevll (... и последние биты станут первыми)

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

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

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

atomicAdd, atomicSub, atomicMin, atomicMax, atomicAnd, atomicOr, atomicXor
atomicInc, atomicDec (с клипированием)
atomicExch (обмен значений), atomicCAS (сравнение и обмен)

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

nvcc

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

булевы (не имеют аргументов)
ключи с одним аргументом (через пробел или знак равенства, за исключением ключей -I, -L для совместимости с gcc)
ключи со списком аргументов (через запятую или используя несколько вхождений ключа)

Ключи могут иметь краткую ("-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"):

--generate-dependencies или -M (генерация зависимостей для make для одного файла типа .c/.cu/.cpp)
--preprocess или -E (препроцессирование)
--cuda (преобразование .cu в .cu.c)
--gpu (генерация .gpu из .cu)
--ptx (генерация .ptx из .cu или .gpu)
--cubin (генерация .cubin из .cu или .gpu или .ptx)
--compile или -c (компиляция .c/.cu/.cpp в .o)
--lib (сборка библиотеки)
--run (сборка и запуск исполняемого файла; автоматически устанавливается путь к динамическим библиотекам CUDA, а также переменные окружения из nvcc.profile)

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

{--output-file | -o} имя-файла (куда выводить результат)
{--pre-include | -include} имя-файла[,...] (этот файл будет вставлен перед обрабатываемым файлом во время препроцессирования)
{--library | -l} имя-библиотеки[,...] (используемые при сборке библиотеки)
{--define-macro | -D} определение-макро[,...]
{--undefine-macro | -U} определение-макро[,...]
{--include-path | -I} имя-каталога[,...] (список поиска для #include)
{--system-include | -isystem} (список поиска для системных #include)
{--library-path | -L} имя-каталога[,...] (список поиска для библиотек)
{--output-directory | -odir} имя-каталога
{-compiler-bindir | -ccbin} имя-каталога (имя каталога, содержащего gcc и g++)

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

--profile | -pg (генерация кода для gprof)
{--debug | -g} [уровень] (код для отладчика, в качестве уровня можно задать "gdb", "dwarf")
{--device-debug | -G} уровень (отладочная информация для кода GPU, также задаёт уровень оптимизации: от 0 до 3)
{--optimize | -O} уровень (уровень оптимизации)
--shared (создавать разделяемую библиотеку)
{--machine | -m} {32 | 64}

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

{--compiler-options | -Xcompiler} ключ[,...]
{--linker-options | -Xlinker} ключ[,...]
{--opencc-options | -Xopencc} ключ[,...] (ключи для nvopencc)
{--cudafe-options | -Xcudafe} ключ[,...]
{--ptxas-options | -Xptxas} ключ[,...]
{--fatbin-options | -Xfatbin} ключ[,...]

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

--dont-use-profile | -noprof (?)
--dryrun
--verbose | -v
--keep | -save-temps (сохранить промежуточные файлы в текущем каталоге)
--clean-targets | -clean (очистить промежуточные файлы, необходимо указать те же самые ключи, что и при использовании -keep)
--run-args аргумент[,...] (аргументы для запуска собранной программы)
--no-align-double (не передавать компилятору "-malign-double" для 32-битной платформы; плохая мысль, будет ошибка сегментации при выполнении; ключ "-malign-double" необходимо указывать также при автономной сборке хостовой подпрограммы с помощью gcc)

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

--device-emulation | -deviceemu (генерировать код для эмулятора GPGPU)
--comic (генерировать код CUDA для многоядерных процессоров, ?)
--use_fast_math | -use_fast_math (использовать библиотеку быстрых, но не точных вычислений)
--host-compilation {C | C++} (режим компиляции хостовых программ, по умолчанию - C++)

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

{--gpu-name | -arch} {sm_10 | sm_11 | sm_13 | sm_14 | sm_20} (версия архитектуры при генерации .ptx; можно также указать версию архитектуры для эмулятора: compute_10, compute_11, compute_13, compute_14, compute_20)
{--gpu-code | -code} {sm_10 | sm_11 | sm_13 | sm_14 | sm_20}[,...] (версия архитектуры при генерации кода из .ptx; можно также указать версию архитектуры для эмулятора: compute_10, compute_11, compute_13, compute_14, compute_20; должен быть совместим с архитектурой, указанной в "--gpu-name"; при указании нескольких архитектур генерируется код для каждой из них)
{--export-dir | -dir} имя-каталога (репозитарий образов для загрузки в GPU)
{--extern-mode | -ext} {all | none | real | virtual} (какие образы записывать в репозитарий)
{--intern-mode | -int} {all | none | real | virtual} (какие образы встраивать в объектный код, .fatbin)
--maxrregcount число (ограничение на число регистров, расходуемых на поток; до 128)

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

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

nvopencc состоит из

gfec (на базе gcc, создаёт WHIRL IR)
inline (все функции являются встроенными)
be (оптимизация (VHO (Very High Optimizer) и WOPT (Whirl OPTimizer, SSA)) и преобразование WHIRL в PTX (CG (Code Generator)))

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

исходный-текст.cpp3.i
-o результат.ptx
-v
-O3 (автоматически добавляется nvcc)
-OPT:Olimit=число (для оптимизации больших программ, 0 - без предела)
-TARG:виртуальная_архитектура
-m64 (-m32 задать не получается, т.к. nvcc вписывает -m64 за ним, но ведь никто не помешает нам оттрасировать все вызовы nvcc и переиграть их, заменив -m64 на -m32 в строках для GPU? у меня это дало выигрыш в 9% и меньшее число регистров, что вообще бесценно)
-show (показать, какие команды выполняются)
-show-defaults (показать параметры по умолчанию)
-std=c99
-std=gnu99

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

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

--verbose (выдаёт для каждой функции выделенное ей число регистров, локальной памяти, разделяемой памяти, констант)
--gpu-name версия-GPU (sm_10; sm_10, sm_11, sm_12, sm_13)
--opt-level уровень-оптимизации (4)
--maxrregcount число (максимальное число регистров, выделяемых на функцию; округляется вверх до 4)
--output-file имя-файла (cubin.bin)
--dont-merge-basicblocks (слияние соседних базовых блоков, производимое по умолчанию в целях оптимизации, мешает отладке)
--debug-info имя-файла (DWARF)
--force-externals (для отладки)
--return-at-end (оставлять в конце функции команду выхода, чтобы было куда поставить break point)
--trace-memory (трассировка памяти, используемой ptxas)
--translation-map имя-файла
--entry имя-функции

Стандартные библиотеки 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).

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

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

сбор информации о выполнении ядра и передаче данных (данные собираются только для одного MP в GPU; счётчики соотносятся с порциями потоков, а не с отдельными потоками)
обработка пролученной информации

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

CUDA_PROFILE=1 (включает сбор информации при выполнении программы)
CUDA_PROFILE_LOG=имя-файла (по умолчанию - cuda_profile.log, времена gputime и cputime в микросекундах)
CUDA_PROFILE_CSV=1 (журнал будет собираться в формате CSV)
CUDA_PROFILE_CONFIG=имя-файла

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

timestamp (записывать в журнал временные отметки)
gld_incoherent (загрузка из глобальной памяти не оптимальна; всегда 0 на GTX 26x)
gld_coherent (загрузка из глобальной памяти оптимальна)
gst_incoherent (запись в глобальную память не оптимальна; всегда 0 на GTX 26x)
gst_coherent (запись в глобальную память оптимальна)
local_load (количество локальных загрузок)
local_store (количество локальных записей)
branch (количество переходов)
divergent_branch (количество расхождений при переходах)
instructions (количество инструкций)
warp_serialize (количество порций, которые пришлось выполнять последовательно из-за конфликта адресов разделяемой или константной памяти)
cta_launched (количество выполненных блоков)

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

timestamp= (временная отметка события, которая поможет в дальнейшем строить графики)
method= (имя программы ядра или memcopy)
gputime= (календарное время, проведённое программой в GPU)
cputime= (календарное время, проведённое программой в хостовой системе, включая ожидания GPU)
occupancy= (доля реально выполнявшихся одновременно порций от максимально возможного числа порций)

Визуализация собранных данных возможна с помощью 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 содержит следующие параметры компиляции:

CharsetLen - задаёт размер алфавита пароля
CharsetInterval - алфавит задаётся с помощью интервала, а не списка допустимых символов, размер интервала м.б. равен 26 (строчные латинские буквы), 75 (буквы и цифры) или 94 (полный набор ASCII); необходимо учитывать, что GPU работает с интервальным алфавитом значительно быстрее
FMUL - определяет быструю операцию умножения 24-битных чисел для GPU; есть шанс, что для будущих моделей GPU обычное умножение станет быстрее
THREADS_PER_GPU_MAX - максимальное число исполнительных процессов на каждый GPU; можно 1 или 2
TASKS_PER_THREAD_MAX - максимальное число задач на исполнительный процесс GPU; можно 1 или 2

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

GPU_SETSIZE - максимальное число GPU в системе (определить их автоматически не получается)
PasswordLen_Min - перебор начинается с паролей указанной длины; не может быть менее 5, если нужно менее, тто пользуйтесь JtR ;)
PasswordLen_Max - перебор заканчивается паролями указанной длины; не может быть более 16 (всё равно терпения не хватит ;)
CPU_NICE - приоритет исполнительных процессов CPU
GPU_NICE - приоритет исполнительных процессов GPU
SYNC_PAUSE - пауза синхронизации GPU (см. ниже)
GPU_CPU - к какому CPU привязываются исполнительные процессы GPU
THREADS_PER_GPU - число исполнительных процессов на каждый GPU
TASKS_PER_THREAD - число задач на исполнительный процесс GPU
QUEUE_LENG - глубина очереди заданий к каждому исполнительному процессу

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

--help
--verbose
--min минимальная-длина-пароля (от 5 до 16)
--max максимальная-длина-пароля (от 5 до 16)
--queue глубина-очереди-заданий
--cpus количество-доступных-CPU (0 - не использовать CPU; по умолчанию используются все CPU, к которым привязан основной процесс (см. taskset.1)
--gpus количество-доступных-GPU (0 - не использовать GPU)
--cpunice приоритет-процессов-CPU (0 - не изменять)
--gpunice приоритет-процессов-GPU (0 - не изменять)
--syncpause пауза-синхронизации-CUDA (микросекунд между cudaStreamQuery, 0 - использовать sched_yield вместо usleep, -1 - использовать cudaThreadSynchronize вместо cudaStreamQuery (бесконечный цикл запросов, обеспечивается максимальная утилизация GPU за счёт напрасного расходования времени CPU)
--nocpua (не привязывать каждый исполнительный процесс CPU к своему процессору)
--gpucpu номер-CPU (к какому CPU привязываются исполнительные процессы GPU; -1 - не привязывать)
--gputhreads число-исполнительных-процессов-на-GPU (не более THREADS_PER_GPU_MAX)
--tasks число-задач-на-исполнительный-процесс-GPU (не более TASKS_PER_THREAD_MAX; позволяет брать из очереди по 2 задачи за раз; полезно для алфавитов маленькой длины

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


	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

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

nvcc по умолчанию компилирует в режиме C++, что несовместимо с моим кодом SSE3
CUDA SDK собирает программы в режиме PIC, что несовместимо с моим кодом SSE3
оптимизатор opencc явно сделан из оптимизатора для обычного процессора, что приводит к логическим ошибкам в коде
в GPU нет команды деления
в SSE3 нет полноценной команды умножения (и вообще, система команд MMX/SSE безумна)

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

Изменения

версия 2.2.1
- только SDK
версия 2.2
- требуется драйвер 185.18.08 (185.18.14)
- добавлена поддержка RHEL 5.3, Fedora 10; убрана поддержка Fedora 8
- поддержка zero-copy, pinned memory (cuMemHostAlloc(), cuMemHostGetDevicePointer(), cudaHostAlloc(), cudaHostGetDevicePointer()) на GT200 (SM (?) предзагружает данные из системной памяти через PCIe) и MCP79 (память доступна GPU из чипсета напрямую!): доступна любому CUDA контексту, отображена на адресное пространство CUDA, write combined (не кешируется, более быстрый доступ со стороны GPU); адресное пространство 32-битно и привязано к процессу
- синхронизация с блокировкой - поток может заснуть в ожидании завершения работы GPU не устраивая бесконечного цикла ожидания
- cudaEventCreateWithFlags() - создание события, использующего блокирующую синхронизацию
- привязка (cudaBindTexture2D(), cuTexRefSetAddress2D()) линейной памяти (cuMemAlloc() и cudaMalloc()) к 2D-текстуре
- nvidia-smi (не заметил) и cudaSetValidDevices() - позволяет задать список допустимых устройств при создании контекста; некоторые устройства можно захватить для эксклюзивного использования для вычислений, а на некоторых запретить создавать CUDA контекст
- cuFuncGetAttribute() - запрос свойств функций
- cudaSetDeviceFlags() - позволяет задать атрибуты отображения памяти хоста и блокирующей синхронизации
- cuDriverGetVersion(), cudaDriverGetVersion(), cudaRuntimeGetVersion()
- можно запросить атрибуты устройства - CU_DEVICE_ATTRIBUTE_INTEGRATED (?), CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, CU_DEVICE_ATTRIBUTE_COMPUTE_MODE (?)
- API драйвера документировано (Doxygen ;)
- cuda-gdb для RHEL5 32 и 64 бит (только на "свободном" от X GPU)
- Visual Profiler 1.2 включён в Toolkit, поддерживает дополнительные счётчики на GT200
версия 2.1
- требуется драйвер 180.22
- добавлена поддержка RHEL 5.2, Fedora 9
- PTX JIT (Just-in-time) API
- cuda-gdb для RHEL5 32 бит
- поддержка шаблонов C++ в ядре
- nvcc -Xopencc -OPT:unroll_size=200000
- CUDA Visual Profiler 1.1

Ссылки

документация по CUDA
CUDA GPU Computing - CUDA on Linux (форум nvidia)
Open64 Sources
Программирование на языке CUDA (форум IXBT)
GPGPU.RU
CUDA, Supercomputing for the Masses
decuda
An Introduction to Parallel Computing: Design and Analysis of Algorithms. Grama, Karypis et al. Second Edition. 2003
Parallel Programming in C with MPI and OpenMP. M.J.Quinn
Parallel Programming in OpenMP. Rohit Chandra and others. Morgan Kaufmann
Scalable Parallel Programming with CUDA (John Nickolls, Ian Buck, Michael Garland, Kevin Skadron. April 28, 2008)
High performance GPU radix sorting in CUDA (удвоение скорости при переходе к Fermi, падение скорости при использовании "программного" ECC)

Карта сайта

News

Автора!

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