Модель CUDA

Размещено на http://www.allbest.ru

1. Модель CUDA

GPU предназначены для выполнения интенсивных расчётов. Задачи с интенсивным обращением к памяти или сложной логикой будут выполняться неэффективно, т.к. GPU обладает слабыми средствами кеширования обращения к памяти (к тому же их требуется настраивать вручную) и "не переносит" ветвлений в программе, особенно если потоки одного блока расходятся по разным веткам.

Модель CUDA предполагает, что программист в начале разбивает задачу на независимые части (блоки), которые могут выполняться параллельно. Затем каждый блок разбивается на множество параллельно выполняющихся потоков (thread), которые могут зависеть друг от друга. CUDA обеспечивает средства расширения языка C для параллельного запуска множества потоков, выполняющих одну и ту же функцию (ядро, kernel). Максимальный размер ядра - 2 миллиона инструкций PTX (если сумеете скомпилировать - мне не удалось превысить 50 тысяч инструкций (блоб в 400KB) из-за падения nvopencc). Потоки объединяются в блоки (до 512 потоков), блоки объединяются в сетки (решётки, grid). Потоки внутри блока запускаются на одном мультипроцессоре (MP), имеют общую разделяемую память и могут (должны) синхронизовать ход выполнения задачи. Каждый поток имеет уникальный идентификатор внутри блока, выражаемый с помощью одномерного, двумерного или трёхмерного индекса (встроенная структурная переменная threadIdx типа dim3). Размерность блока доступна через встроенную переменную blockDim. Максимальные размерности: 512, 512, 64. Решётки могут быть одномерными или двумерными, максимальное значение элемента индекса: 65535. Индекс блока в решётке доступен через встроенную переменную blockIdx. Компоненты индексов нумеруются с нуля.

Порядок выполнения блоков не определён, блоки должны быть независимыми друг от друга. При запуске ядра блоки решётки нумеруются и распределяются по MP, имеющим достаточную свободную ёмкость региcтров, разделяемой памяти и ресурсов планировщика команд. MP состоит из 8 простых процессоров, 2 процессора для сложных операций (например, умножения;), пула регистров, разделяемой памяти и планировщика команд. Планировщик команд последовательно разбивает потоки активного блока на порции (warp), по 4 на каждый простой процессор и выполняет по одной простой команде одновременно для всех потоков порции за 4 цикла. Для исполнения одной команды порции потоков MP должен загрузить операнды для всех потоков порции, выполнить команду (одновременно), записать результат. Если доступ к памяти вызывает задержку (до 600 циклов), то планировщик может перейти к следующей порции. Все потоки порции начинают исполнение программы с одного и того же адреса, но каждый простой процессор имеет собственный счётчик команд и регистр состояния, что позволяет осуществить условное исполнение и ветвление. Однако, каждая ветка условия выполняется всеми потоками порции по очереди. Те потоки, для которых условие ветки не выполняется, "пропускают ход". По завершению расхождения все потоки порции опять одновременно выполняют полезную работу. Таким образом, ветвление при исполнении внутри порции сильно замедляет работу ядра. Потоки из разных порций могут выполнять различные ветки совершенно безнаказанно. По завершению всех потоков блока ресурсы MP освобождаются и на него может быть распределён следующий блок.

Планировщик MP имеет ограничения по максимальному числу одновременно запущенных блоков (до 8), максимальному числу порций (до 24 активных порций в версиях аппаратуры 1.0 и 1.1; до 32 - в версиях аппаратуры 1.2 и 1.3), максимальному числу потоков (до 768 активных потоков в версиях аппаратуры 1.0 и 1.1; до 1024 - в версиях аппаратуры 1.2 и 1.3).

Количество потоков в блоке и количество блоков в решётке выбирается программистом исходя из максимизации загрузки ресурсов MP и с учётом аппаратных ограничений (количество регистров, разделяемой памяти и т.д.). блоков д.б. не менее числа MP, лучше с запасом, чтобы MP не простаивал во время чтения из памяти. Количество потоков в блоке должно быть кратно размеру порции (32).

Ядро выполняется на GPU, в отличие от остальной части программы, выполняемой на CPU хостовой системы. GPU и CPU может быть несколько. GPU должны быть одинаковыми, причём не в режиме SLI. Один хостовый поток (процесс) может использовать только один GPU - тот, который был задействован первым. Хостовая система имеет свою ОП, обмен с глобальной памятью GPU и запуск ядра осуществляется через специальный интерфейс. Запуск ядра является асинхронным, т.е. управление немедленно возвращается хостовой программе. Также асинхронным может являться обмен между хостом и GPU и пересылки внутри GPU. Приложение может использовать либо низкоуровневый интерфейс драйвера CUDA (поставляется вместе с драйвером видеокарты), либо высокоуровневый API приложений CUDA (runtime, использует интерфейс драйвера CUDA). Поверх API приложений реализованы специализированные библиотеки CUDAPP (сортировка и псевдослучайный генератор), FFT (CUFFT, преобразование Фурье) и BLAS (CUBLAS, линейная алгебра).

спецификатор буферный память

2. Среда разработки (CUDA Toolkit)

В составе CUDA Toolkit поставляется компилятор nvcc, документация (в т.ч. ptx_isa_1.4.pdf, nvcc_2.2.pdf), библиотека API приложений CUDA (runtime, libcudart.so), библиотека низкоуровневого интерфейса драйвера CUDA (/usr/include/cuda/cuda.h и /usr/lib/libcuda.so из комплекта драйвера графики), библиотеки CUBLAS и CUFFT, CUDA Profiler 2.2.

Основной программой среды разработки является компилятор nvcc, который управляет всей цепочкой преобразований. Исходная программа может быть на языке C (.c), C++ (.cc, .cxx, .cpp), расширении языка C (.cu; .cup для препроцессированных файлов). В начале, стандартный cpp осуществляет макроподстановки, затем утилита cudafe в 2 прохода разделяет общий код программы, написанный на расширении языка C (.cu; .cup для препроцессированных файлов), на части, предназначенные к исполнению на хосте (.c, компилируется обычным gcc в режиме C++ в .o), и части, предназначенные к исполнению на GPU (.gpu, C). Простенький развёрнутый crypt-des (25 итераций по 16 раундов по 8 S-блоков по 10 переменных) обрабатывался 15 минут. Последние преобразуются nvopencc в код виртуальной машины (.ptx). nvopencc имеет ограничение по умолчанию - не более 26000 регистров каждого типа (а их препроцессор генерирует бесжалостно). После снятия ограничения на оптимизацию потребил (в фазе be) 6GB памяти (реально используются!) и 90 минут CPU, после чего всё равно упал. Код виртуальной машины компилируется оптимизирующим ассемблером ptxas (распределение виртуальных регистров в реальные, имеет свои ограничения) в двоичный блоб (.cubin), который утилита fatbin в сочетании с кодовым именем устройства либо размещает во внешний репозитарий (.fatbin), либо добавляет к коду хостовой части (.cu.c), из которого стандартый компилятор gcc (версии от 3.4 до 4.2; кстати, в Fedora gcc 4.3!) компилирует и собирает программу.

К хостовой программе в дополнение к остальным библиотекам (.a, .so) присоединяется библиотека функций высокого уровня cudart и/или низкого уровня cuda (cudart использует cuda; приложение может использовать только одну из них), обеспечивающих управление и доступ к GPU со стороны хоста. К программе, выполняемой на GPU, присоединяется библиотека функций, специфических для устройства. Имеется также общая часть - подмножество библиотеки C, реализованное как на хосте, так и на GPU. nvcc генерирует код, предназначенный для работы с cudart, при этом инициализация, управление контекстами и модулями скрыты от программиста.

С помощью макроса __CUDACC__ можно определить, компилируется ли программа с помощью nvcc.

Несколько хостовых потоков могут запускать ядра на одном GPU (ресурсы одного потока недоступны другому). Один хостовый поток не может использовать несколько GPU.

Поставляется в виде привычного для NVIDIA исполняемого bash-скрипта, содержащего самораспаковываемый архив (--help, --info, --list, --check). По умолчанию, всё ставится в каталог /usr/local/cuda, но можно задать свой путь вместо /usr/local (я задал /usr/local/cuda2.2, чтобы можно было иметь несколько версий одновременно).

Для автономной компиляции хостовой подпрограммы с помощью gcc необходимо указывать ключ "-malign-double" (nvcc делает это автоматически).

Хостовые подпрограммы компилируются в режиме C++, чтобы компилировать в режиме C требуется указать ключ "--host-compilation=c".

64-битная версия среды разработки генерирует 64-битный код как для хостовой части программы, так и для GPU. Пришлось оттрасировать все вызовы nvcc и переиграть их, заменив -m64 на -m32 в предназначенных для генерации кода для GPU командах. У меня это дало выигрыш в 9% и меньшее число регистров, что вообще бесценно.

3. Расширения языка C

Проrраммы для CUDA (соответствующие файлы обычно имеют расширение .сu) пишутся на «расширенном С и компилируются при помощи команды nvcc.

Вводимые в CUDA расширения языка С состоят из:

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

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

данных переменных;

- директивы, служащей для запуска ядра, задающей как данные, так и иерархию нитей;

- встроенных переменных, содержащих информацию о текущей нити;

- runtime, включающей в себя дополнительные типы данных.

Спецификаторы функций и переменных:

В CUDA используется следующие спецификаторы функций (табл. 2.1).

Таблица 2.1 спецификаторы функций в CUDA

Спецификатор	Функция выполняется на	Функция может вызываться из
_device_	device (GPU)	device (GPU)
_global_	device (GPU)	host (CPU)
_host_	host (CPU)	host (CPU)

Спецификаторы _host_ и _device_ могут быть использованы вместе (это значит, что соответствующая функция может выполняться как на GPU, так и на CPU - соответствующий код для обеих платформ будет автоматически cгeнeрирован компилятором). Спецификаторы global и host не могут быть использованы вместе.

Спецификатор _global_ обозначает ядро, и соответствующая функция должна возвращать значение типа void.

На функции, выполняемые на GPU (device и global), накладываются следующие ограничения:

- нельзя брать их адрес (за исключением _global_ функций);

- не поддерживается рекурсия;

- не поддерживаются stаtiс переменные внутри функции;

- не поддерживается переменное число входных аргументов.

Для задания размещения в памяти GPU переменных используются следующие спецификаторы: _device_, _constant_ и _shared_. На их использование также накладывается ряд ограничений:

- эти спецификаторы не могут быть применены к полям структуры (struct или union);

- соответствующие переменные могут использоваться только в пределах oднoгo файла, их нельзя объявлять как extern;

- запись в переменные типа _constant_ может осуществляться только CPU при помощи специальных функций;

- _shared_ переменые не могут инициализироваться при объявлении.

4. Добавленные типы

В язык добавлены 1/2/3/4- мepныe векторы из базовых типов (сhar, unsi gned, char, short, unsigned short, int, unsigned int, long, unsigned long, longlong, float и double) charl, char2, char3,char4, ucharl, uchar2, uchar3, uchar4, shortl, short2, short3, short4, ushortl, ushort2, ushort3, ushort4, intl, int2, int3, int4, uintl, uint2, uint3, uint4, longl, long2, long3, long4, ulongl, ulong2, ulong3, ulong4, floatl, float2, float3, float4, longlongl, longlong2, doublel и double2.

Обращение к компонентам вектора идет по именам х, у, z и w. Для создания значений векторов заданного типа служит конструкция вида make <typeName >.

int2 а = make_int (1,7); // Созоает вектор (1,7).

float3 и = make_float3 (1, 2, 3.4f ); // Созоает вектор (1.0f, 2.0f, 3.4f ).

Обратите внимание, что для этих типов (в отличие от шейдерных языков

GLSL, Cg и HLSL) не поддерживаются векторные покомпонентные операции, то есть нельзя просто сложить два вектора при помощи оператора «+» это необходима явно делать для каждой компоненты.

Также добавлен тип dim3, используемый для задания размерности. Этот тип основан на типе uint 3, но обладает при этом нормальным конструктором, инициализирующим все незаданные компоненты единицами.

dim3 blocks (16, 16 ); / / Эквивалентно blocks ( 16, 16, 1 ).

dim3 grid (256); / / Эквивалентно grid ( 256, 1, 1 ).

Добавленные переменные:

В язык добавлены следующие специальные переменные:

- gridDim - размер сетки (имеет тип dimз);

- blockDim - размер блока (имеет тип dim3);

- blockldx - индекс текущеrо блока в сетке (имеет тип uint3);

- threadIdx - индекс текущей нити в блоке (имеет тип uint3);

- warpSize - размер warp'a (имеет тип int).

Директива вызова ядра

Для запуска ядра на GPU используется следующая конструкция:

kernelName <<<Dg,Db,Ns,S>>> ( args );

Здесь kemelName это имя (адрес) соответствующей _global_ функции. Через Dg обозначена переменная (или значение) типа dim3, задающая размерность и размер сетки (в блоках). Переменная (или значение) Db - типа dim3, задает размерность и размер блока (в нитях).

Необязательная переменная (или значение) Ns типа size_t задает дополнительный объем разделяемой памяти в байтах, которая должна быть динамически выделена каждому блоку (к уже статически выделенной разделяемой памяти), если не задано, то используется значение О.

Переменная (или значение) Sтипа cudaStream_t задает поток (CUDA stream), в котором должен произойти вызов, по умолчанию используется поток 0.

Через args обозначены арrументы вызова функции kenelName (их может быть несколько ).

Следующий пример запускает ядро с именем mуКеrnel параллельно на n нитях, используя одномерный массив из двумерных (16х16) блоков нитей, и передает на вход ядру два параметра а и n. При этом каждому блоку дополнительно выделяется 512 байт разделяемой памяти и запуск, производится на потоке myStream.

mуКеrnеl <<<dim3(n/256) ,dim3(16,16), 512, myStream>>> ( а, n );

Добавленные функции:

CUDA поддерживает все функции из стандартной библиотеки языка С. Однако при этом следует иметь в виду, что большинство стандартных математических функций используют числа с двойной точностью (double).

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

В таблице 2.2 приведены основные f1оаt - функции и их оптимизированные вepсии пониженной точности.

Для ряда функций можно задать требуемый способ округления. Используемый способ задается при помощи одноrо из следующих суффиксов:

- rn - окруrление к ближайшему;

- rz - окруrление к нулю;

- ru - окруrление вверх;

- rd - окруrление вниз.

Таблица 2.2 Математические floаt - функции в CUDA

Функция	Значение
_fadd_ [rn, rz, ru,rd] (х, у)	Сложение, никоrда не переводимое в команду FMAD
_fmul_ [rn, rz, ru,rd] (х, у)	Умножение, никоrда не переводимое в команду FMAD
_fmaf_ [rn, rz, rU,rd] (х, у, z )	(xЧy)+z
_frcp_ [rn, rz, ru,rd] (х)	1/x
_fsqrt_ [rn, rz, ru,rd] (x)
_fdiv_ [rn, rz, ru,rd] (х, у)	x/y
_fdividef_(х, у)	х/у, но если 2126 < У < 2128, то 0
_expf ( х )	ex
_ехр 10f ( х )	10x
_logf ( х )	log(x)
_log2f ( х )	log2(x)
_log10f ( х )	log10(x)
_sinf ( х )	sinx
_cosf ( х )	cosx
_sincosf ( х, sptr, cptr )	*sptr = sin(x); *cptr = cos(x)
_tanf ( х )	tanx
_powf ( х, у )	xy
_int_as_float ( х )	32 бита, образующие целочисленное значение, интерпретируются как flоаt значение. Так, значение 0хС000000 будет переведено в -2.0f
_float_as_int ( х )	32 бита, образующие flоаt значение, интерпретируются как целочисленное значение. Так, значение 1.0f будет переведено в - 0x3F80000
_ saturate ( х )	min(0,max(1,х))
_float_to_int [rn, rz, ru,rd] (х)	Приведение flоаt значения к целочисленному значению с заданным округлением
_float_to_uint [rn, rz, ru,rd] (х)	Приведение flоаt значения к беззнаковому целочисленному значению с заданным окруrлением
_int_to_float [rn, rz, ru,rd] (х)	При ведение целочисленного значения к flоаt значению с заданным округлением
_uint_to_float [rn, rz, ru,rd] (х)	Приведение беззнаковоrо целочисленноrо значения к flоаt значению с заданным округлением
_float_to_II [rn, rz, ru,rd] (х)	Приведение flоаt значения к 64 битовому целочисленному значению с заданным округлением
_float_to_uII [rn, rz, ru,rd] (х)	При ведение flоаt значения к 64-битовому беззнаковому целочисленному значению с заданным округлением

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



Таблица 2.3 Целочисленныe функции в CUDA

Функция	Значение
_[u]muI24 ( х, у )	Вычисляет произведение младших 24 бит целочисленных параметров х и у, возвращает младшие 32 бита результата. Старшие 8 бит аргументов игнорируются
_[u]mulhi ( х, y )	Значение возвращает старшие 32 бита произведения целочисленных операндов х и y
_[u]mul64hi ( х, y )	Вычисляет произведение 64 битовых целых чисел и возвращает младшие 64 бита этого произведения
_[u]sad ( х, y, z )	Возвращает z + |x - у|
_clz ( х )	Возвращает целое число от 0 до 32 включительно последовательных нулевых битов для целочисленноrо параметра х, начиная со старших бит
_clzll ( х )	Возвращает целое число от 0 до 64 включительно последовательных нулевых битов для целочисленного 64-битовоrо параметра х, начиная со старших бит
_ffs ( х )	Возвращает позицию первоrо (наименее значимоrо) единичнoгo бита для apгумента х. Если х равен нулю, то возвращается нуль
_ffsII ( х )	Возвращает позицию первоrо (наименее значимого) единичного бита для целочисленноrо 64-битового аргумента х. Если х равен нулю, то возвращается нуль
_popc ( х )	Возвращает число бит, которые равны единице в двоичном представлении 32-битовоrо целочисленного aргумента х
_popcII ( х )	Возвращает число бит, которые равны единице в двоичном представлении 64-битовоrо целочисленноrо apгумента х
_brev ( х )	Возвращает число, полученное перестановкой (то есть биты в позициях k и 31-k меняются местами для всех k от 0 до 31) битов исходноrо 32-битовоrо целочисленного apгумента х
_brevll ( х )	Возвращает число, полученное перестановкой (то есть биты в позициях k и 63-k меняются местами для всех k от 0 до 63) битов исходноrо 64-битовоrо целочисленного apгумента х

5. Высокоуровневая библиотека управления GPU

Высокоуровневая библиотека управления GPU (CUDA runtime API, cudart) реализована с использованием низкоуровневой библиотеки (libcuda). Используется интерфейс C++ (именование функций, возможность использовать умолчания). Инициализация происходит неявно при вызове первой функции. Это необходимо учитывать при измерении производительности и разборе ошибок, а также при многозадачно и многопоточной обработке.

Большинство функций API имеет тип cudaError_t, что позволяет обрабатывать ошибки (cudaSuccess - отсутствие ошибок). С помощью функции cudaGetErrorString (cudaError_t) можно получить текстовое сообщение об ошибке. При обработке ошибок необходимо учитывать асинхронность многих функций API. Последнее значение кода возврата можно получить с помощью функции cudaGetLastError.

Управление GPU:

- cudaGetDeviceCount (int *) - узнать число устройств; при отсутствии GPU возвращает число 1 (эмулятор GPU)

- cudaSetDevice (int) - выбрать устройство для использования; необходимо сделать до первого запуска или обращения к устройству; нумерация с 0

- cudaGetDevice (int *) - узнать номер используемого GPU

- cudaGetDeviceProperties (cudaDeviceProp*, dev) - узнать характеристики указанного устройства: имя, объём глобальной памяти, объём разделяемой паяти на MP, число регистров, размер порции, максимальное число потоков на блок, максимальная размерность блока, максимальная размерность сетки, объём памяти констант, версия GPU, частота GPU (непонятно какого функционального блока), количество MP, возможность асинхронной пересылки

- cudaChooseDevice (int*, const struct cudaDeviceProp*) - выбор наиболее подходящего устройства

Запуск ядра:

- cudaConfigureCall (dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, int tokens = 0) - задать размерность сетки и блока процессов для будущего запуска; помещаются в стек для дальнейшего использования в cudaLaunch

- cudaSetupArgument (void* arg, size_t count, size_t offset) - задать аргументы запуска

- cudaLaunch (T entry) - запустить ядро entry, которое задаётся либо функцией __global__, либо текстовой строкой с именем такой функции; размерности и параметры извлекаются из стека, создаваемого cudaConfigureCall

Управление памятью и пересылки (размер в байтах):

- cudaMalloc (void**, size_t) - выделить линейный массив из глобальной памяти

- cudaMallocPitch (void**, size_t* pitch, size_t widthInBytes, size_t height) - выделить двумерный массив из глобальной памяти (заботится о выравниваниях и прочей оптимизации)

- cudaFree (void *)

- cudaMallocHost (void**, size_t) - выделить из ОП хоста в режиме блокировки от подкачки; это сильно ускоряет обмен с GPU; асинхронные операции копирования допустимы только с блокированной памятью

- cudaFreeHost (void *)

- cudaMemset (void*, int, size_t) - заполнить память GPU константой

- cudaMemcpy (void* dst, const void* src, size_t count, enum cudaMemcpyKind kind) - копирование, где cudaMemcpyKind определяет направление пересылки (перекрытие областей запрещено):

- cudaMemcpyHostToHost

- cudaMemcpyHostToDevice

- cudaMemcpyDeviceToHost

- cudaMemcpyDeviceToDevice

- cudaMemcpyAsync (void* dst, const void* src, size_t count, enum cudaMemcpyKind kind, cudaStream_t) - асинхронное копирование, только для блокированной памяти

- cudaMemcpyToSymbol (const T& symbol, const void* src, size_t count, size_t offset, enum cudaMemcpyKind kind) - копирование count байт с адреса src на адрес, определяемый переменной symbol, со смещением offset; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной; cudaMemcpyKind определяет направление пересылки (перекрытие областей запрещено):

- cudaMemcpyHostToDevice

- cudaMemcpyDeviceToDevice

- cudaMemcpyToSymbol (void* dst, const T& symbol, size_t count, size_t offset, enum cudaMemcpyKind kind) - копирование count байт на адрес dst с адреса, определяемого переменной symbol, со смещением offset; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной; cudaMemcpyKind определяет направление пересылки (перекрытие областей запрещено):

- cudaMemcpyDeviceToHost

- cudaMemcpyDeviceToDevice

- cudaGetSymbolAddress (void** devPtr, const T& symbol) - получить указатель на переменную в глобальной памяти GPU; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной

- cudaGetSymbolSize (size_t* size, const T& symbol) - получить размер переменной в глобальной памяти GPU; переменная глобальной памяти или памяти констант задаётся либо переменной с адресом в GPU, либо текстовой строкой с именем переменной

- cudaMallocArray - выделение памяти под CUDA-массив

- cudaFreeArray

- cudaMemset2D

- cudaMemcpy2D, cudaMemcpy2DAsync, cudaMemcpyToArray, cudaMemcpyToArrayAsync, cudaMemcpy2DToArray, cudaMemcpy2DToArrayAsync, cudaMemcpyFromArray, cudaMemcpy2DFromArray, cudaMemcpyArrayToArray, cudaMemcpy2DArrayToArray - копирование матриц

- cudaMalloc3D, cudaMalloc3DArray

- cudaMemset3D

- cudaMemcpy3D

В связи с асинхронностью запуска ядра введены дополнительные средства управления задачами (stream). Задача - это последовательность операций, выполняемых строго по очереди. Относительный порядок операций, принадлежащих различным задачам, не определён. По умолчанию, операция относится к задаче 0. При необходимости организовать выполнение нескольких паралелльных задач необходимо определить задачи и указывать их номера в командах запуска ядра и копирования из памяти в память. Установка переменной окружения CUDA_LAUNCH_BLOCKING в 1 блокирует асинхронное выполнение.

Функции синхронизации:

- cudaStreamCreate (cudaStream_t*)

- cudaStreamQuery (cudaStream_t) - проверить, завершились ли все операции задачи: cudaErrorNotReady (единственный способ синхронизации без 100% загрузки CPU, однако уменьшает эффективность работы)

- cudaStreamSynchronize (cudaStream_t) - ждать завершения всех операций задачи (реализовано в виде непрерывного опроса устройства, что полностью загружает CPU)

- cudaStreamDestroy (cudaStream_t) - ждать завершения всех операций задачи и завершить задачу

- cudaThreadSynchronize (void) - ждать завершения всех операций всех задач (реализовано в виде непрерывного опроса устройства, что полностью загружает CPU)

- cudaThreadExit (void) - ждать завершения всех операций всех задач, освободить все ресурсы; последующее обращение к одной из функций API повторно инициализирует API (глючит, не советую)

Функции работы с событиями (event) позволяют измерять время выполнения операций:

- cudaEventCreate (cudaEvent_t*) - создать описание события

- cudaEventRecord (cudaEvent_t, CUstream) - зафиксировать момент наступления события после выполнения всех операций указанной задачи; сама функция является асинхронной, так что для извлечения точного значения необходимо использовать cudaEventQuery или cudaEventSynchronize; извлечь информацию для не нулевой задачи нечем

- cudaEventQuery (cudaEvent_t) - завершена ли запись наступления события (cudaErrorNotReady)

- cudaEventSynchronize (cudaEvent_t) - ждать окончания записи наступления события

- cudaEventElapsedTime (float*, cudaEvent_t, cudaEvent_t) - вычислить интервал времени между событиями (в мс), точность - 0.5 мкс; не работает для событий из не основной (не нулевой) задачи

- cudaEventDestroy (cudaEvent_t) - удалить описание события

Функции работы с текстурами:

- texture < type, dim, tex_type> g_TexRef; - объявление текстуры

(type - тип хранимых переменных; dim - размерность текстуры (1, 2, 3); tex_type - тип возвращаемых значений)

cudaReadModeNormalizedFloat; - считывание переменной

cudaReadModeElementType. - считывание элемента текстуры

Кроме того, для более полного использования возможностей текстурной памяти можно задать описание канала:

struct cudaChannelFormatDesc {int x, y, z, w;

enum cudaChannelFormatKind f;};

Задает формат возвращаемого значения;

- int x, y, z, w; - число [0,32] проекция исходного значения по битам

- cudaChannelFormatKind - тип возвращаемого значения:

- cudaChannelFormatKindSigned - знаковые;

- intocudaChannelFormatKindUnsigned - беззнаковые ;

- intocudaChannelFormatKindFloat - float.

В CUDA существуют два типа текстур линейная и cudaArray:
Линейная	cudaArray
Можно использовать обычную глобальную память. Ограничения: только для одномерных массивов; нет фильтрации; доступ по целочисленным координатам; обращение по адресу вне допустимого диапазона возвращает ноль. Доступ: tex1Dfetch(tex, int)	Позволяет организовывать данные в1D/ 2D/3D массивы данных вида: 1/2/4 компонентные векторы; 8/16/32 bit signed/unsigned integers; 32 bit float; 16 bit float (driver API). Доступ по семейству функций: tex1D() / tex2D() / tex3D()

6. Линейная текстурная память

Линейная текстурная память не обладает никакими особыми свойствами кроме кэша, однако уже этого иногда достаточно для значительного ускорения работы программы. Как уже сказано выше, линейная текстурная память хранится в обычной глобальной памяти и не требует особых функций копирования. Для включения возможностей линейной текстурной памяти требуется "привязать" обычный массив в глобальной памяти к объявленной текстуре.

Привязывание линейного массива:

cudaBindTexture(size_t shift, texref tex,&src, size_t size));

- shift - смещение при привязки к массиву (к одному массиву можно привязать несколько тектсур)

- tex - объявленная текстура

- src - массив в глобальной памяти, к которому привязывается текстура

- size - размер привязываемой области в байтах

Привязывание "двумерного массива" (в глобальной памяти он все равно хранится как линейный, и обращение к нему идет по одной координате):

- cudaBindTexture2D(size_t shift, texref tex, &src, &channelDesc, int width, int height, int pitch);

- shift - смещение при привязке к массиву (к одному массиву можно привязать несколько текстур);

- tex - объявленная текстура;

- src - массив в глобальной памяти, к которому привязывается текстура;

- channelDesc - описание канала;

- width - ширина двумерного массива;

- height - высота двумерного массива;

- pitch - смещение каждой строки.

После окончания работы с текстурой её надо "отвязать":

- cudaUnbindTexture(texref tex);

Все вышеприведенные функции вызываются с хоста. На устройстве используется функция, которая достает значение из текстуры:

- tex1Dfetch(texRef tex, int index);

tex - объявленная текстура;

index - индекс вынимаемого значения в линейном массиве.

В качестве входного демонстрационного примера в Лабораторной работе №4 создан проект LR4_1, в котором одномерный массив тестовых данных загружается в текстурную память, и ядро GPU вычисляет произведение элемента массива на его индекс, заполняя выходной одномерный массив.

7. CudaArray текстурная память

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

- Нормализация координат (перевод [W, H] => [0,1]).

- Преобразование координат:

- Clamp - координата обрезается по границе;

- Wrap - координата заворачивается.

- Фильтрация (при обращении по float координате):

- Point - возвращается ближайшее заданное значение;

- Linear - производится билинейная интерполяция.

Для использования cudaArray текстурной памяти требуется объявить переменную-указатель на cudaArray:

cudaArray * a;

Затем необходимо выделить память под данные на видеокарте:

- cudaMallocArray(struct cudaArray **arrayPtr, const struct cudaChannelFormatDesc *channelDesc, size_t width, size_t height);

- arrayPtr - указатель на cudaArray;

- channelDesc - описание канала;

- width - ширина массива;

- height - высота массива.

Затем скопировать в выделенную память данные:

- cudaMemcpyToArray(struct cudaArray * dst, size_t wOffset, size_t hOffset, const void * src, size_t count, enum cudaMemcpyKind kind);

- arrayPtr - указатель на cudaArray;

- wOffset - смещение по горизонтали при привязке к массиву;

- hOffset - смещение по вертикали при привязке к массиву;

- src - массив в памяти хоста, к который копируется;

- count - размер данных в байтах;

- kind - направление копирования.

После того как данные скопированы, можно осуществлять привязку cudaArray массива к текстуре:

- cudaBindTextureToArray(const struct textureReference *tex, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc);

- tex - объявленная текстура;

- array - массив в cudaArray, к которому привязывается текстура;

- channelDesc - описание канала.

На устройстве используются функции, которые достают значение из текстуры:

tex1D (texRef tex, float x);

- tex - объявленная текстура;

- x - индекс вынимаемого значения в линейном массиве.

tex2D (texRef tex, float x, float y);

- tex - объявленная текстура;

- x, y - индексы вынимаемого значения в двухмерном массиве.

tex3D (texRef tex, float x, float y, float z);

- tex - объявленная текстура;

- x, y, z - индексы вынимаемого значения в трехмерном массиве.

После использования необходимо отвязать текстуру точно так же как и линейную.

Функции взаимодействия с OpenGL:

Так как ресурсы CUDA физически расположены в памяти графического процессора, то для многих задач было бы удобно создать один объект в графическом API, но иметь возможность отобразить eгo в пространство памяти CUDA.

Создание буферноrо объекта в OpenGL

Расширение vertex_buffer_object добавляет в OpenGL функциональность по использованию буферных объектов. Это расширение определяет интерфейс, который позволяет разного рода данным (прежде вceгo массивам вершинных атрибутов) эффективно храниться в графической памяти устройства.

Данные инкапсулируются внутри буферных объектов, которые по сути являются массивами байтов. Вводится АРI для создания, чтения и записи, используя внутренние механизмы GL или указатели на память CPU.

Последнее получило название отображение буфера (от англ. bufferтapping). Коrда приложение отображает буфер, оно получает указатель на память. Коrда приложение завершило чтение или запись по данному указателю, оно обязано отпустить (в анrл. принято rоворить «иптар») указатель, прежде чем оно вновь начнет использовать отображенный ресурс (это необходимо сделать, чтобы драйвер понял, что ресурс закончил обновление и данные можно копировать из системной памяти в память GPU). Как правило, при отображении приложение может указать флаги для операции, которую собирается производить (чтение, запись, или чтение и запись), что позволяет избегать ненужного копирования ресурсов.

Создание буфера происходит аналоrично созданию друrих ресурсов GL.

берем ero на примере.

При создании буфера сначала создается идентификатор буфера:

void glGenBuffers( GLsizei n, GLuint * buffers ) ;

Идентификаторами ресурсов в GL служат целые числа. Создание объекта caмoгo по себе не означает eгo инициализацию, для этого надо установить этот буфер в качестве текущеrо:



void glBindBuffer( GLenum target, GLuint buffer );

и проинициализировать с помощью следующей функции:

void glBufferData(GLenum target, Gluint size, const GLvoid * data, GLenum usage);

Взаимодействие с OpenGL требует установки СUDА - устройства:

cudaError_t cudaGLSetGLDevice ( int device );

перед любым вызовом.

Прежде чем использовать объект, eгo необходимо зарегистрировать в среде CUDA:

cudaError_t cudaGLRegisterBufferObject(Gluint bufObj);

После регистрации ресурса eгo можно отобразить в пространство памяти

CUDA (получить указатель, который может быть далее передан в ядро) или отпустить ресурс, используя одну из нижеперечисленных функций:

cudaError_t cudaGLMapBufferObject( void ** devPtr, GLuint bufObj);

cudaError_t cudaGLUnmapBufferObject(GLuint bufObj);

По окончании работы стоит снять регистрацию с объекта, чтобы CUDA могла освободить выделенные под нeгo ресурсы.



cudaError_t cudaGLUnregisterBufferObject(GLuint bufObject);

Атомарные операции:

CUDA для GPU, начиная с compute capability 1.1 и выше, поддерживает aтомарные операции над глобальной и разделяемой памятью. Атомарность заключается в том, что гарантируется корректность выполнения операции для случая многих нитей, одновременно пытающихся ее выполнить.

Атомарные операции над 64-битовыми словами и словами в разделяемой памяти поддерживаются только для GPU с compute capability 1.2 и выше.

Обратите внимание, что все атомарные операции (за исключением аtomicExch) работают только с целыми числами.

8. Атомарные арифметические операции

Наиболее часто используемыми атомарными арифметическими операциями являются аtomicAdd и atomicSub, служащие для увеличения или уменьшения вeличины на заданное значение. При этом функция возвращает старое значение.

Обратите внимание, что atomicAdd поддерживает операции над 64-битовыми вeличинами, но только расположенными в глобальной памяти.

int atomicAdd ( int * addr, int value ) ;

unsigned int atomicAdd ( unsigned int * addr, unsigned int value ) ;

unsigned long long atomicAdd ( unsigned long long * addr, unsigned long long value ) ;

int atomicSub. ( int * addr, int value ) ;

unsigned int atomicSub ( unsigned int * addr, unsigned int value ) ;



Операция аtomicExch осуществляет атомарный обмен значениями передаваемое значение записывается по указанному адресу, а предыдущее значение возвращается. При этом подобный обмен происходит как одна транзакция, то есть ни одна нить не может вклиниться между шагами этого обмена. Операции над 64-битовыми значениями поддерживаются только для глобальной памяти.

int atomicExch (int * addr, int value );

unsigned int atomicExch (unsigned int * addr, unsigned int value);

unsigned long long atomicExch (unsigned long long * addr, unsigned long long value );

float atomicExch (float * addr, float value );

Следующие две операции сравнивают значение по адресу с переданным значением, записывают минимум/максимум из этих двух значений по заданному aдресу и возвращают предыдущее значение, находившееся по адресу. Все эти шаги выполняются атомарно, как одна транзакция.

int atomicMin (int * addr, int value);

unsigned int atomicMin (unsigned int * addr, unsigned int value);

int atomicMax (int * addr, int value) ;

unsigned int atomicMax (unsigned int * addr, unsigned int value) ;

Операция atomicInc читает слово по заданному адресу и сравнивает eгo с переданным значением. Если прочитанное слово больше, то по адресу записывается нуль, иначе значение по адресу увеличивается на единицу. Возвращается старое значение.



unsigned int atomicInc ( unsigned int * addr, unsigned int value );

Операция аtomicDec читает слово по переданному адресу. Если прочитанное значение равно нулю или больше переданного значения, то записывает по адресу переданное значение, иначе уменьшает значение по адресу на единицу. Возвращается старое значение.

unsigned int atomicDec ( unsigned int * addr, unsigned int value );

Следующая функция (CAS - Compare And Swap) читает старое 32 или 64-битовое значение по переданному адресу и сравнивает eгo с параметром сompare. В случае совпадения по переданному адресу записывается значение параметра value, иначе значение по адресу не изменяется. Возвращается всегда старое прочитанное значение.

int atomicCAS ( int * addr, int compare, int value );

unsigned int atomicCAS ( unsigned int * addr, unsigned int compare, unsigned int value );

unsigned long long atomicCAS ( unsigned int * addr, unsigned long long compare, unsigned long long value );

9. Атомарные побитовые операции

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

int atomicAnd int * addr, int value ) i

unsigned int atomicAnd unsigned int * addr, unsigned int value ) i

int atomicOr int * addr, int value ) ;

unsigned int atomicOr unsigned int * addr, unsigned int value ) i

int atomicXor int * addr, int value ) ;

unsigned int atomicXor unsigned int * addr, unsigned int value ) ; .

Проверка статуса нитей warp'a

Начиная с compute capability 1.2, поддерживаются следующие две атомарные операции, которые выполняют сравнение переданного значения с нулем и позволяют выяснить, для всех ли нитей warp'a получено истинное значение и есть ли ненулевое значение хотя бы одно для одной нити warp'a.

int_all (int predicate );

int_аnу (int predicate );

file:///C:/Program%20Files/NVIDIA%20GPU%20Computing%20Toolkit/CUDA/v7.0/doc/html/parallel-thread-execution/index.html#ptx-isa-version-4-2

10. Синтаксис

Программы PTX являются набором исходных текстов модулей (файлов). Модули PTX имеют синтаксис в стиле в языке ассемблера с кодами команд и операндами: псевдо-операции, использующие символы и адресацию. В ptxas компилятор оптимизирует и собирает исходные модули PTX, создавая соответствующие двоичные объектные файлы.

Формат источника

Модули источника - текст в формате ASCII. Строки разделяются символом новой строки. Все пробелы эквивалентны; пробелы игнорируются для использования его в качестве разделителя.

C-препроцессор CPP может быть использован для обработки исходных модулей PTX. Строки, начинающиеся с # - директивы препроцессора. Ниже приведены общие директивы препроцессора:

#include, #define, #if, #ifdef, #else, #endif, #line, #file

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

11. Комментарии

Комментарии в PTX используют C / C ++ синтаксис: использование скобок / * и * /, когда комментарии могут занимать несколько строк, и //, когда комментарий простирается до следующего символа новой строки.

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

Команды начинаются с дополнительной метки и заканчиваться точкой с запятой.

Примеры

.reg .b32 r1, r2;

.global .f32 array[N];

start: mov.b32 r1, %tid.x;

shl.b32 r1, r1, 2; // shift thread id by 2 bits

ld.global.b32 r2, array[r1]; // thread[tid] gets array[tid]

add.f32 r2, r2, 0.5; // add 1/2

12. Директивы

Ключевые слова директив начинаются с точки, поэтому невозможен конфликт с пользовательскими идентификаторами. Директивы в PTX приведены в таблице 1.

Таблица 1 Директивы PTX

.address_size	.file	.minnctapersm	.target
.align	.func	.param	.tex
.branchtargets	.global	.pragma	.version
.callprototype	.loc	.reg	.visible
.calltargets	.local	.reqntid	.weak
.const	.maxnctapersm	.section
.entry	.maxnreg	.shared
.extern	.maxntid	.sreg

13. Операторы

Опрераторы состоят из кода операции, разделённый запятой со списком из нуля или более операндов, и завершается точкой с запятой. Операнды могут быть регистровыми переменными, константами, адресными выражениями, или именами меток. Инструкция имеет дополнительный защитный предикат, который управляет условием выполнения.Защитный предикат следует дополнительной метке и предопределяет код операции, записывается как @p, где р - регистр предикат.Защитный предикат может быть отрицательным, записанный как @!С.

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

Ключевые слова операторов перечислены в таблице 2. Все маркеры зарезервированы в PTX.



Таблица 2 Зарезервированные ключевые слова операторов

abs	div	or	sin	vavrg2, vavrg4
add	ex2	pmevent	slct	vmad
addc	exit	popc	sqrt	vmax
and	fma	prefetch	st	vmax2, vmax4
atom	isspacep	prefetchu	sub	vmin
bar	ld	prmt	subc	vmin2, vmin4
bfe	ldu	rcp	suld	vote
bfi	lg2	red	suq	vset
bfind	mad	rem	sured	vset2, vset4
bra	mad24	ret	sust	vshl
brev	madc	rsqrt	testp	vshr
brkpt	max	sad	tex	vsub
call	membar	selp	tld4	vsub2, vsub4
clz	min	set	trap	xor
cnot	mov	setp	txq
copysign	mul	shf	vabsdiff
cos	mul 24	shfl	vabsdiff2, vabsdiff4
cvt	neg	shl	vadd
cvta	not	shr	vadd2, vadd4

14. Идентификаторы

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

followsym: [a-zA-Z0-9_$]

identifier: [a-zA-Z]{followsym}* | {[_$%]{followsym}+



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

Многие языки высокого уровня, такие как C и C ++ придерживаются аналогичных правил для имен идентификаторов, за исключением того знака процента (не используеся). PTX позволяет использовать знак процента в качестве первого символа идентификатора.Знак процента может быть использован, чтобы избежать конфликтов имен, например, между именами переменных, определяемых пользователем и именами сгенерированного компилятором.

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

Таблица 3 Предопределённые идентификаторы

%clock	%laneid	%lanemask_gt	%pm0 ,...,pm7
%clock64	%lanemask_eq	%nctaid	%smid
%ctaid	%lanemask_le	%ntid	%tid
%envreg<32>	%lanemask_lt	%nsmid	%warpid
%gridid	%lanemask_ge	%nwarpid	%WARP_SZ

15. Константы

PTX поддерживает целые и с плавающей точкой константы и константные выражения. Эти константы могут быть использованы при инициализации данных и в качестве операндов инструкции. Тип проверки правил остаётся одинаковым для целого, с плавающей точкой, а также типы битового размера. Для предикат типа данных и инструкций, целочисленные константы могут и интерпретируются как в С, т.е. нулевые значения являются ложными и ненулевые значения True.

16. Целые константы

Целые константы являются 64-разрядными числами знаковыми или беззнаковыми, то есть, каждая целая константа имеет тип .s64 или .u64. Подпись знаковая/беззнаковая необходима, чтобы правильно оценить константные выражения, содержащие такие операции, как разделение и сравнение, где поведение операции зависит от типов операндов.

При использовании в командах или при инициализации данных, каждое целое константное изменяют до соответствующего размера в зависимости от типа данных или команд, использующих её.

Целые литералы могут быть записаны в десятичной, шестнадцатеричной, восьмеричной или двоичной системе счисления. Синтаксис такой же как и в С. За целыми литералами следует знак U, чтобы указать что они беззнаковые

hexadecimal literal: 0[xX]{hexdigit}+U?

octal literal: 0{octal digit}+U?

binary literal: 0[bB]{bit}+U?

decimal literal {nonzero-digit}{digit}*U?

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

17. Константы с плавающей точкой

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

С плавающей точкой литералы могут быть записаны с помощью дополнительной десятичной точки и дополнительного показателя. В отличие от C и C ++, нет суффикса, определяющего размер; литералы всегда представлены в 64-битном формате двойной точности.

PTX включает в себя второе представление константы с плавающей - шестнадцатеричная константа. Постоянная начинается с 0d или 0D затем 16 шестнадцатеричных цифр. Чтобы указать значения одинарной точности с плавающей точкой, постоянная начинается с 0f или 0F затем 8 шестнадцатеричных цифр.

0[fF]{hexdigit}{8} // single-precision floating point

0[dD]{hexdigit}{16} // double-precision floating point

Пример

mov.f32 $f3, 0F3f800000; // 1.0

Предикатные константы

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

18. Константные выражения

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

Постоянные выражения формируются из постоянных литералов, унарного плюса и минуса, основных арифметических операторов (сложение, вычитание, умножение, деление), операторов сравнения, условного тройной оператор (?:) и скобок. Целое постоянные выражения также позволяют унарное логическое отрицание, побитовое дополнение (~), остаток (%), операторы сдвига (<< и >>), операторы бит-типа (&, |, и ^) (!), и логические операторы ( &&, ||).

Постоянные выражения в PTX не поддерживает приведение типов между целыми и с плавающей точкой.

Константные выражения имеют тот же приоритет операций как в C. В таблице 4 приведено старшинство операторов. Приоритет операторов является самым высоким для унарных и уменьшается с каждой строкой в таблице. Операторы в одной и той же строке, имеют одинаковый приоритет.

Таблица 4 Приоритет операции

Тип	Символы операторов	Название опреаторов
Основной	()	скобки
Унарный	+- ! ~	плюс, минус, отрицание, дополнение
	(.s64)(.u64)	типы
Бинарный	*/ %	умножение, деление, остаток
	+-	сложение, вычитание
	>> <<	сдвиги
	< > <= >=	сравнение
	== !=	равно, не равно
	&	битовое И
	^	битовое ИСКЛ-ИЛИ
	|	битовое ИЛИ
	&&	логическое И
	||	логическое ИЛИ
Тройной	?:	условный

19. Пространства состояний, типов и переменных

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

Пространство состояний

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

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

Таблица 6 Пространства состояний

Имя	Значение
.reg	Быстрые регистры
.sreg	Специальные регистры. Только для чтения; предопределенные; платформенные.
.const	Общий, только для чтения памяти.
.global	Глобальная память, общая для всех потоков.
.local	Локальная память, индивидуальная для каждого потока
.param	Параметры ядра, определенные в сетке; или функциональные или локальные параметры, определенные для каждого потока.
.shared	Общая адресуемая память между потоками
.tex	Глобальная текстурная память (не рекомендуется)



Таблица 7 Свойства пространств состояний

Имя	Адресуемость	Инициализуемость	Доступ	Принадлежность
.reg	No	No	R/W	per-thread
.sreg	No	No	RO	per-CTA
.const	Yes	Yes	RO	per-grid
.global	Yes	Yes	R/W	Context
.local	Yes	No	R/W	per-thread
.param (as input to kernel)	Yes2	No	RO	per-grid
.param (used in functions)	Restricted	No	R/W	per-thread
.shared	Yes	No	R/W	per-CTA
.tex	No4	Yes, via driver	RO	Context

20. Типы

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

Таблица 8 Спецификаторы стандартных типов

Базовый тип	Спецификатор
Целочисленный знаковый	.s8, .s16, .s32, .s64
Целочисленный беззнаковый	.u8, .u16, .u32, .u64
С плавающей точкой	.f16, .f32, .f64
Битовый	.b8, .b16, .b32, .b64
Предикатный	.pred

21. Типы текстурных шаблонов и поверхностей

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

- определение переменных в глобальных (модульных) рамках и в списках параметров начального ядра

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