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].
См. также
СсылкиПримечания
|
Portal di Ensiklopedia Dunia
