Спецификаторы доступа CUDA

Спецификатор	Выполняется на	Может вызываться из
__device__	device	device
__global__	device	host
__host__	host	host

 

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

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

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

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

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

 

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

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

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

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

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

 

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

· gridDim - размер grid'а (имеет тип dim3);

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

· blockIdx - индекс текущего блока в grid'е (имеет тип uint3);

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

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

 

В язык добавляются 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>.

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

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

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

Здесь kernelName это имя (адрес) соответствующей __global__ функции, Dg - переменная (или значение) типа dim3, задающая размерность и размер grid'a (в блоках), Db - переменная (или значение) типа dim3, задающая размерность и размер блока (в нитях), Ns - переменная (или значение) типа size_t, задающая дополнительный объем shared-памяти, которая должна быть динамически выделена (к уже статически выделенной shared-памяти), S - переменная (или значение) типа cudaStream_t задает поток (CUDA stream), в котором должен произойти вызов, по умолчанию используется поток 0. Через args обозначены аргументы вызова функции kernelName.

Также в язык С добавлена функция __syncthreads, осуществляющая синхронизацию всех нитей блока. Управление из нее будет возвращено только тогда, когда все нити данного блока вызовут эту функцию. Т.е. когда весь код, идущий перед этим вызовом, уже выполнен (и, значит, на его результаты можно смело рассчитывать). Эта функция очень удобная для организации безконфликтной работы с shared-памятью.

Также CUDA поддерживает все математические функции из стандартной библиотеки С, однако с точки зрения быстродействия лучше использовать их float-аналоги (а не double) - например sinf. Кроме этого CUDA предоставляет дополнительный набор математических функций (__sinf, __powf и т.д.) обеспечивающие более низкую точность, но заметно более высокое быстродействие чем sinf, powf и т.п.

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

На данный момент самой новой версией является CUDA Runtime API 8.0. нужно сказать, что ведется разработка в сторону поддержки все большего числа языков программирования, например для Java существует JCUDA.

 

⇐ Предыдущая1234567Следующая ⇒

Поделиться: