3.3 Расширение языка Си

Технология CUDA вводит ряд дополнительных расширений для языка C, которые необходимы для написания кода для GPU:

1.
Спецификаторы функций, которые показывают, как и откуда будут выполняться функции.
2.
Спецификаторы переменных, которые служат для указания типа используемой памяти GPU.
3.
Спецификаторы запуска ядра GPU.
4.
Встроенные переменные для идентификации нитей, блоков и др. параметров при исполнении кода в ядре GPU.
5.
Дополнительные типы переменных.

Спецификаторы функций определяют, как и откуда буду вызываться функции. Всего в CUDA 3 таких спецификатора:

__host__ — выполнятся на CPU, вызывается с CPU.

__global__ — выполняется на GPU, вызывается с CPU.

__device__ — выполняется на GPU, вызывается с GPU.

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

myKernelFunc<<<gridSize, blockSize, sharedMemSize, cudaStream>>>(float* param1,float * param2), где

1.
myKernelFunc — функция ядра (спецификатор __global__)
2.
gridSize — размерность сетки блоков (dim3), выделенную для расчетов,
3.
blockSize — размер блока (dim3), выделенного для расчетов,
4.
sharedMemSize — размер дополнительной памяти, выделяемой при запуске ядра,
5.
cudaStream — переменная cudaStream_t, задающая поток, в котором будет произведен вызов.

Некоторые переменные при вызове ядра можно опускать, например sharedMemSize и cudaStream.

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

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

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

В язык добавляются 1/2/3/4-мерные вектора из базовых типов — char1, char2, char3, char4, uchar1, uchar2, uchar3, uchar4, short1, short2, short3, short4, ushort1, ushort2, ushort3, ushort4, int1, int2, int3, int4, uint1, uint2, uint3, uint4, long1, long2, long3, long4, ulong1, ulong2, ulong3, ulong4, float1, float2, float3, float2, и double2.

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

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

CUDA поддерживает все математические функции из стандартной библиотеки C. Однако при этом следует иметь в виду, что большинство стандартных математических функций использует число в двойной точностью (double). Но поскольку для GPU операции с double-числами выполняются медленнее, чем операции с float-числами, то предпочтительнее там, где это возможно, использовать float-аналоги стандартных функций. Так, float-аналогом функции sin является функция sinf.

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

В таблице 1 приведены float – функции и их оптимизированные версии пониженной точности.

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

Таблица 3.2: Математические float–функции в CUDA


Функция

Значение



__fadd_[rn, rz, ru, rd](x, y)

Сложение, никогда не переводимое в команду FMAD



__fmul_[rn, rz, ru, rd](x, y)

Умножение, никогда не переводимое в команду FMAD



__fmaf_[rn, rz, ru, rd](x, y, z)

(x * y) + z



__frcp_[rn, rz, ru, rd](x)

1∕x



__fsqrt_[rn, rz, ru, rd](x)

√x--



__fdiv_[rn, rz, ru, rd](x, y)

x∕y



__fdividef(x, y)

x∕y, но если 2126 < y < 2128, то 0



__expf(x)

ex



__exp10f(x)

10x



__logf(x)

log x



__log2f(x)

log 2x



__log10f(x)

log 10x



__sinf(x)

sinx



__cosf(x)

cosx



__sincosf(x, sptr, cptr)

sptr = sinx;cptr = cosx



__tanf(x)

tanx



__powf(x, y)

xy



__int_as_float(x)

32 бита, образующее целочисленное значение, интерпретируются как float – значение. Так, значение 0xС000000 будет переведено в -2.0f



__float_as_int(x)

32 бита, образующие float-значение, интерпретируются как целочисленное значение.



__saturate(x)

Min(0, max(1, x))



__float_to_int_[rn, rz, ru, rd](x)

Приведение float — значения к целочисленному значению с заданным округлением



__float_to_uint_[rn, rz, ru, rd](x)

Приведение float — значения к без знаковому целочисленному значению с заданным округлением



__int_to_float_[rn, rz, ru, rd](x)

Приведение целочисленного значения к float – значению с заданным округлением



__uint_to_float_[rn, rz, ru, rd](x)

Приведение беззнакового целочисленного значения к float – значению с заданным округлением



__float_to_ll_[rn, rz, ru, rd](x)

Приведением float — значения к 64-битовому целочисленному значению с заданным округлением



__float_to_ull_[rn, rz, ru, rd](x)

Приведение float — значения к 64-битовому беззнаковому целочисленному значению с заданным округлением



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

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


Функция

Значение



lstinline__[u]mul24(x, y)

Вычисляет произведение младших 24 бит целочисленных параметров x и y, возвращает младшие 32 бита результата. Старшие 8 бит аргументов игнорируются



lstinline__[u]mulhi(x, y)

Возвращает старшие 32 бита произведения целочисленных операндов x и y



lstinline__[u]mul64hi(x, y)

Вычисляет произведение 64-битовых целых чисел и возвращает младшие 64 бита этого произведения



lstinline__[u]sad(x, y, z)

Возвращает z + |x - y|



lstinline__clz(x)

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



lstinline__clzll(x)

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



lstinline__ffs(x)

Возвращает позицию первого (наименее значимого) единичного бита для аргумента x. Если x равен нулю, то возвращается нуль



lstinline__ffsll(x)

Возвращает позицию первого (наименее значимого) единичного бита для целочисленного 64-битового аргумента x. Если x равен нулю, то возвращается нуль



lstinline__popc(c)

Возвращает число бит, которые равны единице в двоичном представлении 32 – битового целочисленного аргумента x



lstinline__popcll(x)

Возвращает число бит, которые равны единице в двоичном представлении 64-битового целочисленного аргумента x



lstinline__brev(x)

Возвращает число, полученное перестановкой (то есть биты в позициях k и 31-k меняются местами для всех k от 0 до 31) битов исходного 32-битового целочисленного аргумента x



lstinline__brevll(x)

Возвращает число, полученное перестановкой (то есть биты в позициях k и 63-k меняются местами для всех k от 0 до 63) битов исходного 64-битового целочисленного аргумента x



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

Основы CUDA host API

CUDA предоставляет в распоряжение программиста ряд функций, которые могут быть использованы только CPU (называемый CUDA host API). Эти функции отвечают за:

CUDA API для CPU (host API) выступает в 2 формах:

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

CUDA driver API

Низкоуровневый API, дающий больше возможностей программисту, но и требующий большего объема кода и необходимость явных настроек, требование явной инициализации и отсутствие поддержки режима эмуляции (позволяющего компилировать, запускать и отлаживать коды на CUDA с CPU).

CUDA runtime API

Это высокоуровневый API, к тому же CUDA runtime API не требует явной инициализации — она происходит автоматически при первом вызове какой-либо функции. Данный API поддерживает эмуляцию. Плюсом является возможность использования дополнительных библиотек (CUFFT, CUBLAS, CUDPP).

Следует так же отметить, что у каждой функции CUDA runtime API есть прямой аналог в CUDA driver API, то есть переход с CUDA runtime API CUDA runtime API на CUDA driver API не очень сложен, обратное в общем случае не верно.