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 файле посредством строк в начале следующего вида:
.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]

См. также

[править | править код]

Ссылки

[править | править код]

Примечания

[править | править код]
  1. ↑ [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»]
  2. ↑ 1. Overview — cuda-binary-utilities 12.3 documentation. docs.nvidia.com. Дата обращения: 27 декабря 2023. Архивировано 1 января 2024 года.
  3. ↑ nvptx - GCC Wiki. gcc.gnu.org. Дата обращения: 27 декабря 2023. Архивировано 30 сентября 2023 года.
  4. ↑ 1. Using Inline PTX Assembly in CUDA — Inline PTX Assembly in CUDA 12.3 documentation. docs.nvidia.com. Дата обращения: 27 декабря 2023. Архивировано 27 декабря 2023 года.
  5. ↑ CUDA Toolkit Documentation 12.3 Update 1. docs.nvidia.com. Дата обращения: 27 декабря 2023. Архивировано 27 декабря 2023 года.
  6. ↑ gtcasl/gpuocelot. — 2023-12-27. Архивировано 24 сентября 2023 года.