М. В. Ломоносова 2015 Антонюк Валерий Алексеевич Программирование на видеокартах (gpgpu). Спецкурс кафедры мми. М. Физический факультет мгу им. М. В. Ломоносова, 2015. 48 с. Учебное пособие
 Скачать 1.54 Mb.
|
|
CPU (в терминологии CUDA именуемого host) и для GPU (в терминологии CUDA — device), причём оба типа исходного кода используют язык C/C++, только в device-коде пока употребляется некоторое подмножество C99, а вот host-часть программного кода вполне может быть написана на C++. 7 Подобный симбиоз позволяет программисту оставаться в привычной среде одного языка, правда, «разбавленного» некоторыми дополнительными ключевыми словами и типами величин. Основная работа по анализу такой «смеси» ложится на компилятор nvcc. Иерархия нитей Вызов на исполнение отдельных копий кода ядра сопровождается указанием, сколько нужно запустить экземпляров кода (фактически — нитей исполнения), а также как они должны быть организованы для решения конкретной задачи, поскольку эти нити группируются в так называемые блоки (blocks), а те, в свою очередь, формируют сетку блоков (grid). В зависимости от выбранной при запуске конфигурации все нити, исполняющие код ядра, имеют одно-, двух- или трёхмерную организацию (как правило, соответствующую характеру решаемой задачи), причём каждая нить «имеет информацию» о том, где она расположена в выбранной конфигурации. Виды памяти Для работы программного кода на графическом процессоре видеокарты необходима память, доступная этому процессору, поэтому в host-коде понадобится выполнять дополнительные операции по выделению device-памяти, копированию данных с host на device, а после обработки данных — возвращение результата с device на host. Чаще всего используется так называемая глобальная память графической карты, однако нити в пределах одного и того же блока могут использовать и так называемую разделяемую память с быстрым доступом. Кроме этого возможна работа с так называемой константной памятью, доступной всем нитям сетки, — но только на чтение. Компиляция программ При установке пакета CUDA создаются специальные переменные окружения: базовая (с именем CUDA_PATH ) и переменная конкретной версии (например, CUDA_PATH_V4_1 ) — видимо, для возможности одновременного использования нескольких установленных версий CUDA; на их основе определяются переменные окружения, которые помогают указывать в Visual Studio пути поиска включаемых заголовочных файлов ( CUDA_INC_PATH=%CUDA_PATH%include ), статических библиотек ( CUDA_LIB_PATH=%CUDA_PATH%lib\Win32 ), а также исполняемых файлов и динамических библиотек ( CUDA_BIN_PATH=%CUDA_PATH%bin ). Значения этих (да и любых других) переменных окружения могут быть использованы в свойствах проекта Visual Studio для успешной компиляции программ наряду с указанием конкретных путей к файлам, для этого имя переменной окружения заключается в круглые скобки и предваряется символом доллара, например: $(CUDA_INC_PATH) . По сути, можно сказать, что в рамках проекта Visual Studio заданы макроопределения с именами переменных окружения. Далее иногда мы будем использовать такие макросы для обобщённого указания конкретных каталогов. Поскольку какие-то проекты могут использовать также NVIDIA GPU Computing SDK, то иногда необходимо добавлять и пути поиска включаемых файлов и библиотек из местоположения этого SDK (соответствующая переменная окружения называется NVSDKCOMPUTE_ROOT ), например, указывая путь $(NVSDKCOMPUTE_ROOT)\C\common\inc для включаемых заголовочных файлов и путь $(NVSDKCOMPUTE_ROOT)\C\common\lib\Win32 для статических библиотек. 8 Расширения языка C в архитектуре CUDA Программы в архитектуре CUDA пишутся на «расширенном» языке C, а для компиляции их используется, как уже говорилось, специальный компилятор nvcc. Расширения языка включают спецификаторы функций (они показывают, где будет исполняться функция, и откуда может быть вызвана), дополнительные типы данных, спецификаторы переменных (они определяют, какая память будет использоваться для их размещения), синтаксис запуска ядер, встроенные переменные, хранящие информацию о каждой нити, а также математические функции. Спецификаторы функций и переменных Для обозначения того, что функция является ядром (т.е., должна исполняться на GPU и может быть запущена на исполнение с CPU), используется спецификатор __global__ . Спецификаторы __host__ и __device__ помечают функции, исполняемые на CPU и GPU соответственно, причём и вызваны они могут быть только оттуда, где исполняются. Для задания размещения в памяти GPU переменных употребляются спецификаторы __device__ (переменная находится в глобальной памяти и доступна всем нитям), __constant__ (размещена в так называемой константной памяти, откуда она может быть только прочитана любой из нитей), __shared__ (переменная — в разделяемой памяти, где доступна только всем нитям «своего» блока). Дополнительные типы данных CUDA Это 1/2/3/4-мерные векторы из базовых типов языка C ( char , short , int , long — unsigned и signed , — longlong , float и double ), имена их образуются добавлением цифры с числом компонент: 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, float4, longlong1, longlong2, double1, double2. Обращение к компонентам осуществляется по именам ( x , y , z , w ), а создание значений- векторов — с помощью вызова функций make_ , где — один из вышеприведённых типов данных, например: int2 a = make_int2(1,4); float3 u = make_float3(1,2.72f,3.14f); Для задания размерности имеется тип dim3 , он обладает также нормальным конструктором: dim3 Blocks(16,16); // <=> Blocks(16,16); dim3 Grid(256); // <=> Grid(256,1,1); Полный синтаксис запуска ядер на исполнение Имея в своём распоряжении две переменные типа dim3 , можно задавать предпочтительную конфигурацию запуска ядер на исполнение. В простейших случаях, когда достаточно одномерной организации данных, можно обойтись и двумя целочисленными значениями. 2 9 Более сложные случаи запуска могут иметь также и дополнительные параметры; вот как выглядит полный синтаксис запуска ядер на исполнение: kernelName<< Здесь kernelName — имя __global__ -функции, Dg — переменная типа dim3 (параметры сетки блоков), Db — переменная типа dim3 (параметры блока нитей), Ns — дополнительный объём разделяемой памяти в байтах для каждого блока (необязательный параметр, он по умолчанию равен нулю), S — задаёт поток (CUDAstream), в котором должен произойти вызов (поток 0 по умолчанию), args — параметры функции-ядра (их может быть несколько, а общий размер этих параметров пока ограничен 256 байтами). Дополнительные математические функции Поскольку ядра исполняются на GPU, а не на CPU, им необходимы альтернативные реализации как минимум всех тех математических функций, которые были доступны в программах ранее в рамках стандартной математической библиотеки. Кроме double -версий в CUDA также имеются float -аналоги стандартных функций (например, для sin — sinf ). Есть ещё набор функций пониженной точности, но более «быстрых» (для sinf , например, __sinf ). Для функций, реализующих арифметические операции и преобразование типов, можно также задавать с помощью суффиксов их имени способ округления результата: rn (к ближайшему), rz (к нулю), ru (вверх), rd (вниз). Разделяемая (__shared__) память Основная часть динамической памяти GPU доступна как глобальная __device__ -память — всем нитям и блокам, а также CPU. Несмотря на то, что пересылка из памяти CPU в память GPU осуществляется довольно быстро, пересылки в рамках динамической памяти ещё быстрее. Разделяемая ( __shared__ ) память располагается в GPU и потому обладает гораздо большим быстродействием, чем глобальная, но доступна для чтения и записи нитям из одного блока. Первый способ выделения разделяемой памяти в ядре — явное указание размера массива со спецификатором __shared__ Второй способ — при запуске ядра задать дополнительный объём разделяемой памяти в байтах, который необходимо выделить каждому блоку при запуске ядра (третий параметр конфигурации функции-ядра, см. выше). Для доступа в ядрах к такой памяти используется описание массива в функции ядра без явно заданного размера, например: __shared__ float buf[]; При этом предполагается, что сам вызов ядра (с названием kernel ) выглядит так: kernel<<< ... , ... , k*sizeof(float) >>>( ... ); Тогда каждый блок ядер будет иметь доступ к массиву из k вещественных значений в разделяемой памяти (см. реализацию алгоритма умножения матриц далее). 10 Взаимодействие нитей в блоках Нити в рамках блока могут взаимодействовать между собой с помощью общей памяти, а потому появляется необходимость синхронизировать выполнение кода — чтобы чтение результатов из памяти не происходило ранее их формирования. Для этого программист определяет в коде точки синхронизации, используя вызов функции __syncthreads() . Она действует как барьер, перед которым все нити блока должны остановиться, ожидая завершения работы остальных нитей блока. Адресация памяти в нитях Используемая в NVIDIA GPU архитектура относится к типу SIMT (Single Instruction Multiple Threads, «одна инструкция, много нитей»): один и тот же код исполняется большим количеством нитей. Для эффективного распараллеливания таких вычислений, где данные могут обрабатываться одинаково и независимо, нужно, чтобы каждая нить обрабатывала "свою" часть данных, поэтому при программной реализации подобных вычислений возникает необходимость "правильного" распределения данных по нитям. Добиться этого можно такой процедурой сопоставления нитей и участков памяти для них, где разным нитям соответствуют разные участки памяти. При этом часто удобно, чтобы нитям одного блока сопоставлялся один компактный кусок памяти (без пропусков). Как оказывается, каждая копия функции ядра после запуска на исполнение имеет доступ к специальным переменным: к размерностям блоков и сетки (они называются blockDim и gridDim соответственно, обе имеют тип dim3 ), а также к положениям конкретной копии ядра в блоке и этого блока во всей сетке при исполнении кода ( threadIdx , blockIdx , обе переменные имеют тип int3 ). Используя эти переменные, каждая копия ядра может сформировать уникальные индексы для доступа к данным. Можно, например, последовательно перебирать (т.е. нумеровать) нити в блоке по каждому возможному измерению по очереди, при этом количество использованных индексов по одному измерению будет равно размеру блока по этому измерению, а общее количество использованных номеров-индексов будет равно произведению размеров блока по всем возможным измерениям. Аналогично можно поступить и с блоками в сетке. Порядок перебора измерений при этом часто не важен, хотя традиционно всё начинается с компоненты .x , затем — если необходимо — используется компонента .y , потом — если надо — используется компонента .z . Объяснение этому простое: разные измерения в конфигурации запуска ядер, как правило, неравноправны по предельным возможностям, поэтому «традиционный» выбор порядка следования измерений может упрощать последующее масштабирование программы. Подобная схема адресации часто реализуется в программах. Рассмотрим несколько примеров формирования индексов для доступа к данным в зависимости от объявленной конфигурации исполнения ядер. Например, в программе bitreverse.cu , осуществляющей "перевёртывание" битов в каждом из байтов последовательности (организованной в четвёрки целых "слов"), конфигурация содержит всего один блок с одномерной организацией нитей, поэтому для доступа к "словам" достаточно одной компоненты индекса нити ( threadIdx.x ). В программе pi.cu , вычисляющей приближение к числу π суммированием ряда, конфигурация нитей тоже 11 одномерна, равно как и сетка блоков, поэтому для доступа к местам подсчёта отдельных слагаемых ряда используется уникальный индекс, составленный из индекса нити по единственному измерению в блоке ( threadIdx.x ) и индекса блока по единственному измерению в сетке ( blockIdx.x ): int idx = blockIdx.x*blockDim.x+threadIdx.x; Выбор blockDim.x в качестве сомножителя в этом выражении обеспечивает "неразрывность" изменения уникального индекса ( threadIdx.x изменяется от 0 до blockDim.x-1 ). В программе MatrixAdd.cu , предназначенной для суммирования квадратных матриц NxN , конфигурация запуска ядер двумерна и по нитям в блоках, и по блокам в сетке — сообразно смыслу задачи, — а сама матрица разбита на квадратные подблоки такого размера (16x16), чтобы каждый подблок обрабатывался одним блоком нитей. Поэтому формирование уникальных индексов здесь организовано по-другому: создаются индексы сначала по каждому измерению (на основе нужных компонент индекса нити в блоке и индекса блока в сетке), а затем они комбинируются в один общий индекс для адресации элементов матрицы, расположенной в линейной памяти. int i = blockIdx.x * blockDim.x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; int index = i + j*N; В программе compute_pi.cu (ещё один вариант вычисления частичной суммы ряда для π) используется трёхмерная конфигурация нитей в единственном блоке, поэтому формула для уникального индекса использует только компоненты положения нити в блоке, причём перебор измерений здесь производится немного непривычным (хотя и вполне допустимым) образом, т.е., сначала — компонента .z , потом — .y , а затем только — .x : myid = threadIdx.z + threadDim*threadIdx.y+threadDim*threadIdx.y*threadIdx.x; Теперь уже понятно, что общая формула адресации нити в блоке (получение порядкового индекса нити tid ) может быть такова: tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y Она годится для любой возможной размерности конфигурации нитей в блоке, потому что, если конфигурация двумерна ( blockDim.z = 1 ), то единственным возможным значением для threadIdx.z будет 0 ; если же она одномерна (т.е., ещё и blockDim.y = 1 ), то и threadIdx.y будет всегда равно 0 , при этом реальный вид формулы для двумерной или одномерной организации нитей будет проще (и совпадёт с формулами в рассмотренных примерах программ). Аналогично, формула адресации блока в сетке (получение порядкового индекса блока bid ) может быть подобной: bid = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y 12 Она тоже годится для любой возможной размерности конфигурации блоков в сетке (если gridDim.z = 1 , то blockIdx.z = 0 ; если же ещё и gridDim.y = 1 , то и blockIdx.y = 0 ). Вводя вспомогательную величину nthreads для общего числа нитей в каждом из блоков: nthreads = blockDim.x * blockDim.y * blockDim.z мы можем получить линейный индекс id (поскольку память адресуется линейно) для доступа к данным из произвольной нити произвольного блока: id = tid + nthreads * bid Но иногда (как показывает пример с матрицей) удобно отступить от вышеизложенной схемы, организуя как "сквозную" адресацию нитей по каждому измерению, так и "блочную" — чтобы индексы нитей по каждому из двух измерений соответствовали индексам элементов матрицы в отдельном матричном блоке, индексы блоков нитей соответствовали индексам блоков матрицы, а формируемые по каждому измерению комбинированные индексы — индексам элементов матрицы в целом. Суммирование компонент вектора — пример ядра Рассмотрим в качестве примера один из вариантов реализации суммирования компонент вектора (обобщённо реализуемая операция называется в англоязычной литературе reduction и может использовать не только сложение, но и другие бинарные коммутативные операции, например, умножение, логическое сложение и т.п., приводя к формированию произведения и т.д.): __global__ void plus_reduce(int *g_idata, int N, int *g_odata) { __shared__ int sdata[BLOCKSIZE]; // Каждая нить загружает одно значение из глобальной памяти unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; sdata[tid] = (i < N)? g_idata[i] : 0; __syncthreads(); // Суммирование осуществляется в разделяемой памяти for (unsigned int s=blockDim.x/2; s>0; s>>=1) { if (tid < s) sdata[tid] += sdata[tid + s]; __syncthreads(); } // Запись результата для данного блока в глобальную память if (tid == 0) g_odata[blockIdx.x] = sdata[0]; } Здесь предполагается, что окончательное получение суммы осуществляется CPU. Второй вариант: добавление сразу к общей сумме, находящейся в первом элементе. if (tid == 0) atomicAdd(g_odata, sdata[0]); Т.е., может использоваться одна из так называемых атомарных операций, гарантирующая корректность выполнения операции в том случае, когда её пытаются выполнить многие нити одновременно. 13 |