Методичка (CUDA) - черновик. Введение cuda (Compute Unified Device Architecture)Ошибка! Источник ссылки не найден
 Скачать 1.06 Mb.
|
|
3.6. Задание 1. Оценить с помощью NVIDIA Visual Profiler качество реализации алго- ритма из подраздела 2.8. 2. Оптимизировать алгоритм префиксной суммы с учетом рекомендаций из подраздела 3.5. 39 4. Обработка изображения Обработка графической информации является одним из тех направлений, для которых технология CUDA может обеспечить существенное ускорение. Реализация алгоритмов в области обработки изображений связана с реше- нием ряда существенных проблем, актуальных для текущих архитектур графи- ческого процессора: эффективное распределение нагрузки между мультипроцессорами и вычислительными ядрами; формирование транзакций при чтении и записи данных из/в глобальную память (подраздел 2.1); специфические особенности конкретных алгоритмов (например, про- блема обработки краев при фильтрации изображения); 4.1. Проблема загрузки изображения Одна из проблем, с которой сталкивается разработчик при реализации лю- бого алгоритма обработки изображений, - формирование транзакций при до- ступе к глобальной памяти. При этом можно выделить две случая, которые мо- гут привести к нарушению шаблона доступа: один пиксель представляет собой 3 байта данных (каналы RGB); отсутствие выравнивания в памяти для второй и последующих строк исходного изображения. При чтении данных в ядре из глобальной памяти разработчик должен ре- шить проблему эффективной загрузки в разделяемую память. Проблема заклю- чается в том, что пиксель для большинства алгоритмов представляет собой 3- байтную последовательность данных. Простейший способ загрузки – побайтное чтение данных – приводит к нарушению шаблона доступа к памяти и в резуль- тате транзакции не формируются, т.к. требуется последовательное обращение нитей к 4-байтным данным (подробнее в подразделе 2.1). На современных ар- хитектурах за счет наличия кэша частично данная проблема сглаживается, од- нако это все равно приводит к существенному замедлению работы алгоритма. Эффективное решение данной проблемы сводится к одному из двух реше- ний. Первое решение заключается в добавлении канала альфа к каждому пиксе- лю изображения, однако, с одной стороны, это приводит к увеличению требуе- мой глобальной памяти, а с другой стороны – требуется реализация алгоритма конверсии из 3 байтного формата в 4 байтный и обратно, что увеличивает пол- ное время работы программы. 40 Второй способ эффективного решения проблемы заключается в том, чтобы одна нить загружала не один пиксель, а сразу 4 (рисунок 4.1). Как видно из ри- сунка 4.1, для этого потребуется выполнить 3 загрузки данных типа int R G B R G B R G B R G B 1 пиксель 1 пиксель 1 пиксель 1 пиксель 4 байта 4 байта 4 байта Рис. 4.1. Демонстрация загрузки и сохранения данных Фрагмент кода, выполняющий загрузку в соответствии с рисунком 4.1, по- казан ниже: __global__ void filter( unsigned *input, const unsigned inputWidth, …) { __shared__ unsigned smem[(BLOCK_DIM_X*3)*BLOCK_DIM_Y]; const int baseX = (blockIdx.x*blockDim.x+threadIdx.x)*3; const int baseY = blockIdx.y * blockDim.y + threadIdx.y; smem[threadIdx.y * BLOCK_DIM_X * 3 + threadIdx.x] = input[baseY * inputWidth + baseX]; smem[threadIdx.y*BLOCK_DIM_X*3+threadIdx.x+blockDim.x] = input[baseY * inputWidth + baseX + blockDim.x]; smem[threadIdx.y*BLOCK_DIM_X*3+threadIdx.x+blockDim.x*2] = input[baseY*inputWidth+baseX + blockDim.x * 2]; __syncthreads(); … } В представленном фрагменте кода предполагается, что BLOCK_DIM_X и BLOCK_DIM_Y – число нитей по осям X и Y соответственно конфигурации ядра, inputWidth – ширина изображения в 4-байтных словах. Проблема с негарантированным формированием транзакций при чтении второй и последующих строк заключается в том, что данные строки должны быть выравнены на границу 128 байт. В общем случае при произвольной ши- 41 рине изображения данное условие гарантируется только для первой строки, т.к. функция cudaMalloc осуществляет выравнивание стартового адреса. Решением данной проблемы является искусственное увеличение ширины изображения до границы 128 байт при инициализации памяти устройства (ри- сунок 4.2). В рамках технологии CUDA предлагается 2 способа решения: ручное выравнивание и расширение изображения; использование CUDA API. x x x width h ei g h t pitch Рис. 4.2. Выравнивание строк изображения Фрагмент кода при ручном выравнивании представлен ниже: void cudaInit( char *inputImg, // входное изображение const int width, // ширина в байтах const int height, // высота изображения … ) { … int *d_inputImg = NULL; // выравниванием ширину изображения const int pitch = ( ( width + 127 ) / 128 ) * 128; // выделяем память с учетом новой ширины cudaMalloc( (void **)&d_inputImg, pitch * height ); // копируем данные с расширением на новую ширину for (int i = 0; i < height; i++) { cudaMemcpy( &d_inputImg[i * pitch], //адрес приемника &inputImg[i * width], // адрес источника width, // число байт для копирования 42 cudaMemcpyHostToDevice ); } … } При использовании CUDA API достаточно воспользоваться функцией вы- деления памяти: cudaError_t cudaMallocPitch ( void** devPtr, // указатель на выделенную память size_t* pitch, // фактически сформированная ширина в байтах size_t width, // реальная ширина изображения в байтах size_t height // реальная высота изображения ) В результате выполнения данной функции будет выделено как минимум width * height байт. При этом в переменной pitch будет возвращено факти- чески выделенная ширина изображения в байтах с учетом выравнивания. Для расширения массива до pitch байт CUDA API предлагает использовать следующую функцию: cudaError_t cudaMemcpy2D ( void* dst, // приемник size_t dpitch, // фактическая ширина приемника в байтах const void* src, // источник size_t spitch, // фактическая ширина источника в байтах size_t width, // реальная ширина копируемых данных в байтах size_t height, // реальная высота копируемых данных cudaMemcpyKind kind // направление копирования ) В результате будет выполнено копирование height строк по width байт каждая из src в dst . При этом смещение для строки i приемника будет опреде- ляться как i * dpitch , а для источника – i * spitch Фрагмент кода, демонстрирующий работу с CUDA API, показан ниже: void cudaInitAPI( char *inputImg, // входное изображение char *outputImg, // выходное изображение const int width, // ширина в байтах const int height, // высота изображения … 43 ) { … int *d_inputImg = NULL; int pitch = 0; cudaMalloc((void **)&d_inputImg, &pitch, width, height); cudaMemcpy2D( d_ inputImg, pitch, inputImg, width, width, height, cudaMemcpyHostToDevice ); … // можем работать с d_ inputImg, при этом ширина – pitch байт cudaMemcpy2D( outputImg, width, d_ inputImg, pitch, width, height, cudaMemcpyDeviceToHost ); // в outputImg ширина изображения снова width байт … } 4.2. Проблема границ. Разделение ядер на краевые и центральные Еще одна проблема, с которой сталкивается разработчик, - обработка гра- ницы изображения. В особенности это касается правой и нижней частей изоб- ражения, т.к. требуется обработать возможный выход за пределы обрабатывае- мой области. Несмотря на то, что изображение расширяется до границы 128 байт, тратить ресурсы и обрабатывать участок, содержащий «мусор», появля- ющийся за счет расширения, не всегда имеет смысл. Поэтому в код ядра добав- ляется условный оператор, ограничивающий участок обработки. Для центральных блоков данный оператор не имеет смысла, т.к. все нити блока «примут» одинаковое решение. В результате дивергенция потоков не сформируется, однако ресурсы на обработку условия затрачены будут. Для того чтобы этого избежать, ядро разбивается на две ветки: первая – обрабатывает центральные блоки изображения и исключает все проверочные условные операторы; вторая – обрабатывает только правый и нижний края изображения. 44 4.3. Разбиение на блоки Формирование конфигурации исполнения ядра является важной задачей, т.к. даже идеально оптимизированное ядро при неправильной конфигурации способно существенно замедлить работу. Исходя из практики разработки ядер для обработки цветных изображений, наиболее оптимальным числом нитей по оси X изображения для большинства задач является 32. Данное число нитей позволяет использовать свойства варпа для синхронизации нитей. Обрабатываемое число пикселей в этом случае – 128, что, с одной стороны, позволяет обрабатывать сравнительно небольшие по ши- рине изображения, а с другой стороны увеличить число блоков для более каче- ственной нагрузки имеющихся потоковых мультипроцессоров (Streaming Mul- tiprocessor). Четких критериев для определения числа нитей по оси Y нет и зависит от решаемой задачи и вычислительной сложности ядра. 4.4. Конкатенация ядер 4.5. Задание 1. Свертка изображения с фильтром 3х3 45 5. Оптимизация доступа к памяти и технология CUDA Stream 5.1. Pinned память Как уже отмечалось, операции копирования между ОЗУ и памятью ви- деокарты обладают существенной латентностью. Для оптимизации передачи данных рекомендуется обращение с формированием транзакций, как механиз- ма, позволяющего максимально эффективно задействовать шину памяти. Все современные операционные системы работают в режиме с виртуаль- ной адресацией памяти. Для каждого запущенного процесса выделяется свое адресное пространство, в рамках которого процесс может осуществлять чтение и запись данных. Сторонний процесс не может получить доступ на чте- ние/запись в области памяти текущего процесса без специального разрешения со стороны ОС. В большинстве операционных систем виртуальное адресное пространство делится на страницы по как минимум 4096 байт. Для доступа к физическому адресу данных в процессор встроена таблица PTE (Page Table Entry), осуществ- ляющая трансляцию между виртуальным и физическим адресами. В случае ес- ли запрашиваемая страница не загружена, генерируется исключение, которое должна обработать ОС. Схема работы показана на рисунке 5.1. Страница Инструкции Адрес Таблица страниц Страница (4096 байт) Физическая память Рис. 5.1. Виртуальное адресное пространство 46 Для организации доступа к глобальной памяти на видеокарте используется упрощенный механизм виртуальной адресации – сохраняется безопасность и локальность в памяти для конкретной CUDA программы, однако убрана стра- ничная адресация. Различия в механизме адресации для CPU и GPU привели к тому, что доступ к памяти является сравнительно сложной задачей: определяется номер страницы, содержащей необходимые данные в ОЗУ, и ее наличие в памяти; осуществляется копирование страницы в специальный pinned-буфер; выполняется синхронное копирование страницы в память GPU с ис- пользованием DMA; освобождаются запрошенные ресурсы. Данный алгоритм повторяется для каждой последующей страницы. Очевидно, что из-за большого числа операций приблизиться к пиковой пропускной способности шины PCI Express не представляется возможным. GPU CPU Страничная адресация памяти pinned-память Глобальная память Постраничная адресация памяти GPU CPU pinned-память Глобальная память Pinned-режим Рис. 5.2. Сравнение передачи данных в постраничном и pinned-режимах Для решения данной проблемы в CUDA API предложен способ выделения памяти в ОЗУ непосредственно в pinned-буфере в обход постраничной адреса- ции. С одной стороны, разработчик теряет в объеме доступно памяти ОЗУ и возможность кэшировать данные на диск, а с другой – существенно упрощает механизм копирования из ОЗУ в память GPU и приближается к пиковой про- пускной способности PCI Express. Для выделения памяти в pinned-режиме до- статочно выделить память в ОЗУ с применением следующей функции: cudaError_t cudaMallocHost ( void** ptr, // указатель на выделяемую память в ОЗУ size_t size // объем выделяемой памяти 47 ) ; Заранее выделенная системная память может быть также зарегистрирована в pinned-режиме с использованием CUDA API: cudaError_t cudaHostRegister ( void * ptr, // указатель на память ОЗУ size_t size, // размер регистрируемой памяти в байтах unsigned int flags ) К параметрам управления (к флагам) относятся следующие допустимые значения: cudaHostRegisterPortable – зарегистрированная память переводится в pinned-режим для всех контекстов, а не только для текущего; cudaHostRegisterMapped – проецирует выделенную память в адресное пространство CUDA. Указатель на память в GPU может быть получен с помощью cudaHostGetDevicePointer() Однако следует учитывать, что передача данных может быть ускорена, но в целом система станет работать медленнее, особенно если задействовать в ал- горитмах заранее выделенную системную память и осуществлять ее регистра- цию с использованием cudaHostRegister() 5.2. Унифицированная память (Unified Memory) Данная технология была введена в CUDA 6.0 и поддерживается архитек- турой Kepler и более новыми. Основная задача заключается в том, чтобы пол- ностью скрыть от разработчика различия в CPU и GPU памяти. Рассмотрим следующий пример: CPU CUDA 6 void sampleFunc(FILE *fp, int N) { char *data; data = (char *)malloc(N); fread(data, 1, N, fp); testFunction(data, N); free(data); void sampleFunc(FILE *fp, int N) { char *data; cudaMallocManaged(&data, N); fread(data, 1, N, fp); testFunction<<<…>>>(data, N); cudaFree(data); 48 } } Код слева – код для исполнения на CPU, код справа – для GPU. Основное отличие – в выделении памяти. Для этого используется функция cudaMal- locManaged() By migrating data on demand between the CPU and GPU, Unified Memory can offer the performance of local data on the GPU, while providing the ease of use of globally shared data. The complexity of this functionality is kept under the covers of the CUDA driver and runtime, ensuring that application code is simpler to write. The point of migration is to achieve full bandwidth from each processor; the 250 GB/s of GDDR5 memory is vital to feeding the compute throughput of a Kepler GPU. 5.3. CUDA Stream Достижение максимальной производительности возможно только при мак- симальной утилизации всех имеющихся ресурсов (в случае с технологией CUDA – CPU и GPU). Данный принцип лежит в основе управления нитями на GPU: как только нити варпа переходят в состояние ожидания получения ре- зультата, планировщик пытается найти варп, который может начать либо про- должить выполнение следующих операций, и переключается на него. В резуль- тате при правильном проектировании конфигурации системы возможно до- стигнуть идеального баланса производительности для GPU, когда задержки из- за ожидания выполнения операции сводятся к минимуму. Кроме того асинхронный запуск ядра позволяет одновременно с расчетами задействовать ресурсы CPU. Однако данный подход не позволяет полностью перекрыть операции копирования с хоста на устройство и наоборот. Для реше- ния данной проблемы NVIDIA© предлагает использовать CUDA Stream. CUDA Stream представляет собой очередь команд для GPU, для которого справедливы следующие правила: команды в рамках одного stream выполняются строго последовательно; команды из разных stream могут выполняться по мере освобождения необходимых ресурсов, т.е. возможно параллельное исполнение команд из разных stream. По-умолчанию все команды выполняются в 0 stream’е, который создается автоматически системой. Стандартная последовательность команд (копирова- ние с хоста на устройство, выполнение ядра, копирование результата с устрой- ства на хост) при использовании только одного stream’а показана на рис.5.3. Используемый при этом код выглядит следующим образом: Примечание [D5]: Как их назвать на русском? 49 cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice); kernel<<<1, N>>>(d_a); cpuFunction(); cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost); Копирование с хоста на устройство Исполнение ядра Копирование с устройства на хост Время Рис. 5.3. Исполнение при использовании stream’а по-умолчанию Основная проблема в данной ситуации – простаивание GPU при копирова- нии данных и простаивание шины при исполнении ядра. Для использования различных stream’ов требуется реорганизовать код сле- дующим образом: ядро kernel для каждого stream’а должно обрабатывать свой локальных участок данных полностью независимо по отношению к остальным stream’ам (например, при обработке изображений – каждый stream об- рабатывает свой рисунок); функции копирования требуется заменить на их асинхронные версии; инициализировать необходимое число stream’ов. Инициализация stream’а выглядит следующим образом (все представлен- ные команды будут относиться к одному stream’у): cudaStream_t stream1; // создаем stream1 cudaStreamCreate(&stream1); // асинхронное копирование на устройство cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1); // запуск ядра kernelStream<<<1, N, 0, stream1>>>(d_a); … // дальнейшая работа со stream1, CPU // разрушаем определенный stream1 cudaStreamDestroy(stream1); 50 Для определения того, что использование CUDA Stream может обеспечить повышение производительности, на этапе инициализации рекомендуется вы- полнить следующие действия: проверить флаг deviceOverlap в структуре cudaDeviceProp . В соот- ветствии с документации все устройства с Compute Capability 1.1 и но- вее данные флаг установлен в истинное значение; память на хосте должна быть определена как pinned. Для синхронизации, например, со stream1 введена функция cudaStreamSynchronize(stream1) , которая блокирует хост до тех пор, пока не будет полностью выполнены все команды для stream1 . Для того чтобы узнать завершены команды stream’а или нет без блокировки хоста используется функ- ция cudaStreamQuery(stream) Новая версия кода с использование nStreams stream’ов показана ниже: for (int i = 0; i < nStreams; ++i) { int offset = i * streamSize; cudaMemcpyAsync( &d_a[offset], &a[offset], streamBytes, cudaMemcpyHostToDevice, stream[i] ); kernel<<>>(d_a, offset); cudaMemcpyAsync( &a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i] ); } Копирование с хоста на устройство stream 0 Исполнение ядра stream 0 Копирование с устройства на хост stream 0 Время Копирование с хоста на устройство stream 1 Исполнение ядра stream 1 Копирование с устройства на хост stream 1 Рис. 5.4. Перекрытие операций с использование различных stream’ов 51 Для представленного кода возможно обеспечить перекрытие, показанное на рис. 5.4, начиная с архитектуры NVIDIA© Fermi и более новых. Это связано с аппаратной организацией системы очередей. Изначально в архитектуре NVIDIA© Tesla содержалось две очереди – одна для копирования, одна для ис- полнения ядер. Команды на исполнение поступают в очередь в порядке их сле- дования по коду, в результате в очередь на копирование попадает вначале ко- манда копирования с хоста на устройство для stream[0] , затем с устройства на хост для этого же stream’а и так далее для всех nStream stream’ов. Система вы- полняет команды последовательно, в результате очередь блокируется из-за необходимости копирования с устройства на хост, а для этого требуется вы- полнить ядро. В результате для архитектуры NVIDIA© Tesla данный код не даст никакого повышения производительности. Начиная с архитектуры NVID- IA© Fermi, число буферов копирования было расширено до двух устройств (одно для копирования с хоста на устройство, второе – в обратном направле- нии), что определяет эффективность работы алгоритма. Использование stream’ов будет заблокировано или прервано, если: переменная окружения CUDA_LAUNCH_BLOCKING = 1; выделение page-locked памяти; выделение и инициализация памяти на устройстве; любая команда из CUDA API, работающая с нулевым stream’ом (stream’ом по-умолчанию); переключение между режимами, управляющими L1кэшем и разделяе- мой памятью для Compute Capability 2.x и Compute Capability 3.x. 5.4. Задание Оптимизировать алгоритм из задания раздела 4 для использования техно- логии CUDA Stream. Примечание [D6]: Перед Fermi ??? 52 6. Список литературы [1] CUDA C Programming Guide [Электронный ресурс]. – 2014. – Ре- жим доступа: http://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_ Guide.pdf. [2] BrookGPU [Электронный ресурс]. – 2014. – Режим доступа: http://graphics.stanford.edu/projects/brookgpu/. [3] The open standard for parallel programming of heterogeneous sys- tems[Электронный ресурс]. – 2014. – Режим доступа: http://www.khronos.org/opencl/. [4] Параллельные вычисления на GPU. Архитектура и программная модель CUDA / А. Боресков [и др.]. – М. : Изд-во МГУ, 2012. – 336 с. [5] Wilt, N. The CUDA Handbook: A Comprehensive Guide to GPU Pro- gramming / N. Wilt. – Addison Wesley, 2013. – 528 p. [6] Cook, S. CUDA Programming (Applications of GPU Computing Series) / S. Cook. – Morgan Kaufmann, 2012. – 560 p. [7] Farber, R. CUDA Application Design and Development / R. Farber. – Morgan Kaufmann, 2011. – 336 p. [8] Kirk, D. Programming Massively Parallel Processors: A Hands-on Ap- proach / D. Kirk, Wen-Mei Hwu ; 2nd edition. – Morgan Kaufmann, 2012. – 514 p. |