hoo2/PDS
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

HW3: RC4 - V0-V1 parts of the report

Browse Source
This commit is contained in:
Christos Choutouridis 2025-02-20 17:46:49 +02:00
parent 1fa62a85ce
commit c21c1edeab
8 changed files with 2420 additions and 139 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

BIN
homework_3/report/homework_3_report.pdf
View File

Binary file not shown.

84
homework_3/report/homework_3_report.tex
View File

@ -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 στον πίνακα η υλοποίηση εκτελεί το διπλό βρόχο που περιγράψαμε παραπάνω και αναθέτει το σώμα το οποίο εκτελεί τις ανταλλαγές στη κάρτα γραφικών.
Το κάθε thread εκτελεί μια ανταλλαγή με την άνω διευθυνσιοδότηση και ο αλγόριθμος τερματίζει στο τέλος του διπλού βρόχου.
Το πλεονέκτημα αυτής της μεθόδου είναι η απλότητά της.
Κάτι που προφανώς πληρώνουμε με έναν αρκετά μεγάλο αριθμό κλήσεων προς την κάρτα γραφικών που προσθέτει overhead.
Έχοντας κατανέμει λοιπόν τα 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)}
Από την εκφώνηση της εργασίας ήδη υπάρχουν οδηγίες για τις πρώτες βελτιστοποιήσεις.
Η πρώτη βελτιστοποίηση λοιπόν αφορά ακριβώς αυτές τις κλήσεις των συναρτήσεων που επιταχύνονται από την κάρτα γραφικών.
Γίνεται δηλαδή η \textbf{θεώρηση} από την εκφώνηση, ότι\textbf{ η ελαχιστοποίηση αυτών των κλήσεων θα βελτιστοποιήσει τη διαδικασία}.
Φυσικά αυτό πρέπει πάντα να αποδειχτεί στην πράξη.
\subsection{Δεύτερη έκδοση (V1)} \label{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$ είναι και το μέγιστο βήμα από το οποίο μπορεί να ξεκινήσει το 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}

261
homework_3/report/img/V0Addresing.svg
View File

@ -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" />
<defs
id="defs1">
<rect
x="120.30141"
y="341.30875"
width="234.15368"
height="40.183152"
width="234.15369"
height="40.183151"
id="rect5" />
<rect
x="171.64655"
@ -67,8 +67,8 @@
id="path135" />
</marker>
<rect
x="75.770079"
y="296.06457"
x="75.770081"
y="296.06458"
width="207.66614"
height="78.576378"
id="rect2" />
@ -130,7 +130,7 @@
<rect
x="171.64655"
y="128.23883"
width="423.76601"
width="423.76602"
height="42.273273"
id="rect4-1" />
<rect
@ -309,8 +309,8 @@
x="20.343727"
y="60.406105"
inkscape:export-filename="V0Addressing.png"
inkscape:export-xdpi="96"
inkscape:export-ydpi="96" />
inkscape:export-xdpi="200"
inkscape:export-ydpi="200" />
<text
xml:space="preserve"
transform="matrix(0.26458333,0,0,0.26458333,-24.631251,38.437251)"
@ -318,9 +318,9 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-3);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan14"><tspan
id="tspan19"><tspan
style="fill:#000000"
id="tspan13">0</tspan></tspan></text>
id="tspan2">0</tspan></tspan></text>
<text
xml:space="preserve"
transform="matrix(0.26458333,0,0,0.26458333,96.45285,38.454477)"
@ -328,9 +328,9 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-3-8);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan16"><tspan
id="tspan23"><tspan
style="fill:#000000"
id="tspan15">N/2</tspan></tspan></text>
id="tspan22">N/2</tspan></tspan></text>
<text
xml:space="preserve"
transform="matrix(0.26458333,0,0,0.26458333,215.16783,38.453616)"
@ -338,9 +338,9 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-3-8-9);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan18"><tspan
id="tspan25"><tspan
style="fill:#000000"
id="tspan17">N</tspan></tspan></text>
id="tspan24">N</tspan></tspan></text>
<path
style="fill:none;stroke:#000000;stroke-width:0.629912;stroke-dasharray:none;stroke-opacity:1;marker-start:url(#marker4-6);marker-end:url(#Dot-0)"
d="m 49.749862,65.166268 -7.310658,-0.09543 -6.876337,-0.08975"
@ -353,40 +353,44 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan22"><tspan
dx="0 0 0 1.0199995 -1.02 1.0200014 -1.02 0 0 0 0 0 0 0 0 0 0 1.0199914 -1.02 1.0199995 -1.02 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1.0200067"
id="tspan29"><tspan
style="fill:#000000"
id="tspan19">Tid = threadIdx.x + blockIdx.x*blockDim.x
id="tspan26">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan27">k</tspan><tspan
dx="1.0200005 -1.02 1.0199995 -1.02 0 0 0 0 0 0 0 0 0 0 1.0200028 -1.0199962 1.0200033 -1.0200038 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 1.0200028"
style="fill:#000000"
id="tspan28"> = threadIdx.x + blockIdx.x*blockDim.x
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan24"><tspan
y="133.24558"
id="tspan31"><tspan
style="fill:#000000"
id="tspan23">
id="tspan30">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan26"><tspan
y="133.24558"
id="tspan33"><tspan
style="fill:#000000"
id="tspan25">
id="tspan32">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan28"><tspan
y="133.24558"
id="tspan35"><tspan
style="fill:#000000"
id="tspan27">
id="tspan34">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan30"><tspan
y="133.24558"
id="tspan37"><tspan
style="fill:#000000"
id="tspan29">
id="tspan36">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan32"><tspan
y="133.24558"
id="tspan39"><tspan
style="fill:#000000"
id="tspan31">
id="tspan38">
</tspan></tspan></text>
<text
xml:space="preserve"
@ -400,40 +404,40 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1-5);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan34"><tspan
id="tspan41"><tspan
dx="0 0 0 1.0200005"
style="fill:#000000"
id="tspan33">Pid
id="tspan40">Pid
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan36"><tspan
id="tspan43"><tspan
style="fill:#000000"
id="tspan35">
id="tspan42">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan38"><tspan
id="tspan45"><tspan
style="fill:#000000"
id="tspan37">
id="tspan44">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan40"><tspan
id="tspan47"><tspan
style="fill:#000000"
id="tspan39">
id="tspan46">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan42"><tspan
id="tspan49"><tspan
style="fill:#000000"
id="tspan41">
id="tspan48">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan44"><tspan
id="tspan51"><tspan
style="fill:#000000"
id="tspan43">
id="tspan50">
</tspan></tspan></text>
<rect
style="fill:#ffffff;fill-opacity:1;stroke:#000000;stroke-width:0.687452;stroke-dasharray:none;stroke-opacity:1"
@ -457,15 +461,15 @@
sodipodi:nodetypes="ccc" />
<text
xml:space="preserve"
transform="matrix(0.26458333,0,0,0.26458333,55.513176,2.606821)"
transform="matrix(0.26458333,0,0,0.26458333,50.750224,2.606821)"
id="text4-44"
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-30);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan46"><tspan
dx="0 0 0 1.02 -1.0200005"
y="132.39453"
id="tspan53"><tspan
dx="0 0 0 0 0 0 0 1.02 -1.0199966 1.02 -1.0200043"
style="font-style:normal;font-size:16px;font-family:Candara;-inkscape-font-specification:'Candara, Normal';fill:#000000"
id="tspan45">N/2 Threads</tspan></tspan></text>
id="tspan52">Threads = N/2</tspan></tspan></text>
</g>
<text
xml:space="preserve"
@ -474,40 +478,40 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1-5-4);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan48"><tspan
id="tspan55"><tspan
dx="0 0 0 0 1.02"
style="fill:#000000"
id="tspan47">Pid'
id="tspan54">Pid'
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan50"><tspan
id="tspan57"><tspan
style="fill:#000000"
id="tspan49">
id="tspan56">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan52"><tspan
id="tspan59"><tspan
style="fill:#000000"
id="tspan51">
id="tspan58">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan54"><tspan
id="tspan61"><tspan
style="fill:#000000"
id="tspan53">
id="tspan60">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan56"><tspan
id="tspan63"><tspan
style="fill:#000000"
id="tspan55">
id="tspan62">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan58"><tspan
id="tspan65"><tspan
style="fill:#000000"
id="tspan57">
id="tspan64">
</tspan></tspan></text>
<text
xml:space="preserve"
@ -516,44 +520,44 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1-5-7-3);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan62"><tspan
id="tspan69"><tspan
style="fill:#000000"
id="tspan59">Tid</tspan><tspan
id="tspan66">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan60">1</tspan><tspan
id="tspan67">1</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan61">
id="tspan68">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan64"><tspan
id="tspan71"><tspan
style="fill:#000000"
id="tspan63">
id="tspan70">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan66"><tspan
id="tspan73"><tspan
style="fill:#000000"
id="tspan65">
id="tspan72">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan68"><tspan
id="tspan75"><tspan
style="fill:#000000"
id="tspan67">
id="tspan74">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan70"><tspan
id="tspan77"><tspan
style="fill:#000000"
id="tspan69">
id="tspan76">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan72"><tspan
id="tspan79"><tspan
style="fill:#000000"
id="tspan71">
id="tspan78">
</tspan></tspan></text>
<path
style="fill:none;stroke:#000000;stroke-width:0.629912;stroke-dasharray:none;stroke-opacity:1;marker-start:url(#marker4-6-9)"
@ -567,44 +571,44 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1-5-7-3-4);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan76"><tspan
id="tspan83"><tspan
style="fill:#000000"
id="tspan73">Tid</tspan><tspan
id="tspan80">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan74">2</tspan><tspan
id="tspan81">2</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan75">
id="tspan82">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan78"><tspan
id="tspan85"><tspan
style="fill:#000000"
id="tspan77">
id="tspan84">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan80"><tspan
id="tspan87"><tspan
style="fill:#000000"
id="tspan79">
id="tspan86">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan82"><tspan
id="tspan89"><tspan
style="fill:#000000"
id="tspan81">
id="tspan88">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan84"><tspan
id="tspan91"><tspan
style="fill:#000000"
id="tspan83">
id="tspan90">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan86"><tspan
id="tspan93"><tspan
style="fill:#000000"
id="tspan85">
id="tspan92">
</tspan></tspan></text>
<path
style="fill:none;stroke:#000000;stroke-width:0.629912;stroke-dasharray:none;stroke-opacity:1;marker-start:url(#marker4-6-9-6)"
@ -618,44 +622,44 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1-5-7-3-4-6);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan90"><tspan
id="tspan97"><tspan
style="fill:#000000"
id="tspan87">Tid'</tspan><tspan
id="tspan94">Tid'</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan88">2</tspan><tspan
id="tspan95">2</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan89">
id="tspan96">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan92"><tspan
id="tspan99"><tspan
style="fill:#000000"
id="tspan91">
id="tspan98">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan94"><tspan
id="tspan101"><tspan
style="fill:#000000"
id="tspan93">
id="tspan100">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan96"><tspan
id="tspan103"><tspan
style="fill:#000000"
id="tspan95">
id="tspan102">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan98"><tspan
id="tspan105"><tspan
style="fill:#000000"
id="tspan97">
id="tspan104">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan100"><tspan
id="tspan107"><tspan
style="fill:#000000"
id="tspan99">
id="tspan106">
</tspan></tspan></text>
<path
style="fill:none;stroke:#000000;stroke-width:0.629912;stroke-dasharray:none;stroke-opacity:1;marker-start:url(#marker4-6-9-6-2)"
@ -669,40 +673,49 @@
style="font-style:italic;font-variant:normal;font-weight:normal;font-stretch:normal;font-size:13.3333px;line-height:0;font-family:Arial;-inkscape-font-specification:'Arial Italic';text-align:start;letter-spacing:2px;word-spacing:-1.02px;writing-mode:lr-tb;direction:ltr;white-space:pre;shape-inside:url(#rect4-1-3);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan102"><tspan
dx="0 0 0 0 1.02 -1.0199981 1.0199995 -1.02 0 0 1.0199986 -1.02 1.0199995 -1.02 0 0 1.0200062"
id="tspan113"><tspan
style="fill:#000000"
id="tspan101">Tid' = Tid + N/2
id="tspan108">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan109">k</tspan><tspan
dx="0 1.02 -1.02 1.0199995 -1.02"
style="fill:#000000"
id="tspan110">' = Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan111">k</tspan><tspan
dx="1.0200062 -1.02 1.0199995 -1.02 0 0 1.0199986"
style="fill:#000000"
id="tspan112"> + N/2
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan104"><tspan
y="133.24558"
id="tspan115"><tspan
style="fill:#000000"
id="tspan103">
id="tspan114">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan106"><tspan
y="133.24558"
id="tspan117"><tspan
style="fill:#000000"
id="tspan105">
id="tspan116">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan108"><tspan
y="133.24558"
id="tspan119"><tspan
style="fill:#000000"
id="tspan107">
id="tspan118">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan110"><tspan
y="133.24558"
id="tspan121"><tspan
style="fill:#000000"
id="tspan109">
id="tspan120">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan112"><tspan
y="133.24558"
id="tspan123"><tspan
style="fill:#000000"
id="tspan111">
id="tspan122">
</tspan></tspan></text>
</g>
</svg>

Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 25 KiB

After

Width:  |  Height:  |  Size: 26 KiB

BIN
homework_3/report/img/V0Addressing.png
View File

Binary file not shown.
Side by Side Swipe Overlay

Before

Width:  |  Height:  |  Size: 16 KiB

After

Width:  |  Height:  |  Size: 37 KiB

BIN
homework_3/report/img/V1Addressing.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 55 KiB

1146
homework_3/report/img/V1Addressing.svg Normal file
View File

File diff suppressed because it is too large Load Diff
Side by Side

After

Width:  |  Height:  |  Size: 40 KiB

BIN
homework_3/report/img/V1Unrolling.png Normal file
View File

Binary file not shown.
Side by Side

After

Width:  |  Height:  |  Size: 45 KiB

1068
homework_3/report/img/V1Unrolling.svg Normal file
View File

File diff suppressed because it is too large Load Diff
Side by Side

After

Width:  |  Height:  |  Size: 37 KiB