Parallel Thread Execution
 | В этой статье использованы только первичные либо аффилированные источники. |
Parallel Thread Execution (PTX или NVPTX[1]) — это низкоуровневая виртуальная машина и архитектура набора инструкций для параллельного выполнения потоков, используемая в программной среде CUDA от Nvidia. Компилятор NVCC переводит код, написанный на CUDA, языке, похожем на C++, в инструкции PTX (язык ассемблера, представленный в виде ASCII текста, а графический драйвер содержит компилятор, который переводит инструкции PTX в исполняемый двоичный код[2], который может выполняться на процессорных ядрах GPU от Nvidia. В коллекции компиляторов GNU также есть базовая возможность генерации PTX в контексте оффлоадинга[3] OpenMP. Встроенная сборка 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Инструкция setp.cc.type устанавливает предикатный регистр равным результату сравнения двух регистров соответствующего типа. Также существует инструкция set, где set.le.u32.u64 %r101, %rd12, %rd28 устанавливает 32-битный регистр %r101 в значение 0xffffffff, если 64-битный регистр %rd12 меньше или равен 64-битному регистру %rd28. В противном случае %r101 устанавливается в 0x00000000.
Есть несколько предопределенных идентификаторов, которые обозначают псевдорегистры. В частности, %tid, %ntid, %ctaid и %nctaid содержат, соответственно, индексы потоков, размеры блоков, индексы блоков и размеры сетки.
Пространстве состояний[править | править код]
Команды загрузки (ld) и сохранения (st) относятся к одному из нескольких различных пространств состояний (банков памяти), например ld.param. Существует восемь пространств состояний:[5].
.reg- регистры
.sreg- специальные, только для чтения, специфичные для платформы регистры
.const- общая, только для чтения память
.global- глобальная память, общая для всех потоков
.local- локальная память, приватная для каждого потока
.param- параметры, передаваемые в ядро
.shared- память, общая между потоками в блоке
.tex- глобальная текстурная память (устарела) Общая память объявляется в PTX файле посредством строк в начале следующего вида:Написание ядер на PTX требует явной регистрации модулей PTX через API драйвера CUDA, что обычно бывает более громоздким, чем использование API времени выполнения CUDA и компилятора CUDA от Nvidia, nvcc. Проект GPU Ocelot предоставлял API для регистрации модулей PTX наряду с вызовами ядер API времени выполнения CUDA, хотя проект GPU Ocelot более не поддерживается активно.[6]
.shared .align 8 .b8 pbatch_cache[15744]; // define 15,744 bytes, aligned to an 8-byte boundary
См. также[править | править код]
- Standard Portable Intermediate Representation (SPIR)
- CUDA binary (cubin) — a type of fat binary
Ссылки[править | править код]
Примечания[править | править код]
- ↑ [Cite web|url=https://llvm.org/docs/NVPTXUsage.html%7Ctitle=User Guide for NVPTX Back-end — LLVM 18.0.0git documentation|website=llvm.org|access-date=2023-12-27 «User Guide for NVPTX Back-end — LLVM 7 documentation»]
- ↑ 1. Overview — cuda-binary-utilities 12.3 documentation. docs.nvidia.com. Дата обращения: 27 декабря 2023. Архивировано 1 января 2024 года.
- ↑ nvptx - GCC Wiki. gcc.gnu.org. Дата обращения: 27 декабря 2023. Архивировано 30 сентября 2023 года.
- ↑ 1. Using Inline PTX Assembly in CUDA — Inline PTX Assembly in CUDA 12.3 documentation. docs.nvidia.com. Дата обращения: 27 декабря 2023. Архивировано 27 декабря 2023 года.
- ↑ CUDA Toolkit Documentation 12.3 Update 1. docs.nvidia.com. Дата обращения: 27 декабря 2023. Архивировано 27 декабря 2023 года.
- ↑ gtcasl/gpuocelot. — 2023-12-27. Архивировано 24 сентября 2023 года.
 | На эту статью не ссылаются другие статьи Википедии. |