Методичка (CUDA) - черновик. Введение cuda (Compute Unified Device Architecture)Ошибка! Источник ссылки не найден
 Скачать 1.06 Mb.
|
|
3.1. Функции синхронизации нитей внутри блока К функциям синхронизации нитей внутри блока относятся следующие 4 функции – __syncthreads , __syncthreads_count , int __syncthreads_and , __syncthreads_or . Описание функций сведено в таблицу 3.1. Таблица 3.1 Функции синхронизации нитей внутри блока Функция Compute Capability Возвращаемое значение 1 2 3 void __syncthreads() 1.0 и новее – int __syncthreads_count(int pred- icate) 2.0 и новее Число нитей, для кото- рых выражение predicate не равно 0 int __syncthreads_and(int predi- cate) 2.0 и новее Не нулевое значение только в том случае, если 25 значение predicate для всех нитей блока не рав- но 0 int __syncthreads_or(int predi- cate) 2.0 и новее Не нулевое значение только в том случае, если значение predicate хотя бы для одной нити блока не равно 0 3.2. Варп Понятие варпа (warp) в архитектуре CUDA является одним из ключевых наравне с такими понятиями как нить, блок и решетка блоков. Варп представ- ляет собой блок из 32 нитей мультипроцессора GPU, которые выполняют одну и ту же инструкцию, но каждый над своим блоком данных, т.е. фактически варп определяет минимальное число нитей, которые будут работать в один момент времени. Тот факт, что 32 нити работают синхронно, при правильном проекти- ровании алгоритма (с учетом разбиения блоков записи и чтения в/из глобаль- ной памяти по 32 нити) позволяет избавиться от операции __syncthreads() , обеспечивающей барьерную синхронизацию внутри блока потоков. Рассмотрим следующий пример. Допустим, имеется некоторый массив длиной elementCount (тип данных – int ). Требуется написать ядро, которое вычисляет среднее арифметическое двух соседних элементов и сохраняет ре- зультат в целочисленном виде. Схема стандартного подхода представлена на рисунке 3.1, цифрами на рисунке обозначены индексы элементов, которые за- действуются. 0 1 2 3 4 5 6 7 8 9 10 11 0 avg avg avg avg avg avg 1 2 3 4 5 Рис. 3.1. Схема вычисления среднего арифметического В представленном ниже алгоритме предполагается, что elementCount де- лится на 2. Код ядра, вычисляющего среднее арифметическое, показан ниже: 26 __global__ void average( int *inElements, int *outElements, unsigned elementCount ) { // разделяемая память, используемая для сохранения данных __shared__ int elementsBlock[512*2]; // вычисляем смещение для текущего блока const int baseOffsetIn = blockIdx.x * blockDim.x * 2; const int baseOffsetOut = blockIdx.x * blockDim.x; // сохраняем входной массив if ( ( baseOffsetIn + threadIdx.x ) < elementCount ) elementsBlock[threadIdx.x] = inElements[baseOffsetIn + threadIdx.x]; if ((baseOffsetIn+threadIdx.x+blockDim.x) < elementCount ) elementsBlock[threadIdx.x + blockDim.x] = inElements[baseOffsetIn + threadIdx.x + blockDim.x]; // синхронизация данных для гарантии того, что все элементы // были загружены __syncthreads(); // вычисляем среднее арифметическое и сохраняем данные if ( (baseOffsetOut + threadIdx.x) < (elementCount >> 1) ) outElements[baseOffsetOut + threadIdx.x] = ( elementsBlock[threadIdx.x * 2] + elementsBlock[threadIdx.x * 2 + 1] ) >> 1; } Код запуска ядра выглядит следующим образом: dim3 threads(512); dim3 blocks( ( elementCount + 1023 ) / ( 2 * 512 ) ); average<< ); В конфигурации запуска ядра задается 512 нитей на блок (переменная threads ) и динамически вычисляется число блоков (переменная blocks ). При этом выражение ( elementCount + 1023 ) / ( 2 * 512 ) гарантирует пра- вильное число блоков при условии, что остаток от деления elementCount на 1024 не равен нулю. Таким образом, весь массив разбивается на блоки по 1024 элемента в каждом, при этом каждая нить обрабатывает по 2 соседних элемен- 27 та. Для соблюдения условия формирования транзакций (см. подраздел 2.2) чте- ние 1024 элементов разбивается на 2 этапа – последовательно по 512 элементов с сохранением результата в разделяемую память (переменная elementsBlock ). Вызов функции __syncthreads() гарантирует, что все 1024 элемента бы- ли загружены и размещены в разделяемой памяти. Если убрать функцию син- хронизации, существует не нулевая вероятность получить некорректный ре- зультат. После синхронизации осуществляется расчет среднего арифметическо- го и сохранение результата в заданную позицию. Если задействовать свойства варпа, можно реорганизовать ядро таким об- разом, чтобы функция __syncthreads() стала лишней. Преобразуем загрузку элементов следующим образом: вместо блоков 2 по 512 элементов, будем за- гружать блоками 2 по 32 элемента (в соответствии с размером варпа). То есть первых 32 нити загружает последовательно 64 элемента данных, необходимых для вычисления среднего арифметического, следующие 32 нити еще 64 элемен- та, которые необходимы, и так далее. Разбиение на блоки по 32 в соответствии со свойством синхронности нитей варпа гарантирует, что все необходимые данные будут доступны без явной синхронизации. Получившийся код ядра представлен ниже: __global__ void average(int *inElements, int *outElements, unsigned elementCount) { // разделяемая память, используемая для сохранения данных __shared__ int elementsBlock[512*2]; // индекс нити в реорганизованной структуре const int threadId = threadIdx.x & 31; // индекс варпа const int warpId = threadIdx.x >> 5; // вычисляем глобальное смещение для текущего блока const int baseOffsetIn = blockIdx.x * blockDim.x * 2; const int baseOffsetOut = blockIdx.x * blockDim.x; // вычисляем локальное смещение внутри блока const int localOffsetIn = warpId * 32 * 2 + threadId; // в данном примере выходное локальное смещение // совпадает с threadIdx.x, // однако в общем случае это может быть не справедливо const int localOffsetOut = warpId * 32 + threadId; // сохраняем входной массив 28 if ( ( baseOffsetIn + localOffsetIn ) < elementCount ) elementsBlock[localOffsetIn] = inElements[baseOffsetIn + localOffsetIn]; if ( (baseOffsetIn + localOffsetIn + 32) < elementCount ) elementsBlock[localOffsetIn + 32] = inElements[baseOffsetIn + localOffsetIn + 32]; // вычисляем среднее арифметическое и сохраняем данные if ((baseOffsetOut + localOffsetOut) < (elementCount>>1)) outElements[baseOffsetOut + localOffsetOut] = (elementsBlock[localOffsetOut*2] + elementsBlock[localOffsetOut*2+1])>>1; } 3.3. Атомарные операции Атомарные операции представляют собой набор функций, обеспечиваю- щий гарантированное чтение – модификацию – сохранение одного 32- либо 64- битного слова в глобальной или разделяемой памяти. Сохранение при этом происходит по тому же адресу, из которого было произведено чтение данных. В отличие от стандартных операций загрузки, модификации и сохранения, в ато- марных операциях гарантируется: полное выполнение операции. Если, например, при чтении будет сгене- рировано исключение, операция модификации и сохранения выполнена не будет; отсутствие влияния других нитей на результат работы, т.е. другие нити будут вынуждены ожидать полного завершения выполнения данной операции, чтобы прочитать из нее данные. При этом накладываются следующие ограничения: атомарные операции доступны только при исполнении на устройстве (обязательна метка __device__ либо __global__ ); не применимы к памяти выделенной как Mapped page-locked memory. На современных архитектурах графического процессора поддерживаются операции сложения, вычитания, операции над битами. Актуальный список ато- марных операций и их описание, а также версию Compute Capability, начиная с которой данная операция поддерживается, можно найти в официальной доку- ментации [1]. 29 3.4. NVIDIA© Visual Profiler Профилировщик NVIDIA© Visual Profiler является основным средством для анализа производительности разработанного приложения с применением технологии CUDA. Данное приложение доступно для всех основных операци- онных системам (Windows, Linux, Max OS) и входит в стандартный набор NVIDIA© CUDA Toolkit. Настройка приложения (настройка сессии в терминологии CUDA) для под- готовки к анализу сводится лишь к указанию приложения, которое разработчик хочет протестировать, указанию рабочей директории и параметров командной строки при их наличии (рисунок 3.2). В примере на рисунке 3.2 осуществляется запуск приложения SampleTest.exe для тестирования и указана рабочая папка «K:\!Projects\». Результат работы профилировщика с настройками по-умолчанию для те- стового приложения (пример из раздела 3.2) показан на рисунке 3.3. Централь- ную часть экрана приложения занимает линия времени (timeline), на которой отображается последовательность запусков CUDA-функций и ядер с привязкой ко времени. Это позволяет быстро оценить, на какие действия тратится макси- мум времени, чтобы разработчик в дальнейшем мог оптимизировать это. Спра- ва при наведении курсора мыши на ядро на линии времени появляются полу- ченные свойства ядра. Данные свойства позволяют узнать полученную произ- водительность ядра (теоретическую и практическую), число регистров, прихо- дящихся на одну нить и другие параметры. Снизу экрана отображаются все функции-ядра, вызываемые в приложении и их краткая характеристика и, что самое важное, отображается реальное время работы ядра. 30 Рис. 3.2. Создание новой сессии в профилировщике Рис. 3.3. Результат работы профилировщика с настройками по-умолчанию После профилирования с настройками по-умолчанию через контекстное меню Run можно вызвать диалоговое окно, в котором указать характеристики, которые требуется измерить (рисунок 3.4). 31 Рис. 3.4. Диалоговое окно выбора метрик и событий для анализа Все метрики сгруппированы в следующие группы: память, инструкции, мультипроцессор, кэш, текстуры. Данные метрики позволяют оценить эффек- тивность работы с глобальной памятью (например, путем оценки теоретической и эффективной пропускной способности шины), работу мультипроцессора и другие характеристики. Более подробное описание всех метрик можно найти в официальной документации []. Важной составляющей профилировщика является встроенная система ана- лиза производительности приложения (вкладка Analysis). При выборе ядра, ко- торое требуется проанализировать, становится доступными следующие режимы анализа: «Kernel Performance Limiter» – определяет, является ли вычислительное устройство ограничителем производительности; 32 «Kernel Latency» – определяет, насколько качественно осуществляется сокрытие латентности операций работы с памятью и арифметических операций; «Kernel Compute» – оценка утилизации мультипроцессоров ядра ви- деокарты; «Kernel Memory» – анализ эффективной пропускной способности памя- ти; «Memory Access Pattern» – анализ шаблона обращения к памяти; «Divergent Execution» – анализ наличия дивергенции потоков; «Data Movement and Concurrency» – оценка эффективности перекрытия вычислительной нагрузки с функциями копирования данных (cudaMemcpy); «Compute Utilization» – анализ вычислительной нагрузки; «Kernel Performance» – оценка производительности ядра. После запуска анализа выводятся результаты работы. При этом в случае, если система обнаружила критические проблемы, выдается предупреждение об этом и рекомендации, как устранить данную проблему. На рисунке 3.5 пред- ставлен результат анализа ядра average после удаления из кода __syncthreads() в режиме «Kernel Latency». 33 Рис. 3.5. Результат анализа в режиме «Kernel Latency» Как видно из рисунка 3.5, используемая в тесте конфигурация запуска яв- ляется не эффективной, т.к. не позволяет качественно нагрузить имеющиеся мультипроцессоры видеокарты. При этом система рекомендовала увеличить число блоков, используемых в CUDA-ядре. 3.5. Базовые алгоритмы с использованием CUDA К базовым относятся следующие алгоритмы: редукция и операция вычис- ления префиксной суммы (scan). Редукция – сумма значений всех элементов массива. Графически схема простейшего решения данной задачи представлена на рисунке 3.6. В качестве входного массива для алгоритма используется разделяемая память. 0 1 2 3 4 5 6 7 0 1 2 3 Входной массив: Индекс нити: 1 1 5 3 9 5 13 7 Входной массив: Синхронизация 0 1 Индекс нити: 6 1 5 3 22 5 13 7 Входной массив: Синхронизация 0 Индекс нити: 28 1 5 3 22 5 13 7 Входной массив: Синхронизация Результат работы Рис. 3.6. Алгоритм редукции. Базовая версия Код алгоритма редукции показан ниже: __global__ void reduceBase(int* inputData, int* outputData) { extern __shared__ int smemData[]; // загружаем данные в разделяемую память const unsigned tid = threadIdx.x; smemData[tid] = inputData[blockIdx.x * blockDim.x+ threadIdx.x]; 34 // синхронизируемся __syncthreads(); // выполняем редукцию над элементами массива for( unsigned s = 1; s < blockDim.x; s <<= 1 ) { if( tid % ( s << 1 ) == 0 ) { smemData[tid] += smemData[tid + s]; } // синхронизируемся, чтобы гарантировать // корректность входных данных // на следующей итерации цикла __syncthreads(); } // сохраняем результат вычислений if ( tid == 0 ) outputData[blockIdx.x] = smemData[0]; } Анализ кода на производительность указывает на несколько ключевых проблем в ядре: дивергенция потоков при выполнении условия: tid%(s<<1) == 0 ; возможно нарушение доступа по банкам разделяемой памяти: smemData[tid] += smemData[tid + s] ; большое число вызовов функции синхронизации __syncthreads() Для решения описанных проблем представим схему вычисления как пока- зано на рисунке 3.7. 35 0 1 2 3 4 5 6 7 0 1 2 3 Входной массив: Индекс нити: 4 6 8 10 4 5 6 7 Входной массив: Синхронизация Индекс нити: 12 16 8 10 4 5 6 7 Входной массив: Синхронизация Индекс нити: 28 1 5 3 4 5 6 7 Входной массив: Синхронизация Результат работы 0 1 0 Рис. 3.7. Алгоритм редукции. Оптимизированная версия Последовательное обращение нитей к памяти позволяет решить проблему доступа к банкам разделяемой памяти. Кроме того появляется возможность сразу провести первое суммирование при чтении данных из глобальной памяти, тем самым убрав одну синхронизацию. Для этого достаточно правильной настройки конфигурации и чтение данных из глобальной памяти должно осу- ществляться с шагом blockDim.x , т.е. sum = inputData[threadIdx.x] + inputData[threadIdx.x + blockDim.x] . Данную оптимизацию читателю предлагается реализовать самостоятельно. Более детальное описание алгоритма и сравнение производительности можно найти в книге [4]. Префиксной суммой называется массив следующего вида: }} { },..., { }, { , , 0 { 2 2 1 0 2 1 0 1 0 0 n a a a a a a a a a a , где – любое арифметическое действие. Простейшая версия последовательной реализации алгоритма пре- фиксной суммы сводится к следующему: sum[0] = 0; for (int i = 0; i < size; i++) sum[i] = sum[i – 1] a[i – 1]; В качестве примера в дальнейшем выбрано арифметическое действие – сумма. Параллельная реализация является не тривиальной задачей и сводится к 36 двум подзадачам – построение дерева сумм (рисунок 3.8) и на основе получен- ного дерева построение префиксных сумм (рисунок 3.9). 0 1 2 3 4 5 6 7 + + + + Входной массив: 0 1 2 5 4 9 6 13 Входной массив: Синхронизация 0 1 2 6 4 9 6 22 Входной массив: Синхронизация d = 2 0 1 2 6 4 9 6 28 Входной массив: Синхронизация Результат работы + + + d = 0 d = 1 Рис. 3.8. Построение дерева сумм 0 1 2 6 4 9 6 28 Входной массив: 0 1 2 6 4 9 6 0 Входной массив: Синхронизация 0 1 2 0 4 9 6 6 Входной массив: Синхронизация d = 1 0 0 2 1 4 6 6 15 Входной массив: Синхронизация + + zero d = 0 + d = 2 0 0 1 3 6 10 15 21 Входной массив: Синхронизация + + + + 37 Рис. 3.9. Построение префиксных сумм по дереву сумм Задачей первой части алгоритма является нахождение сумм элементов массива (по-аналогии с редукцией). После этого обнуляется последний элемент, содержащий сумму всех элементов. Далее начинается обработка пар, начиная с пары элементов, удаленных на половину длины блока и заканчивая парами со- седних элементов. На каждом шаге один из элементов пары копируется на ме- сто второго, а на место первого записывается сумма исходных элементов. Код функции-ядра, реализующий данный алгоритм, представлен ниже: const int blockSize = 256; __global__ void scan( float *inData, float *outData, int size) { __shared__ float temp[2 * blockSize]; int tid = threadIdx.x; int offset = 1; // чтение данных temp[tid] = inData[tid]; temp[tid + blockSize] = inData[tid + blockSize]; // построение дерева сумм for (int d = n >> 1; d > 0; d >>= 1) { __syncthreads(); if ( tid < d ) { int ai = offset * ( 2 * tid + 1 ) – 1; int bi = offset * ( 2 * tid + 2 ) – 1; temp[bi] += temp[ai]; } offset <<= 1; } // обнуление одного элемента if ( tid == 0 ) { temp[n – 1] = 0; } // построение префиксных сумм по дереву сумм for ( int d = 1; d < n; d <<= 1 ) { 38 offset >>= 1; __syncthreads(); if ( tid < d ) { int ai = offset * ( 2 * tid + 1 ) – 1; int bi = offset * ( 2 * tid + 2 ) – 1; float t = temp[ai]; temp[ai] = temp[bi]; temp[bi] += t; } } __syncthreads(); outData[2 * tid] = temp[2 * tid]; outData[2 * tid + 1] = temp[2 * tid + 1]; } Основной проблемой данного алгоритма являются постоянные конфликты доступа к банкам памяти вплоть до 16 порядка. Читателю предлагается само- стоятельно подумать над тем, как устранить данную проблему. |