Compare commits

...

4 Commits

View File

@@ -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*{Заключение}