Модель CUDA

- ссылки на текстуры, пробники, или поверхности с помощью текстуры и поверхности инструкций загрузки / сохранения (tex, suld, sust, sured).

- получение значения именованного члена по инструкции запроса (txq, suq)

- cоздание ссылки на непрозрачные переменные с помощью MOV, например, mov.u64 reg, opaque_var;.

- непрозрачные переменные не могут появиться в инициализаторах, например, для инициализации указателя на непрозрачной переменной.

Косвенный доступ к текстурам поддерживается только в едином режиме текстур. Три встроенных типа - это .texref, .samplerref и .surfref. Для работы с текстурами и шаблонами PTX имеет два режима работы. В едином режиме, текстурная и шаблонная информация доступна через одну комманду .texref.

В автономном режиме, текстурная и шаблонная информации имеют свои собственные источники, что позволяет им быть определенными отдельно и в сочетании в зависимости от использования в программе. В автономном режиме, поля типа .texref, которые описывают свойства шаблона игнорируются, так как эти свойства определяются .samplerref - переменными.

Таблица 9 и Таблица 10 представляют перечень членов каждого типа для унифицированных и независимых текстурных модов. Эти члены и их значения имеют точные сопоставления методов и значений, определенных в текстурах класса HW, а также выставленных значений через API.

Таблица 9 Непрозрачный тип поля в едином режиме работы

Член	Значения .texref	Значения .surfref
width	в элементах
height	в элементах
depth	в элементах
channel_data_type	enum тип, соответствующий языку API
channel_order	enum тип, соответствующий языку API
normalized_coords	0, 1	N/A
filter_mode	nearest, linear	N/A
addr_mode_0, addr_mode_1, addr_mode_2	wrap,mirror, clamp_ogl, clamp_to_edge, clamp_to_border	N/A
array_size	как номер в текстурах текстурного массива	as number of surfaces in a surface array
num_mipmap_levels	как номер уровня в массиве MIP-типа	N/A
num_samples	как номер шаблона в мультишаблонной текстуре	N/A
memory_layout	N/A	1 for linear memory layout; 0 otherwise

Таблица 10 Непрозрачный тип поля в раздельном режиме работы

Member	.samplerref values	.texref values	.surfref values
width	N/A	в элементах
height	N/A	в элементах
depth	N/A	в элементах
channel_data_type	N/A	enum тип, соответствующий языку API
channel_order	N/A	enum тип, соответствующий языку API
normalized_coords	N/A	0, 1	N/A
force_unnormalized_coords	0, 1	N/A	N/A
filter_mode	nearest, linear	ignored	N/A
addr_mode_0, addr_mode_1, addr_mode_2	wrap,mirror, clamp_ogl, clamp_to_edge, clamp_to_border	N/A
array_size	N/A	как номер тектуры в массиве текстур	как номер шаблона в массиве шаблонов
num_mipmap_levels	N/A	как номер уровня в массиве MIP-типа	N/A
num_samples	N/A	как номер шаблона в мультишаблонной текстуре	N/A
memory_layout	N/A	N/A	1 for linear memory layout; 0 otherwise

Тип channel_data_type и channel_order

Channel_data_type и channel_order поля имеют типы перечисления соответствующие исходному языку API. В настоящее время, OpenCL является только язык программирования, который определяет эти поля. В таблице 12 и таблице 11 показаны перечисления значений, определенные в OpenCL версии 1.0 для типа данных канала и порядка каналов.

CL_SNORM_INT8	0x10D0
CL_SNORM_INT16	0x10D1
CL_UNORM_INT8	0x10D2
CL_UNORM_INT16	0x10D3
CL_UNORM_SHORT_565	0x10D4
CL_UNORM_SHORT_555	0x10D5
CL_UNORM_INT_101010	0x10D6
CL_SIGNED_INT8	0x10D7
CL_SIGNED_INT16	0x10D8
CL_SIGNED_INT32	0x10D9
CL_UNSIGNED_INT8	0x10DA
CL_UNSIGNED_INT16	0x10DB
CL_UNSIGNED_INT32	0x10DC
CL_HALF_FLOAT	0x10DD
CL_FLOAT	0x10DE

Таблица 12 Типы каналов комманд OpenCL

CL_R	0x10B0
CL_A	0x10B1
CL_RG	0x10B2
CL_RA	0x10B3
CL_RGB	0x10B4
CL_RGBA	0x10B5
CL_BGRA	0x10B6
CL_ARGB	0x10B7
CL_INTENSITY	0x10B8
CL_LUMINANCE	0x10B9

22. Переменные

В PTX, объявление переменной описывает как тип переменной и ее пространство состояний. В дополнение к базовым типам, PTX поддерживает типы для простых совокупных объектов, таких как вектора и массивы.

Объявления переменных

Все для хранения данных задается с объявлениями переменных. Каждая переменная должна находиться в одном из пространств состояний, перечисленных в предыдущем разделе.

А имена переменных - декларации пространства, в котором переменная находится, его тип и размер, его имя, по желанию размер массива, по желанию инициализатор, и дополнительный фиксированный адрес переменной.

Предикатные переменные могут быть объявлены только в реестре пространства состояний.

Examples

.global .u32 loc;

.reg .s32 i;

.const .f32 bias[] = {-1.0, 1.0};

.global .u8 bg[4] = {0, 0, 0, 0};

.reg .v4 .f32 accel;

.reg .pred p, q, r;

23. Векторы

Поддерживаются векторные типы ограниченной длины. Векторы длиной 2 и 4 любого непредикатного фундаментального типа может быть объявлен с помощью префикса типа с .х2 или .v4. Векторы должны быть основаны на фундаментальном типе, и они могут находиться в регистре места. Вектор не может превышать 128-бит в длину; например, .v4.f64 не допускается. Трехмерный векторы может быть обработан с помощью .v4 вектор, в котором четвертый элемент обеспечивает заполнение. Это общий случай для трехмерных сеток, текстур и т.д.

Examples

.global .v4 .f32 V; // a length-4 vector of floats

.shared .v2 .u16 uv; // a length-2 vector of unsigned ints

.global .v4 .b8 v; // a length-4 vector of bytes

24. Объявление массива

Объявление массива небходимо для того, чтобы позволить программисту выделить резервное пространство. Чтобы объявить массив: имя переменной следует из одномерных объявлений, подобных объявлениям фиксированного размера массива в C, размер каждого измерения постоянное выражение.



Examples

.local .u16 kernel[19][19];

.shared .u8 mailbox[128];

25. Выравнивание

Выравнивание хранения для всех адресуемых переменных может быть указано в объявлении переменной. Выравнивание задается с помощью спецификатора дополнительного подсчета .alignbyte сразу после пространства состояний спецификатора.

Переменные будут выровнены по адресу, который является целым числом, кратным байту-количеству. Значение выравнивания байт-счетчик должен быть степенью двойки. Для массивов, выравнивание определяет выравнивание адресов для начального адреса всего массива, а не для отдельных элементов.

Выравнивание по умолчанию для скалярных и матричных переменных кратно размеру базового типа. Выравнивание по умолчанию для векторных переменных является кратным общему размеру вектора.

Examples

// allocate array at 4-byte aligned address. Elements are bytes.

.const .align 4 .b8 bar[8] = {0,0,0,0,2,0,0,0};

26. Параметризованные имена переменных

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

Например, предположим, что программа использует большое количество, скажем, сто, из .b32 переменных, названных% r0,% r1, ...,% R99. Эти 100 регистровых переменных могут быть объявлены следующим образом:

.reg .b32 %r<100>; // declare %r0, %r1, ..., %r99

27. Атрибуты переменных

Переменные могут быть объявлены с дополнительной директивой .attribute, которая позволяет указать специальные атрибуты переменных. Ключевое слово .attribute следует перед спецификацией атрибута в скобках. Несколько атрибутов разделены запятой.

Поддерживаются следующие атрибуты:

.managed - атрибут указывает, что переменная будет выделена в единой виртуальной среде памяти, где хозяин и другие устройства в системе может ссылаться на переменную непосредственно. Этот атрибут может быть использован только с переменными в .global пространстве состояний.PTX ISA

Examples

.global .attribute(.managed) .s32 g;

.global .attribute(.managed) .u64 x;

28. Операнды команд

Тип данных операнда

Все операнды в инструкции - известный тип из своих деклараций. Каждый тип операнда должен быть совместим с типом определяемого шаблоном инструкции и типа команды. Там нет автоматического преобразования между типами.

Тип бит совместим с каждым типом, имеющий тот же размер. Целые типы общего размера совместимы друг с другом

Операнды - источники

Исходные операнды обозначаются в описаниях команд именами a, b и с. Операнды для инструкций ALU должны быть переменными .reg, объявленными в регистровом пространстве состояний. Для большинства операций, размеры операндов должны быть совместимыми.

cvt (конвертировать) команда проводится над различными типами операндов, так как его работа заключается в преобразовании практически из любого типа данных в любой другой тип данных (и размер).

ld, st, mov, и cvt - инструкции, чтобы скопировать данные из одного места в другое. Инструкциb ld и st перемещают данные из / в адресные пространства состояний в / из регистров. Инструкция mov копирует Данные между регистрами.

Большинство инструкций имеет дополнительный защитный предикат, который контролирует условное выполнение. Предикатные операнды обозначаются именами p, q, r, s.

Операнды назначения

Инструкции PTX, которые сохраняют результат в поле, обозначаются именем d (для назначения) в описаниях команд. Операнд назначения является скалярной или векторной переменной в регистровом пространстве состояний.

Адреса в качестве операндов

Адресная арифметика выполняется с использованием целочисленной арифметики и логических команд. Примеры включают арифметику указателей и указатели сравнения. Все адреса и адресные расчёты на основе байта; нет поддержки C-стиля арифметики указателей.

Инструкция mov может быть использован для перемещения адрес переменной в указатель. Адрес - смещения в пространстве состояний, в котором объявляется переменная. Загрузки и сохранения операции перемещения данных между регистрами и местах - в адресных пространствах состояний. Синтаксис такой же, что используется во многих языках ассамблера.

Адрес выражения включают имена переменных, адресные регистры, регистр адреса плюс смещение байта, а также немедленные адресные выражения, которые имеют значения во время компиляции к постоянному адресу.

Вот несколько примеров:

.shared .u16 x;

.reg .u16 r0;

.global .v4 .f32 V;

.reg .v4 .f32 W;

.const .s32 tbl[256];

.reg .b32 p;

.reg .s32 q;

ld.shared.u16 r0,[x];

ld.global.v4.f32 W, [V];

ld.const.s32 q, [tbl+12];

mov.u32 p, tbl;

29. Массивы в качестве операндов

Массивы всех типов могут быть объявлены и идентификатор становится адресом постоянной в пространстве, где объявлен массив. Размер массива является постоянным в программе.

Элементы массива можно получить с помощью рассчитанного адреса байта или по индексу в массиве с использованием квадратных скобок обозначения. Выражение в квадратных скобках является либо постоянным числом, регистровой переменной или просто регистр с постоянным смещением выражения, где смещение является постоянным выражением, которое прибавляется или вычитается из регистровой переменной.

Examples:

ld.global.u32 s, a[0];

ld.global.u32 s, a[N-1];

mov.u32 s, a[1]; // move address of a[1] into s

Вектор как операнд

Операнды вектора поддерживают ограниченное число команд: mov, ld, st, and tex.

Векторы могут быть также аргументами, поддерживаемыми в функциях.

Элементы вектора могут быть выделены при помощи суффиксов .x, .y, .z and .w, или как в цветовых полях .r, .g, .b and .a.

Список в скобках используется для того, чтобы распределить вектор.

.reg .v4 .f32 V;

.reg .f32 a, b, c, d;

mov.v4.f32 {a,b,c,d}, V;

Векторы могут быть использованы для реализации загрузки и хранения, которые могут улучшить производительность памяти.

В регистрах операций загрузки / сохранения может быть вектор или скобки, заключающие список типизированных скаляров.

Examples:

ld.global.v4.f32 {a,b,c,d}, [addr+offset];

ld.global.v2.u32 V2, [addr+offset2];

Элементы в фигурных скобках - вектора, скажем, {Ra, Rb, Rc, Rd}, соответствуют извлеченным элементам следующим образом:

Ra = V.x = V.r

Rb = V.y = V.g

Rc = V.z = V.b

Rd = V.w = V.a

30. Метки и имена функций в качестве операндов

Этикетки и имена функций могут быть использованы только во вложенных и вызываемых инструкциях и в инструкциях перемещения, чтобы получить адрес метки или функции в регистр для того, чтобы вызвать.

Функциональные объявления и определения

В PTX, функции декларируются и определяются с помощью .func директивы. Объявление функции определяет дополнительный список возвращаемых параметров, имя функции, необязательный список входных параметров; вместе это задает интерфейс функции или прототип. Определение функции определяется как интерфейс и тело функции. Функция должна быть объявлена или определена до того, как называется.

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

Example

.func foo

{

...

ret;

}

...

call foo;

...

Скалярные и векторные базового типа входные и возвращаемые параметры могут быть представлены просто как регистровые переменные. По вызову аргументы могут быть зарегистрированы как переменные и константы, а возвращаемые значения могут быть помещены непосредственно в регистровые переменные. Аргументы и возвращаемые переменные, должны иметь тип и размер, которые соответствуют формальным параметрам.

Example

.func (.reg .u32 %res) inc_ptr ( .reg .u32 %ptr, .reg .u32 %inc )

{

add.u32 %res, %ptr, %inc;

ret;

}

...

call (%r1), inc_ptr, (%r1,4);

...

31. Набор комманд

Инструкции PTX, как правило, имеют от нуля до четырех операндов плюс, по желанию, защитный предикат после символа @ слева от opcode:

- @p opcode;

- @p opcode a;

- @p opcode d, a;

- @p opcode d, a, b;

- @p opcode d, a, b, c;

Для инструкций, которые создают результирующее значение, d операнд является операндом назначения, a, b, c являются исходными операндами.

Инструкция setp пишет два регистра назначения. Мы используем символ | для разделения нескольких регистров назначения.

setp.lt.s32 p|q, a, b; // p = (a < b); q = !(a < b);

Для некоторых инструкций операнд назначения является обязательным. Битоприемник операнд обозначается символом подчеркивания (_) может быть использован вместо регистра назначения.

Предикатное выполнение

В PTX, предикатные регистры являются виртуальными и имеют .pred как спецификатор типа. Так, предикат регистры могут быть объявлены как .reg .pred p, q, r;

Все инструкции имеют дополнительный защитный предикат, который управляет условным выполнением инструкции.Синтаксисом, указывающим условное выполнение, является префикс - инструкция с @ {!} p, где р - предикатная переменная, возможно отрицательная. Инструкции без предикат выполняются безоговорочно.

Предикаты обычно устанавливается в результате сравнения, выполненного setp- инструкции.

В качестве примера, рассмотрим код на высоком уровне:

if (i < n)

j = j + 1;

Который может быть записан на PTX как:

setp.lt.s32 p, i, n; // p = (i < n)

@p add.s32 j, j, 1; // if i < n, add 1 to j

Сравнения

Конец формы

Таблица 17 Операции сравнения для целочисленного и битового типов

Значение	Знаковая операция	Беззнаковая операция	Битовая операция
a == b	eq	eq	eq
a != b	ne	ne	ne
a < b	lt	lo	n/a
a <= b	le	ls	n/a
a > b	gt	hi	n/a
a >= b	ge	hs	n/a

Таблица 18 Операции сравнения для чисел с плавающей точкой

Значение	Оператор с плавающей точкой
a == b && !isNaN(a) && !isNaN(b)	eq
a != b && !isNaN(a) && !isNaN(b)	ne
a < b && !isNaN(a) && !isNaN(b)	lt
a <= b && !isNaN(a) && !isNaN(b)	le
a > b && !isNaN(a) && !isNaN(b)	gt
a >= b && !isNaN(a) && !isNaN(b)	ge

Таблица 19 Операции сравнения для чисел с плавающей точкой с применением NaN

Значение	Оператор
a == b || isNaN(a) || isNaN(b)	equ
a != b || isNaN(a) || isNaN(b)	neu
a < b || isNaN(a) || isNaN(b)	ltu
a <= b || isNaN(a) || isNaN(b)	leu
a > b || isNaN(a) || isNaN(b)	gtu
a >= b || isNaN(a) || isNaN(b)	geu

Управление предикатами

Предикатные значения могут быть вычислены и управляемы, используя следующие инструкции: and, or, xor, not, и mov.

Там нет прямого преобразования между предикатами и целочисленными значениями, и никакого прямого способа загрузки или хранения предикатных значений регистров. Тем не менее, setp может быть использован для создания предикат из целого числа, и выбор (selp -инструкция на основе предиката) может быть использован для создания целочисленного значение в зависимости от значения предиката; например:

selp.u32 %r1,1,0,%p; // convert predicate to 32-bit value

Арифметические операции для целых чисел

Операция	Значение
add	Складывает два значения
sub	Вычитает два значения
mul	Умножает два значения
mad	Умножает два значения, извлекает высокую или низкую половину результата, а также добавляет третье значение.
mul24	Вычисляет произведение двух 24-битных целочисленных значений, проведенных в 32-разрядных регистрах источника и вернуть либо высокие или низкие 32-бит 48-разрядного результата.
mad24	Вычисляет произведение двух 24-битных целочисленных значений, проведенных в 32-разрядных регистрах источников, а также добавить третье 32-битное значение или к высоким или к низким 32-битам 48-разрядного результата. Вернуться либо высокие или низкие 32-бит 48-разрядного результата.
sad	Добавляет абсолютное значение AB-С и записывает полученное значение в D
div	Делит одно значение на другое
rem	Находит остаток от деления одного числа на другое
abs	d = |a|;
neg	d = -a;
min	d = (a < b) ? a : b;
max	d = (a > b) ? a : b;
popc	Подсчитывает количество единичных битов и помещает в результат в 32-разрядный регистр назначения
clz	Подсчитывает количество ведущих нулей, начиная с наиболее значимого бита и поместить результат в 32-разрядный регистр назначения
bfind	Найти номер самого значительного бита из беззнаковых битов в A и поместите результат в D. Для целых чисел без знака, bfind возвращает битовую позицию самого значительного 1. Для целых чисел, bfind возвращает битовую позицию самого значительного 0 для отрицательных входов и самым значительным 1 для неотрицательных входов.
brev	Выполняет побитовый разворот ввода
bfe	Извлечение битового поля из A и поместить нуль или знаковый результат в D. Источник B дает положение отправного бита битового поля и C дает битовую длину поля в битах.
bfi	Совместите и вставьте битное поле из А в В, и поместить результат в F. C дает исходное положение бита для вставки, и источник D дает битовую длину поля в битах.

Арифметические операции для чисел с пласающей точкой

Операция	Значение
testp	testp проверяет общие свойства чисел с плавающей точкой и возвращает значение предиката 1, если это правда и 0, если False. testp.finite Правда, если вход не бесконечен или NaN testp.infinite Правда, если вход положителен или отрицателен бесконечен testp.number Правда, если вход не NaN testp.notanumber Правда, если вход NaN testp.normal Правда, если вход обычный номер (не NaN, не бесконечность) testp.subnormal Правда, если вход субнормальная номер (не NaN, не бесконечность)
copysign	Скопируйте знаковый бит в значении B, и возвращает результат в виде D.
add	Складывает два значения
sub	Вычитает два значения
mul	Умножает два значения
fma	Выполнение конденсированное умножения-сложения без потери точности в промежуточном продукте и того.
mad	Умножение два значения и добавляет третий, а затем записывает полученное значение в регистр назначения.
div	Делит одно значение на другое
abs	d = |a|;
neg	d = -a;
min	d = (a < b) ? a : b;
max	d = (a > b) ? a : b;
rcp	Выводит обратное значение числа
sqrt	Вычисляет значение квадратного корня
rsqrt	Выводит обратное значение квадратного корня числа
sin	Выводит значение синуса числа
cos	Выводит значение косинуса числа
lg2	Выводит значение двоичного логарифма числа
ex2	Выводит степень двойки

Операции сравнения

Операция	Значение
set	Сравнивает два числовых значения и, возможно, сочетает в себе результат с другим значением предиката, применяя логический оператор. Если этот результат верен, 1.0f написано с плавающей точкой типов назначения, а 0xffffffff написана для целочисленных типов назначения. В противном случае, 0x00000000 написано.
setp	Сравнивает два значения и сочетает в себе результат с другим значением предиката, применяя логический оператор. Это результат записывается в первый операнд назначения.Связанное значение, рассчитанное с помощью дополнения к результатам сравнения записывается во второй операнд назначения.
selp	Условный оператор
slct	Условный оператор. Если с ? 0, A хранится в D, иначе B хранится в D.

Операции преобразования и движения данных

Операция	Значение
mov	Запись в регистр D значения А. Операнд может быть регистром, специальным регистром, содержащая дополнительный смещенный в адресное пространство памяти, этикетки, или имя функции. Для переменных, объявленных в .const, .global, .local и .shared, mov помещает адрес переменной (например, адрес переменной в ее пространстве состояний) в регистр назначения.Общий адрес переменной в константном, глобальном, локальном или общем пространстве состояний может быть получен сначала взятием адреса в пространстве состояний с mov, а затем превращением его в универсальный адрес с помощью cvta инструкции; с другой стороны, общий адрес переменной, объявленной в константном, глобальном, локальном или общем пространстве состояний быть взяты непосредственно с помощью cvta инструкции.
shfl	Обмен регистрами данных между нитями
prmt	Выбирает четыре произвольных байта из двух 32-разрядных регистров, и собрать их в 32-разрядный регистр назначения.
ld	Загружает регистровые переменные D от места указанного адреса операнда источника А в указанном пространстве состояний. Если состояние пространства не дается, выполнить загрузку, используя общую адресацию
ldu	Загружает данные только для чтения с адреса, который является общим для всех нитей по основе
st	Сохраняет регистровую переменную в переменную адресного пространства состояний.
prefetch, prefetchu	Выводит строку, содержащую общий адрес на определенном уровне иерархии памяти, в указанном пространстве состояний.
isspacep	Определяет, соответствует ли адрес данному пространству состояний
cvta	Преобразование адреса из константного, глобального, локального, или общего пространства состояния общим, или наоборот.
cvt	Преобразует типы

Операции для работы с текстурами

- Статическая инициализация текстурных и шаблонных дескрипторов

- Модуль-сфера, а также записи зарезервированных текстурных и шаблонных дескрипторов.

- Возможность запрашивать поля внутри текстуры и шаблонных дескрипторов.

Режимы текстурирования

Для работы с текстурами и шаблонами PTX имеет два режима работы. В едином режиме, информация о текстуре и шаблоне доступна через однин обработчик .texref. В автономном режиме, информация о текстуре и шаблоне имеют свои собственные обработчики, что позволяет им быть определенными отдельно и в сочетании. Преимущество единого режима в том, что поддерживает 128 шаблонов на ядро, с тем ограничением, что они соответствуют 1-к-1 с 128 возможными текстурами на ядро. Преимущество автономного режима в том, что текстуры и шаблоны могут быть скомбинированы, однако количество шаблонов значительно ограничена до 16 на ядро.

Режим текстурирования выбирается с помощью .target с опцией texmode_unified и texmode_independent. Модуль PTX может объявить только один режим текстурирования. По умолчанию предполагается использовать единый режим.



32. Мипмапы

Мипмап последовательность текстур, каждая из которых представляет собой прогрессивно меньше представление разрешение и того же изображения. Высота и ширина каждого изображения, или уровень детализации (LOD), в MIPMAP является мощностью в два раза меньшей, чем в предыдущем уровне. Мипмапов используются в графических приложениях для повышения скорости рендеринга и уменьшить артефакты сглаживания.

Операция	Значение
tex	Выполняет поиск тектурной памяти. tex.{1D,2D,3D} Эта функция загружает данные из текстуры c именем в операнде А в точке с координатами, заданными операндом В в конечном операнде D. Операнд C является скалярным или одиночным кортежем 1D текстур; вектором из двух элементов для 2D текстур; и вектором из четырех элементов для 3D текстур, где четвертый элемент игнорируется. tex.{a1d,a2d} Далее следует выбор текстуры из массива. Инструкция сначала выбирает текстуру из массива текстур названного операнда с помощью индекса. Инструкция загружает данные из выбранной текстуры в точке с координатами, заданными остальными элементами операнда tex.cube Эта функция загружает данные из cubemap текстуры в операндом А в точке с координатами, заданных операндом C в конечном D. Cubemap - текстуры специальные двумерные слоистые текстуры, состоящие из шести слоев, которые представляют граней куба. Все слои в cubemap одного и того же размера и площади (т.е., ширина равна высоте). tex.2dms Инструкция сначала выбирает текстуру из массива текстур, определённых в операнде A с помощью индекса, заданного первым элементом массива координат вектора C. Инструкция загружает данные из выбранного мульти-образца текстуры, заданного во втором элементе операнда С, в точку с координатами, заданными остальными элементами операнда С в конечном операнде D.
tld4	Выборка текстуры по 4 параметрам
txq	Атрибуты текстур и шаблонов (.width, .height, .depth, .channel_data_type)
istypep	Запрос указывает ли регистр на непрозрачную переменную указанного типа.

Операции для работы с поверхностями

Операция	Значение
suld	Загрузка из памяти поверхностей
sust	Запись в память поверхностей
sured	Преобразование поверхности
suq	Определение атрибутов поверхности

Операции управления потоками

Операция	Значение
{}	Группировка инструкций
@	Предикатное выражение
bra	Переход к цели и продолжение исполнения
call	Вызов функции
ret	Возвращение значения из функции
exit	Завершение нити

Операции параллельной синхронизации и коммуникации

Операция	Значение
bar	Функция синхронизации потоков для обмена данными между ними
membar	Функция ожидания, пока все потоки выполнят обращение к памяти
atom	Атомарные операции преобразования в общении нить-нить
red	Операции приведения на уровнях глобальной и разделяемой памяти
vote	Голосование различных нитей группы

Видео - операции

Операция	Значение
vadd	Сложение слов
vsub	Вычитание слов
vabsdiff	Абсолютное значение разности слов
vmin	Слово-минимум
vmax	Слово -максимум
vshl	Сдвиг слова влево
vshr	Сдвиг слова вправо
vmad	Расчёт (a*b) + c
vset	Сравнение двух слов
vadd2	Целое двойное SIMD сложение
vsub2	Целое двойное SIMD вычитание.
vavrg2	Целое двойное SIMD среднее
vabsdiff2	Целое двойное SIMD абсолютное значение разности.
vmin2	Целый двойной SIMD минимум
vmax2	Целый двойной SIMD максимум
vset2	Целое двойное SIMD сравнение
vadd4	Целое четверное SIMD сложение.
vsub4	Целое четверное SIMD вычитание.
vavrg4	Целое четверное SIMD среднее.
vabsdiff4	Целое четверное SIMD абсолютное значение разницы.
vmin4	Целый четверной SIMD минимум
vmax4	Целый четверной SIMD максимум.
vset4	Целое четверное SIMD сравнение

Другие инструкции

Операция	Значение
trap	Прервать выполнение и генерировать прерывание для процессора.
brkpt	Приостанавливает выполнение
pmevent	Инициировать одно или более высоко производительное событие монитора.

Размещено на Allbest.ru