Compare commits
4 Commits
850ae49a14
...
5b0f9f1904
| Author | SHA1 | Date | |
|---|---|---|---|
| 5b0f9f1904 | |||
| 447ad9c0b8 | |||
| d0b593fa92 | |||
| 4ee74c099f |
@@ -539,6 +539,190 @@ K40. Ниже приведены вычислительные возможнос
|
|||||||
|
|
||||||
Волновой алгоритм либо находит кратчайший путь от начальной к конечной точке, либо информирует о неудаче, если путь к конечной точке блокируется препятствиями.
|
Волновой алгоритм либо находит кратчайший путь от начальной к конечной точке, либо информирует о неудаче, если путь к конечной точке блокируется препятствиями.
|
||||||
|
|
||||||
|
\subsection{Метод распараллеливания алгоритма}
|
||||||
|
Рассмотрим два подхода к распараллеливанию этой задачи на CUDA: с использованием глобальной памяти и с использованием разделяемой (shared) памяти.
|
||||||
|
|
||||||
|
В обоих подходах восстановление пути происходит на CPU и не подвергается распараллеливанию.
|
||||||
|
|
||||||
|
\subsubsection{Глобальная память}
|
||||||
|
Используется одномерная сетка блоков и одномерная сетка потоков внутри блоков. Таким образом уникальный идентификатор потока вычисляется по формуле:
|
||||||
|
|
||||||
|
\texttt{tid = threadIdx.x + blockIdx.x * blockDim.x}
|
||||||
|
|
||||||
|
Каждый поток обрабатывает несколько ячеек исходной матрицы с шагом \texttt{blockDim.x * gridDim.x} (общее количество потоков во всех блоках). Каждая клетка исходной матрицы обрабатывается только одним потоком.
|
||||||
|
|
||||||
|
На каждом шаге алгоритма каждый поток пытается уменьшить значения своих клеток. Поток обращается к глобальной памяти, чтобы получить значения соседних клеток. Значение текущей клетки обновляется, если значение одной из соседних клеток, увеличенное на единицу, меньше текущего значения клетки. То есть по следующей формуле:
|
||||||
|
|
||||||
|
$$
|
||||||
|
cell = \min(cell, left + 1, right + 1, top + 1, bottom + 1),
|
||||||
|
$$
|
||||||
|
где $cell$ -- значение в текущей клетке, $left$ -- значение в соседней слева, $right$ -- значение в соседней справа, $top$ -- значение в соседней сверху, $bottom$ -- значение в соседней снизу. Клетки с препятствиями игнорируются и не участвуют в формуле.
|
||||||
|
|
||||||
|
Следующий шаг алгоритма запускается, если хотя бы одному потоку удалось обновить значение в какой-нибудь своей клетке. Алгоритм завершается, если больше не осталось незаполненных клеток и дальнейшее обновление расстояний невозможно. Для отслеживания изменений используется флаг, который также хранится в глобальной памяти.
|
||||||
|
|
||||||
|
\subsubsection{Разделяемая память}
|
||||||
|
Используется одномерная сетка блоков. Потоки внутри блоков организованы в двумерную сетку.
|
||||||
|
|
||||||
|
Каждый блок потоков обрабатывает несколько соразмерных ему подматриц исходной матрицы с шагом, равным количеству используемых блоков. Каждая подматрица обрабатывается только одним блоком. Посколько блок и поматрица имеют равные размеры, каждому потоку блока в соответствие ставится одна ячейка поматрицы. Таким образом каждая ячейка исходной матрицы обрабатывается только одним потоком.
|
||||||
|
|
||||||
|
На каждом шаге алгоритма каждый блок потоков копирует соответствующую подматрицу исходной матриц из глобальной памяти в разделяемую. Затем на подматрице запускается паралелльная версия волнового алгоритма, описанная в предыдущем разделе, с двумя отличиями:
|
||||||
|
\begin{itemize}
|
||||||
|
\item К глобальной памяти обращаются только те потоки, которые обрабатывают клетки на границах подматрицы. Потому что некоторые соседи граничных клеток находятся вне скопированной в разделяемую память подматрицы. Такой подход позволяет значительно сократить обращения к глобальной памяти при использовании больших размеров блоков.
|
||||||
|
\item Локальный флаг для отслеживания изменений хранится в разделяемой памяти каждого потока.
|
||||||
|
\end{itemize}
|
||||||
|
После выполнения волнового алгоритма на подматрице в разделяемой памяти, её значения копируются в соответствующую подматрицу исходной матрицы в глобальной памяти.
|
||||||
|
|
||||||
|
Таким образом, на каждом шаге алгоритма к волне прибавляются не просто отдельные ячейки, а целые подматрицы, соразмерные используемым блокам потоков.
|
||||||
|
|
||||||
|
Следующий шаг алгоритма запускается, если хотя бы одному блоку потоков удалось обновить значения в своей подматрице. Алгоритм завершается, если больше не осталось незаполненных подматриц и дальнейшее обновление расстояний невозможно. Для отслеживания изменений в подматрицах используется дополнительный флаг, хранящийся в глобальной памяти.
|
||||||
|
|
||||||
|
|
||||||
|
\newpage
|
||||||
|
\section{Описание эксперимента}
|
||||||
|
В этом разделе выполняется исследование времени решения задачи при изменении следующих параметров:
|
||||||
|
|
||||||
|
\begin{itemize}
|
||||||
|
\item Размеры матрицы: 500 × 500, 1000 × 1000, 1500 × 1500;
|
||||||
|
\item Количество блоков: 1, 10, 100, 1000, 10000;
|
||||||
|
\item Количество потоков: 1, 9, 100, 1024;
|
||||||
|
\item Используемая память: глобальная, глобальная и разделяемая.
|
||||||
|
\end{itemize}
|
||||||
|
|
||||||
|
Для каждой комбинации параметров было проведено 100 измерений с использованием событий CUDA, после чего вычислялось среднее значение.
|
||||||
|
|
||||||
|
\newpage
|
||||||
|
\section{Анализ результатов}
|
||||||
|
В таблицах 1, 2, 3 приведены результаты измерения времени в миллисекундах для глобальной памяти для размеров матрицы 500 × 500, 1000 × 1000, 1500 × 1500 соответственно.
|
||||||
|
|
||||||
|
\begin{table}[h!]
|
||||||
|
\centering
|
||||||
|
\caption{Результаты измерения времени исполнения программы для матрицы 100 × 100 и глобальной памяти. Время указано в миллисекундах.}
|
||||||
|
\footnotesize
|
||||||
|
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||||||
|
\hline
|
||||||
|
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||||||
|
\hline
|
||||||
|
& 1 & 10 & 100 & 1000 & 10000 \\
|
||||||
|
\hline
|
||||||
|
1 & 171.85 & 120.98 & 19.41 & 12.84 & 23.68 \\
|
||||||
|
\hline
|
||||||
|
9 & 233.35 & 29.13 & 8.14 & 5.68 & 10.55 \\
|
||||||
|
\hline
|
||||||
|
100 & 27.57 & 7.79 & 3.90 & \textbf{3.79} & 8.97 \\
|
||||||
|
\hline
|
||||||
|
1024 & 7.96 & 4.04 & 4.16 & 6.47 & 32.24 \\
|
||||||
|
\hline
|
||||||
|
\end{tabularx}
|
||||||
|
\end{table}
|
||||||
|
|
||||||
|
\begin{table}[h!]
|
||||||
|
\centering
|
||||||
|
\caption{Результаты измерения времени исполнения программы для матрицы 500 × 500 и глобальной памяти. Время указано в миллисекундах.}
|
||||||
|
\footnotesize
|
||||||
|
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||||||
|
\hline
|
||||||
|
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||||||
|
\hline
|
||||||
|
& 1 & 10 & 100 & 1000 & 10000 \\
|
||||||
|
\hline
|
||||||
|
1 & 13124 & 15096 & 1694 & 1027 & 1140 \\
|
||||||
|
\hline
|
||||||
|
9 & 33257 & 4075 & 554 & 291 & 223 \\
|
||||||
|
\hline
|
||||||
|
100 & 3643 & 536 & 93 & 55 & 77 \\
|
||||||
|
\hline
|
||||||
|
1024 & 596 & 88 & \textbf{49} & 64 & 196 \\
|
||||||
|
\hline
|
||||||
|
\end{tabularx}
|
||||||
|
\end{table}
|
||||||
|
|
||||||
|
\begin{table}[h!]
|
||||||
|
\centering
|
||||||
|
\caption{Результаты измерения времени исполнения программы для матрицы 1000 × 1000 и глобальной памяти. Время указано в миллисекундах.}
|
||||||
|
\footnotesize
|
||||||
|
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||||||
|
\hline
|
||||||
|
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||||||
|
\hline
|
||||||
|
& 1 & 10 & 100 & 1000 & 10000 \\
|
||||||
|
\hline
|
||||||
|
1 & 89324 & 121341 & 12760 & 7010 & 6472 \\
|
||||||
|
\hline
|
||||||
|
9 & 264810 & 29369 & 4435 & 2208 & 1356 \\
|
||||||
|
\hline
|
||||||
|
100 & 28946 & 2985 & 628 & 329 & 360 \\
|
||||||
|
\hline
|
||||||
|
1024 & 3701 & 609 & \textbf{277} & 322 & 579 \\
|
||||||
|
\hline
|
||||||
|
\end{tabularx}
|
||||||
|
\end{table}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
|
\begin{table}[h!]
|
||||||
|
\centering
|
||||||
|
\caption{Результаты измерения времени исполнения программы для матрицы 100 × 100 и разделяемой памяти. Время указано в миллисекундах.}
|
||||||
|
\footnotesize
|
||||||
|
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||||||
|
\hline
|
||||||
|
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||||||
|
\hline
|
||||||
|
& 1 & 10 & 100 & 1000 & 10000 \\
|
||||||
|
\hline
|
||||||
|
1 & 260.60 & 230.96 & 47.04 & 18.13 & 32.70 \\
|
||||||
|
\hline
|
||||||
|
9 & 28.65 & 19.92 & 7.27 & 3.10 & 4.38 \\
|
||||||
|
\hline
|
||||||
|
100 & 5.46 & 4.71 & \textbf{1.35} & 1.36 & 1.45 \\
|
||||||
|
\hline
|
||||||
|
1024 & 4.39 & 2.68 & 2.65 & 2.67 & 3.00 \\
|
||||||
|
\hline
|
||||||
|
\end{tabularx}
|
||||||
|
\end{table}
|
||||||
|
|
||||||
|
\begin{table}[h!]
|
||||||
|
\centering
|
||||||
|
\caption{Результаты измерения времени исполнения программы для матрицы 500 × 500 и разделяемой памяти. Время указано в миллисекундах.}
|
||||||
|
\footnotesize
|
||||||
|
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||||||
|
\hline
|
||||||
|
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||||||
|
\hline
|
||||||
|
& 1 & 10 & 100 & 1000 & 10000 \\
|
||||||
|
\hline
|
||||||
|
1 & 18837 & 31345 & 3566 & 1786 & 1724 \\
|
||||||
|
\hline
|
||||||
|
9 & 1120 & 2109 & 347 & 163 & 148 \\
|
||||||
|
\hline
|
||||||
|
100 & 166 & 166 & 65 & 37 & \textbf{17} \\
|
||||||
|
\hline
|
||||||
|
1024 & 99 & 64 & 34 & 112 & 114 \\
|
||||||
|
\hline
|
||||||
|
\end{tabularx}
|
||||||
|
\end{table}
|
||||||
|
|
||||||
|
\begin{table}[h!]
|
||||||
|
\centering
|
||||||
|
\caption{Результаты измерения времени исполнения программы для матрицы 1000 × 1000 и разделяемой памяти. Время указано в миллисекундах.}
|
||||||
|
\footnotesize
|
||||||
|
\begin{tabularx}{\textwidth}{|X|X|X|X|X|X|}
|
||||||
|
\hline
|
||||||
|
\textbf{Число потоков в блоков} & \multicolumn{5}{c|}{\textbf{Число блоков}} \\
|
||||||
|
\hline
|
||||||
|
& 1 & 10 & 100 & 1000 & 10000 \\
|
||||||
|
\hline
|
||||||
|
1 & 127512 & 48471 & 24677 & 14065 & 12445 \\
|
||||||
|
\hline
|
||||||
|
9 & 5775 & 16595 & 2062 & 1138 & 727 \\
|
||||||
|
\hline
|
||||||
|
100 & 656 & 961 & 491 & 124 & \textbf{103} \\
|
||||||
|
\hline
|
||||||
|
1024 & 398 & 190 & 222 & 384 & 618 \\
|
||||||
|
\hline
|
||||||
|
\end{tabularx}
|
||||||
|
\end{table}
|
||||||
|
|
||||||
|
|
||||||
|
|
||||||
\newpage
|
\newpage
|
||||||
\section*{Заключение}
|
\section*{Заключение}
|
||||||
|
|||||||
Reference in New Issue
Block a user