% !TEX TS-program = xelatex
% !TEX encoding   = UTF-8 Unicode
% !TEX spellcheck = el-GR
%
% PDS homework_3 report
%
% Requires compilation with pdfLaTeX or XeLaTeX
%
% authors:
%   Χρήστος Χουτουρίδης ΑΕΜ 8997
%   cchoutou@ece.auth.gr

%
% Options:
%
% 1) mainlang=<language>
%    Default:  english
%    Set the default language of the document which affects hyphenations,
%    localization (section, dates, etc...)
%
%    example: \documentclass[mainlang=greek]{AUThReport}
%
% 2) <language>
%    Add hyphenation and typesetting support for other languages
%    Currently supports: english, greek, german, frenc
%
%    example: \documentclass[english, greek]{AUThReport}
%
% 3) short: Requests a shorter title for the document
%    Default: no short
%
%    example: \documentclass[short]{AUThReport}
%
\documentclass[a4paper, 11pt, mainlang=greek, english]{AUThReport}

\CurrentDate{\today}

% Document setup
%---------------------------------

% \WorkGroup{Ομάδα Χ}

\AuthorName{Χρήστος Χουτουρίδης}
\AuthorMail{cchoutou@ece.auth.gr}
\AuthorAEM{8997}

%\CoAuthorName{Όνομα Επίθετο}
%\CoAuthorAEM{AEM}
%\CoAuthorMail{xxx@ece.auth.gr}

\DocTitle{Εργασία 3: Bitonic sort with CUDA}
\DocSubTitle{Ανάλυση και υλοποίηση του αλγόριθμου bitonic sort για υποσύστημα γραφικών με χρήση CUDA}

\Department{Τμήμα ΗΜΜΥ. Τομέας Ηλεκτρονικής}
\ClassName{Παράλληλα και Διανεμημένα Συστήματα}
%
\InstructorName{Νικόλαος Πιτσιάνης}
\InstructorMail{nikos.pitsianis@eng.auth.gr}

%\CoInstructorName{Όνομα Επίθετο}
%\CoInstructorMail{xxx@ece.auth.gr}


% Local package requirements
%---------------------------------

\usepackage{enumitem}
\usepackage{tabularx}
\usepackage{array}
\usepackage{multirow}
\usepackage{float}
\usepackage{xcolor}
\usepackage{soul}
\usepackage{amsmath}
\usepackage{footnote}
\usepackage{footmisc}


%\definecolor{red-highlight}{rgb}{1.0, 0.11, 0.0}
\definecolor{red-highlight}{rgb}{1.0, 0.31, 0.0}
\definecolor{green-highlight}{rgb}{0.0, 0.62, 0.42}

\newcommand{\bxrd}[1]{\colorbox{red-highlight}{#1}}
\newcommand{\bxgr}[1]{\colorbox{green-highlight}{#1}}

\newcommand{\trd}[1]{\textcolor{red-highlight}{\textbf{#1}}}
\newcommand{\tgr}[1]{\textcolor{green-highlight}{\textbf{#1}}}

% “virtual” cites
% ----------------------------------------
\newcommand{\citeKnuthACM}{%
    Donald E. Knuth, \emph{Structured Programming with go to Statements}, ACM Computing Surveys, vol. 6, no. 4, 1974
}

% Links
\newcommand{\hwTwo}{https://git.hoo2.net/hoo2/PDS/src/branch/master/homework_2}
\newcommand{\hwThree}{https://git.hoo2.net/hoo2/PDS/src/branch/master/homework_3}
\newcommand{\repoAnalyse}{https://git.hoo2.net/hoo2/PDS/src/branch/master/homework_3/analyse}
\newcommand{\repoProf}{https://git.hoo2.net/hoo2/PDS/src/branch/master/homework_3/analyse/prof.sh}

\newcommand{\bitonicSortWikipedia}{https://en.wikipedia.org/wiki/Bitonic_sorter}

\begin{document}

\InsertTitle

%\tableofcontents


\section{Εισαγωγή}

Η παρούσα εργασία αποτελεί συνέχεια της \href{\hwTwo}{προηγούμενης υλοποίησης} του ίδιου αλγόριθμου σε κατανεμημένα συστήματα με τη χρήση MPI.
Στην τρέχουσα έκδοση ο αλγόριθμος καλείται να εκτελεστεί σε υποσυστήματα γραφικών (GPU) με χρήση CUDA.
Τα υποσυστήματα γραφικών μας δίνουν την δυνατότητα για πολύ μεγάλο αριθμό παραλληλοποίησης, κάτι που αποτελεί σημαντική διαφοροποίηση σε σχέση με την MPI έκδοση.
Πρακτικά ο τρόπος με τον οποίο χρησιμοποιείται το υποσύστημα γραφικών σε αυτή την εργασία είναι σαν accelerator συγκεκριμένων συναρτήσεων.
Η υλοποίηση του αλγόριθμου έγινε σε τρεις φάσεις ακολουθώντας τις υποδείξεις της εκφώνησης.
Στην πρώτη φάση, το υποσύστημα γραφικών  απλώς επιταχύνει το κύριο σώμα του αλγόριθμου, ενώ στις επόμενες έγινε προσπάθεια τροποποίησης αυτού ώστε να εκμεταλλευτούμε τα χαρακτηριστικά και τους πόρους του υποσυστήματος γραφικών, με αποδοτικότερο τρόπο.


\subsection{Παραδοτέα}
Τα παραδοτέα της εργασίας αποτελούνται από:
\begin{itemize}
    \item Την παρούσα αναφορά.
    \item Το \href{\hwThree}{σύνδεσμο με το αποθετήριο} που περιέχει τον κώδικα για την παραγωγή των εκτελέσιμων, της αναφοράς και τις μετρήσεις.
\end{itemize}

\section {Υλοποίηση}
Πριν ξεκινήσουμε με τις λεπτομέρειες του αλγόριθμου, καλό θα ήταν να περιγράψουμε τις βασικές συναρτήσεις και δομές που χρησιμοποιούνται στην εργασία, ώστε να βοηθήσουμε στην καλύτερη κατανόηση της υλοποίησης.
\begin{itemize}
    \item \texttt{\textbf{Log}}:
    Πρόκειται για ένα τύπο ο οποίος μας δίνει logging δυνατότητες τις οποίες μπορεί ο χρήστης να ενεργοποιήσει από τη γραμμή εντολών με την παράμετρο \texttt{-v} ή \texttt{--verbose}.
    \item \texttt{\textbf{Timing}}:
    Η τάξη αυτή προσφέρει δυνατότητες χρονομέτρησης μίας ή και πολλαπλών κλήσεων αθροίζοντας τους χρόνους (accumulation).
    Επίσης προσφέρει δυνατότητες για πολλαπλές εκτελέσεις ολόκληρης της ταξινόμησης όπου οι χρόνοι μπορούν να ταξινομηθούν και να επιστραφεί ο ενδιάμεσος.
    Ο χρήστης μπορεί να επιλέξει την εμφάνιση των αποτελεσμάτων αλλά και τον αριθμό των εκτελέσεων από τη γραμμή εντολών με την παράμετρο \texttt{--perf}.
    \item \texttt{\textbf{bitonicSort()}}:
    Πρόκειται για την βασική συνάρτηση που υλοποιεί τον αλγόριθμο της εργασίας και είναι σε μορφή template για διαφορετικούς τύπους δεδομένων προς ταξινόμηση.
    Έχει υλοποιηθεί 3 φορές, μία για κάθε έκδοση της εκφώνησης V0, V1, V2 στο αρχείο \texttt{bitonicsort.hpp}.
    \item \texttt{\textbf{inBlockStep()/interBlockStep()}}:
    Ειδικά για τις εκδόσεις V1 και V2 όπου χρησιμοποιείται loop unrolling, το κύριο σώμα του αλγόριθμου bitonic sort έχει χωριστεί σε δύο συναρτήσεις.
    Μία που εκτελείτε εκτός του unrolling (\textit{interBlockStep()}) και μία που εκτελεί το ίδιο το unrolling (\textit{inBlockStep()}).
\end{itemize}
\par
Όλοι οι παραπάνω τύποι είναι templates, και βρίσκονται στα αντίστοιχα headers (.hpp) αντί για αρχεία .cpp.
Η συνάρτηση \textit{main()} που υλοποιεί τον αλγόριθμο μαζί με τον validator και το command line interface βρίσκονται στο αρχείο main.cpp.
Στο αρχείο config.h υπάρχουν οι compile time ρυθμίσεις της υλοποίησης όπου ο χρήστης μπορεί να επιλέξει για παράδειγμα τον τύπο δεδομένων προς ταξινόμηση.

\subsection{Πρώτη έκδοση (V0)} \label{sec:V0}
Η πρώτη έκδοση υλοποιεί τον βασικό αλγόριθμο επιταχύνοντας απλώς το κύριο σώμα του στην κάρτα γραφικών.
Καθώς θεωρούμε πως ο αναγνώστης είναι ήδη εξοικειωμένος με την διτονική ταξινόμηση, δεν θα αναλύσουμε την λειτουργία της σε μεγάλο βάθος.
Για πίνακα μεγέθους $2^N$ ο αλγόριθμος συνοψίζεται στα εξής:
\begin{itemize}
    \item Εκτελούνται $N$ ακολουθίες ανταλλαγών με αύξον αριθμό: $m = 1, 2,\dots,N$.
    \item Η κάθε ακολουθία $m$ αποτελείται από ανταλλαγές μεταξύ γειτόνων, των οποίων η απόσταση ξεκινάει από $2^{m-1}$ και μειώνεται με διαδοχικές ακέραιες διαιρέσεις με το 2, εωσότου γίνει $1$ ($2^{m-1}, 2^{m-2}, ..., 1$).
    \item Οι ανταλλαγές χωρίζουν τα μεγαλύτερα και τα μικρότερα στοιχεία με στόχο στο τέλος της κάθε ακολουθίας τα στοιχεία του πίνακα να αποτελούν διαδοχικές διτονικές ακολουθίες.
    Μέτά την πρώτη ακολουθία ανταλλαγών να έχουμε $\frac{N}{2}$ διτονικές ακολουθίες, μετά την δεύτερη $\frac{N}{4}$ ακολουθίες και ούτω κάθε εξής, έως ότου στην τελευταία να έχουμε $\frac{N}{2N} = \frac{1}{2}$ διτονική, δηλαδή μια πλήρως ταξινομημένη λίστα.
\end{itemize}
\par
\InsertFigure{!ht}{0.8}{fig:bitonicSort}{img/BitonicSort.png}{4 ακολουθίες ανταλλαγών για ταξινόμηση 16 στοιχείων.}
Η εικόνα \ref{fig:bitonicSort} παρουσιάζει αυτή τη λογική\footnote{Πηγή \href{\bitonicSortWikipedia}{wikipedia}}.

\subsubsection{Πλήθος απαιτούμενων διεργασιών (threads) και καταμερισμός στον πίνακα} \label{sec:V0addressing}
Στην παρούσα εργασία, επιθυμούμε οι ανταλλαγές να λάβουν χώρα εντός του υποσυστήματος γραφικών και να εκτελεστούν από τις διεργασίες (threads) του υποσυστήματος γραφικών.
Σε αντίθεση όμως με την προηγούμενη εργασία, όπου για $M$ υποσύνολα του πίνακα είχαμε $M$ MPI διεργασίες και όλες απαιτούνταν να εκτελέσουν διαχωρισμό μεγίστων -- ελαχίστων, \textbf{εδώ η κάθε διεργασία μπορεί να διαχωρίσει 2 στοιχεία}.
\textbf{Δηλαδή για ταξινόμηση ενός πίνακα $N$ στοιχείων απαιτούνται $\frac{N}{2}$ διεργασίες}.
\par
Βλέπουμε δηλαδή πως δεν είναι απαραίτητο να αναθέσουμε μία διεργασία σε κάθε στοιχείο του πίνακα.
Με αυτό σαν βάση προβήκαμε στον πρώτο καταμερισμό των διεργασιών στα στοιχεία του πίνακα.
Έτσι χωρίσαμε τον πίνακα σε δύο μέρη.
Στο πρώτο μέρος η διευθυνσιοδότηση έγινε ακολουθώντας την αρίθμηση των threads, ενώ για το δεύτερο το διπλάσιο αυτής.
\par
Όπως γίνεται φανερό \textit{έπρεπε να βρεθεί ένας τρόπος ώστε οι μισές διεργασίες να διευθυνσιοδοτήσουν το πρώτο μέρος και οι άλλες μισές το δεύτερο}.
Ο διαχωρισμός έγινε ελέγχοντας τον γείτονα (partner) για την κάθε ανταλλαγή.

\InsertFigure{!ht}{0.85}{fig:V0Addressing}{img/V0Addressing.png}{Διευθυνσιοδότηση των threads (V0).}

Η λογική πίσω από αυτό τον έλεγχο είναι ότι αν ένα thread με διεύθυνση $Tid_1$ έχει να ανταλλάξει μέ έναν γείτονα με διεύθυνση $Pid=Tid_2$, τότε το thread που αρχικά πέφτει στη διεύθυνση $Tid_2$ έχει να ανταλλάξει με το $Tid_1$.
Η ανταλλαγή όμως αυτή έχει δρομολογηθεί στο $Tid_1$ και έτσι το συγκεκριμένο thread μπορεί να μεταφερθεί στο δεύτερο μισό του πίνακα στην αντίστοιχη θέση $Tid_2'$ και να ανταλλάξει με τον γείτονα που προκύπτει σε εκείνη τη θέση $Pid'$.
Ο γείτονας σε εκείνη τη θέση φυσικά είναι ο αντίστοιχος του $Tid_1$.
Το διάγραμμα \ref{fig:V0Addressing} παρουσιάζει αυτόν τον καταμερισμό.
\par
Έχοντας κατανέμει λοιπόν τα threads στον πίνακα, η υλοποίηση (\textit{bitonicSort()}) εκτελεί το διπλό βρόχο που περιγράψαμε παραπάνω και \textbf{αναθέτει το κάθε βήμα το οποίο εκτελεί τις ανταλλαγές στη κάρτα γραφικών}.
Το βήμα αυτό αντιστοιχεί στις \textbf{κάθετες γραμμές} του σχήματος \ref{fig:bitonicSort}.
Αυτό γίνεται με την κλήση του CUDA kernel \textit{bitonicStep\texttt{<<<}Nbl, Nth\texttt{>>>}()}, όπου δημιουργούμε \boldmath $N_{bl}$ blocks από $N_{th}$, τέτοια ώστε $N_{bl} \cdot N_{th} = \frac{N}{2}$  \unboldmath threads, τα οποία εκτελούν το σώμα της συνάρτησης.
Ο χρήστης μπορεί να επιλέξει τον αριθμό των threads ανά block (ή διαφορετικά blocksize) από τη γραμμή εντολών με το όρισμα \texttt{-b} ή \texttt{--block-size}.
Έτσι το κάθε thread εκτελεί μια ανταλλαγή με την άνω διευθυνσιοδότηση και ο αλγόριθμος τερματίζει στο τέλος του διπλού βρόχου.
Το πλεονέκτημα αυτής της μεθόδου είναι η απλότητά της, κάτι που προφανώς πληρώνουμε με έναν αρκετά μεγάλο αριθμό κλήσεων προς την κάρτα γραφικών το οποίο προσθέτει overhead, όπως θα δούμε και παρακάτω.

\subsection{Δεύτερη έκδοση (V1)} \label{sec:V1}
Από την εκφώνηση της εργασίας ήδη υπάρχουν οδηγίες για βελτιστοποιήσεις.
Η πρώτη βελτιστοποίηση λοιπόν αφορά αυτές ακριβώς τις κλήσεις των kernels, των συναρτήσεων δηλαδή που επιταχύνονται από την κάρτα γραφικών.
Γίνεται δηλαδή η \textbf{θεώρηση} από την εκφώνηση, ότι \textbf{η ελαχιστοποίηση αυτών των κλήσεων θα βελτιστοποιήσει τη διαδικασία}.
Φυσικά, αυτό μένει πάντα να αποδειχτεί στην πράξη.
\par
Ο τρόπος που προτείνεται ανήκει στην κατηγορία βελτιστοποιήσεων που λέγεται loop unrolling, όπου ένα μέρος του βρόχου αφαιρείται από τον έλεγχο του βρόχου και καλείται "με το χέρι".
Στη δική μας περίπτωση το μέρος που αφαιρείται μπορεί να μπει μέσα στην κλήση kernel αξιοποιώντας καλύτερα τους πόρους που μπορεί να χρησιμοποιήσει, μειώνοντας έτσι και τον αριθμό κλήσεων.
\par
Ενώ όμως στην έκδοση V0 η κάθε κλήση kernel αφορούσε ένα στιγμιότυπο (μια κάθετη γραμμή) του αλγόριθμου, εδώ οι κλήσεις αφορούν περισσότερα.
Ο κάθε kernel δηλαδή θα εκτελέσει τμήματα περισσοτέρων της μίας γραμμής του σχήματος \ref{fig:bitonicSort}.
Αυτό όμως δημιουργεί data race, καθώς οι τιμές αυτών των στοιχείων διαμοιράζονται μεταξύ των διαφορετικών κλήσεων.
\par
Για να λύσουμε αυτό το πρόβλημα εργαστήκαμε έτσι ώστε:
\begin{itemize}
	\item Να \textbf{περιορίσουμε} το εύρος των δεδομένων που "βλέπει" ο κάθε kernel.
	\item Να \textbf{συγχρονίσουμε} τις κλήσεις εσωτερικά του kernel στο κάθε στιγμιότυπο.
\end{itemize}
\par
Για το λόγο αυτό αρχικά αλλάξαμε τον τρόπο της διευθυνσιοδότησης από την έκδοση V0.
Η διευθυνσιοδότηση του πίνακα πλέον δεν ακολουθεί την αρίθμηση των thread, αλλά αντίθετα τα threads δείχνουν μόνο στα ζυγά blocks αφήνοντας ελεύθερο "ένα -- παρά -- ένα" block.
Δεδομένου όμως ότι ο αριθμός των threads ενός block αρκεί για να ανταλλάξει 2 blocks του πίνακα, μεταθέτουμε τα πλεονάζοντα threads απόσταση ένα block.
Με αυτό τον τρόπο έχουμε περιορίσει τα threads μέσα στις δυάδες των block.
Έτσι το εύρος διευθύνσεων που διευθυνσιοδοτείται από τα threads του block είναι όσο πιο κοντά γίνεται.
Το σχήμα \ref{fig:V1Addressing} παρουσιάζει αυτή τη μεθοδολογία.

\InsertFigure{H}{0.85}{fig:V1Addressing}{img/V1Addressing.png}{Διευθυνσιοδότηση των threads (V1).}

\par
Επομένως, για να λύσουμε το πρόβλημα του data race πρέπει να βεβαιωθούμε πως η κάθε κλήση του kernel αφορά αποστάσεις που βρίσκονται εντός των δυάδων των blocks.
Όπως είδαμε όμως και στην παράγραφο \ref{sec:V0} για την m-οστή ακολουθία, η απόσταση των ανταλλαγών μεταξύ των γειτόνων ξεκινάει από $2^{m-1}$, όπου $m-1$ είναι το αρχικό βήμα του εσωτερικού βρόχου (έστω $S_0$).
Άρα για το βήμα $S_i$ του εσωτερικού βρόχου από το οποίο οι ανταλλαγές βρίσκονται εσωτερικά των δυάδων των block αρκεί:
\[
	BlockSize = N_{th} \ge 2^{m_i-1} = 2^{S_i} \Rightarrow \log_2(N_{th}) \ge S_i
\]

\WrapFigure{0.35}{r}{fig:V1Unrolling}{img/V1Unrolling.png}{Βήματα διτονικής ταξινόμησης πίνακα 128 θέσεων (block size = 16).}
Όπου το $S_i=\log_2(N_{th})$ είναι και το μέγιστο βήμα από το οποίο μπορεί να ξεκινήσει το loop unrolling.
\par
Στο σχήμα \ref{fig:V1Unrolling} φαίνεται ένα παράδειγμα ταξινόμησης πίνακα 128 θέσεων και block size = 16, όπου τα βήματα με σκιαγράφηση μπορούν να υλοποιηθούν μέσα στη συνάρτηση kernel.
Μπορεί κανείς να παρατηρήσει πως τα βήματα με αποστάσεις $16, 8, 4, 2, 1$ είναι τα βήματα $4, 3, 2, 1, 0$ όπου $\log_2(N_{th}) = \log_2(16) = 4$.
\par
Για να υλοποιήσουμε αυτή την προσέγγιση, σπάσαμε την \textit{bitonicStep()} σε δύο εκδόσεις.
Η μία (\textit{interBlockStep()}) καλείται για αποστάσεις μεγαλύτερες από δύο block size και η άλλη (\textit{inBlockStep()})για ίσες και μικρότερες.
Επίσης, η δεύτερη υλοποιεί και το μέρος του βρόχου που αναλογεί στα βήματα εντός kernel.
Ο παρακάτω κώδικας υλοποιεί αυτή τη προσέγγιση:
\begin{verbatim}
size_t Nth = config.blockSize;
size_t Nbl = NBlocks(size);

auto Stages          = static_cast<size_t>(log2(size));
auto InnerBlockSteps = static_cast<size_t>(log2(Nth));
for (size_t stage = 1; stage <= Stages; ++stage) {
    size_t step = stage - 1;
    for ( ; step > InnerBlockSteps; --step) {
        interBlockStep<<<Nbl, Nth>>>(dev_data, size, step, stage);
        cudaDeviceSynchronize();
    }
    inBlockStep<<<Nbl, Nth>>>(dev_data, size, step, stage);
    cudaDeviceSynchronize();
}
\end{verbatim}

\par
Το μόνο που μένει πλέον είναι το πρόβλημα του \textbf{συγχρονισμού}.
Εδώ χρειάζεται να κρατήσουμε τον συγχρονισμό που είχαμε σε κάθε κάθετη γραμμή του σχήματος \ref{fig:bitonicSort}, απλώς αυτός θα λάβει χώρα εσωτερικά του kernel.
Η συνάρτηση που μας δίνει αυτή τη δυνατότητα είναι η \textit{\_\_syncthreads()}, την οποία και καλούμε μετά από την ολοκλήρωση του κάθε βήματος.
\par
Όπως είναι φανερό ο διαχωρισμός μεταξύ των βημάτων που μπορούν να γίνουν εσωτερικά και αυτών εξωτερικά του kernel έχει άνω όριο, αλλά όχι κάτω.
Στην υλοποίησή μας παρόλα αυτά δεν δίνουμε τη δυνατότητα ρύθμισης αυτού του ορίου.
Αυτό γιατί ούτως ή άλλως υπάρχει ρύθμιση για το μέγεθος του block, το οποίο και χρησιμοποιούμε για να κάνουμε και αυτό το διαχωρισμό.
Έτσι \textbf{αν ο χρήστης επιλέξει block size τότε όλες οι ανταλλαγές που χωρούν σε αυτό το μέγεθος θα λάβουν χώρα εντός του kernel}.

\subsection{Τρίτη έκδοση (V2)} \label{sec:V2}
Η δεύτερη βελτιστοποίηση που προτείνεται από την εκφώνηση αφορά τη \textbf{χρήση της κοινής μνήμης} μεταξύ των threads στο ίδιο block.
Η μνήμη αυτή (shared memory) μπορεί, ανάλογα πάντα και με την αρχιτεκτονική, να είναι μία ή και δύο τάξεις μεγέθους γρηγορότερη από την global.
Ενώ βέβαια αυτό φαντάζει \textbf{ειδυλλιακό}, θα πρέπει να τονίσουμε ότι για την συγκεκριμένη βελτιστοποίηση θα χρειαστούν αρκετές παραχωρήσεις.
Για αυτό το λόγο είδαμε αυτή την λύση με \textbf{αρκετό σκεπτικισμό}.

Τα βασικά βήματα που πρέπει να ακολουθήσουμε ώστε να χρησιμοποιήσουμε την κοινή μνήμη είναι:
\begin{itemize}
    \item \textbf{Αντιγραφή} των στοιχείων που ανταλλάσσονται από τα threads του block στη shared μνήμη.
    \item \textbf{Προσαρμογή της διευθυνσιοδότησης} των threads εντός του block ώστε να αντιστοιχούν στα όρια της shared μνήμης.
    \item \textbf{Εκτέλεση} του αλγόριθμου της έκδοσης \textbf{V1} στην κοινή μνήμη με τις προσαρμοσμένες διευθύνσεις
    \item \textbf{Αντιγραφή} των στοιχείων πίσω στη global μνήμη.
\end{itemize}
Το σχήμα \ref{fig:V2Addressing} παρουσιάζει αυτή την προσέγγιση.
\InsertFigure{H}{0.85}{fig:V2Addressing}{img/V2Addressing.png}{Διευθυνσιοδότηση των threads (V2).}
\par
Βλέπουμε λοιπόν πως \textbf{ο αλγόριθμος είναι πρακτικά ο ίδιος}, απλώς \textbf{οι ανταλλαγές} στα βήματα εντός του block \textbf{λαμβάνουν χώρα στη shared μνήμη} αντί της global.
Επίσης πολύ εύκολα μπορούμε να αντιληφθούμε πως οι απαιτήσεις σε shared μνήμη είναι δυο φορές το μέγεθος των blocks.
Τη μνήμη αυτή τη δεσμεύουμε δυναμικά μέσω ορίσματος στον kernel \textit{inBlockStep\texttt{<<<}Nbl, Nth, kernelMemSize\texttt{>>>}()}.
\par
Η αλλαγή για την δεύτερη έκδοση βρίσκεται επί της ουσίας στη συνάρτηση \textit{inBlockStep()}.
Η συνάρτηση \textit{interBlockStep()} παραμένει ως έχει, καθώς εργάζεται με αποστάσεις μεγαλύτερες αυτών που χωρούν στη shared μνήμη που έχουμε δεσμεύσει για τα blocks.

\section{Επιδόσεις και βελτιστοποίηση} \label{sec:performance}

Όπως φαίνεται και από το σχήμα \ref{fig:V2Addressing}, στην δεύτερη έκδοση, ενώ χρησιμοποιούμε την shared μνήμη για τις ανταλλαγές εντός block, αυτό \textbf{δεν φαίνεται να περιορίζει την πρόσβαση στη global μνήμη}.
Μάλιστα η πρόσβαση είναι πλέον ντετερμινιστική.
Το κάθε thread κάνει 2 αναγνώσεις και 2 εγγραφές, ενώ στην έκδοση V1 οι εγγραφές για παράδειγμα γίνονται μόνο όταν χρειάζεται.
Επίσης η λογική που εκτελείται εντός thread είναι περισσότερη καθώς πρέπει να γίνουν οι αντιγραφές αλλά και η προσαρμογή των διευθύνσεων.
Θα πρέπει λοιπόν το κέρδος από τις γρήγορες ανταλλαγές στη shared μνήμη να είναι πολύ μεγάλο, ούτως ώστε να αντισταθμίσουμε αυτό το κόστος.
Αυτός είναι και ο λόγος που όπως αναφέραμε βλέπουμε την έκδοση V2 με σκεπτικισμό.

\par
Αρχικά για να διερευνήσουμε τις επιδόσεις εκτελέσαμε τοπικά σε μία GTX 1650 τις εκδόσεις και είδαμε ότι η έκδοση V2 υστερεί κατά τι.
Έπειτα επιβεβαιώσαμε τις ανησυχίες μας στη συστοιχία όπου στην ampere A100 πάλι η V2 ήταν πιο αργή και μόνο στην Tesla P100 πιο γρήγορη.
Για να διαπιστώσουμε όμως επακριβώς τι συμβαίνει χρησιμοποιήσαμε το εργαλείο \href{\repoProf}{ncu} ώστε να πάρουμε το προφίλ των εκτελέσεων.
Οι έξοδοι του εργαλείου είναι διαθέσιμοι στους καταλόγους RC1 και RC2 στο αποθετήριο της εργασίας \href{\repoAnalyse}{εδώ}.
\par
Ενδιαφέρον παρουσιάζουν οι εξής μετρήσεις για την \textit{inBlockStep()} που αποτελεί και την διαφορά μεταξύ των εκδόσεων:
\begin{table}[H]
    \centering
    \renewcommand{\arraystretch}{1.25}
    \begin{tabular}{l | r r c}
                                         \textbf{Metric} & \textbf{V1} & \textbf{V2} & \textbf{diff}  \\ \hline
        gpu time duration.sum                            &      184 μs & \trd{230 μs}& +25\%   \\
        average sectors per request-mem global ld.ratio  &        3.47 &        4.00 & - \\
        average sectors per request-mem global st.ratio  &        4.71 &        4.00 & - \\
        mem shared cmd read.sum                          &           0 &        525k & - \\
        mem shared cmd write.sum                         &           0 &        314k & - \\
        smps average warp latency stalled barrier.pct    &       159\% & \tgr{116\%} &  -27.0\% \\
        smsp cycles active.avg                           &        253k &  \trd{315k} &  +24.5\% \\
        smsp cycles active.sum                           &       16.2M & \trd{20.2M} &  +24.7\% \\
        smsp inst executed.avg                           &        132k &  \trd{189k} &  +43.2\% \\
        smsp inst executed.sum                           &       8.46M & \trd{12.1M} &  +43.0\% \\
    \end{tabular}
    \caption{Σύγκριση της inBlockStep() των εκδόσεων V1 και V2.}
    \label{tab:profiling_v1v2}
\end{table}

Πέραν του ότι βλέπουμε ξεκάθαρα ότι ο χρόνος της συνάρτησης είναι μεγαλύτερος δικαιολογώντας τις μετρήσεις, παρατηρούμε πως:
\begin{itemize}
    \item Οι \textbf{αναγνώσεις/εγγραφές από/προς την global μνήμη} ναι μεν έγιναν ντετερμινιστικές (ratio 4 ακριβώς, όσες και οι προσβάσεις που αναλογούν στο thread), αλλά η έκδοση V1 είχε ήδη πολύ καλή συμπεριφορά, καθώς είχε γίνει προσπάθεια ώστε η πρόσβαση στη μνήμη να είναι όσο μειωμένη γίνεται.
    Μάλιστα το read ήταν και καλύτερο.
    Έτσι ουσιαστικά δεν μπορέσαμε να κερδίσουμε κάτι από εδώ.
    \item Εκτός από την παραπάνω πρόσβαση στη μνήμη, \textbf{προστέθηκε η πρόσβαση στη shared μνήμη} η οποία είναι σε σημαντικό ποσοστό ($525k$ αναγνώσεις και $314k$ εγγραφές).
    \item \textbf{Ο αριθμός των κύκλων αλλά και των εντολών} που εκτελούνται από τα τους streaming multiprocessors (SMs) έχει εκτοξευτεί κατά \boldmath $25\%$ και $43\%$ \unboldmath αντίστοιχα.
    Αυτό θεωρούμε ότι συμβαίνει πάλι γιατί η έκδοση V1 είχε αρκετά προσεγμένο κώδικα με σκοπό τη μείωση των εντολών που εκτελούνται.
    Έτσι η έστω και μικρή αλλαγή για να γίνουν οι αντιγραφές από και προς την global αλλά και η προσαρμογή των διευθύνσεων αποδεικνύεται σημαντική.
\end{itemize}

Με βάση τα παραπάνω αποτελέσματα συμπεραίνουμε πως για να βελτιωθεί η έκδοσή μας χρειάζεται:
\begin{itemize}
    \item Να καταφέρουμε με \textit{κάποιο τρόπο} να \textbf{μειώσουμε τις εντολές}.
    \item Να κάνουμε την \textbf{εκτέλεση των εντολών} που αφορά τις ανταλλαγές στο σώμα της \textit{inBitonicStep()} \textbf{ακόμα πιο γρήγορη}.
    \item Να αναζητήσουμε κάποιου \textbf{άλλου είδους βελτιστοποίηση} που δεν έχει να κάνει με το shared memory optimization.
\end{itemize}
Αν και καταβάλαμε φιλότιμη προσπάθεια να μειώσουμε τις εντολές, τελικά αυτή η κατεύθυνση δεν μας οδήγησε πουθενά.
Αντίθετά για τα άλλα δύο σημεία προσπαθήσαμε κάποιες λύσεις.

\subsection{Ελαχιστοποίηση των αρχικών κλήσεων της \textit{inBlockStep()} - Prephase} \label{sec:Opt_prephase}

Μια βελτιστοποίηση που σκεφτήκαμε και δεν έχει να κάνει με το shared memory optimization είναι ο αριθμός των \textit{inBlockStep()} στην αρχή του αλγόριθμου.
Αν προσέξει κανείς το σχήμα \ref{fig:V1Unrolling} εύκολα θα παρατηρήσει πως στην αρχή του αλγόριθμου ο διπλός βρόχος περιέχει βήματα που χωράνε εξολοκλήρου στο block.
Παρόλα αυτά γίνονται ξεχωριστές κλήσεις στην \textit{inBlockStep()} χωρίς να μεσολαβεί κάποια κλήση στην \textit{interBlockStep()}.
Μπορούμε έτσι να \textbf{συνενώσουμε όλες αυτές τις πρώτες ξεχωριστές κλήσεις της \textit{inBlockStep()} σε μία συνάρτηση} η οποία δεν χρειάζεται να επιστρέφει παρά μόνο όταν τα βήματα του βρόχου φτάσουν να χρειαστεί να καλέσουν την πρώτη \textit{interBlockStep()}.
Η συνάρτηση αυτή είναι η \textit{prephase()}.
\par
Πριν αποδεχτούμε βέβαια αυτή τη βελτιστοποίηση έπρεπε πρώτα να την μετρήσουμε.
Η έκδοση στην οποία δοκιμάστηκε αυτή η προσέγγιση είναι η RC2 και ο αναγνώστης μπορεί να τη βρει στο αποθετήριο της εργασίας.
Ο πίνακας παρακάτω δείχνει τα αποτελέσματα που πήραμε από το profiling εργαλείο ncu, για prephase που αντικαθιστά 28 κλήσεις \textit{inBlockStep()}, αλλά και κάποιους ενδεικτικούς χρόνους εκτέλεσης από τη συστοιχία:
\begin{table}[H]
    \centering
    \renewcommand{\arraystretch}{1.25}
    \begin{tabular}{l | r r c}
                                         \textbf{Metric} & \textbf{InBlockStep()} & \textbf{prephase()} \\ \hline
        gpu time duration.sum                            &                 184 μs &            1.06 ms  \\
        average sectors per request-mem global ld.ratio  &                   3.47 &               3.91  \\
        average sectors per request-mem global st.ratio  &                   4.71 &               5.48  \\
        smsp inst executed.avg                           &                   132k &               770k  \\
        smsp inst executed.sum                           &                  8.46M &              49.3M  \\ \hline
           \textbf{Sorting Time} & \textbf{No-Opt (RC1)} & \textbf{Prephase-Opt (RC2)} \\ \hline
        V1Q27 - A100             &            156 [msec] &  153 [msec] \\
        V1Q30 - A100             &           1580 [msec] & 1566 [msec] \\
    \end{tabular}
    \caption{Σύγκριση των εκδόσεων \textit{InBlockStep() - RC1} και \textit{prephase() - RC2}.}
    \label{tab:profiling_prephase}
\end{table}
Από τον παραπάνω πίνακα συμπεραίνουμε πως η αλλαγή βελτιώνει τόσο τη συμπεριφορά όσο και τους χρόνους εκτέλεσης.
Στο \href{\repoAnalyse}{αποθετήριο} της εργασίας υπάρχουν όλες οι μετρήσεις και τα αποτελέσματα του profiling τόσο για την έκδοση RC1 (πριν τη βελτιστοποίηση) όσο και για την RC2 (μετά τη βελτιστοποίηση).

\subsection{Ελαχιστοποίηση του χρόνου εκτέλεσης της \textit{inBlockStep()}} \label{sec:Opt_registers}

Παρατηρώντας τον αλγόριθμο της έκδοσης V2 μπορεί κάποιος να παρατηρήσει ότι η πρόσβαση στη global μνήμη ακολουθεί τα εξής.
Το κάθε thread:
\begin{enumerate}
    \item Διαβάζει δύο θέσεις μνήμης(πίνακα) από τη global
    \item Γράφει δύο θέσεις πίνακα στη shared
    \item Διαβάζει δύο θέσεις από τη shared
    \item Γράφει με μία πιθανότητα δύο θέσεις στη shared
    \item Διαβάζει δύο θέσεις από τη shared
    \item Γράφει δύο θέσεις στη global
\end{enumerate}
Συνολικά global: 2RD/2WR και shared: 4RD/(2 ή 4)WR.
Αν κοιτάξει κάποιος τα αποτελέσματα από το profiling στον πίνακα \ref{tab:profiling_v1v2} όπου οι αναγνώσεις/εγγραφές στη shared είναι RD:525k/WR:314k θα δει ότι επιβεβαιώνεται αυτός ο συλλογισμός.
\par
Η ιδέα αφορά ουσιαστικά το βήμα 3.
Αντί να διαβάζουμε από τη θέση του shared πίνακα που μόλις γράψαμε, \textbf{θα μπορούσαμε να κρατάμε την τιμή του τρέχοντος thread που διαβάσαμε από την global, σε ένα register}.
Έπειτα η σύγκριση με τον κάθε γείτονα θα γίνεται μεταξύ register και shared και μόνο όταν γίνεται ανταλλαγή, η τιμή του register θα επαναγράφεται στη shared.
Αυτό το trick μας \textbf{μειώνει τον αριθμό των προσβάσεων στη shared σε: 3RD/(2 ή 4)WR} κάτι που αναμένεται να μειώσει τις αναγνώσεις από \boldmath $525k$ σε $393k$ \unboldmath.
Απαιτεί βέβαια προσοχή γιατί δημιουργούμε πλεονασμό δεδομένων μεταξύ του register και της θέσης στον πίνακα share που αυτός αντιστοιχεί.
\par
Με προσεκτική αλλαγή στην \textit{exchange()} ώστε να επιστρέφει bool για το αν έγινε ανταλλαγή, δεν προστίθεται κάποιο branch στη ροή ελέγχου κάτι που μας επιτρέπει απλώς να προσθέσουμε δύο εντολές για ανάγνωση και εγγραφή στη register μεταβλητή.
\par
Δυστυχώς όμως δεν καταφέραμε ποτέ να πετύχουμε μια λειτουργική υλοποίηση.
Είτε η μεταβλητή που θα έπρεπε να είναι στους registers ήταν στη local μνήμη, είτε δημιουργούνταν κάποιο mismatch μεταξύ της μεταβλητής του register και της αντίστοιχης θέσης στον shared πίνακα.
Ο αναγνώστης μπορεί να βρει ένα draft στο branch RC3 του αποθετηρίου αν θέλει να πειραματιστεί\footnote{Βέβαια δεν μπορώ να σκεφτώ ένα καλό λόγο για να το κάνει κάποιος αυτό στον εαυτό του.}!

\section{Τελική έκδοση} \label{sec:Final_version}

Έχοντας λοιπόν προηγηθεί η παραπάνω ανάλυση, για την τελική έκδοση της εφαρμογής επιλέξαμε απλά να κρατήσουμε την βελτιστοποίηση με το prephase.
Η επαλήθευση του αλγόριθμου έγινε μέσω ενός validator που ενσωματώθηκε στο εκτελέσιμο και μπορεί να καλεστεί από το command line argument \texttt{--validation}, ο οποίος στο τέλος της ταξινόμησης ελέγχει αν ο πίνακας είναι ορθά ταξινομημένος.
\par
Στον κατάλογο hpc, βρίσκονται scripts που δημιουργούν τα batch scripts για τη συστοιχία και τα οποία χρησιμοποιήθηκαν για τις τελικές μετρήσεις.
Ο αναγνώστης μπορεί να δοκιμάσει τον αλγόριθμο στη συστοιχία για παράδειγμα στην ampere, δίνοντας από το root directory:
\begin{verbatim}
    $ module load gcc/9.2.0 cuda/11.1.0
    $ make -j hpc-build
    $ cd hpc && ./makeSlurmScripts.sh
    $ cd ..  && ./hpc/submitJobs.sh hpc ampere
\end{verbatim}
Ο αναγνώστης μπορεί επίσης τοπικά να μεταγλωττίσει και να εκτελέσει τον αλγόριθμο, δίνοντας από το root directory:
\begin{verbatim}
    $ make -j hpc-build
    $ ./out/v1/bitonicCUDA -v --validation --perf 5 -q 24
    $ # or
    $ ./out/v2/bitonicCUDA -v --validation --perf 5 -b 512 -q 24
\end{verbatim}
Το οποίο για πίνακα μεγέθους q=24 θα εκτελέσει την ταξινόμηση 5 φορές, θα ελέγξει την εγκυρότητά της και θα εκτυπώσει τον ενδιάμεσο χρόνο, χρησιμοποιώντας στην πρώτη εντολή την έκδοση v1 και στη δεύτερη την έκδοση v2 αλλά και block size 512 threads.
Δίνοντας το όρισμα \texttt{-h} ο χρήστης μπορεί να δει όλες τι επιλογές εκτέλεσης και μια μικρή τεκμηρίωση ώστε να πειραματιστεί.
\par
Παρακάτω παραθέτουμε τα γραφήματα και τα αποτελέσματα στους τελικούς χρόνους από τη συστοιχία ampere, gpu(tesla), αλλά και από μια κάρτα GTX1650, όπου κρατήσαμε τον ενδιάμεσο χρόνο από 7 εκτελέσεις.

\InsertFigure{H}{1.025}{fig:TotalTime}{../matlab/total_time_plot}{Συνολικός χρόνος εκτέλεσης της ταξινόμησης (Με μεταφορά δεδομένων).}
\InsertFigure{H}{1.025}{fig:SortTime}{../matlab/sort_time_plot}{Χρόνος εκτέλεσης της ταξινόμησης (Χωρίς μεταφορά δεδομένων).}
\InsertFigure{H}{1.025}{fig:SpeedUp}{../matlab/speedup_plot}{Επιτάχυνση σε σχέση με τη σειριακή έκδοση (Συστοιχία AMD 7742).}

\begin{table}[H]
    \centering
    \renewcommand{\arraystretch}{1.25}

    \begin{tabular}{ll|rrrrrr}
                     &       &       Q20  &        Q22 &       Q24 &        Q26 &        Q28 &        Q30  \\ \hline
    Serial - AMD7742 & Total &    74 msec &   328 msec & 1441 msec &  6293 msec & 27.17 sec  & 115.58 sec \\ \hline
        V0 - ampere  & Total &  8733 usec &    59 msec &   56 msec &   229 msec &   976 msec &  3971 msec \\
                     & Sort  &  3134 usec &  6074 usec &   29 msec &   120 msec &   554 msec &  2453 msec \\
        V0 - tesla   & Total &  6559 usec &    23 msec &  101 msec &   442 msec &  1986 msec &  9071 msec \\
                     & Sort  &  5419 usec &    19 msec &   85 msec &   385 msec &  1762 msec &  8034 msec \\
        V0 - gtx1650 & Total &    16 msec &    71 msec &  327 msec &  1493 msec &  6782 msec &        N/A \\
                     & Sort  &    15 msec &    66 msec &  305 msec &  1408 msec &  6445 msec &        N/A \\ \hline
        V1 - ampere  & Total &  2676 usec &    10 msec &   34 msec &   134 msec &   664 msec &  2938 msec \\
                     & Sort  &  1417 usec &  3699 usec &   17 msec &    71 msec &   351 msec &  1556 msec \\
        V1 - tesla   & Total &  4559 usec &    18 msec &   80 msec &   358 msec &  1618 msec &  7326 msec \\
                     & Sort  &  3418 usec &    14 msec &   64 msec &   300 msec &  1394 msec &  6411 msec \\
        V1 - gtx1650 & Total &  8271 usec &    37 msec &  161 msec &   774 msec &  3701 msec &        N/A \\
                     & Sort  &  6666 usec &    31 msec &  139 msec &   689 msec &  3358 msec &        N/A \\ \hline
        V2 - ampere  & Total &  2883 usec &    12 msec &   34 msec &   135 msec &   749 msec &  3063 msec \\
                     & Sort  &  1574 usec &  4134 usec &   19 msec &    76 msec &   372 msec &  1650 msec \\
        V2 - tesla   & Total &  4407 usec &    17 msec &   75 msec &   338 msec &  1533 msec &  6932 msec \\
                     & Sort  &  3264 usec &    13 msec &   60 msec &   280 msec &  1308 msec &  6048 msec \\
        V2 - gtx1650 & Total &  8985 usec &    38 msec &  169 msec &   810 msec &  3864 msec &        N/A \\
                     & Sort  &  7294 usec &    33 msec &  148 msec &   726 msec &  3527 msec &        N/A \\
    \end{tabular}
    \caption{Μέρος των τελικών μετρήσεων στη συστοιχία.}
    \label{tab:Final_Reults}
\end{table}

\subsection{Συμπεράσματα}
Από τα διαγράμματα και τις μετρήσεις βλέπουμε πως:
\begin{itemize}
    \item Η έκδοση που επιταχύνεται από την GPU σε σχέση με την εκτέλεση στη CPU του ίδιου μηχανήματος είναι \textbf{γρηγορότερη ακόμα και για μικρό μέγεθος δεδομένων}.
    Φυσικά, εδώ θα μπορούσαμε να συγκρίνουμε μια σειριακή εκτέλεση εντός GPU, αλλά αυτό θα εμφάνιζε κυρίως το θεωρητικό και μαθηματικό πλεονέκτημα της παραλληλοποίησης και θα εξαρτώνταν κυρίως από των αριθμό των παράλληλων threads.
    Αντίθετα η σύγκριση με τη CPU έγινε για να αποκτηθεί μια αίσθηση από ένα πιο πιθανό σενάριο ανάγκης για επιτάχυνση από GPU.
    \item Η κάρτα \textbf{Tesla P100-PCIE-12GB} της συστοιχίας gpu παρουσιάζει \textbf{διαφορετική συμπεριφορά για την έκδοση V2}, σε σχέση με την NVIDIA A100-SXM4-40GB της ampere, αλλά και της GTX1650.
    \item Επίσης η \textbf{υπεροχή} της Α100 στις επιδόσεις \textbf{στη μεταφορά δεδομένων} φαίνονται έντονα στο σχήμα \ref{fig:SpeedUp}, όπου η κάρτα κερδίζει αρκετό έδαφος καθώς το μέγεθος προς ταξινόμηση μεγαλώνει.
\end{itemize}

\section{Σύνοψη}
Κλείνοντας να πούμε πως προσπαθήσαμε να παρουσιάσουμε ένα σχετικά αποδοτικό αλγόριθμο δίνοντας βάση στην βελτιστοποίηση μέσω ανάλυσης.
Ξεκινώντας από την έκδοση V0 η οποία χρησιμοποιεί απλή παραλληλοποίηση, περάσαμε στη V1 που μειώνει τις kernel κλήσεις μέσω loop unrolling και στη V2 που αξιοποιεί shared memory.
Μέσα από προφίλ εκτελέσεων και ανάλυση επιδόσεων, προέκυψε αρχικά το συμπέρασμα ότι η V2 παρουσιάζει δικαιολογημένες αποκλίσεις, και στη συνέχεια η βελτιστοποίηση του prephase.
Οι πειραματικές μετρήσεις που πραγματοποιήθηκαν σε διάφορες αρχιτεκτονικές GPU (Ampere A100, Tesla P100, GTX 1650) επιβεβαίωσαν την αποτελεσματικότητα της παραλληλοποίησης έναντι μιας σειριακής εκτέλεσης σε CPU (AMD 7742), με σημαντικό speedup, ειδικά για μεγάλα dataset sizes.
\par
Τέλος, αν και δεν καταφέραμε να υλοποιήσουμε όλες μας τις ιδέες, θεωρούμε πως ασχοληθήκαμε με κάποιες βελτιστοποιήσεις που ίσως να αξίζουν περισσότερης διερεύνησης, πέρα από το πλαίσιο της εργασίας.

\end{document}
