Выполнение параллельного потока - Parallel Thread Execution
Параллельное выполнение резьбы ( PTX или NVPTX ) является низкоуровневым параллельным потоком выполнение виртуальной машины и набор инструкций архитектуры используется в Nvidia «s CUDA среды программирования. NVCC компилятор транслирует код , написанный на CUDA, на C ++ -как языка, в инструкции PTX, и графический драйвер содержит компилятор , который транслирует инструкции PTX в двоичный код , который может быть запущен на обрабатывающих ядер графических процессоров NVIDIA . GNU Compiler Collection также имеет базовые возможности для генерации 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
инструкция, где устанавливает 32-битовый регистр , чтобы , если 64-битный регистр меньше или равно 64-битовый регистр . В противном случае установлено значение .
set.le.u32.u64 %r101, %rd12, %rd28
%r101
0xffffffff
%rd12
%rd28
%r101
0x00000000
Есть несколько предопределенных идентификаторов, которые обозначают псевдорегистраторы. Среди прочего, %tid, %ntid, %ctaid
и %nctaid
содержат, соответственно, индексы резьбы, размеры блоков, индексы блоков и размеры сетки.
Государственные пространства
Команды Load ( ld
) и store ( 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, что обычно более громоздко, чем использование CUDA Runtime API и компилятора Nvidia CUDA, nvcc. Проект GPU Ocelot предоставил API для регистрации модулей PTX вместе с вызовами ядра CUDA Runtime API, хотя GPU Ocelot больше не поддерживается активно.
Смотрите также
использованная литература
внешние ссылки
- PTX ISA версии 1.4 NVIDIA, 31 марта 2009 г.
- PTX ISA версии 2.3 NVIDIA, 03.11.2011
- PTX ISA версии 3.2 NVIDIA, 2013-07-19
- PTX ISA версии 4.0 NVIDIA, 12 апреля 2014 г.
- PTX ISA версии 4.3 NVIDIA, 2015-08-15
- PTX ISA версии 5.0 NVIDIA, 2017-06-xx
- PTX ISA версии 6.0 NVIDIA, 2017-09-xx
- PTX ISA версии 6.1 NVIDIA, 2018-03-xx
- PTX ISA версии 6.2 NVIDIA, 2018-08-xx
- PTX ISA версии 6.3 NVIDIA, 2018-10-xx
- PTX ISA версии 6.4 NVIDIA, 2019-08-xx
- PTX ISA версии 6.5 NVIDIA, 2019-11-xx
- PTX ISA версии 7.0 NVIDIA, 2020-08-xx
- PTX ISA версии 7.1 NVIDIA, 2020-09-xx
- Страница PTX ISA в NVIDIA Developer Zone