Parallel Thread Execution

Parallel Thread Execution (PTX или NVPTX) — это низкоуровневая виртуальная машина и архитектура набора инструкций для параллельного выполнения потоков, используемая в программной среде CUDA от Nvidia. Компилятор NVCC переводит код, написанный на CUDA, языке, похожем на C++, в инструкции PTX (язык ассемблера, представленный в виде ASCII текста, а графический драйвер содержит компилятор, который переводит инструкции PTX в исполняемый двоичный код, который может выполняться на процессорных ядрах GPU от Nvidia. В коллекции компиляторов GNU также есть базовая возможность генерации PTX в контексте оффлоадинга OpenMP. Встроенная сборка PTX может использоваться в CUDA.

Регистры

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. Существует восемь пространств состояний:.

.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 более не поддерживается активно.

См. также

• Standard Portable Intermediate Representation (SPIR)
• CUDA binary (cubin) — a type of fat binary

Ссылки

• PTX ISA page on NVIDIA Developer Zone