Описание алгоритмов

report/report.tex

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