Методичка (CUDA) - черновик. Введение cuda (Compute Unified Device Architecture)Ошибка! Источник ссылки не найден
 Скачать 1.06 Mb.
|
|
1.5. Задание Реализовать CPU и GPU версии программ в соответствии с вариантом за- дания. Результат работы на GPU проверяется с помощью реализации на CPU, которая является эталонной. Для обеих версий измерить время работы алго- ритма. Варианты задания: 1. транспонирование матрицы; 2. сложение двух матриц в соответствии с вариантом, выданным препо- давателем; 3. перемножение двух матриц; 4. выполнить обнуление элементов выше (ниже) главной диагонали мат- рицы. 13 2. Иерархия памяти в CUDA Технология CUDA использует следующие типы памяти: регистры, локаль- ная, глобальная, разделяемая, константная и текстурная память (рисунок 2.1). Однако разработчику недоступны к управлению регистры и локальная память. Все остальные типы он вправе использовать по своему усмотрению. Решетка блоков Блок (0,0) Разделяемая память Регистры Регистры Нить (0,0) Локальная память Локальная память Глобальная память Локальная память Текстурная память Нить (0,1) Блок (0,0) Разделяемая память Регистры Регистры Нить (0,0) Локальная память Локальная память Нить (0,1) Рис. 2.1. Иерархия памяти CUDA Характеристики каждого типа памяти представлены в таблице 2.1. Для ло- кальной и глобальной типов памяти был добавлен кэш, начиная с Compute Ca- pability 2.0 (обозначено символом «*»). Таблица 2.1 Типы памяти CUDA Тип памяти Расположение Кэш Доступ Видимость Время жизни Регистры Мультипроцессор Нет R/W Нить Нить Локальная DRAM GPU * R/W Нить Нить Разделяемая Мультипроцессор Нет R/W Блок Блок 14 Глобальная DRAM GPU * R/W Сетка Константная DRAM GPU Есть R Сетка Текстурная DRAM GPU Есть R Сетка 2.1. Глобальная память Глобальная память (global memory) – тип памяти с самой высокой латент- ностью, из доступных на GPU. Переменные в данном типе памяти можно выде- лить с помощью спецификатора __global__ , а так же динамически, с помощью функций из семейства cudaMalloc . Глобальная память в основном служит для хранения больших объемов данных, над которыми осуществляется обработка, и для сохранения результата работы. Данные перемещения осуществляются с ис- пользованием функций cudaMemcpy . В алгоритмах, требующих высокой произ- водительности, количество операций с глобальной памятью необходимо свести к минимуму. Пример сложения двух векторов из подраздела 1.2.1 демонстрирует ис- пользование глобальной памяти. Для ускорения передачи данных между ОЗУ и памятью устройства ис- пользуется механизм объединения запросов (coalescing). Шаблон доступа к гло- бальной памяти отличается в зависимости от версии архитектуры графического процессора (Compute Capability), на котором осуществляется запуск. Самые жесткие требования накладывает версия 1.х. Для Compute Capability 1.0 и 1.1 накладываются следующие ограничения для полуварпа (16 нитей из варпа): размер слова, к которому осуществляется доступ нитью, должен быть 4, 8 или 16 байтным. При этом если работа осуществляется с 4 либо 8 байтными словами, тогда все 16 слов должны находиться в одном 64- либо 128-байтном сегменте соответственно. При 16 байтных словах первых 8 слов должны лежать в одном 128-байтном сегменте, следую- щие 8 слов следующим за первым 128-байтном сегменте; нити должны обращаться последовательно, т. е. k-я нить полуварпа должна обращаться к k-ому слову. Если данные условия выполняются, формируются 64-, 128- или 2 по 128 байт транзакции соответственно для 4, 8 и 16 байтных словах. Для Compute Capability 1.2 и 1.3 шаблон доступа несколько усовершен- ствован (при этом доступ по-прежнему осуществляется полу-варпами): 15 осуществляется поиск сегмента в памяти, содержащего адрес, запраши- ваемый активной нитью с минимальным индексом thread ID. Размер сегмента зависит от размера слова, к которому осуществляется доступ: o 32 байта для 1-байтных слов; o 64 байта для 2-байтных слов; o 128 байт для 4-, 8- и 16-байтных слов; ищутся все активные нити, запрашиваемый адрес которых находится в одном сегменте; анализируется возможность понижения размера транзакции: o если размер транзакции 128 байт и только старшая или младшая его части загружаются, происходит уменьшение размера тран- закции до 64 байт: o по аналогии для 64 байт возможно понижение размера транзак- ции до 32 байт; выполняется транзакция и обработанные нити отмечаются как не ак- тивные; алгоритм повторяется для всех нитей полуварпа. Для Compute Capability 2.x шаблон доступа существенно упростился. В ос- новном это связано с вводом кэша. При компиляции появляется дополнитель- ный флаг, позволяющий задействовать L1 и L2 кэши (опция -Xptxas -dlcm=ca, включенная по-умолчанию) либо только L2 (опция -Xptxas -dlcm=cg). Размер кэш-линии — 128 байт, на которые проецируется соответственно 128-байтный выравненный сегмент в глобальной памяти. При этом если задей- ствуется L1 и L2 кэши, формируются 128-байтные транзакции, в случае, если только L2 кэш, тогда формируются только 32-байтные транзакции (в зависимо- сти от местоположения запрашиваемых данных (например, в произвольном по- рядке) 32-байтные транзакции могут оказаться более эффективными). В случае если размер слова, запрашиваемого нитью, больше 4 байт, запра- шиваемый варпом, блок данных разделяется на 128-байтные запросы, обраба- тываемые независимо. Транзакции по 128 байт будут формироваться до тех пор, пока все нити варпа не получат необходимые данные. В отличие от Compute Capability 1.x шаблон доступа в архитектуре 2.х до- пускает доступ нитью к словам в произвольном порядке. Для Compute Capability 3.x убрана возможность управления кешированием данных в L1, т.к. данный уровень кэша зарезервирован для работы с локальной памятью. При этом схема работы для L2 сохраняется такой же, как и для Com- pute Capability 2.x. Для Compute Capability 3.5 добавлена поддержка работы с 16 кэшем, доступным из ядра только на чтение. Более подробно этот режим рас- смотрен в [1]. На рисунке 2.2 показан пример эффективного обращения в глобальную память, в результате такого чтения будет сформирована одна транзакция. При этом предполагается, что каждая нить обращается к 4-байтному слову. Следует учитывать, что некоторые нити могут не обращаться в свою ячейку памяти (например, нить 2 на рисунке 2.2), при этом транзакция нарушена не будет. За- гружаемый блок данных выделен с помощью точек. Адрес 128 Адрес 132 Адрес 136 Адрес 140 Адрес 244 Адрес 248 Адрес 252 Нить 0 Нить 1 Нить 3 Нить 29 Нить 30 Нить 31 Адрес 124 Адрес 120 Адрес 256 Адрес 260 Сегемент 2 Рис. 2.2. Пример выравненного обращения к глобальной памяти На рисунке 2.3 показан пример обращения к памяти со смещением. В ре- зультате такого обращения будут сформированы 2 транзакции и загружены два 128-байтных сегмента данных. Загружаемый блок данных выделен с помо- щью точек. 17 Адрес 128 Адрес 132 Адрес 136 Адрес 140 Адрес 244 Адрес 248 Адрес 252 Нить 0 Нить 1 Нить 3 Нить 29 Нить 30 Нить 31 Адрес 124 Адрес 120 Адрес 256 Адрес 260 Адрес 238 Сегемент 1 Сегемент 2 Сегемент 3 Адрес 96 Адрес 0 Адрес 255 Рис. 2.3. Пример не выравненного обращения к глобальной памяти В тоже время отключение L1 (рисунок 2.4) позволяет уменьшить объем данных, который будет передан при выполнении. При этом будут загружены данные, начиная с адреса 96. Загружаемый блок данных выделен с помощью точек. 18 Адрес 128 Адрес 132 Адрес 136 Адрес 140 Адрес 244 Адрес 248 Адрес 252 Нить 0 Нить 1 Нить 3 Нить 29 Нить 30 Нить 31 Адрес 124 Адрес 120 Адрес 256 Адрес 256 Адрес 260 Адрес 260 Адрес 238 Сегемент 1 Сегемент 2 Сегемент 3 Адрес 96 Адрес 0 Адрес 0 Адрес 255 Адрес 255 Адрес 92 Адрес 92 Рис. 2.4. Пример не выравненного обращения к глобальной памяти с отключенным кэшем L1 2.2. Разделяемая память Разделяемая память (shared memory) относится к типу памяти с низкой латентностью. Данный тип памяти рекомендуется использовать для минимиза- ции обращения к глобальной памяти, а так же для хранения локальных пере- менных функций. Адресация разделяемой памяти осуществляется между нитя- ми одного блока, что может быть использовано для обмена данными между по- Примечание [D2]: ???? 19 токами в пределах одного блока. Для размещения данных в разделяемой памяти используется спецификатор __shared__ Разделяемая память делится на блоки фиксированного размера, доступ к которым осуществляется одновременно. Таким образом, если нити обращаются к различным банкам, латентность доступа будет минимальна. Если хотя бы од- на из нитей не получила данные из-за конфликта доступа к банку, осуществля- ется повторный доступ к памяти и данные догружаются. Данная процедура бу- дет повторяться до тех пор, пока не будут разрешены все конфликты и загру- жены или сохранены все данные. В результате время доступа будет произведе- нием порядка конфликта на латентность доступа к разделяемой памяти. Рассмотрим следующий пример. Имеется массив на N элементов типа int Требуется выполнить сложение трех подряд идущих элементов и сохранить ре- зультат в качестве одного элемента выходного массива (рисунок 2.5). 0 1 2 3 4 5 N-3 N-2 N-1 + + + 0 1 N/3-1 Рис. 2.5. Постановка задачи для примера При формировании конфигурации исходим из следующих предположений: одна нить формирует один выходной элемент; число нитей в блоке – 512; число элементов N гарантированно делится на 512*3=1536. Простейшее решение – каждая нить читает последовательно 3 элемента из глобальной памяти. Однако такое решение приводит к тому, что возникает конфликт третьего порядка при чтении из глобальной памяти и, как результат, нарушение условия формирования транзакций. Для того чтобы устранить данный недостаток рекомендуется задейство- вать разделяемую память как промежуточный буфер. При этом запись в разде- ляемую память осуществляется блоками по 512 последовательных элементов. Пример эффективного решения данной задачи показан ниже: 20 __global__ void Sum( const int *input, const int *output, const int N ) { __shared__ smem[512 * 3]; // копируем входные данные в разделяемую память smem[threadIdx.x] = input[blockIdx.x * 512 * 3 + threadIdx.x]; smem[threadIdx.x + 512] = input[blockIdx.x * 512 * 3 + threadIdx.x + 512]; smem[threadIdx.x + 1024] = input[blockIdx.x * 512 * 3 + threadIdx.x + 1024]; // синхронизируем нити, чтобы гарантировать, // что в разделяемую память записаны все данные __syncthreads(); // находим сумму и сохраняем результат работы output[blockIdx.x * 512 + threadIdx.x] = smem[threadIdx.x * 3] + smem[threadIdx.x * 3 + 1] + smem[threadIdx.x * 3 + 2]; } Функция __syncthreads() гарантирует, что все нити блока выполнили запись в разделяемую память. Разрешение конфликта 3-го порядка, возникаю- щего при чтении данных из разделяемой памяти, является гораздо менее за- тратным по сравнению с конфликтом при чтении из глобальной памяти, поэто- му в данном случае этим можно пренебречь. Запись в глобальную память осуществляется с формированием транзакций, поэтому промежуточного сохранения в разделяемую память не требуется. 2.3. Константная память Константная память (constant memory) является одной из самых быстрых из доступных на GPU. Отличительной особенностью данного типа памяти яв- ляется возможность записи данных с хоста, но при этом в пределах GPU воз- можно лишь чтение из этой памяти, что и обуславливает её название. Кон- стантная память обладает собственным кэшем, что обеспечивало выигрыш на платформах с Compute Capability 1.x, однако на более новых архитектурах из-за ввода кэшей L1 и L2 константная память частично потеряла свою актуальность. 21 Для размещения данных в константной памяти предусмотрен специфика- тор __constant__ . Если необходимо использовать массив в константной памя- ти, то его размер необходимо указать на этапе компиляции, так как динамиче- ское выделение в отличие от глобальной памяти в константной не поддержива- ется. Для записи с хоста в константную память используется функция cudaMemcpyToSymbol 2.4. Текстурная память Текстурная память (texture memory) входит в состав текстурных блоков, используемых в графических задачах для формирования текстур. В текстурном блоке аппаратно реализована фильтрация текстурных координат, интерполя- ция, нормализация текстурных координат в случаях, когда они выходят за до- пустимые пределы. Наличие кэша делала эффективным использование данного типа памяти на архитектурах с Compute Capability 1.x, однако с появлением ар- хитектуры Fermi значимость данного фактора потеряла актуальность, т.к. по- явился собственный кэш для глобальной памяти. Текстурная память выделяется с помощью функции cudaMallocArray и освобождается с использованием cudaFreeArray : cudaError_t cudaMallocArray ( cudaArray_t* array, const cuda- ChannelFormatDesc* desc, size_t width, size_t height = 0, unsigned int flags = 0 ); cudaError_t cudaFreeArray ( cudaArray_t array ); Для того чтобы получить доступ к выделенной памяти в CUDA-ядре тре- буется ассоциировать указатель с текстурной ссылкой: cudaError_t cudaBindTexture ( size_t* offset, const tex- tureReference* texref, const void* devPtr, const cudaChannelFor- matDesc* desc, size_t size = UINT_MAX ); Для получения доступа к данным из CUDA-ядра используются специаль- ные функции tex1D , tex2D , tex3D В связи с практически полной потерей актуальности текстурной памяти в задачах вычислений общего назначения, подробно указанные функции не рас- сматриваются. Примечание [D3]: 22 2.5. Регистровая память Регистровая память (register) является самой быстрой из всех типов па- мяти. Определить общее количество регистров, доступных GPU, можно с по- мощью функции cudaGetDeviceProperties Количество регистров, доступных одной нити, определяется следующим образом: где – общее количество регистров, доступных GPU; – число нитей в блоке; – число блоков в ре- шетке блоков. Все регистры GPU 32 разрядные. При этом в технологии CUDA нет явных способов использования регистровой памяти, всю работу по размещению дан- ных в регистрах берет на себя компилятор. 2.6. Локальная память Локальная память (local memory) может быть использована компилято- ром, если все локальные переменные не могут быть размещены в регистровой памяти. По скоростным характеристикам локальная память значительно мед- леннее, чем регистровая. Для проверки, используется ли локальная память, ре- комендуется при компиляции включить опцию PTXAS Output в свойствах про- екта в Microsoft© Visual Studio (опция --ptxas-options=-v при компиляции через командную строку) и проанализировать окно вывода информации о сборке проекта: Число для ___, отличное от нуля, говорит о том, что компилятор не смог использовать только регистры и задействовал локальную память. В данной си- туации рекомендуется проанализировать код и уменьшить число используемых переменных. 2.7. Разделяемая память как управляемый программный кэш, регистро- вый кэш Латентность доступа к глобальной памяти крайне высока, поэтому в пра- вильно спроектированном алгоритме чтение и запись элемента данных должно Примечание [D4]: Не слишком ли сложно??? 23 осуществляться только один раз, при этом должны формироваться транзакции. Поэтому, как правило, ядро начинается с загрузки данных из глобальной в раз- деляемую память и заканчивается сохранением результата работы из разделяе- мой в глобальную память. Латентность доступа к разделяемой памяти в сравне- нии с глобальной гораздо ниже, однако она не нулевая. Поэтому при использо- вании одной нитью значений из разделяемой памяти более одного раза может оказаться эффективной программная реализация дополнительного кэша, по- строенного на уровне регистров (локальных переменных). В данные перемен- ные размещаются наиболее часто используемые данные и в дальнейшем могут использоваться для аккумулирования промежуточного результата. Однако работа с массивами переменных в качестве кэша данных является задачей, требующей осторожности и внимательности. Если компилятор не сможет распределить переменные по доступным регистрам, будет задействова- на локальная память, что приведет к существенному замедлению работы ядра. Поэтому для контроля используемого числа регистров одной нитью и количе- ства байт задействованной локальной памяти рекомендуется воспользоваться рекомендациями из подраздела 2.7. 2.8. Задание 1. Выполнить оптимизацию доступа к глобальной памяти для задания из подраздела 1.5. 2. Проанализировать возможность использования регистрового кэша для задания из подраздела 1.5. 24 3. Синхронизация данных Архитектура CUDA предлагает 2 основных способа синхронизации дан- ных: синхронизация устройства, когда графический процессор гарантиро- ванно завершает поставленные на выполнение задачи, а работа цен- трального процессора для данного алгоритма блокируется до окончания работы GPU. При этом синхронизация может происходить как в явном виде (функция cudaDeviceSynchronize ), так и в не явном виде (напри- мер, при использовании не асинхронной версии функции копирования данных cudaMemcpy ); синхронизация нитей внутри блока. Межблочная синхронизация нитей на момент написания методического пособия не предусмотрена технологией и осуществляется самостоятельно раз- работчиком собственными силами. Кроме описанных выше способов можно выделить дополнительный спо- соб синхронизации, который при правильном проектировании алгоритма поз- воляет избавиться синхронизации нитей внутри блока в явном виде, – исполь- зование свойств варпа. Данный способ будет описан ниже. |