Алматы 2015 Almaty
жүктеу/скачать 19,96 Mb. Pdf көрінісі
|
- Бұл бет үшін навигация:
- PTX – ассемблер
Расширения языка С++ Основной программой среды разработки является компилятор nvcc, который управляет всей цепочкой преобразований. Исходная программа может быть на языке C (.c), C++ (.cc, .cxx, .cpp), расширении языка C (.cu; .cup для препроцессированных файлов) [3]. Проrраммы для CUDA (соответствующие файлы обычно имеют расширение .сu) пишутся на «расширенном С и компилируются при помощи команды nvcc. Вводимые в CUDA расширения языка С состоят из: – спецификаторов функций, показывающих, где будет выполняться функция и откуда она может быть вызвана; – спецификаторов переменных, задающих тип памяти, используемый для данных переменных; – директивы, служащей для запуска ядра, задающей как данные, так и иерархию нитей; – встроенных переменных, содержащих информацию о текущей нити; – runtime, включающей в себя дополнительные типы данных. В CUDA используется следующие спецификаторы функций (таблица 1).
Таблица 1 Спецификаторы функций и переменных:
Спецификатор Функция выполняется на
_device_ device (GPU) device (GPU) _global_ device (GPU) host (CPU) _host_ host (CPU) host (CPU)
В язык добавлены 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. Также добавлен тип dim3, используемый для задания размерности. В язык добавлены следующие специальные переменные: – gridDim – размер сетки (имеет тип dimз); – blockDim – размер блока (имеет тип dim3); – blockldx – индекс текущеrо блока в сетке (имеет тип uint3); – threadIdx – индекс текущей нити в блоке (имеет тип uint3); – warpSize – размер warp'a (имеет тип int). Для запуска ядра на GPU используется следующая конструкция: kernelName << Здесь kemelName это имя (адрес) соответствующей _global_ функции. Через Dg обозначена переменная (или значение) типа dim3, задающая размерность и размер сетки (в блоках). Переменная (или значение) Db – типа dim3, задает размерность и размер блока (в нитях). Необязательная переменная (или значение) Ns типа size_t задает дополнительный объем разделяемой памяти в байтах, которая должна быть динамически выделена каждому блоку (к уже статически выделенной разделяемой памяти), если не задано, то используется значение О. Переменная (или значение) S типа cudaStream_t задает поток (CUDA stream), в котором должен произойти вызов, по умолчанию используется поток 0. Через args обозначены арrументы вызова функции kenelName (их может быть несколько ). CUDA поддерживает все функции из стандартной библиотеки языка С. Однако при этом следует иметь в виду, что большинство стандартных математических функций используют числа с двойной точностью (double). CUDA предоставляет также специальный набор функций пониженной точности, но обеспечивающих большее быстродействие. В таблице 2 приведены основные f1оаt – функции и их оптимизированные вepсии пониженной точности. Для ряда функций можно задать требуемый способ округления. Используемый способ задается при помощи одноrо из следующих суффиксов: – rn – окруrление к ближайшему; – rz – окруrление к нулю; – ru – окруrление вверх; – rd – окруrление вниз. 744 Таблица 2 Функции для работы с числами с плавающей точкой
Функция Значение _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_(х, у) х/у, но если 2 126
< У < 2 128
, то 0 _expf ( х ) e x
_ехр 10f ( х ) 10 x _logf ( х ) log(x)
_log2f ( х ) log
2 (x)
_log10f ( х ) log
10 (x)
_sinf ( х ) sinx
_cosf ( х ) cosx
_sincosf ( х, sptr, cptr ) *sptr = sin(x); *cptr = cos(x) _tanf ( х ) tanx
_powf ( х, у ) x y _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-битовому беззнаковому целочисленному значению с заданным округлением
Кроме ряда оптимизированных функций для работы с числами с плавающей точкой, также есть ряд быстрых функций для работы с целыми числами, приводимый в таблице 3.
745 Таблица 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гумента х Для того, чтобы подключить к работе графический процессор и использовать его, существуют – cudaGetDeviceCount (int *) - узнать число устройств; при отсутствии GPU возвращает число 1 (эмулятор GPU); – cudaSetDevice (int) - выбрать устройство для использования; необходимо сделать до первого запуска или обращения к устройству; нумерация с 0. – cudaGetDevice (int *) - узнать номер используемого GPU; – cudaGetDeviceProperties (cudaDeviceProp*, dev) - узнать характеристики указанного устройства: имя, объём глобальной памяти, объём разделяемой паяти на MP, число регистров, размер порции, максимальное число потоков на блок, максимальная размерность блока, максимальная размерность сетки, объём памяти констант, версия GPU, частота GPU (непонятно какого функционального блока), количество MP, возможность асинхронной пересылки; – cudaChooseDevice (int*, const struct cudaDeviceProp*) - выбор наиболее подходящего устройства; В состав программы входит ядро, реализуемая через центральный процессор(CPU) и необходимая для распределения потоков в GPU и имеет собственные функции управления: – cudaConfigureCall (dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, int tokens = 0) - задать 746 размерность сетки и блока процессов для будущего запуска; помещаются в стек для дальнейшего использования в 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. При необходимости организовать выполнение нескольких паралелльных задач 747 необходимо определить задачи и указывать их номера в командах запуска ядра и копирования из памяти в память. Установка переменной окружения 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 (глючит, не советую)
Для анализа программы и написания аппаратно–зависимых программ используется специальный язык PTX – ассемблер. Программы PTX являются набором исходных текстов модулей (файлов). Модули PTX имеют синтаксис в стиле в языке ассемблера с кодами команд и операндами: псевдо- операции, использующие символы и адресацию [2]. Модули источника – текст в формате ASCII. Строки разделяются символом новой строки. Все пробелы эквивалентны; пробелы игнорируются для использования его в качестве разделителя. C-препроцессор CPP может быть использован для обработки исходных модулей PTX. Строки, начинающиеся с # – директивы препроцессора. PTX является чувствительным к регистру и использует нижний регистр для ключевых слов. Каждый модуль PTX должен начинаться с .version директивы с указанием языковой версии PTX, с последующим .target директивы с указанием целевой предполагаемой архитектуры . Опрераторы состоят из кода операции, разделённый запятой со списком из нуля или более операндов, и завершается точкой с запятой. Операнды могут быть регистровыми переменными, константами, адресными выражениями, или именами меток. Инструкция имеет дополнительный защитный предикат, который управляет условием выполнения. Защитный предикат следует дополнительной метке и предопределяет код операции, записывается как @p, где р – регистр предикат. Защитный предикат может быть отрицательным, записанный как @!С. Пользовательские идентификаторы используют расширенным правилам языка C++: они либо начинаться с буквы, за которой следует нуль или более буквы, цифры, подчеркивание или символ доллара; или они начинаются с символа подчеркивания, доллара, или процента, за которым следуют один или более буквы, цифры, подчеркивание или знак доллара. PTX поддерживает целые и с плавающей точкой константы и константные выражения. Эти константы могут быть использованы при инициализации данных и в качестве операндов инструкции. Тип проверки правил остаётся одинаковым для целого, с плавающей точкой, а также типы битового размера. Для предикат типа данных и инструкций, целочисленные константы могут и интерпретируются как в С, т.е. нулевые значения являются ложными и ненулевые значения True. Хотя конкретные ресурсы, имеющиеся в данной целевой GPU будет варьироваться, виды ресурсов будет общим для всех платформ, и эти ресурсы абстрагируются в PTX через пространства состояний и типов данных. Пространство состояний – набор с конкретными характеристиками. Все переменные находятся в каком-то пространстве состояний. Характеристики пространства состояний включают его размер, адресацию, скорость доступа, права доступа, и уровень обмена между потоками. Пространства состояний, определенные в PTX являются побочным продуктом параллельного программирования и графического программирования. В PTX основные типы отражают собственные типы данных, поддерживаемые целевыми архитектурами. Фундаментальный тип определяет как базовый тип и размер. Регистровые переменные всегда имеют фундаментальный тип, и команды действуют для этих типов. PTX включает в себя встроенные непрозрачные типы для определения текстуры и поверхностные дескрипторные переменные. Эти типы названы как и поля, аналогичные структурам, но вся информация о макете, заказ поля, базовый адрес, и общий размер скрыты в программе PTX,
748 отсюда и термин непрозрачный. Использование этих непрозрачных типов ограничивается: – определение переменных в глобальных (модульных) рамках и в списках параметров начального ядра – статическая инициализация модуля – область видимости переменных, используя разделённые запятыми выражений статического назначения для названных членов типа – ссылки на текстуры, пробники, или поверхности с помощью текстуры и поверхности инструкций загрузки / сохранения (tex, suld, sust, sured). – получение значения именованного члена по инструкции запроса (txq, suq) – cоздание ссылки на непрозрачные переменные с помощью MOV, например, mov.u64 reg, opaque_var;. – непрозрачные переменные не могут появиться в инициализаторах, например, для инициализации указателя на непрозрачной переменной. Косвенный доступ к текстурам поддерживается только в едином режиме текстур. Три встроенных типа – это .texref, .samplerref и .surfref. Для работы с текстурами и шаблонами PTX имеет два режима работы. В едином режиме, текстурная и шаблонная информация доступна через одну комманду .texref. В автономном режиме, текстурная и шаблонная информации имеют свои собственные источники, что позволяет им быть определенными отдельно и в сочетании в зависимости от использования в программе. В автономном режиме, поля типа .texref, которые описывают свойства шаблона игнорируются, так как эти свойства определяются .samplerref - переменными. В PTX, объявление переменной описывает как тип переменной и ее пространство состояний. В дополнение к базовым типам, PTX поддерживает типы для простых совокупных объектов, таких как вектора и массивы. Поддерживаются векторные типы ограниченной длины. Векторы длиной 2 и 4 любого непредикатного фундаментального типа может быть объявлен с помощью префикса типа с .х2 или .v4. Векторы должны быть основаны на фундаментальном типе, и они могут находиться в регистре места. Вектор не может превышать 128-бит в длину; например, .v4.f64 не допускается. Трехмерный векторы может быть обработан с помощью .v4 вектор, в котором четвертый элемент обеспечивает заполнение. Это общий случай для трехмерных сеток, текстур и т.д. Объявление массива небходимо для того, чтобы позволить программисту выделить резервное пространство. Чтобы объявить массив: имя переменной следует из одномерных объявлений, подобных объявлениям фиксированного размера массива в C, размер каждого измерения постоянное выражение. Выравнивание хранения для всех адресуемых переменных может быть указано в объявлении переменной. Выравнивание задается с помощью спецификатора дополнительного подсчета .alignbyte сразу после пространства состояний спецификатора. Переменные будут выровнены по адресу, который является целым числом, кратным байту- количеству. Значение выравнивания байт-счетчик должен быть степенью двойки. Для массивов, выравнивание определяет выравнивание адресов для начального адреса всего массива, а не для отдельных элементов. Выравнивание по умолчанию для скалярных и матричных переменных кратно размеру базового типа. Выравнивание по умолчанию для векторных переменных является кратным общему размеру вектора. Все операнды в инструкции – известный тип из своих деклараций. Каждый тип операнда должен быть совместим с типом определяемого шаблоном инструкции и типа команды. Там нет автоматического преобразования между типами.Тип бит совместим с каждым типом, имеющий тот же размер. Целые типы общего размера совместимы друг с другом. Исходные операнды обозначаются в описаниях команд именами a, b и с. Операнды для инструкций ALU должны быть переменными .reg, объявленными в регистровом пространстве состояний. Для большинства операций, размеры операндов должны быть совместимыми. cvt (конвертировать) команда проводится над различными типами операндов, так как его работа заключается в преобразовании практически из любого типа данных в любой другой тип данных (и размер). ld, st, mov, и cvt – инструкции, чтобы скопировать данные из одного места в другое. Инструкци b ld и st перемещают данные из / в адресные пространства состояний в / из регистров. Инструкция mov копирует Данные между регистрами. Большинство инструкций имеет дополнительный защитный предикат, который контролирует условное выполнение. Предикатные операнды обозначаются именами p, q, r, s. 749 Инструкции PTX, которые сохраняют результат в поле, обозначаются именем d (для назначения) в описаниях команд. Операнд назначения является скалярной или векторной переменной в регистровом пространстве состояний. Инструкции PTX, как правило, имеют от нуля до четырех операндов плюс, по желанию, защитный предикат после символа @ слева от opcode:Для инструкций, которые создают результирующее значение, d операнд является операндом назначения, a, b, c являются исходными операндами.Инструкция setp пишет два регистра назначения. Мы используем символ | для разделения нескольких регистров назначения.Для некоторых инструкций операнд назначения является обязательным. Операнд – битоприемник обозначается символом подчеркивания (_) и может быть использован вместо регистра назначения. В PTX, предикатные регистры являются виртуальными и имеют .pred как спецификатор типа. Так, предикат регистры могут быть объявлены как .reg .pred p, q, r; Все инструкции имеют дополнительный защитный предикат, который управляет условным выполнением инструкции.Синтаксисом, указывающим условное выполнение, является префикс – инструкция с @ {!} p, где р – предикатная переменная, возможно отрицательная. Инструкции без предикат выполняются безоговорочно. Предикаты обычно устанавливается в результате сравнения, выполненного setp- инструкции. Предикатные значения могут быть вычислены и управляемы, используя следующие инструкции: and, or, xor, not, и mov. Там нет прямого преобразования между предикатами и целочисленными значениями, и никакого прямого способа загрузки или хранения предикатных значений регистров. Тем не менее, setp может быть использован для создания предикат из целого числа, и выбор (selp –инструкция на основе предиката) может быть использован для создания целочисленного значение в зависимости от значения предиката. жүктеу/скачать 19,96 Mb. Достарыңызбен бөлісу:
|