hoo2
/
PDS
1
0
Форк 0
Код Проблеми 0 Запити на злиття 0 Релізи 3 Вікі Активність
Переглянути джерело

HW3: RC4 - A first report without the final plots

tags/v3.0
Christos Choutouridis 2 дні тому
джерело
c21c1edeab
коміт
dac00ed194
9 змінених файлів з 2033 додано та 193 видалено
Розділений перегляд
Параметри порівняння
Показати статистику Завантажити патч Завантажити файл різниці
  1. +45
    -0
      homework_3/analyse/RC4-8f44cbf/ampere-sum.txt
  2. +45
    -0
      homework_3/analyse/RC4-8f44cbf/gtx1650-sum.txt
  3. +45
    -0
      homework_3/analyse/RC4-8f44cbf/tesla-sum.txt
  4. BIN
      homework_3/report/homework_3_report.pdf
  5. +215
    -6
      homework_3/report/homework_3_report.tex
  6. BIN
      homework_3/report/img/V1Addressing.png
  7. +187
    -187
      homework_3/report/img/V1Addressing.svg
  8. BIN
      homework_3/report/img/V2Addressing.png
  9. +1496
    -0
      homework_3/report/img/V2Addressing.svg

+ 45
- 0
homework_3/analyse/RC4-8f44cbf/ampere-sum.txt Переглянути файл

@@ -0,0 +1,45 @@
GPU: NVIDIA A100-SXM4-40GB
Block size: 512
--------------------------------------

Code version: V0
Total Mem-xch Sorting
Q20: 8733 [usec] 5649 [usec] 3134 [usec]
Q21: 12 [msec] 8329 [usec] 4105 [usec]
Q22: 59 [msec] 53 [msec] 6074 [usec]
Q23: 37 [msec] 22 [msec] 14 [msec]
Q24: 56 [msec] 26 [msec] 29 [msec]
Q25: 98 [msec] 39 [msec] 59 [msec]
Q26: 229 [msec] 102 [msec] 120 [msec]
Q27: 420 [msec] 164 [msec] 256 [msec]
Q28: 976 [msec] 429 [msec] 554 [msec]
Q29: 1910 [msec] 745 [msec] 1164 [msec]
Q30: 3971 [msec] 1521 [msec] 2453 [msec]

Code version: V1
Total Mem-xch Sorting
Q20: 2676 [usec] 1254 [usec] 1417 [usec]
Q21: 6308 [usec] 4098 [usec] 2193 [usec]
Q22: 10 [msec] 6454 [usec] 3699 [usec]
Q23: 18 [msec] 10 [msec] 8329 [usec]
Q24: 34 [msec] 17 [msec] 17 [msec]
Q25: 83 [msec] 45 [msec] 37 [msec]
Q26: 134 [msec] 63 [msec] 71 [msec]
Q27: 331 [msec] 177 [msec] 153 [msec]
Q28: 664 [msec] 318 [msec] 351 [msec]
Q29: 1364 [msec] 621 [msec] 742 [msec]
Q30: 2938 [msec] 1384 [msec] 1556 [msec]

Code version: V2
Total Mem-xch Sorting
Q20: 2883 [usec] 1307 [usec] 1574 [usec]
Q21: 5861 [usec] 3444 [usec] 2411 [usec]
Q22: 12 [msec] 8167 [usec] 4134 [usec]
Q23: 21 [msec] 12 [msec] 8940 [usec]
Q24: 34 [msec] 15 [msec] 19 [msec]
Q25: 86 [msec] 45 [msec] 40 [msec]
Q26: 135 [msec] 59 [msec] 76 [msec]
Q27: 331 [msec] 167 [msec] 163 [msec]
Q28: 749 [msec] 371 [msec] 372 [msec]
Q29: 1459 [msec] 666 [msec] 788 [msec]
Q30: 3063 [msec] 1409 [msec] 1650 [msec]

+ 45
- 0
homework_3/analyse/RC4-8f44cbf/gtx1650-sum.txt Переглянути файл

@@ -0,0 +1,45 @@
GPU: NVIDIA GeForce GTX 1650
Block size: 512
--------------------------------------

Code version: V0
Total Mem-xch Sorting
Q20: 16 [msec] 1627 [usec] 15 [msec]
Q21: 33 [msec] 2906 [usec] 30 [msec]
Q22: 71 [msec] 5624 [usec] 66 [msec]
Q23: 153 [msec] 10 [msec] 142 [msec]
Q24: 327 [msec] 21 [msec] 305 [msec]
Q25: 696 [msec] 42 [msec] 654 [msec]
Q26: 1493 [msec] 84 [msec] 1408 [msec]
Q27: 3177 [msec] 168 [msec] 3009 [msec]
Q28: 6782 [msec] 336 [msec] 6445 [msec]
Q29: 14.42 [sec] 673 [msec] 13.75 [sec]
Q30: N/A N/A N/A

Code version: V1
Total Mem-xch Sorting
Q20: 8271 [usec] 1613 [usec] 6666 [usec]
Q21: 17 [msec] 2987 [usec] 14 [msec]
Q22: 37 [msec] 5752 [usec] 31 [msec]
Q23: 78 [msec] 10 [msec] 67 [msec]
Q24: 161 [msec] 21 [msec] 139 [msec]
Q25: 353 [msec] 42 [msec] 311 [msec]
Q26: 774 [msec] 84 [msec] 689 [msec]
Q27: 1709 [msec] 168 [msec] 1539 [msec]
Q28: 3701 [msec] 338 [msec] 3358 [msec]
Q29: 8106 [msec] 676 [msec] 7428 [msec]
Q30: N/A N/A N/A

Code version: V2
Total Mem-xch Sorting
Q20: 8985 [usec] 1665 [usec] 7294 [usec]
Q21: 18 [msec] 2849 [usec] 15 [msec]
Q22: 38 [msec] 5518 [usec] 33 [msec]
Q23: 81 [msec] 10 [msec] 70 [msec]
Q24: 169 [msec] 21 [msec] 148 [msec]
Q25: 370 [msec] 42 [msec] 328 [msec]
Q26: 810 [msec] 84 [msec] 726 [msec]
Q27: 1779 [msec] 168 [msec] 1609 [msec]
Q28: 3864 [msec] 337 [msec] 3527 [msec]
Q29: 8424 [msec] 672 [msec] 7749 [msec]
Q30: N/A N/A N/A

+ 45
- 0
homework_3/analyse/RC4-8f44cbf/tesla-sum.txt Переглянути файл

@@ -0,0 +1,45 @@
GPU: Tesla P100-PCIE-12GB
Block size: 512
--------------------------------------

Code version: V0
Total Mem-xch Sorting
Q20: 6559 [usec] 1147 [usec] 5419 [usec]
Q21: 12 [msec] 2031 [usec] 10 [msec]
Q22: 23 [msec] 3801 [usec] 19 [msec]
Q23: 49 [msec] 8534 [usec] 40 [msec]
Q24: 101 [msec] 15 [msec] 85 [msec]
Q25: 210 [msec] 29 [msec] 181 [msec]
Q26: 442 [msec] 57 [msec] 385 [msec]
Q27: 939 [msec] 112 [msec] 826 [msec]
Q28: 1986 [msec] 223 [msec] 1762 [msec]
Q29: 4211 [msec] 446 [msec] 3765 [msec]
Q30: 9071 [msec] 1037 [msec] 8034 [msec]

Code version: V1
Total Mem-xch Sorting
Q20: 4559 [usec] 1144 [usec] 3418 [usec]
Q21: 9014 [usec] 2018 [usec] 6995 [usec]
Q22: 18 [msec] 3821 [usec] 14 [msec]
Q23: 39 [msec] 8564 [usec] 30 [msec]
Q24: 80 [msec] 15 [msec] 64 [msec]
Q25: 169 [msec] 29 [msec] 139 [msec]
Q26: 358 [msec] 57 [msec] 300 [msec]
Q27: 761 [msec] 113 [msec] 648 [msec]
Q28: 1618 [msec] 224 [msec] 1394 [msec]
Q29: 3440 [msec] 447 [msec] 2992 [msec]
Q30: 7326 [msec] 914 [msec] 6411 [msec]

Code version: V2
Total Mem-xch Sorting
Q20: 4407 [usec] 1146 [usec] 3264 [usec]
Q21: 8751 [usec] 2025 [usec] 6728 [usec]
Q22: 17 [msec] 3818 [usec] 13 [msec]
Q23: 38 [msec] 8515 [usec] 29 [msec]
Q24: 75 [msec] 15 [msec] 60 [msec]
Q25: 159 [msec] 29 [msec] 129 [msec]
Q26: 338 [msec] 58 [msec] 280 [msec]
Q27: 719 [msec] 113 [msec] 606 [msec]
Q28: 1533 [msec] 225 [msec] 1308 [msec]
Q29: 3262 [msec] 447 [msec] 2815 [msec]
Q30: 6932 [msec] 883 [msec] 6048 [msec]

BIN
homework_3/report/homework_3_report.pdf Переглянути файл


+ 215
- 6
homework_3/report/homework_3_report.tex Переглянути файл

@@ -95,6 +95,8 @@
% 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}

@@ -184,7 +186,7 @@
Έτσι το κάθε thread εκτελεί μια ανταλλαγή με την άνω διευθυνσιοδότηση και ο αλγόριθμος τερματίζει στο τέλος του διπλού βρόχου.
Το πλεονέκτημα αυτής της μεθόδου είναι η απλότητά της, κάτι που προφανώς πληρώνουμε με έναν αρκετά μεγάλο αριθμό κλήσεων προς την κάρτα γραφικών το οποίο προσθέτει overhead, όπως θα δούμε και παρακάτω.

\subsection{Δεύτερη έκδοση (V1)} \label{V1}
\subsection{Δεύτερη έκδοση (V1)} \label{sec:V1}
Από την εκφώνηση της εργασίας ήδη υπάρχουν οδηγίες για βελτιστοποιήσεις.
Η πρώτη βελτιστοποίηση λοιπόν αφορά αυτές ακριβώς τις κλήσεις των kernels, των συναρτήσεων δηλαδή που επιταχύνονται από την κάρτα γραφικών.
Γίνεται δηλαδή η \textbf{θεώρηση} από την εκφώνηση, ότι \textbf{η ελαχιστοποίηση αυτών των κλήσεων θα βελτιστοποιήσει τη διαδικασία}.
@@ -221,24 +223,231 @@
\]

\WrapFigure{0.35}{r}{fig:V1Unrolling}{img/V1Unrolling.png}{Βήματα διτονικής ταξινόμησης πίνακα 128 θέσεων (block size = 16).}
Όπου το $S_i$ είναι και το μέγιστο βήμα από το οποίο μπορεί να ξεκινήσει το loop unrolling.
Όπου το $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()})για μικρότερες.
Η μία (\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{συγχρονισμού}.
Το μόνο που μένει πλέον είναι το πρόβλημα του \textbf{συγχρονισμού}.
Εδώ χρειάζεται να κρατήσουμε τον συγχρονισμό που είχαμε σε κάθε κάθετη γραμμή του σχήματος \ref{fig:bitonicSort}, απλώς αυτός θα λάβει χώρα εσωτερικά του kernel.
Η συνάρτηση που μας δίνει αυτή τη δυνατότητα είναι η \textit{\_\_syncthreads()}, την οποία και καλούμε μετά από την ολοκλήρωση της κάθε ακολουθίας.
Η συνάρτηση που μας δίνει αυτή τη δυνατότητα είναι η \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<<<Nbl, Nth, kernelMemSize>>>()}.
\par
Η αλλαγή για την δεύτερη έκδοση βρίσκεται επί της ουσίας στη συνάρτηση \textit{inBlockStep()}.
Η συνάρτηση \textit{interBlockStep()} παραμένει ως έχει, καθώς εργάζεται με αποστάσεις μεγαλύτερες αυτών που χωρούν στη shared μνήμη που έχουμε δεσμεύσει για τα blocks.

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

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

\par
Αρχικά για να διερευνήσουμε τις επιδόσεις εκτελέσαμε τοπικά σε μία GTX 1650 τις εκδόσεις και είδαμε ότι η 2η έκδοση υστερεί κατά τι.
Έπειτα επιβεβαιώσαμε τις ανησυχίες μας στη συστοιχία όπου στην 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{Σύγκριση των εκδόσεων 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{Τρίτη έκδοση (V2)} \label{V2}
\subsection{Ελαχιστοποίηση των αρχικών \textit{inBlockStep()} κλήσεων} \label{sec:Opt_prephase}

Μια βελτιστοποίηση που σκεφτήκαμε και δεν έχει να κάνει με το shared memory optimization είναι ο αριθμός των \textit{inBlockStep()} στην αρχή του αλγόριθμου.
Αν προσέξει κανείς το σχήμα \ref{fig:V1Unrolling} εύκολα θα παρατηρήσει πως στην αρχή του αλγόριθμου ο διπλός βρόχος περιέχει βήματα που χωράνε εξολοκλήρου στο block.
Παρόλα αυτά γίνονται ξεχωριστές κλήσεις στην \textit{inBlockStep()} χωρίς να μεσολαβεί κάποια κλήση στην \textit{interBlockStep()}.
Μπορούμε έτσι να \textbf{συνενώσουμε όλες αυτές τις πρώτες ξεχωριστές κλήσεις της \textit{inBlockStep()} σε μία συνάρτηση} η οποία δεν χρειάζεται να επιστρέφει παρά μόνο όταν τα βήματα του βρόχου φτάσουν να χρειαστεί να καλέσουν την πρώτη \textit{interBlockStep()}.
Η συνάρτηση αυτή είναι η \textit{prephace()}.
\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 εκτελέσεις.

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

\begin{tabular}{ll|rrrrrr}
& & Q20 & Q22 & Q24 & Q26 & Q28 & Q30 \\ \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}

\end{document}

BIN
homework_3/report/img/V1Addressing.png Переглянути файл

Перед Після
Ширина: 1916  |  Висота: 488  |  Розмір: 55 KiB Ширина: 1916  |  Висота: 488  |  Розмір: 56 KiB

+ 187
- 187
homework_3/report/img/V1Addressing.svg Переглянути файл

@@ -7,7 +7,7 @@
viewBox="0 0 297 210"
version="1.1"
id="svg1"
inkscape:version="1.4 (1:1.4+202410161351+e7c3feb100)"
inkscape:version="1.4 (e7c3feb100, 2024-10-09)"
sodipodi:docname="V1Addressing.svg"
xmlns:inkscape="http://www.inkscape.org/namespaces/inkscape"
xmlns:sodipodi="http://sodipodi.sourceforge.net/DTD/sodipodi-0.dtd"
@@ -24,9 +24,9 @@
inkscape:deskcolor="#d1d1d1"
inkscape:document-units="mm"
showguides="true"
inkscape:zoom="0.60177118"
inkscape:cx="422.91822"
inkscape:cy="571.64585"
inkscape:zoom="0.85103296"
inkscape:cx="536.99448"
inkscape:cy="440.64098"
inkscape:window-width="1920"
inkscape:window-height="1011"
inkscape:window-x="0"
@@ -427,9 +427,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="tspan3"><tspan
id="tspan17"><tspan
style="fill:#000000"
id="tspan2">0</tspan></tspan></text>
id="tspan1">0</tspan></tspan></text>
<text
xml:space="preserve"
transform="matrix(0.26458333,0,0,0.26458333,215.16783,38.453616)"
@@ -437,9 +437,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="tspan5"><tspan
id="tspan19"><tspan
style="fill:#000000"
id="tspan4">N</tspan></tspan></text>
id="tspan18">N</tspan></tspan></text>
<text
xml:space="preserve"
transform="matrix(0.26458333,0,0,0.26458333,-24.658218,67.876598)"
@@ -447,44 +447,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="tspan9"><tspan
id="tspan23"><tspan
style="fill:#000000"
id="tspan6">Tid</tspan><tspan
id="tspan20">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan7">k</tspan><tspan
id="tspan21">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 0 0 1.0199723"
style="fill:#000000"
id="tspan8"> = threadIdx.x + 2*blockIdx.x*blockDim.x
id="tspan22"> = threadIdx.x + 2*blockIdx.x*blockDim.x
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan11"><tspan
id="tspan25"><tspan
style="fill:#000000"
id="tspan10">
id="tspan24">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan13"><tspan
id="tspan27"><tspan
style="fill:#000000"
id="tspan12">
id="tspan26">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan15"><tspan
id="tspan29"><tspan
style="fill:#000000"
id="tspan14">
id="tspan28">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan191"><tspan
id="tspan31"><tspan
style="fill:#000000"
id="tspan16">
id="tspan30">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan193"><tspan
id="tspan33"><tspan
style="fill:#000000"
id="tspan192">
id="tspan32">
</tspan></tspan></text>
<text
xml:space="preserve"
@@ -498,49 +498,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="tspan199"><tspan
id="tspan39"><tspan
style="fill:#000000"
id="tspan194">Tid</tspan><tspan
id="tspan34">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan195">k</tspan><tspan
id="tspan35">k</tspan><tspan
dx="0 1.02 -1.02 1.0199995 -1.02"
style="fill:#000000"
id="tspan196">' = Tid</tspan><tspan
id="tspan36">' = Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan197">k</tspan><tspan
id="tspan37">k</tspan><tspan
dx="1.0200062 -1.02 1.0199995 -1.02 0 0 0 0 0 0 0 0 0 1.019999"
style="fill:#000000"
id="tspan198"> + blockDim.x
id="tspan38"> + blockDim.x
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan201"><tspan
id="tspan41"><tspan
style="fill:#000000"
id="tspan200">
id="tspan40">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan203"><tspan
id="tspan43"><tspan
style="fill:#000000"
id="tspan202">
id="tspan42">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan205"><tspan
id="tspan45"><tspan
style="fill:#000000"
id="tspan204">
id="tspan44">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan207"><tspan
id="tspan47"><tspan
style="fill:#000000"
id="tspan206">
id="tspan46">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan209"><tspan
id="tspan49"><tspan
style="fill:#000000"
id="tspan208">
id="tspan48">
</tspan></tspan></text>
<path
style="fill:none;stroke:#000000;stroke-width:0.629912;stroke-dasharray:1.25982, 0.629912;stroke-dashoffset:0;stroke-opacity:1"
@@ -559,40 +559,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="tspan211"><tspan
id="tspan51"><tspan
dx="0 0 0 0 1.02"
style="fill:#000000"
id="tspan210">Pid'
id="tspan50">Pid'
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan213"><tspan
id="tspan53"><tspan
style="fill:#000000"
id="tspan212">
id="tspan52">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan215"><tspan
id="tspan55"><tspan
style="fill:#000000"
id="tspan214">
id="tspan54">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan217"><tspan
id="tspan57"><tspan
style="fill:#000000"
id="tspan216">
id="tspan56">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan219"><tspan
id="tspan59"><tspan
style="fill:#000000"
id="tspan218">
id="tspan58">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan221"><tspan
id="tspan61"><tspan
style="fill:#000000"
id="tspan220">
id="tspan60">
</tspan></tspan></text>
<path
style="fill:none;stroke:#000000;stroke-width:0.613089;stroke-dasharray:none;stroke-opacity:1;marker-start:url(#marker4-6);marker-end:url(#Dot-0)"
@@ -606,44 +606,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="tspan225"><tspan
id="tspan65"><tspan
style="fill:#000000"
id="tspan222">Tid</tspan><tspan
id="tspan62">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan223">1</tspan><tspan
id="tspan63">1</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan224">
id="tspan64">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan227"><tspan
id="tspan67"><tspan
style="fill:#000000"
id="tspan226">
id="tspan66">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan229"><tspan
id="tspan69"><tspan
style="fill:#000000"
id="tspan228">
id="tspan68">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan231"><tspan
id="tspan71"><tspan
style="fill:#000000"
id="tspan230">
id="tspan70">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan233"><tspan
id="tspan73"><tspan
style="fill:#000000"
id="tspan232">
id="tspan72">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan235"><tspan
id="tspan75"><tspan
style="fill:#000000"
id="tspan234">
id="tspan74">
</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)"
@@ -657,44 +657,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="tspan239"><tspan
id="tspan79"><tspan
style="fill:#000000"
id="tspan236">Tid</tspan><tspan
id="tspan76">Tid</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan237">2</tspan><tspan
id="tspan77">2</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan238">
id="tspan78">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan241"><tspan
id="tspan81"><tspan
style="fill:#000000"
id="tspan240">
id="tspan80">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan243"><tspan
id="tspan83"><tspan
style="fill:#000000"
id="tspan242">
id="tspan82">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan245"><tspan
id="tspan85"><tspan
style="fill:#000000"
id="tspan244">
id="tspan84">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan247"><tspan
id="tspan87"><tspan
style="fill:#000000"
id="tspan246">
id="tspan86">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan249"><tspan
id="tspan89"><tspan
style="fill:#000000"
id="tspan248">
id="tspan88">
</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)"
@@ -708,44 +708,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="tspan253"><tspan
id="tspan93"><tspan
style="fill:#000000"
id="tspan250">Tid'</tspan><tspan
id="tspan90">Tid'</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan251">2</tspan><tspan
id="tspan91">2</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan252">
id="tspan92">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan255"><tspan
id="tspan95"><tspan
style="fill:#000000"
id="tspan254">
id="tspan94">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan257"><tspan
id="tspan97"><tspan
style="fill:#000000"
id="tspan256">
id="tspan96">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan259"><tspan
id="tspan99"><tspan
style="fill:#000000"
id="tspan258">
id="tspan98">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan261"><tspan
id="tspan101"><tspan
style="fill:#000000"
id="tspan260">
id="tspan100">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan263"><tspan
id="tspan103"><tspan
style="fill:#000000"
id="tspan262">
id="tspan102">
</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)"
@@ -779,44 +779,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-3-2);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan267"><tspan
id="tspan107"><tspan
style="fill:#000000"
id="tspan264">Bl</tspan><tspan
id="tspan104">Bl</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan265">k</tspan><tspan
id="tspan105">2k</tspan><tspan
dx="1.0200005"
style="fill:#000000"
id="tspan266">
id="tspan106">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan269"><tspan
id="tspan109"><tspan
style="fill:#000000"
id="tspan268">
id="tspan108">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan271"><tspan
id="tspan111"><tspan
style="fill:#000000"
id="tspan270">
id="tspan110">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan273"><tspan
id="tspan113"><tspan
style="fill:#000000"
id="tspan272">
id="tspan112">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan275"><tspan
id="tspan115"><tspan
style="fill:#000000"
id="tspan274">
id="tspan114">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan277"><tspan
id="tspan117"><tspan
style="fill:#000000"
id="tspan276">
id="tspan116">
</tspan></tspan></text>
<text
xml:space="preserve"
@@ -825,44 +825,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-3-2-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="tspan281"><tspan
id="tspan121"><tspan
style="fill:#000000"
id="tspan278">Bl</tspan><tspan
id="tspan118">Bl</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan279">k+1</tspan><tspan
dx="1.0200005"
id="tspan119">2k+1</tspan><tspan
dx="1.0199986"
style="fill:#000000"
id="tspan280">
id="tspan120">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan283"><tspan
id="tspan123"><tspan
style="fill:#000000"
id="tspan282">
id="tspan122">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan285"><tspan
id="tspan125"><tspan
style="fill:#000000"
id="tspan284">
id="tspan124">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan287"><tspan
id="tspan127"><tspan
style="fill:#000000"
id="tspan286">
id="tspan126">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan289"><tspan
id="tspan129"><tspan
style="fill:#000000"
id="tspan288">
id="tspan128">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan291"><tspan
id="tspan131"><tspan
style="fill:#000000"
id="tspan290">
id="tspan130">
</tspan></tspan></text>
<text
xml:space="preserve"
@@ -871,40 +871,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-3-7);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan293"><tspan
id="tspan133"><tspan
dx="0 0 0 0 0 0 0 0 0 0 1.019999"
style="fill:#000000"
id="tspan292">blockDim.x
id="tspan132">blockDim.x
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan295"><tspan
id="tspan135"><tspan
style="fill:#000000"
id="tspan294">
id="tspan134">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan297"><tspan
id="tspan137"><tspan
style="fill:#000000"
id="tspan296">
id="tspan136">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan299"><tspan
id="tspan139"><tspan
style="fill:#000000"
id="tspan298">
id="tspan138">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan301"><tspan
id="tspan141"><tspan
style="fill:#000000"
id="tspan300">
id="tspan140">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan303"><tspan
id="tspan143"><tspan
style="fill:#000000"
id="tspan302">
id="tspan142">
</tspan></tspan></text>
<text
xml:space="preserve"
@@ -913,44 +913,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-3-2-1-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="tspan307"><tspan
id="tspan147"><tspan
style="fill:#000000"
id="tspan304">Bl</tspan><tspan
id="tspan144">Bl</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan305">k+2</tspan><tspan
dx="1.0200005"
id="tspan145">2k+2</tspan><tspan
dx="1.0199986"
style="fill:#000000"
id="tspan306">
id="tspan146">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan309"><tspan
id="tspan149"><tspan
style="fill:#000000"
id="tspan308">
id="tspan148">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan311"><tspan
id="tspan151"><tspan
style="fill:#000000"
id="tspan310">
id="tspan150">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan313"><tspan
id="tspan153"><tspan
style="fill:#000000"
id="tspan312">
id="tspan152">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan315"><tspan
id="tspan155"><tspan
style="fill:#000000"
id="tspan314">
id="tspan154">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan317"><tspan
id="tspan157"><tspan
style="fill:#000000"
id="tspan316">
id="tspan156">
</tspan></tspan></text>
<g
id="g16">
@@ -982,44 +982,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-3-2-1-9-7);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan321"><tspan
id="tspan161"><tspan
style="fill:#000000"
id="tspan318">Bl</tspan><tspan
id="tspan158">Bl</tspan><tspan
style="font-size:65%;baseline-shift:sub"
id="tspan319">k+3</tspan><tspan
dx="1.0200005"
id="tspan159">2k+3</tspan><tspan
dx="1.0199986"
style="fill:#000000"
id="tspan320">
id="tspan160">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan323"><tspan
id="tspan163"><tspan
style="fill:#000000"
id="tspan322">
id="tspan162">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan325"><tspan
id="tspan165"><tspan
style="fill:#000000"
id="tspan324">
id="tspan164">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan327"><tspan
id="tspan167"><tspan
style="fill:#000000"
id="tspan326">
id="tspan166">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan329"><tspan
id="tspan169"><tspan
style="fill:#000000"
id="tspan328">
id="tspan168">
</tspan></tspan><tspan
x="171.64648"
y="133.24558"
id="tspan331"><tspan
id="tspan171"><tspan
style="fill:#000000"
id="tspan330">
id="tspan170">
</tspan></tspan></text>
<g
id="g15">
@@ -1060,40 +1060,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-3-7-7);display:inline;fill:none;stroke:#000000;stroke-width:0.865512;stroke-dasharray:none;stroke-opacity:1"><tspan
x="171.64648"
y="131.94739"
id="tspan333"><tspan
id="tspan173"><tspan
dx="0 0 0 0 1.0200005 -1.02 0 0 0 1.0200009 -1.02 0 1.0199966 -1.02 0 0 0 0 0 1.0199914"
style="fill:#000000"
id="tspan332">Next Pair of Blocks
id="tspan172">Next Pair of Blocks
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan335"><tspan
id="tspan175"><tspan
style="fill:#000000"
id="tspan334">
id="tspan174">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan337"><tspan
id="tspan177"><tspan
style="fill:#000000"
id="tspan336">
id="tspan176">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan339"><tspan
id="tspan179"><tspan
style="fill:#000000"
id="tspan338">
id="tspan178">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan341"><tspan
id="tspan181"><tspan
style="fill:#000000"
id="tspan340">
id="tspan180">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan343"><tspan
id="tspan183"><tspan
style="fill:#000000"
id="tspan342">
id="tspan182">
</tspan></tspan></text>
<path
style="fill:none;fill-opacity:1;stroke:#000000;stroke-width:0.5;stroke-dasharray:none;stroke-dashoffset:0;stroke-opacity:1;marker-start:url(#marker4-6-7);marker-end:url(#marker4-6-7)"
@@ -1107,40 +1107,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-3-7-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="tspan345"><tspan
id="tspan185"><tspan
dx="0 0 0 0 0 0 0 1.0200005 -1.02 0 0 0 1.0200009 -1.02 0 1.0199966 -1.02 0 0 0 0 0 1.0200067"
style="fill:#000000"
id="tspan344">Current Pair of Blocks
id="tspan184">Current Pair of Blocks
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan347"><tspan
id="tspan187"><tspan
style="fill:#000000"
id="tspan346">
id="tspan186">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan349"><tspan
id="tspan189"><tspan
style="fill:#000000"
id="tspan348">
id="tspan188">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan351"><tspan
id="tspan191"><tspan
style="fill:#000000"
id="tspan350">
id="tspan190">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan353"><tspan
id="tspan193"><tspan
style="fill:#000000"
id="tspan352">
id="tspan192">
</tspan></tspan><tspan
x="171.64648"
y="131.94739"
id="tspan355"><tspan
id="tspan195"><tspan
style="fill:#000000"
id="tspan354">
id="tspan194">
</tspan></tspan></text>
</g>
</svg>

BIN
homework_3/report/img/V2Addressing.png Переглянути файл

Перед Після
Ширина: 1916  |  Висота: 564  |  Розмір: 66 KiB

+ 1496
- 0
homework_3/report/img/V2Addressing.svg
Різницю між файлами не показано, бо вона завелика
Переглянути файл


Писати Попередній перегляд
Завантаження…
Відмінити
Зберегти