Некоторые возможности CUDA, часть 1

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

Поддержка функции printf на стороне GPU

Одной из крайне простых и полезных возможностей CUDA является поддержка функции printf на стороне GPU. Каждая нить сетки может вызвать printf для форматированного вывода данных в любой момент времени.

Однако при этом есть некоторые отличия от стандартной функции printf. Во-первых, функция возвращает число обработанных аргументов, не считая строки формата.

Поддерживаются почти все основные спецификаторы формата. Может выведено за один вызов до 32 аргументов, не считая самой строки формата. Однако обратите внимание, что вызов printf на стороне GPU просто помещает данные в специальный буфер. Окончательное форматирование выполняется уже на стороне CPU.

Вывод данных при помощи printf идет в специальный кольцевой буфер, размер которого можно задать до запуска ядра. По умолчанию размер этого буфера равен 1 Мб. Для задания размера этого буфера можно использовать следующий вызов:

Обратите внимание, что содержимое этого буфера не выводится автоматически перед завершением работы приложения. Автоматический вывод содержимого буфера и его очистка производятся в следующих случаях:

• запуск ядра;
• явная синхронизация через вызов cudaDeviceSynchronize, cudaStreamSynchronize и cudaEventSynchronize;
• любое блокирующее (т.е. не асинхронное) копирование памяти через cudaMemcpy*;
• Уничтожение текущего контекста CUDA через вызов функции cudaDeviceReset;
• Перед вызовом callback-функции, установленного через вызов cudaStreamAddCallback;

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

__global__ void hello ( float a )
{
    printf ( "Hello from thread %d, a = %f\n", threadIdx.x, a );
}

int main ()
{
    hello<<<1,5>>> ( 3.1415926f );
    cudaDeviceSynchronize ();       // force buffer flush

    return 0;
}

Выделение и освобождение памяти на стороне GPU

Стандартный CUDA API для выделения и освобождения памяти (функции cudaMalloc и cudaFree) могут вызываться только на стороне CPU. Однако иногда возникает необходимость иметь возможность выделять и освобождать глобальную память GPU прямо на стороне GPU. С этой целью были введены GPU-версии таких функций, как malloc, free, memcpy и memset.

Функция malloc служит для выделения заданного числа байт глобальной памяти из кучи (heap). При этом данная память остается выделенной и сохраняет свое содержимое и после завершения работы ядра, выделившего эту память. При этом гарантируется, что выделяемая та4ким образом память всегда будет выровнена как минимум по 16 байтам.

Для освобождения выделенной таким образом памяти используется функция free. Для копирования и записи фиксированного значения во все байты выделенной таким образом области памяти служат функции memcpy и memset. Прототипы всех этих четырех функций полностью совпадают со стандартными функциями библиотеки языка С.

Однако в отличии от вызов cudaMalloc функция malloc выделяет память из кучи (heap) фиксированного размера. По умолчанию этот размер равен 1 Мб, но его можно явно задать при помощи вызова cudaDeviceSetLimit как это показано ниже.

cudaDeviceSetLimit ( cudaLimitMallocHeapSize, size_t size );

Обратите внимание, что нельзя динамически менять размер кучи. При этом вся память, выделенная через malloc может быть освобождена только через вызов free на стороне GPU. Аналогично, вся память, выделенная через cudaMalloc* не может быть освобождена через вызов free. Кроме того, память, выделенная через malloc не может использоваться на стороне CPU в вызовах API и в функциях cudaMemcpy* и cudaMemset.

Описатель __launch_bounds

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

__global__void _launch_bounds__  ( maxThreadsPerBlock, minBlockPerMultiProcessor ) kernel ()
{
    . . .
}

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

Обычно эти параметры зависят от поколения GPU (точнее, от compute capability). Поэтому на этапе компиляции можно использовать макрос __CUDA_ARCH__ для определения того, под какую конкретную архитектуру в данный момент компилируется ядро.

#if __CUDA_ARCH >= 600       // compute capability 6.0 or higher
. . .
#else
#if __CUDA_ARCH >= 350      // compute capability 3.5 or higher
. . .
#endf

Опция #pragma unroll

По умолчанию nvcc автоматически разворачивает небольшие циклы, когда число итерация известно на этапе компиляции. Однако есть и способ явного управления разворачиванием циклов при помощи директивы #pragma unroll.

Это директива может быть поставлено только непосредственно перед циклом и действует только на этот цикл. В ней можно указать необязательное значение, управляющее степенью разворачивания. Если это число не указано, то цикл разворачивается полностью (если число итерация известно на этапе компиляции). Задание этого параметра равным единице фактически запрещает разворачивание данного цикла.

#pragma unroll
for ( int i = 0; i < 12; i++ )
    p1 [i] += p2 [i];

#pragma unroll 1
for ( int i = 0; i < 12; i++ )
    p1 [i] += p2 [i];

#pragma unroll 4
for ( int i = 0; i < 12; i++ )
    p1 [i] += p2 [i];