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

Имеется четыре типа расширений языка Си в CUDA. Это:

1. Квалификаторы типа функции, определяющие, где эта функция будет выполняться - в центральном процессоре или на видеокарте, и возможно ли вызывать эту функцию из центрального процессора или видеокарты.
2. Квалификаторы типа переменных, определяющие распределение памяти для этих переменных на видеокарте.
3. Директива, определяющая как ядро выполняется на видеокарте.
4. Четыре встроенных переменных, определяющих размерность грида, размерность блока, индекс блока и индекс нити.

Рассмотрим эти расширения подробней.

I.Квалификаторы типа функций

1. __device__ квалификатор говорит, что функция выполняется на видеокарте и может быть вызвана только из видеокарты.
2. __global__ квалификатор объявляет функцию выполняемой в  центральном процессоре и вызываемой только процессором. Квалифицирование функции только __host__ аналогично объявлению функции без какого-то ни было квалификатора. Однако функцию можно квалифицировать комбинацией __host__ __device__. Тогда функция будет откомпилирована и для центрального процессора и для видеокарты.  

 Ограничения:

•  функции квалифицированные __device__ и __global__ не поддерживают рекурсии.
•  функции квалифицированные __device__ и __global__ не могут декларировать внутри себя статические переменные.
• функции квалифицированные __device__ и __global__ не могут иметь переменное количество аргументов.
• функции квалифицированные __device__ не могут вызываться по их адресу.
• функции не могут быть одновременно квалифицированы как __global__ и __host__.
• функции квалифицированные как __global__ обязаны возвращать пустой тип.
• функции квалифицированные как __global__ должны специфицировать  свою конфигурацию выполнения.
• функции квалифицированные как __global__ выполняются синхронно.
• параметры функций квалифицированных как __global__ передаются через общую память в видеокарту и ограничены 256-ю байтами.

II. Квалификаторы переменных.

1. __device__ квалификатор говорит, что переменная размещается на видеокарте. Может использоваться совместно с другими квалификаторами типа переменных. Если таковых нет, то это означает, что переменная размещена в глобальной памяти, имеет время жизни равное времени жизни приложения и доступна из нитей в гриде и с центрального процессора через библиотеку времени исполнения.
2. __constant__ квалификатор опционально используемый предыдущим квалификатором объявляет, что переменная размещается в области постоянной памяти, имеет время жизни равное времени жизни приложения и доступна из нитей в гриде и с центрального процессора через библиотеку времени исполнения.
3. __shared__ квалификатор опционально используемый с  __device__ квалификатором объявляет, что перменная размещается в общей памяти блока нитей, имеет время жизни равное времени жизни блока и доступна только из нитей данного блока.

Ограничения:

• Эти квалификаторы не допустимы для членов struct и union, в формальных параметрах функций и локальныхперменных внутри функций, что выполняются на центральном процессоре.
• __shared__ и __shared__ не могут использоваться вместе.
• __shared__ и __shared__ переменные подразумеваются статическими.
• __constant__ переменные не могут использоваться внутри видеокарты, только на центральном процессоре. Поэтому они имеют область видимости файла.
• __shared__ переменные не могут иметь инициализацию как часть их декларации.

III. Конфигурация выполнения.

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

Описание аргументов в угловых скобках следующее:

1. Dg - имеет тип dim3 и определяет размерность и длину грида так что выражение Dg.x*Dg.y равно числу запущенных блоков.
2. Db - имеет тип dim3 и определяет размерность и длину каждого блока, так что выражение Db.x*Db.y*Db.z равно числу нитей в каждом блоке.
3. Ns - имеет тип size_t и определяет количество байтов в общей памяти, которая выделяется динамически при данном вызове функции дополнительно к ее статической памяти. Ns - необязательный аргумент со значением по умолчанию раным нулю.

Аргументы конфигурации выполнения вычисляются раньше  "обычных" аргументов функции.

IV. Встроенные переменные

1. gridDim - переменная типа dim3. Содержит размерность грида.
2. blockIdx - переменная типа uint3. Содержит индекс блока в гриде.
3. blockDim - переменная типа dim3. Содержит размерность блока.
4. threadIdx - переменная типа uint3. Содержит индекс нити в блоке.

Ограничения:

• не допустима операция получения адреса от любой встроенной переменной.
• не допустима операция присваивания для любой встроенной переменной.