Практикум по программированию

Размещено на http://allbest.ru

Размещено на http://allbest.ru

Федеральное государственное автономное образовательное учреждение

высшего образования

Санкт-Петербургский политехнический университет Петра Великого

Институт компьютерных наук и технологий

Высшая школа «Киберфизические системы и управление»

Реферат

по дисциплине «Практикум по Программированию»

Выполнил:

студент гр. з3530902/00001

Кузнецов Р.О.

Руководитель:

ст. преподаватель

Жиленков А.А.

Санкт-Петербург

2022

N-тело

Любая симуляция N-body - это симуляция динамической системы, развивающейся под воздействием физических сил. Численная аппроксимация выполняется по мере того, как тела непрерывно взаимодействуют друг с другом. Моделирование N-body широко применяется в физике и астрономии, например, для того, чтобы ученые могли понять динамику частиц во Вселенной. Моделирование N-тел используется во многих других областях, включая вычислительную гидродинамику для понимания моделирования турбулентных потоков жидкости.

Относительно простой метод решения моделирования N-тел заключается в использовании метода грубой силы, техника, которая имеет O() сложность. Этот подход неудобен по своей параллельной природе. Существуют различные оптимизации на уровне алгоритмов, которые могут уменьшить сложность вычислений. сложность. Вместо того чтобы применять метод всех пар ко всей симуляции, его можно использовать для определения сил во взаимодействиях на близких расстояниях. Даже в этом случае создание ядра для решения сил на CUDA очень полезно, так как это также улучшит производительность компонентов дальнего поля. Ускорение одного компонента разгружает работу других компонентов, поэтому все приложение выигрывает от ускорения одного ядра.

Реализация моделирования N тел на GPU

Алгоритм в основном представляет собой алгоритм расчета силы, , для сетки из тел. Общая сила/ускорение, , на тело, i, является результатом суммирования всех записей в строке i. С точки зрения параллелизма, это неловко параллельная задача O().

С точки зрения производительности, приложение ограничено памятью и будет ограничено пропускной способностью памяти. Хорошо то, что большая часть данных может быть использована повторно и храниться в памяти с высокой пропускной способностью и низкой задержкой, такой как общая память. Повторное использование и хранение данных в общей памяти снижает нагрузку на глобальную память и, следовательно, помогает достичь пиковой производительности вычислений.

На следующей диаграмме показана стратегия, которую мы будем использовать:

Вместо того чтобы снова и снова загружать память из глобальной памяти, мы используем плитку. На предыдущей диаграмме показано, что каждая строка вычисляется параллельно. Размер тайла определяется максимальным количеством элементов, которые можно хранить в общей памяти, что не влияет на загрузку ядра. Каждый блок загружает данные в общую память, после чего выполняется синхронизация. Как только данные были загружены в общую память, расчет силы/ускорения выполняется в каждом блоке. Видно, что хотя отдельный ряд вычисляется параллельно, для достижения оптимального повторного использования данных, взаимодействие в каждом ряду выполняется последовательно.

Обзор реализации симуляции N-body

Рассмотрим реализацию этого в формате псевдокода, а затем объясним ее логику. В этом примере мы используем гравитационный потенциал, чтобы проиллюстрировать основную форму вычислений в моделировании всех пар N-тел. Реализованный код можно найти в 07_parallel_programming_pattern/05_n-body. Для начала работы выполняем следующие шаги:

1. Инициализируем n-пространство случайными переменными

Data [i] = 2.0f * (rand()/max ) - 1.0f

2. Хранение данных в промежуточном разделяемом пространстве памяти для эффективного повторного использования. Синхронизируем его, чтобы гарантировать, что все потоки внутри блока видят обновленные значения в общей памяти:

for (int tile = 0; tille < gridDim.x; tile++) {

..._shared_ float3 shared_position[blockDim.x];

float4 temp_position = p[tile * blockDim.x + threadIdx.x];

shared_position [threadIdx.x]

= make_float3(temp_position.x, temp_position.y, temp_position.z);

_syncthreads();

...

}

3. Рассчитываем силу путем итерации каждого блока:

for (int j = 0; j < BLOCK_SIZE; j++) {

// Calculate Force

_synccthreads();

}

4. Наконец, скомпилируем приложение с помощью компилятора OWDD с помощью следующей командой:

$nvcc -run --gpu-architecture=sm_70 -o n-body n_body.cu

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

Вычисление гистограммы

В неловко параллельной работе, в идеале, вы должны поручить вычисления каждому потоку работающим с независимыми данными, что исключает гонки данных. К настоящему времени мы поняли, что некоторые модели не подходят под эту категорию. Один из таких примеров - когда мы вычисляем гистограмму. Шаблон гистограммы отображает частоту элемента данных, например, количество раз, когда мы использовали слово CUDA в каждой главе, количество раз, когда каждая буква встречалась в этой главе, и так далее. Гистограмма имеет следующий вид:

симуляция динамическая система гистограмма вычисление атомарные функции

В этом разделе использую атомарные операции для сериализации доступа к данным, чтобы получить правильные результаты.

Этапы компиляции и выполнения

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

1. Подготавливаю приложение для GPU. Этот код можно найти

07_parallel_programming_pattern/08_histogram.

2. Скомпилирую приложение с помощью компилятора OWDD с помощью следующей команды:

$ nvcc -c scrImagePgmPpmPackage.cpp

$ nvcc -c image_histogram.cu

$ nvcc -run -o image_histogram image_histogram.o

scrImagePgmPpmPackage.o

ScrImagePgmPpmPackage.cpp предоставляет исходный код, который мы можем использовать для чтения и записи изображений с расширениями QHN. Код расчета гистограммы можно найти в image_histogram.cu.

Понимание параллельной гистограммы

Такие паттерны, как гистограмма, требуют атомарных операций, что означает обновление значения по определенному адресу в последовательном режиме, чтобы устранить разногласия между несколькими потоками, тем самым обновляя один и тот же адрес. Это требует координации между несколькими потоками.

Приватизация - это техника, которая использует память с низкой задержкой, такую как общая память, для снижения пропускной способности и уменьшения задержки, как показано на следующей диаграмме:

По сути, вместо использования атомарных операций в глобальной памяти, мы используем атомарные операции над общей памятью. Причина очевидна. Атомарные операции в глобальной памяти более дорогостоящие по сравнению с выполнением тех же операций в общей памяти/кэша L1. Начиная с архитектуры Maxwell и далее, атомарные операции поддерживаются аппаратно. Реализация приватизированной общей памяти в идеале должна давать 2х-кратную производительность, начиная с архитектуры Maxwell и далее. Однако атомарные операции ограничены определенными функциями и размерами данных.

Вычисление гистограммы с помощью атомарных функций

CUDA функции

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

1. Выделяю общую память на блок, равный размеру гистограммы на блок. Поскольку это изображение в формате char элементы будут находиться в диапазоне 0-255:

_shared_ unsigned int histo_private [256] ;

2. Инициализируйте массив общей памяти для каждого блока:

if (localId <256)

histo_private [localId] = 0 ;

3. Синхронизируйте это, чтобы убедиться, что все потоки в блоке видят инициализированный массив:

_syncthreads () ;

4. Считывание данных изображения из глобальной памяти текстур:

unsigned char imageData = text2D<unsigned

char> (texObj, (float) (tidX) , (float) (tidY) ) ;

5. Выполните операцию AtomicAdd () операцию над общей памятью:

atomicAdd (&(histo_private [imageData] ), 1) ;

6. Синхронизация по блоку перед записью в глобальную память:

_syncthreads ();

7. Запишите гистограмму на блок в глобальную память:

if (localId <256)

imageHistogram [histStartIndex+localId] =

histo_private [localId] ;

Теперь мы закончили реализацию вычисления гистограммы на GPU. Подводя итог, можно сказать, что гистограммы легко реализовать с помощью общей атомарной памяти. Этот подход может достичь высокой производительности на картах Maxwell и далее благодаря встроенной аппаратной поддержке разделяемой атомарной памяти.

Переборка в CUDA с использованием динамического параллелизм

Одним из ключевых алгоритмов, который является фундаментальным строительным блоком для любого приложения, является сортировка. Существует множество алгоритмов сортировки, которые были подробно изучены. Некоторые алгоритмы сортировки относятся к категории "разделяй и властвуй". Эти алгоритмы подходят для параллелизма и подходят для таких архитектур, как GPU где данные, подлежащие сортировке, могут быть разделены для сортировки. Одним из таких алгоритмов является Quicksort. Он представляет собой трехэтапный подход, который заключается в следующем:

1. Выберите элемент из массива, который необходимо отсортировать. Этот элемент действует как поворотный элемент.

2. Вторым шагом является разбиение на разделы, куда попадают все элементы. Все элементы, которые меньше поворотного элемента, сдвигаются влево, а все элементы, которые больше или равны стержню, сдвигаются вправо от стержневого элемента. Этот шаг также известен как разбиение на разделы.

3. Рекурсивно выполняйте шаги 1 и 2, пока все подмассивы не будут отсортированы.

Наихудшая сложность Quicksort составляет O(), что может показаться не идеальным по сравнению с другими процессами сортировки. процессами сортировки, сложность которых в наихудшем случае равна O() ), такими как сортировка слиянием и y heap sort). Однако на практике видно, что Quicksort эффективен. Выбор поворотного элемент может быть выбран с учетом, а иногда и случайно, так что сложность в худшем случае сложность практически не возникает. Кроме того, видно, что Quicksort меньше загружает память и имеет меньшие по сравнению с другими алгоритмами сортировки, такими как сортировка слиянием, которая требует дополнительного хранилища. В более практических реализациях Quicksort используется рандомизированная версия. На сайте рандомизированная версия имеет ожидаемую временную сложность O( ). В худшем случае сложность также возможна в рандомизированной версии, но она не возникает для определенного шаблона (например, отсортированного массива), и рандомизированная версия Quicksort хорошо работает на практике. Хотя можно написать целую главу об особенностях алгоритма сортировки, я планирую рассмотреть только те возможности CUDA, которые помогут эффективно реализовать Quicksort на GPU. В этом разделе мы будем использовать динамический параллелизм, который был введен начиная с CUDA 6.0 и GPU с архитектурой 3.5 и выше. Итак, рассмотрим, какой вклад вносит динамический параллелизм в алгоритм сортировки.

Quicksort и динамический параллелизм CUDA

Алгоритм Quicksort требует рекурсивного запуска ядер. До сих пор алгоритмы, которые мы рассматривали, вызывали ядро один раз через CPU. После того как ядро завершает выполнение, мы возвращаемся к потоку CPU и снова запускаем его. Это приводит к передаче управления CPU, а также может привести к передаче данных между CPU и GPU, что является дорогостоящей операцией. Раньше было очень сложно эффективно реализовать такие алгоритмы, как Quicksort на графических процессорах, которые требуют таких функций, как рекурсия. С архитектурой GPU 3.5 и CUDA 5.0 была введена новая функция, называемая динамическим параллелизмом. Динамический параллелизм позволяет потокам внутри ядра запускать новые ядра с GPU без возврата управления обратно. GPU без возврата управления обратно на CPU. Слово динамический происходит от того, что что он динамически основывается на данных времени выполнения. Несколько ядер могут быть запущены потоками одновременно. Следующая диаграмма упрощает это объяснение:

Если бы мы перевели эту концепцию на то, как выполняется Quicksort, то это выглядело бы примерно так примерно так:

Глубина 0 - это вызов от центрального процессора. Для каждого подмассива мы запускаем два ядра: одно для левого и одно для правого массивов. массива и одно для правого массива. Рекурсия останавливается после того, как максимальная глубина ядра будет или количество элементов становится меньше 32, что является размером искривления. Для того чтобы ядро запуск ядра должен быть в ненулевом потоке и асинхронным, чтобы ядро подмассива запускалось независимо, нам нужно создать поток перед каждым запуском ядра:

cudaStream_t s ;

cudaStreamCreatWithFlags (&s, cudaSteamNonBlocking ) ;

cpd_simple_quicksort <<< 1, 1, 0, s>>> (data, left, nright, depth+1);

cudaStreamDestroy ( s ) ;

Это действительно важный шаг, потому что в противном случае запуск ядра может быть сериализован. Для получения более подробной информации о потоках обратитесь к ядру multi-GPU.

Размещено на Allbest.ru