Параллельное выполнение потоков
 |
Параллельное выполнение потоков ( PTX или NVPTX) [1] ) — это низкоуровневая с параллельным потоков выполнением виртуальная машина и архитектура набора команд, используемая в среде программирования Nvidia Compute Unified Device Architecture ( CUDA ). Компилятор Nvidia CUDA (NVCC) преобразует код, написанный на CUDA, языке, похожем на C++ , в инструкции PTX ( язык ассемблера, представленный в виде текста американского стандартного кода для обмена информацией ( ASCII )), а графический драйвер содержит компилятор , который преобразует PTX. инструкции в исполняемый двоичный код, [2] который может работать на вычислительных ядрах (GPU) Nvidia графических процессоров . Коллекция компиляторов GNU также имеет базовые возможности генерации PTX в контексте разгрузки OpenMP . [3] Встроенную сборку PTX можно использовать в CUDA. [4]
Регистры
[ редактировать ]PTX использует произвольно большой набор регистров процессора ; Вывод компилятора представляет собой почти чистую статическую форму с одним присвоением , где последовательные строки обычно относятся к последовательным регистрам. Программы начинаются с объявлений вида
.reg .u32 %r<335>; // declare 335 registers %r0, %r1, ..., %r334 of type unsigned 32-bit integer
Это язык ассемблера с тремя аргументами, и почти во всех инструкциях явно указан тип данных (по знаку и ширине), с которыми они работают. Имена регистров начинаются с символа %, а константы являются буквальными, например:
shr.u64 %rd14, %rd12, 32; // shift right an unsigned 64-bit integer from %rd12 by 32 positions, result in %rd14
cvt.u64.u32 %rd142, %r112; // convert an unsigned 32-bit integer to 64-bit
Существуют регистры предикатов, но скомпилированный код в модели шейдера 1.0 использует их только в сочетании с командами ветвления; условная ветвь
@%p14 bra $label; // branch to $label
The setp.cc.type Инструкция устанавливает предикатный регистр в результат сравнения двух регистров соответствующего типа, также существует set инструкция, где set.le.u32.u64 %r101, %rd12, %rd28 устанавливает 32-битный регистр %r101 к 0xffffffff если 64-битный регистр %rd12 меньше или равно 64-битному регистру %rd28. В противном случае %r101 установлено на 0x00000000.
Существует несколько предопределенных идентификаторов, обозначающих псевдорегистры. Среди прочего, %tid, %ntid, %ctaid, и %nctaid содержат соответственно индексы потоков, размеры блоков, индексы блоков и размеры сетки. [5]
Государственные пространства
[ редактировать ]Нагрузка ( ld) и хранить ( st) команды относятся к одному из нескольких отдельных пространств состояний (банков памяти), например ld.param.
Есть восемь государственных пространств: [5]
.reg- регистрирует
.sreg- специальные регистры, доступные только для чтения и зависящие от платформы
.const- общая память, доступная только для чтения
.global- глобальная память, общая для всех потоков
.local- локальная память, частная для каждого потока
.param- параметры, передаваемые в ядро
.shared- память, разделяемая между потоками в блоке
.tex- глобальная память текстур (устарела)
Общая память объявляется в файле PTX с помощью строк в начале формы:
.shared .align 8 .b8 pbatch_cache[15744]; // define 15,744 bytes, aligned to an 8-byte boundary
Написание ядер в PTX требует явной регистрации модулей PTX через API драйвера CUDA, что обычно более громоздко, чем использование API среды выполнения CUDA и компилятора CUDA от Nvidia, nvcc. Проект GPU Ocelot предоставил API для регистрации модулей PTX наряду с вызовами ядра API среды выполнения CUDA, хотя GPU Ocelot больше не поддерживается активно. [6]
См. также
[ редактировать ]- Стандартное переносимое промежуточное представление (SPIR)
- Бинарный файл CUDA (кубин) – разновидность жирного бинарного файла.
Ссылки
[ редактировать ]- ^ «Руководство пользователя для серверной части NVPTX — документация LLVM 7» . llvm.org .
- ^ «Двоичные утилиты CUDA» . docs.nvidia.com . Проверено 19 октября 2019 г.
- ^ "нвпткс" . GCC Wiki .
- ^ «Встроенная сборка PTX в CUDA» . docs.nvidia.com . Проверено 3 ноября 2019 г.
- ^ Jump up to: а б «PTX ISA Версия 2.3» (PDF) .
- ^ «GPUOCelot: среда динамической компиляции для PTX» . github.com . 7 ноября 2022 г.