diff --git a/homework_3/report/homework_3_report.pdf b/homework_3/report/homework_3_report.pdf index 042b9a5..8e1eb4a 100644 Binary files a/homework_3/report/homework_3_report.pdf and b/homework_3/report/homework_3_report.pdf differ diff --git a/homework_3/report/homework_3_report.tex b/homework_3/report/homework_3_report.tex index 930d530..bf7edd9 100644 --- a/homework_3/report/homework_3_report.tex +++ b/homework_3/report/homework_3_report.tex @@ -143,22 +143,21 @@ Η συνάρτηση \textit{main()} που υλοποιεί τον αλγόριθμο μαζί με τον validator και το command line interface βρίσκονται στο αρχείο main.cpp. Στο αρχείο config.h υπάρχουν οι compile time ρυθμίσεις της υλοποίησης όπου ο χρήστης μπορεί να επιλέξει για παράδειγμα τον τύπο δεδομένων προς ταξινόμηση. -\subsection{Πρώτη έκδοση (V0)} +\subsection{Πρώτη έκδοση (V0)} \label{sec:V0} Η πρώτη έκδοση υλοποιεί τον βασικό αλγόριθμο επιταχύνοντας απλώς το κύριο σώμα του στην κάρτα γραφικών. Καθώς θεωρούμε πως ο αναγνώστης είναι ήδη εξοικειωμένος με την διτονική ταξινόμηση, δεν θα αναλύσουμε την λειτουργία της σε μεγάλο βάθος. Για πίνακα μεγέθους $2^N$ ο αλγόριθμος συνοψίζεται στα εξής: \begin{itemize} \item Εκτελούνται $N$ ακολουθίες ανταλλαγών με αύξον αριθμό: $m = 1, 2,\dots,N$. - \item Η κάθε ακολουθία $m$ αποτελείται από ανταλλαγές μεταξύ γειτόνων, των οποίων η απόσταση ξεκινάει από $2^{m-1}$ και μειώνεται με διαδοχικές ακέραιες διαιρέσεις με το 2, εωσότου γίνει ένα: $2^{m-1}, 2^{m-2}, ..., 1$. + \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{!h}{0.8}{fig:lala}{lala.png}{\eng{Makes lala}} \InsertFigure{!ht}{0.8}{fig:bitonicSort}{img/BitonicSort.png}{4 ακολουθίες ανταλλαγών για ταξινόμηση 16 στοιχείων.} Η εικόνα \ref{fig:bitonicSort} παρουσιάζει αυτή τη λογική\footnote{Πηγή \href{\bitonicSortWikipedia}{wikipedia}}. -\subsubsection{Πλήθος απαιτούμενων διεργασιών (threads) και καταμερισμός στον πίνακα} +\subsubsection{Πλήθος απαιτούμενων διεργασιών (threads) και καταμερισμός στον πίνακα} \label{sec:V0addressing} Στην παρούσα εργασία, επιθυμούμε οι ανταλλαγές να λάβουν χώρα εντός του υποσυστήματος γραφικών και να εκτελεστούν από τις διεργασίες (threads) του υποσυστήματος γραφικών. Σε αντίθεση όμως με την προηγούμενη εργασία, όπου για $M$ υποσύνολα του πίνακα είχαμε $M$ MPI διεργασίες και όλες απαιτούνταν να εκτελέσουν διαχωρισμό μεγίστων -- ελαχίστων, \textbf{εδώ η κάθε διεργασία μπορεί να διαχωρίσει 2 στοιχεία}. \textbf{Δηλαδή για ταξινόμηση ενός πίνακα $N$ στοιχείων απαιτούνται $\frac{N}{2}$ διεργασίες}. @@ -170,21 +169,76 @@ \par Όπως γίνεται φανερό \textit{έπρεπε να βρεθεί ένας τρόπος ώστε οι μισές διεργασίες να διευθυνσιοδοτήσουν το πρώτο μέρος και οι άλλες μισές το δεύτερο}. Ο διαχωρισμός έγινε ελέγχοντας τον γείτονα (partner) για την κάθε ανταλλαγή. -\InsertFigure{!ht}{0.85}{fig:threadDivision}{img/V0Addressing.png}{Διευθυνσιοδότηση των threads (V0).} + +\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:threadDivision} παρουσιάζει αυτόν τον καταμερισμό. +Το διάγραμμα \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{V1} +Από την εκφώνηση της εργασίας ήδη υπάρχουν οδηγίες για βελτιστοποιήσεις. +Η πρώτη βελτιστοποίηση λοιπόν αφορά αυτές ακριβώς τις κλήσεις των kernels, των συναρτήσεων δηλαδή που επιταχύνονται από την κάρτα γραφικών. +Γίνεται δηλαδή η \textbf{θεώρηση} από την εκφώνηση, ότι \textbf{η ελαχιστοποίηση αυτών των κλήσεων θα βελτιστοποιήσει τη διαδικασία}. +Φυσικά, αυτό μένει πάντα να αποδειχτεί στην πράξη. +\par +Ο τρόπος που προτείνεται ανήκει στην κατηγορία βελτιστοποιήσεων που λέγεται loop unrolling, όπου ένα μέρος του βρόχου αφαιρείται από τον έλεγχο του βρόχου και καλείται "με το χέρι". +Στη δική μας περίπτωση το μέρος που αφαιρείται μπορεί να μπει μέσα στην κλήση kernel αξιοποιώντας καλύτερα τους πόρους που μπορεί να χρησιμοποιήσει, μειώνοντας έτσι και τον αριθμό κλήσεων. +\par +Ενώ όμως στην έκδοση V0 η κάθε κλήση kernel αφορούσε ένα στιγμιότυπο (μια κάθετη γραμμή) του αλγόριθμου, εδώ οι κλήσεις αφορούν περισσότερα. +Ο κάθε kernel δηλαδή θα εκτελέσει τμήματα περισσοτέρων της μίας γραμμής του σχήματος \ref{fig:bitonicSort}. +Αυτό όμως δημιουργεί data race, καθώς οι τιμές αυτών των στοιχείων διαμοιράζονται μεταξύ των διαφορετικών κλήσεων. \par -Έχοντας κατανέμει λοιπόν τα threads στον πίνακα η υλοποίηση εκτελεί το διπλό βρόχο που περιγράψαμε παραπάνω και αναθέτει το σώμα το οποίο εκτελεί τις ανταλλαγές στη κάρτα γραφικών. -Το κάθε thread εκτελεί μια ανταλλαγή με την άνω διευθυνσιοδότηση και ο αλγόριθμος τερματίζει στο τέλος του διπλού βρόχου. -Το πλεονέκτημα αυτής της μεθόδου είναι η απλότητά της. -Κάτι που προφανώς πληρώνουμε με έναν αρκετά μεγάλο αριθμό κλήσεων προς την κάρτα γραφικών που προσθέτει overhead. - -\subsection{Δεύτερη έκδοση (V1)} -Από την εκφώνηση της εργασίας ήδη υπάρχουν οδηγίες για τις πρώτες βελτιστοποιήσεις. -Η πρώτη βελτιστοποίηση λοιπόν αφορά ακριβώς αυτές τις κλήσεις των συναρτήσεων που επιταχύνονται από την κάρτα γραφικών. -Γίνεται δηλαδή η \textbf{θεώρηση} από την εκφώνηση, ότι\textbf{ η ελαχιστοποίηση αυτών των κλήσεων θα βελτιστοποιήσει τη διαδικασία}. -Φυσικά αυτό πρέπει πάντα να αποδειχτεί στην πράξη. +Για να λύσουμε αυτό το πρόβλημα εργαστήκαμε έτσι ώστε: +\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$ είναι και το μέγιστο βήμα από το οποίο μπορεί να ξεκινήσει το 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. + +\par +Το μόνο που μένει είναι πλέον το πρόβλημα του \textbf{συγχρονισμού}. +Εδώ χρειάζεται να κρατήσουμε τον συγχρονισμό που είχαμε σε κάθε κάθετη γραμμή του σχήματος \ref{fig:bitonicSort}, απλώς αυτός θα λάβει χώρα εσωτερικά του kernel. +Η συνάρτηση που μας δίνει αυτή τη δυνατότητα είναι η \textit{\_\_syncthreads()}, την οποία και καλούμε μετά από την ολοκλήρωση της κάθε ακολουθίας. +\par +Όπως είναι φανερό ο διαχωρισμός μεταξύ των βημάτων που μπορούν να γίνουν εσωτερικά και αυτών εξωτερικά του kernel έχει άνω όριο, αλλά όχι κάτω. +Στην υλοποίησή μας παρόλα αυτά δεν δίνουμε τη δυνατότητα ρύθμισης αυτού του ορίου. +Αυτό γιατί ούτως ή άλλως υπάρχει ρύθμιση για το μέγεθος του block, το οποίο και χρησιμοποιούμε για να κάνουμε και αυτό το διαχωρισμό. + +\subsection{Τρίτη έκδοση (V2)} \label{V2} \end{document} diff --git a/homework_3/report/img/V0Addresing.svg b/homework_3/report/img/V0Addresing.svg index cc97677..9375088 100644 --- a/homework_3/report/img/V0Addresing.svg +++ b/homework_3/report/img/V0Addresing.svg @@ -7,7 +7,7 @@ viewBox="0 0 297 210" version="1.1" id="svg1" - inkscape:version="1.4 (e7c3feb100, 2024-10-09)" + inkscape:version="1.4 (1:1.4+202410161351+e7c3feb100)" sodipodi:docname="V0Addresing.svg" xmlns:inkscape="http://www.inkscape.org/namespaces/inkscape" xmlns:sodipodi="http://sodipodi.sourceforge.net/DTD/sodipodi-0.dtd" @@ -25,21 +25,21 @@ inkscape:document-units="mm" showguides="true" inkscape:zoom="0.70710678" - inkscape:cx="439.82042" - inkscape:cy="613.76869" + inkscape:cx="689.42911" + inkscape:cy="606.69762" inkscape:window-width="1920" inkscape:window-height="1011" inkscape:window-x="0" inkscape:window-y="32" inkscape:window-maximized="1" - inkscape:current-layer="layer1" /> + inkscape:current-layer="g6" /> @@ -130,7 +130,7 @@ + inkscape:export-xdpi="200" + inkscape:export-ydpi="200" /> 0 + id="tspan2">0 N/2 + id="tspan22">N/2 N + id="tspan24">N TidkTid = threadIdx.x + blockIdx.x*blockDim.x + id="tspan28"> = threadIdx.x + blockIdx.x*blockDim.x + id="tspan30"> + id="tspan32"> + id="tspan34"> + id="tspan36"> + id="tspan38"> Pid + id="tspan40">Pid + id="tspan42"> + id="tspan44"> + id="tspan46"> + id="tspan48"> + id="tspan50"> N/2 Threads + id="tspan52">Threads = N/2 Pid' + id="tspan54">Pid' + id="tspan56"> + id="tspan58"> + id="tspan60"> + id="tspan62"> + id="tspan64"> TidTid11 + id="tspan68"> + id="tspan70"> + id="tspan72"> + id="tspan74"> + id="tspan76"> + id="tspan78"> TidTid22 + id="tspan82"> + id="tspan84"> + id="tspan86"> + id="tspan88"> + id="tspan90"> + id="tspan92"> Tid'Tid'22 + id="tspan96"> + id="tspan98"> + id="tspan100"> + id="tspan102"> + id="tspan104"> + id="tspan106"> Tid' = Tid + N/2 + id="tspan108">Tidk' = Tidk + N/2 + id="tspan114"> + id="tspan116"> + id="tspan118"> + id="tspan120"> + id="tspan122"> diff --git a/homework_3/report/img/V0Addressing.png b/homework_3/report/img/V0Addressing.png index 80405a4..e762dce 100644 Binary files a/homework_3/report/img/V0Addressing.png and b/homework_3/report/img/V0Addressing.png differ diff --git a/homework_3/report/img/V1Addressing.png b/homework_3/report/img/V1Addressing.png new file mode 100644 index 0000000..732ee69 Binary files /dev/null and b/homework_3/report/img/V1Addressing.png differ diff --git a/homework_3/report/img/V1Addressing.svg b/homework_3/report/img/V1Addressing.svg new file mode 100644 index 0000000..fe9c87b --- /dev/null +++ b/homework_3/report/img/V1Addressing.svg @@ -0,0 +1,1146 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + 0 + N + Tidk = threadIdx.x + 2*blockIdx.x*blockDim.x + + + + + + + + Tidk' = Tidk + blockDim.x + + + + + + + + + Pid' + + + + + + + + Tid1 + + + + + + + + Tid2 + + + + + + + + Tid'2 + + + + + + + + + + + + Blk + + + + + + + Blk+1 + + + + + + + blockDim.x + + + + + + + Blk+2 + + + + + + + + + + + Blk+3 + + + + + + + + + + + + + + + + Next Pair of Blocks + + + + + + + + Current Pair of Blocks + + + + + + + + diff --git a/homework_3/report/img/V1Unrolling.png b/homework_3/report/img/V1Unrolling.png new file mode 100644 index 0000000..7b8e649 Binary files /dev/null and b/homework_3/report/img/V1Unrolling.png differ diff --git a/homework_3/report/img/V1Unrolling.svg b/homework_3/report/img/V1Unrolling.svg new file mode 100644 index 0000000..8276e79 --- /dev/null +++ b/homework_3/report/img/V1Unrolling.svg @@ -0,0 +1,1068 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + 1 + 2 + 3 + 4 + 5 + 6 + 7 + + 1 + 2 + 1 + 4 + 2 + 1 + 8 + 4 + 2 + 1 + 16 + 8 + 4 + 2 + 1 + 32 + 64 + 16 + 8 + 4 + 2 + 1 + 32 + 16 + 8 + 4 + 2 + 1 + + Seq- + + + + + + + Distances + + + + + + + + Inside kernel + + + + + + + + + Step: + + + + + + + 6 + 0 + 1 + 2 + 3 + 4 + 5 + uence + + + + + + + +