Параллельное выполнение потока - Parallel Thread Execution

Параллельное выполнение потока (PTX, или же NVPTX[1]) является низкоуровневым параллельно нить исполнение виртуальная машина и архитектура набора команд используется в Nvidia с CUDA среда программирования. В nvcc компилятор переводит код, написанный на CUDA, C ++ -подобный язык, в инструкции PTX, а графический драйвер содержит компилятор, который переводит инструкции PTX в двоичный код[2] которые могут быть запущены на процессорных ядрах Графические процессоры Nvidia. В Коллекция компиляторов GNU также имеет базовую возможность для генерации PTX в контексте OpenMP разгрузка.[3] Встроенная сборка PTX может использоваться в CUDA.[4]

Регистры

PTX использует произвольно большой набор регистров; вывод компилятора почти чистый форма разового назначения, причем последовательные строки обычно относятся к последовательным регистрам. Программы начинаются с деклараций формы

.reg .u32 %р<335>; // объявляем 335 регистров% r0,% r1, ...,% r334 типа 32-битного целого числа без знака

Это язык ассемблера с тремя аргументами, и почти все инструкции явно указывают тип данных (с точки зрения знака и ширины), с которым они работают. Имена регистров начинаются с символа%, а константы являются буквальными, например:

shr.u64 % rd14, % rd12, 32; // сдвигаем вправо 64-битное целое число без знака из% rd12 на 32 позиции, получаем% rd14cvt.u64.u32 % rd142, % r112; // преобразовываем 32-битное целое число без знака в 64-битное

Существуют регистры предикатов, но скомпилированный код в модели шейдеров 1.0 использует их только в сочетании с командами ветвления; условная ветвь

@%стр. 14 бюстгальтер $ label; // переход к $ label

В setp.cc.type инструкция устанавливает регистр предиката на результат сравнения двух регистров соответствующего типа, также есть набор инструкция, где set.le.u32.u64 % r101, % rd12, % rd28 устанавливает 32-битный регистр % r101 к 0xffffffff если 64-битный регистр % rd12 меньше или равен 64-битному регистру % rd28. Иначе % r101 установлен на 0x00000000.

Есть несколько предопределенных идентификаторов, которые обозначают псевдорегистры. Среди прочего, % tid,% ntid,% ctaid, и % nctaid содержат, соответственно, индексы резьбы, размеры блоков, индексы блоков и размеры сетки.[5]

Государственные пространства

Нагрузка (ld) и сохраните (ул) команды относятся к одному из нескольких отдельных пространств состояний (банков памяти), например ld.param.Существует восемь пространств состояний:[5]

  • .reg : регистры
  • .sreg : специальные регистры только для чтения, зависящие от платформы
  • .const : общая, постоянная память
  • .Глобальный : глобальная память, общая для всех потоков
  • .местный : локальная память, частная для каждого потока
  • .param : параметры, переданные ядру
  • .общий : память, разделяемая между потоками в блоке
  • .tex : глобальная память текстур (не рекомендуется)

Общая память объявляется в файле PTX через строки в начале формы:

.общий .выровнять 8 .b8 pbatch_cache[15744]; // определение 15744 байта, выровненного по 8-байтовой границе

Написание ядер в PTX требует явной регистрации модулей PTX через API драйвера CUDA, что обычно более громоздко, чем использование CUDA Runtime API и компилятора NVIDIA CUDA, nvcc. Проект GPU Ocelot предоставил API для регистрации модулей PTX вместе с вызовами ядра CUDA Runtime API, хотя GPU Ocelot больше не поддерживается активно.[6]

Смотрите также

Рекомендации

  1. ^ "Руководство пользователя серверной части NVPTX - документация LLVM 7". llvm.org.
  2. ^ «Бинарные утилиты CUDA». docs.nvidia.com. Получено 2019-10-19.
  3. ^ "nvptx". GCC Wiki.
  4. ^ «Встроенная сборка PTX на CUDA». docs.nvidia.com. Получено 2019-11-03.
  5. ^ а б "PTX ISA версии 2.3" (PDF).
  6. ^ «GPUOCelot: платформа динамической компиляции для PTX». github.com.

внешняя ссылка