From beb4b32a38394d2c847ba5a733a4906781d95881 Mon Sep 17 00:00:00 2001 From: couturie Date: Sun, 3 Mar 2013 11:29:51 +0100 Subject: [PATCH] update ch3 --- BookGPU/Chapters/chapter12/ch12.aux | 196 ++++++++--------- BookGPU/Chapters/chapter13/ch13.aux | 154 ++++++------- BookGPU/Chapters/chapter16/ch16.aux | 126 +++++------ BookGPU/Chapters/chapter18/ch18.aux | 70 +++--- BookGPU/Chapters/chapter3/biblio3.bib | 35 +++ BookGPU/Chapters/chapter3/ch3.aux | 85 ++++---- BookGPU/Chapters/chapter3/ch3.tex | 79 ++++--- .../chapter3/code/kernMedianRegTri9.cu | 2 +- .../chapter3/code/kernMedianRegTri9.cu~ | 25 +-- .../chapter3/code/kernMedianSeparable.cu | 8 +- .../chapter3/code/kernMedianSeparable.cu~ | 15 +- BookGPU/Chapters/chapter3/img/debitPlot1.pdf | Bin 14363 -> 16071 bytes BookGPU/Chapters/chapter3/img/debitPlot2.pdf | Bin 15561 -> 16016 bytes BookGPU/Chapters/chapter6/ch6.aux | 202 +++++++++--------- 14 files changed, 518 insertions(+), 479 deletions(-) diff --git a/BookGPU/Chapters/chapter12/ch12.aux b/BookGPU/Chapters/chapter12/ch12.aux index 26d263a..d3b6ca3 100644 --- a/BookGPU/Chapters/chapter12/ch12.aux +++ b/BookGPU/Chapters/chapter12/ch12.aux @@ -3,107 +3,107 @@ \@writefile{toc}{\author{Rapha\IeC {\"e}l Couturier}{}} \@writefile{toc}{\author{Jacques Bahi}{}} \@writefile{loa}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {chapter}{\numberline {11}Solving sparse linear systems with GMRES and CG methods on GPU clusters}{249}} +\@writefile{toc}{\contentsline {chapter}{\numberline {11}Solving sparse linear systems with GMRES and CG methods on GPU clusters}{251}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\newlabel{ch12}{{11}{249}} -\@writefile{toc}{\contentsline {section}{\numberline {11.1}Introduction}{249}} -\newlabel{ch12:sec:01}{{11.1}{249}} -\@writefile{toc}{\contentsline {section}{\numberline {11.2}Krylov iterative methods}{250}} -\newlabel{ch12:sec:02}{{11.2}{250}} -\newlabel{ch12:eq:01}{{11.1}{250}} -\newlabel{ch12:eq:02}{{11.2}{250}} -\newlabel{ch12:eq:03}{{11.3}{250}} -\newlabel{ch12:eq:11}{{11.4}{251}} -\@writefile{toc}{\contentsline {subsection}{\numberline {11.2.1}CG method}{251}} -\newlabel{ch12:sec:02.01}{{11.2.1}{251}} -\newlabel{ch12:eq:04}{{11.5}{251}} -\newlabel{ch12:eq:05}{{11.6}{251}} -\newlabel{ch12:eq:06}{{11.7}{251}} -\newlabel{ch12:eq:07}{{11.8}{251}} -\newlabel{ch12:eq:08}{{11.9}{251}} -\newlabel{ch12:eq:09}{{11.10}{251}} -\@writefile{loa}{\contentsline {algocf}{\numberline {9}{\ignorespaces Left-preconditioned CG method\relax }}{252}} -\newlabel{ch12:alg:01}{{9}{252}} -\newlabel{ch12:eq:10}{{11.11}{252}} -\@writefile{toc}{\contentsline {subsection}{\numberline {11.2.2}GMRES method}{253}} -\newlabel{ch12:sec:02.02}{{11.2.2}{253}} -\newlabel{ch12:eq:12}{{11.12}{253}} -\newlabel{ch12:eq:13}{{11.13}{253}} -\newlabel{ch12:eq:14}{{11.14}{253}} -\newlabel{ch12:eq:15}{{11.15}{253}} -\newlabel{ch12:eq:16}{{11.16}{253}} -\newlabel{ch12:eq:17}{{11.17}{253}} -\newlabel{ch12:eq:18}{{11.18}{253}} -\newlabel{ch12:eq:19}{{11.19}{253}} -\@writefile{loa}{\contentsline {algocf}{\numberline {10}{\ignorespaces Left-preconditioned GMRES method with restarts\relax }}{254}} -\newlabel{ch12:alg:02}{{10}{254}} -\@writefile{toc}{\contentsline {section}{\numberline {11.3}Parallel implementation on a GPU cluster}{255}} -\newlabel{ch12:sec:03}{{11.3}{255}} -\@writefile{toc}{\contentsline {subsection}{\numberline {11.3.1}Data partitioning}{255}} -\newlabel{ch12:sec:03.01}{{11.3.1}{255}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.1}{\ignorespaces A data partitioning of the sparse matrix $A$, the solution vector $x$ and the right-hand side $b$ into four portions.\relax }}{256}} -\newlabel{ch12:fig:01}{{11.1}{256}} -\@writefile{toc}{\contentsline {subsection}{\numberline {11.3.2}GPU computing}{256}} -\newlabel{ch12:sec:03.02}{{11.3.2}{256}} -\@writefile{toc}{\contentsline {subsection}{\numberline {11.3.3}Data communications}{257}} -\newlabel{ch12:sec:03.03}{{11.3.3}{257}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.2}{\ignorespaces Data exchanges between \textit {Node 1} and its neighbors \textit {Node 0}, \textit {Node 2} and \textit {Node 3}.\relax }}{258}} -\newlabel{ch12:fig:02}{{11.2}{258}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.3}{\ignorespaces Columns reordering of a sparse sub-matrix.\relax }}{259}} -\newlabel{ch12:fig:03}{{11.3}{259}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.4}{\ignorespaces General scheme of the GPU cluster of tests composed of six machines, each with two GPUs.\relax }}{260}} -\newlabel{ch12:fig:04}{{11.4}{260}} -\@writefile{toc}{\contentsline {section}{\numberline {11.4}Experimental results}{260}} -\newlabel{ch12:sec:04}{{11.4}{260}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.5}{\ignorespaces Sketches of sparse matrices chosen from the Davis's collection.\relax }}{261}} -\newlabel{ch12:fig:05}{{11.5}{261}} -\@writefile{lot}{\contentsline {table}{\numberline {11.1}{\ignorespaces Main characteristics of sparse matrices chosen from the Davis's collection.\relax }}{262}} -\newlabel{ch12:tab:01}{{11.1}{262}} -\@writefile{lot}{\contentsline {table}{\numberline {11.2}{\ignorespaces Performances of the parallel CG method on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs.\relax }}{262}} -\newlabel{ch12:tab:02}{{11.2}{262}} -\@writefile{lot}{\contentsline {table}{\numberline {11.3}{\ignorespaces Performances of the parallel GMRES method on a cluster 24 CPU cores vs. on cluster of 12 GPUs.\relax }}{263}} -\newlabel{ch12:tab:03}{{11.3}{263}} -\newlabel{ch12:eq:20}{{11.20}{263}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.6}{\ignorespaces Parallel generation of a large sparse matrix by four computing nodes.\relax }}{264}} -\newlabel{ch12:fig:06}{{11.6}{264}} -\@writefile{lot}{\contentsline {table}{\numberline {11.4}{\ignorespaces Main characteristics of sparse banded matrices generated from those of the Davis's collection.\relax }}{265}} -\newlabel{ch12:tab:04}{{11.4}{265}} -\@writefile{lot}{\contentsline {table}{\numberline {11.5}{\ignorespaces Performances of the parallel CG method for solving linear systems associated to sparse banded matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs.\relax }}{265}} -\newlabel{ch12:tab:05}{{11.5}{265}} -\@writefile{toc}{\contentsline {section}{\numberline {11.5}Hypergraph partitioning}{265}} -\newlabel{ch12:sec:05}{{11.5}{265}} -\@writefile{lot}{\contentsline {table}{\numberline {11.6}{\ignorespaces Performances of the parallel GMRES method for solving linear systems associated to sparse banded matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs.\relax }}{266}} -\newlabel{ch12:tab:06}{{11.6}{266}} -\@writefile{lot}{\contentsline {table}{\numberline {11.7}{\ignorespaces Main characteristics of sparse five-bands matrices generated from those of the Davis's collection.\relax }}{266}} -\newlabel{ch12:tab:07}{{11.7}{266}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.7}{\ignorespaces Parallel generation of a large sparse five-bands matrix by four computing nodes.\relax }}{267}} -\newlabel{ch12:fig:07}{{11.7}{267}} -\@writefile{lot}{\contentsline {table}{\numberline {11.8}{\ignorespaces Performances of parallel CG solver for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs\relax }}{267}} -\newlabel{ch12:tab:08}{{11.8}{267}} -\@writefile{lot}{\contentsline {table}{\numberline {11.9}{\ignorespaces Performances of parallel GMRES solver for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs\relax }}{268}} -\newlabel{ch12:tab:09}{{11.9}{268}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.8}{\ignorespaces An example of the hypergraph partitioning of a sparse matrix decomposed between three computing nodes.\relax }}{269}} -\newlabel{ch12:fig:08}{{11.8}{269}} -\@writefile{lot}{\contentsline {table}{\numberline {11.10}{\ignorespaces Performances of the parallel CG solver using hypergraph partitioning for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPU.\relax }}{270}} -\newlabel{ch12:tab:10}{{11.10}{270}} -\@writefile{lot}{\contentsline {table}{\numberline {11.11}{\ignorespaces Performances of the parallel GMRES solver using hypergraph partitioning for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPU.\relax }}{271}} -\newlabel{ch12:tab:11}{{11.11}{271}} -\@writefile{lot}{\contentsline {table}{\numberline {11.12}{\ignorespaces The total communication volume between 12 GPU computing nodes without and with the hypergraph partitioning method.\relax }}{272}} -\newlabel{ch12:tab:12}{{11.12}{272}} -\newlabel{ch12:fig:09.01}{{11.9(a)}{273}} -\newlabel{sub@ch12:fig:09.01}{{(a)}{273}} -\newlabel{ch12:fig:09.02}{{11.9(b)}{273}} -\newlabel{sub@ch12:fig:09.02}{{(b)}{273}} -\@writefile{lof}{\contentsline {figure}{\numberline {11.9}{\ignorespaces Weak-scaling of the parallel CG and GMRES solvers on a GPU cluster for solving large sparse linear systems.\relax }}{273}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Sparse band matrices}}}{273}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Sparse five-bands matrices}}}{273}} -\newlabel{ch12:fig:09}{{11.9}{273}} -\@writefile{toc}{\contentsline {section}{\numberline {11.6}Conclusion}{273}} -\newlabel{ch12:sec:06}{{11.6}{273}} -\@writefile{toc}{\contentsline {section}{Bibliography}{274}} +\newlabel{ch12}{{11}{251}} +\@writefile{toc}{\contentsline {section}{\numberline {11.1}Introduction}{251}} +\newlabel{ch12:sec:01}{{11.1}{251}} +\@writefile{toc}{\contentsline {section}{\numberline {11.2}Krylov iterative methods}{252}} +\newlabel{ch12:sec:02}{{11.2}{252}} +\newlabel{ch12:eq:01}{{11.1}{252}} +\newlabel{ch12:eq:02}{{11.2}{252}} +\newlabel{ch12:eq:03}{{11.3}{252}} +\newlabel{ch12:eq:11}{{11.4}{253}} +\@writefile{toc}{\contentsline {subsection}{\numberline {11.2.1}CG method}{253}} +\newlabel{ch12:sec:02.01}{{11.2.1}{253}} +\newlabel{ch12:eq:04}{{11.5}{253}} +\newlabel{ch12:eq:05}{{11.6}{253}} +\newlabel{ch12:eq:06}{{11.7}{253}} +\newlabel{ch12:eq:07}{{11.8}{253}} +\newlabel{ch12:eq:08}{{11.9}{253}} +\newlabel{ch12:eq:09}{{11.10}{253}} +\@writefile{loa}{\contentsline {algocf}{\numberline {9}{\ignorespaces Left-preconditioned CG method\relax }}{254}} +\newlabel{ch12:alg:01}{{9}{254}} +\newlabel{ch12:eq:10}{{11.11}{254}} +\@writefile{toc}{\contentsline {subsection}{\numberline {11.2.2}GMRES method}{255}} +\newlabel{ch12:sec:02.02}{{11.2.2}{255}} +\newlabel{ch12:eq:12}{{11.12}{255}} +\newlabel{ch12:eq:13}{{11.13}{255}} +\newlabel{ch12:eq:14}{{11.14}{255}} +\newlabel{ch12:eq:15}{{11.15}{255}} +\newlabel{ch12:eq:16}{{11.16}{255}} +\newlabel{ch12:eq:17}{{11.17}{255}} +\newlabel{ch12:eq:18}{{11.18}{255}} +\newlabel{ch12:eq:19}{{11.19}{255}} +\@writefile{loa}{\contentsline {algocf}{\numberline {10}{\ignorespaces Left-preconditioned GMRES method with restarts\relax }}{256}} +\newlabel{ch12:alg:02}{{10}{256}} +\@writefile{toc}{\contentsline {section}{\numberline {11.3}Parallel implementation on a GPU cluster}{257}} +\newlabel{ch12:sec:03}{{11.3}{257}} +\@writefile{toc}{\contentsline {subsection}{\numberline {11.3.1}Data partitioning}{257}} +\newlabel{ch12:sec:03.01}{{11.3.1}{257}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.1}{\ignorespaces A data partitioning of the sparse matrix $A$, the solution vector $x$ and the right-hand side $b$ into four portions.\relax }}{258}} +\newlabel{ch12:fig:01}{{11.1}{258}} +\@writefile{toc}{\contentsline {subsection}{\numberline {11.3.2}GPU computing}{258}} +\newlabel{ch12:sec:03.02}{{11.3.2}{258}} +\@writefile{toc}{\contentsline {subsection}{\numberline {11.3.3}Data communications}{259}} +\newlabel{ch12:sec:03.03}{{11.3.3}{259}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.2}{\ignorespaces Data exchanges between \textit {Node 1} and its neighbors \textit {Node 0}, \textit {Node 2} and \textit {Node 3}.\relax }}{260}} +\newlabel{ch12:fig:02}{{11.2}{260}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.3}{\ignorespaces Columns reordering of a sparse sub-matrix.\relax }}{261}} +\newlabel{ch12:fig:03}{{11.3}{261}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.4}{\ignorespaces General scheme of the GPU cluster of tests composed of six machines, each with two GPUs.\relax }}{262}} +\newlabel{ch12:fig:04}{{11.4}{262}} +\@writefile{toc}{\contentsline {section}{\numberline {11.4}Experimental results}{262}} +\newlabel{ch12:sec:04}{{11.4}{262}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.5}{\ignorespaces Sketches of sparse matrices chosen from the Davis's collection.\relax }}{263}} +\newlabel{ch12:fig:05}{{11.5}{263}} +\@writefile{lot}{\contentsline {table}{\numberline {11.1}{\ignorespaces Main characteristics of sparse matrices chosen from the Davis's collection.\relax }}{264}} +\newlabel{ch12:tab:01}{{11.1}{264}} +\@writefile{lot}{\contentsline {table}{\numberline {11.2}{\ignorespaces Performances of the parallel CG method on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs.\relax }}{264}} +\newlabel{ch12:tab:02}{{11.2}{264}} +\@writefile{lot}{\contentsline {table}{\numberline {11.3}{\ignorespaces Performances of the parallel GMRES method on a cluster 24 CPU cores vs. on cluster of 12 GPUs.\relax }}{265}} +\newlabel{ch12:tab:03}{{11.3}{265}} +\newlabel{ch12:eq:20}{{11.20}{265}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.6}{\ignorespaces Parallel generation of a large sparse matrix by four computing nodes.\relax }}{266}} +\newlabel{ch12:fig:06}{{11.6}{266}} +\@writefile{lot}{\contentsline {table}{\numberline {11.4}{\ignorespaces Main characteristics of sparse banded matrices generated from those of the Davis's collection.\relax }}{267}} +\newlabel{ch12:tab:04}{{11.4}{267}} +\@writefile{lot}{\contentsline {table}{\numberline {11.5}{\ignorespaces Performances of the parallel CG method for solving linear systems associated to sparse banded matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs.\relax }}{267}} +\newlabel{ch12:tab:05}{{11.5}{267}} +\@writefile{toc}{\contentsline {section}{\numberline {11.5}Hypergraph partitioning}{267}} +\newlabel{ch12:sec:05}{{11.5}{267}} +\@writefile{lot}{\contentsline {table}{\numberline {11.6}{\ignorespaces Performances of the parallel GMRES method for solving linear systems associated to sparse banded matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs.\relax }}{268}} +\newlabel{ch12:tab:06}{{11.6}{268}} +\@writefile{lot}{\contentsline {table}{\numberline {11.7}{\ignorespaces Main characteristics of sparse five-bands matrices generated from those of the Davis's collection.\relax }}{268}} +\newlabel{ch12:tab:07}{{11.7}{268}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.7}{\ignorespaces Parallel generation of a large sparse five-bands matrix by four computing nodes.\relax }}{269}} +\newlabel{ch12:fig:07}{{11.7}{269}} +\@writefile{lot}{\contentsline {table}{\numberline {11.8}{\ignorespaces Performances of parallel CG solver for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs\relax }}{269}} +\newlabel{ch12:tab:08}{{11.8}{269}} +\@writefile{lot}{\contentsline {table}{\numberline {11.9}{\ignorespaces Performances of parallel GMRES solver for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPUs\relax }}{270}} +\newlabel{ch12:tab:09}{{11.9}{270}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.8}{\ignorespaces An example of the hypergraph partitioning of a sparse matrix decomposed between three computing nodes.\relax }}{271}} +\newlabel{ch12:fig:08}{{11.8}{271}} +\@writefile{lot}{\contentsline {table}{\numberline {11.10}{\ignorespaces Performances of the parallel CG solver using hypergraph partitioning for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPU.\relax }}{272}} +\newlabel{ch12:tab:10}{{11.10}{272}} +\@writefile{lot}{\contentsline {table}{\numberline {11.11}{\ignorespaces Performances of the parallel GMRES solver using hypergraph partitioning for solving linear systems associated to sparse five-bands matrices on a cluster of 24 CPU cores vs. on a cluster of 12 GPU.\relax }}{273}} +\newlabel{ch12:tab:11}{{11.11}{273}} +\@writefile{lot}{\contentsline {table}{\numberline {11.12}{\ignorespaces The total communication volume between 12 GPU computing nodes without and with the hypergraph partitioning method.\relax }}{274}} +\newlabel{ch12:tab:12}{{11.12}{274}} +\newlabel{ch12:fig:09.01}{{11.9(a)}{275}} +\newlabel{sub@ch12:fig:09.01}{{(a)}{275}} +\newlabel{ch12:fig:09.02}{{11.9(b)}{275}} +\newlabel{sub@ch12:fig:09.02}{{(b)}{275}} +\@writefile{lof}{\contentsline {figure}{\numberline {11.9}{\ignorespaces Weak-scaling of the parallel CG and GMRES solvers on a GPU cluster for solving large sparse linear systems.\relax }}{275}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Sparse band matrices}}}{275}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Sparse five-bands matrices}}}{275}} +\newlabel{ch12:fig:09}{{11.9}{275}} +\@writefile{toc}{\contentsline {section}{\numberline {11.6}Conclusion}{275}} +\newlabel{ch12:sec:06}{{11.6}{275}} +\@writefile{toc}{\contentsline {section}{Bibliography}{276}} \@setckpt{Chapters/chapter12/ch12}{ -\setcounter{page}{276} +\setcounter{page}{278} \setcounter{equation}{25} \setcounter{enumi}{4} \setcounter{enumii}{0} diff --git a/BookGPU/Chapters/chapter13/ch13.aux b/BookGPU/Chapters/chapter13/ch13.aux index 9460acd..bb6c472 100644 --- a/BookGPU/Chapters/chapter13/ch13.aux +++ b/BookGPU/Chapters/chapter13/ch13.aux @@ -5,86 +5,86 @@ \@writefile{toc}{\author{Pierre Spit\IeC {\'e}ri}{}} \@writefile{toc}{\author{Jacques Bahi}{}} \@writefile{loa}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {chapter}{\numberline {12}Solving sparse nonlinear systems of obstacle problems on GPU clusters}{277}} +\@writefile{toc}{\contentsline {chapter}{\numberline {12}Solving sparse nonlinear systems of obstacle problems on GPU clusters}{279}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\newlabel{ch13}{{12}{277}} -\@writefile{toc}{\contentsline {section}{\numberline {12.1}Introduction}{277}} -\newlabel{ch13:sec:01}{{12.1}{277}} -\@writefile{toc}{\contentsline {section}{\numberline {12.2}Obstacle problems}{278}} -\newlabel{ch13:sec:02}{{12.2}{278}} -\@writefile{toc}{\contentsline {subsection}{\numberline {12.2.1}Mathematical model}{278}} -\newlabel{ch13:sec:02.01}{{12.2.1}{278}} -\newlabel{ch13:eq:01}{{12.1}{278}} -\newlabel{ch13:eq:02}{{12.2}{278}} -\@writefile{toc}{\contentsline {subsection}{\numberline {12.2.2}Discretization}{279}} -\newlabel{ch13:sec:02.02}{{12.2.2}{279}} -\newlabel{ch13:eq:03}{{12.3}{279}} -\newlabel{ch13:eq:04}{{12.4}{279}} -\newlabel{ch13:eq:05}{{12.5}{279}} -\@writefile{toc}{\contentsline {section}{\numberline {12.3}Parallel iterative method}{280}} -\newlabel{ch13:sec:03}{{12.3}{280}} -\newlabel{ch13:eq:06}{{12.6}{280}} -\newlabel{ch13:eq:07}{{12.7}{280}} -\newlabel{ch13:eq:08}{{12.8}{280}} -\newlabel{ch13:eq:09}{{12.9}{280}} -\newlabel{ch13:eq:10}{{12.10}{281}} -\newlabel{ch13:eq:11}{{12.11}{281}} -\newlabel{ch13:eq:12}{{12.12}{281}} -\newlabel{ch13:eq:13}{{12.13}{282}} -\newlabel{ch13:eq:14}{{12.14}{282}} -\newlabel{ch13:eq:15}{{12.15}{282}} -\newlabel{ch13:eq:16}{{12.16}{282}} -\@writefile{toc}{\contentsline {section}{\numberline {12.4}Parallel implementation on a GPU cluster}{283}} -\newlabel{ch13:sec:04}{{12.4}{283}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.1}{\ignorespaces Data partitioning of a problem to be solved among $S=3\times 4$ computing nodes.\relax }}{283}} -\newlabel{ch13:fig:01}{{12.1}{283}} -\@writefile{loa}{\contentsline {algocf}{\numberline {11}{\ignorespaces Parallel solving of the obstacle problem on a GPU cluster\relax }}{284}} -\newlabel{ch13:alg:01}{{11}{284}} -\newlabel{ch13:eq:18}{{12.17}{284}} -\@writefile{loa}{\contentsline {algocf}{\numberline {12}{\ignorespaces Parallel iterative solving of the nonlinear systems on a GPU cluster ($Solve()$ function)\relax }}{285}} -\newlabel{ch13:alg:02}{{12}{285}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.2}{\ignorespaces Decomposition of a sub-problem in a GPU into $nz$ slices.\relax }}{286}} -\newlabel{ch13:fig:02}{{12.2}{286}} -\newlabel{ch13:list:01}{{12.1}{286}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.1}Skeleton codes of a GPU kernel and a CPU function}{286}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.3}{\ignorespaces Matrix constant coefficients in a three-dimensional domain.\relax }}{288}} -\newlabel{ch13:fig:03}{{12.3}{288}} -\newlabel{ch13:eq:17}{{12.18}{288}} -\newlabel{ch13:list:02}{{12.2}{288}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.2}GPU kernels of the projected Richardson method}{288}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.4}{\ignorespaces Computation of a vector element with the projected Richardson method.\relax }}{290}} -\newlabel{ch13:fig:04}{{12.4}{290}} -\newlabel{ch13:list:03}{{12.3}{290}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.3}Memory access to the cache texture memory}{290}} -\@writefile{toc}{\contentsline {section}{\numberline {12.5}Experimental tests on a GPU cluster}{291}} -\newlabel{ch13:sec:05}{{12.5}{291}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.5}{\ignorespaces GPU cluster of tests composed of 12 computing nodes (six machines, each with two GPUs.\relax }}{293}} -\newlabel{ch13:fig:05}{{12.5}{293}} -\@writefile{lot}{\contentsline {table}{\numberline {12.1}{\ignorespaces Execution times in seconds of the parallel projected Richardson method implemented on a cluster of 24 CPU cores.\relax }}{293}} -\newlabel{ch13:tab:01}{{12.1}{293}} -\@writefile{lot}{\contentsline {table}{\numberline {12.2}{\ignorespaces Execution times in seconds of the parallel projected Richardson method implemented on a cluster of 12 GPUs.\relax }}{294}} -\newlabel{ch13:tab:02}{{12.2}{294}} -\@writefile{toc}{\contentsline {section}{\numberline {12.6}Red-Black ordering technique}{294}} -\newlabel{ch13:sec:06}{{12.6}{294}} -\newlabel{ch13:list:04}{{12.4}{295}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.4}GPU kernels of the projected Richardson method using the red-black technique}{295}} -\newlabel{ch13:fig:06.01}{{12.6(a)}{296}} -\newlabel{sub@ch13:fig:06.01}{{(a)}{296}} -\newlabel{ch13:fig:06.02}{{12.6(b)}{296}} -\newlabel{sub@ch13:fig:06.02}{{(b)}{296}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.6}{\ignorespaces Red-Black ordering for computing the iterate vector elements in a three-dimensional space.\relax }}{296}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Red-Black ordering on x, y and z axises}}}{296}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Red-Black ordering on y axis}}}{296}} -\@writefile{lot}{\contentsline {table}{\numberline {12.3}{\ignorespaces Execution times in seconds of the parallel projected Richardson method using read-black ordering technique implemented on a cluster of 12 GPUs.\relax }}{297}} -\newlabel{ch13:tab:03}{{12.3}{297}} -\@writefile{lof}{\contentsline {figure}{\numberline {12.7}{\ignorespaces Weak scaling of both synchronous and asynchronous algorithms of the projected Richardson method using red-black ordering technique.\relax }}{298}} -\newlabel{ch13:fig:07}{{12.7}{298}} -\@writefile{toc}{\contentsline {section}{\numberline {12.7}Conclusion}{298}} -\newlabel{ch13:sec:07}{{12.7}{298}} -\@writefile{toc}{\contentsline {section}{Bibliography}{299}} +\newlabel{ch13}{{12}{279}} +\@writefile{toc}{\contentsline {section}{\numberline {12.1}Introduction}{279}} +\newlabel{ch13:sec:01}{{12.1}{279}} +\@writefile{toc}{\contentsline {section}{\numberline {12.2}Obstacle problems}{280}} +\newlabel{ch13:sec:02}{{12.2}{280}} +\@writefile{toc}{\contentsline {subsection}{\numberline {12.2.1}Mathematical model}{280}} +\newlabel{ch13:sec:02.01}{{12.2.1}{280}} +\newlabel{ch13:eq:01}{{12.1}{280}} +\newlabel{ch13:eq:02}{{12.2}{280}} +\@writefile{toc}{\contentsline {subsection}{\numberline {12.2.2}Discretization}{281}} +\newlabel{ch13:sec:02.02}{{12.2.2}{281}} +\newlabel{ch13:eq:03}{{12.3}{281}} +\newlabel{ch13:eq:04}{{12.4}{281}} +\newlabel{ch13:eq:05}{{12.5}{281}} +\@writefile{toc}{\contentsline {section}{\numberline {12.3}Parallel iterative method}{282}} +\newlabel{ch13:sec:03}{{12.3}{282}} +\newlabel{ch13:eq:06}{{12.6}{282}} +\newlabel{ch13:eq:07}{{12.7}{282}} +\newlabel{ch13:eq:08}{{12.8}{282}} +\newlabel{ch13:eq:09}{{12.9}{282}} +\newlabel{ch13:eq:10}{{12.10}{283}} +\newlabel{ch13:eq:11}{{12.11}{283}} +\newlabel{ch13:eq:12}{{12.12}{283}} +\newlabel{ch13:eq:13}{{12.13}{284}} +\newlabel{ch13:eq:14}{{12.14}{284}} +\newlabel{ch13:eq:15}{{12.15}{284}} +\newlabel{ch13:eq:16}{{12.16}{284}} +\@writefile{toc}{\contentsline {section}{\numberline {12.4}Parallel implementation on a GPU cluster}{285}} +\newlabel{ch13:sec:04}{{12.4}{285}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.1}{\ignorespaces Data partitioning of a problem to be solved among $S=3\times 4$ computing nodes.\relax }}{285}} +\newlabel{ch13:fig:01}{{12.1}{285}} +\@writefile{loa}{\contentsline {algocf}{\numberline {11}{\ignorespaces Parallel solving of the obstacle problem on a GPU cluster\relax }}{286}} +\newlabel{ch13:alg:01}{{11}{286}} +\newlabel{ch13:eq:18}{{12.17}{286}} +\@writefile{loa}{\contentsline {algocf}{\numberline {12}{\ignorespaces Parallel iterative solving of the nonlinear systems on a GPU cluster ($Solve()$ function)\relax }}{287}} +\newlabel{ch13:alg:02}{{12}{287}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.2}{\ignorespaces Decomposition of a sub-problem in a GPU into $nz$ slices.\relax }}{288}} +\newlabel{ch13:fig:02}{{12.2}{288}} +\newlabel{ch13:list:01}{{12.1}{288}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.1}Skeleton codes of a GPU kernel and a CPU function}{288}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.3}{\ignorespaces Matrix constant coefficients in a three-dimensional domain.\relax }}{290}} +\newlabel{ch13:fig:03}{{12.3}{290}} +\newlabel{ch13:eq:17}{{12.18}{290}} +\newlabel{ch13:list:02}{{12.2}{290}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.2}GPU kernels of the projected Richardson method}{290}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.4}{\ignorespaces Computation of a vector element with the projected Richardson method.\relax }}{292}} +\newlabel{ch13:fig:04}{{12.4}{292}} +\newlabel{ch13:list:03}{{12.3}{292}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.3}Memory access to the cache texture memory}{292}} +\@writefile{toc}{\contentsline {section}{\numberline {12.5}Experimental tests on a GPU cluster}{293}} +\newlabel{ch13:sec:05}{{12.5}{293}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.5}{\ignorespaces GPU cluster of tests composed of 12 computing nodes (six machines, each with two GPUs.\relax }}{295}} +\newlabel{ch13:fig:05}{{12.5}{295}} +\@writefile{lot}{\contentsline {table}{\numberline {12.1}{\ignorespaces Execution times in seconds of the parallel projected Richardson method implemented on a cluster of 24 CPU cores.\relax }}{295}} +\newlabel{ch13:tab:01}{{12.1}{295}} +\@writefile{lot}{\contentsline {table}{\numberline {12.2}{\ignorespaces Execution times in seconds of the parallel projected Richardson method implemented on a cluster of 12 GPUs.\relax }}{296}} +\newlabel{ch13:tab:02}{{12.2}{296}} +\@writefile{toc}{\contentsline {section}{\numberline {12.6}Red-Black ordering technique}{296}} +\newlabel{ch13:sec:06}{{12.6}{296}} +\newlabel{ch13:list:04}{{12.4}{297}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {12.4}GPU kernels of the projected Richardson method using the red-black technique}{297}} +\newlabel{ch13:fig:06.01}{{12.6(a)}{298}} +\newlabel{sub@ch13:fig:06.01}{{(a)}{298}} +\newlabel{ch13:fig:06.02}{{12.6(b)}{298}} +\newlabel{sub@ch13:fig:06.02}{{(b)}{298}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.6}{\ignorespaces Red-Black ordering for computing the iterate vector elements in a three-dimensional space.\relax }}{298}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Red-Black ordering on x, y and z axises}}}{298}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Red-Black ordering on y axis}}}{298}} +\@writefile{lot}{\contentsline {table}{\numberline {12.3}{\ignorespaces Execution times in seconds of the parallel projected Richardson method using read-black ordering technique implemented on a cluster of 12 GPUs.\relax }}{299}} +\newlabel{ch13:tab:03}{{12.3}{299}} +\@writefile{lof}{\contentsline {figure}{\numberline {12.7}{\ignorespaces Weak scaling of both synchronous and asynchronous algorithms of the projected Richardson method using red-black ordering technique.\relax }}{300}} +\newlabel{ch13:fig:07}{{12.7}{300}} +\@writefile{toc}{\contentsline {section}{\numberline {12.7}Conclusion}{300}} +\newlabel{ch13:sec:07}{{12.7}{300}} +\@writefile{toc}{\contentsline {section}{Bibliography}{301}} \@setckpt{Chapters/chapter13/ch13}{ -\setcounter{page}{301} +\setcounter{page}{303} \setcounter{equation}{18} \setcounter{enumi}{4} \setcounter{enumii}{0} diff --git a/BookGPU/Chapters/chapter16/ch16.aux b/BookGPU/Chapters/chapter16/ch16.aux index dfac2f4..1b44748 100644 --- a/BookGPU/Chapters/chapter16/ch16.aux +++ b/BookGPU/Chapters/chapter16/ch16.aux @@ -4,72 +4,72 @@ \@writefile{toc}{\author{H. Wang}{}} \@writefile{toc}{\author{H. Yu}{}} \@writefile{loa}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {chapter}{\numberline {15}GPU-Accelerated Envelope-Following Method}{341}} +\@writefile{toc}{\contentsline {chapter}{\numberline {15}GPU-Accelerated Envelope-Following Method}{343}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {15.1}Introduction}{341}} -\newlabel{fig:ef1}{{15.1(a)}{343}} -\newlabel{sub@fig:ef1}{{(a)}{343}} -\newlabel{fig:ef2}{{15.1(b)}{343}} -\newlabel{sub@fig:ef2}{{(b)}{343}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.1}{\ignorespaces Transient envelope-following analysis. (Both two figures reflect backward-Euler style envelope-following.)\relax }}{343}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Illustration of one envelope skip.}}}{343}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {The envelope changes in a slow time scale.}}}{343}} -\newlabel{fig:ef_intro}{{15.1}{343}} -\@writefile{toc}{\contentsline {section}{\numberline {15.2}The envelope-following method in a nutshell}{344}} -\newlabel{sec:ef}{{15.2}{344}} -\newlabel{eq:dae}{{15.1}{344}} -\newlabel{eq:Newton}{{15.2}{345}} -\newlabel{eq:A}{{15.3}{345}} -\@writefile{toc}{\contentsline {section}{\numberline {15.3}New parallel envelope-following method}{346}} -\newlabel{sec:gmres}{{15.3}{346}} -\@writefile{toc}{\contentsline {subsection}{\numberline {15.3.1}GMRES solver for Newton update equation}{346}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.2}{\ignorespaces The flow of envelope-following method.\relax }}{347}} -\newlabel{fig:ef_flow}{{15.2}{347}} -\@writefile{loa}{\contentsline {algocf}{\numberline {14}{\ignorespaces Standard GMRES algorithm.\relax }}{348}} -\newlabel{alg:GMRES}{{14}{348}} -\newlabel{line:mvp}{{5}{348}} -\newlabel{line:newnorm}{{11}{348}} -\@writefile{toc}{\contentsline {subsection}{\numberline {15.3.2}Parallelization on GPU platforms}{348}} -\newlabel{sec:gpu}{{15.3.2}{348}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.3}{\ignorespaces GPU parallel solver for envelope-following update.\relax }}{349}} -\newlabel{fig:gmres}{{15.3}{349}} -\@writefile{toc}{\contentsline {subsection}{\numberline {15.3.3}Gear-2 based sensitivity calculation}{350}} -\newlabel{sec:gear}{{15.3.3}{350}} -\newlabel{eq:BE}{{15.4}{350}} -\newlabel{eq:sens1}{{15.5}{350}} -\newlabel{eq:Gear_t2}{{15.6}{351}} -\newlabel{eq:sens2}{{15.7}{351}} -\newlabel{eq:Gear_t3}{{15.8}{351}} -\newlabel{eq:sensM}{{15.9}{351}} -\@writefile{loa}{\contentsline {algocf}{\numberline {15}{\ignorespaces The matrix-free method for Krylov subspace construction.\relax }}{352}} -\newlabel{alg:mf_Gear}{{15}{352}} -\newlabel{line:mf_Gear_loop}{{4}{352}} -\newlabel{line:shift}{{8}{352}} -\@writefile{toc}{\contentsline {section}{\numberline {15.4}Numerical examples}{352}} -\newlabel{sec:exp}{{15.4}{352}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.4}{\ignorespaces Diagram of a zero-voltage quasi-resonant flyback converter.\relax }}{353}} -\newlabel{fig:flyback}{{15.4}{353}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.5}{\ignorespaces Illustration of power/ground network model.\relax }}{353}} -\newlabel{fig:pg}{{15.5}{353}} -\newlabel{fig:flybackWhole}{{15.6(a)}{354}} -\newlabel{sub@fig:flybackWhole}{{(a)}{354}} -\newlabel{fig:flybackZoom}{{15.6(b)}{354}} -\newlabel{sub@fig:flybackZoom}{{(b)}{354}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.6}{\ignorespaces Flyback converter solution calculated by envelope-following. The red curve is traditional SPICE simulation result, and the back curve is the envelope-following output with simulation points marked.\relax }}{354}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {The whole plot}}}{354}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Detail of one EF simulation period}}}{354}} -\newlabel{fig:flyback_wave}{{15.6}{354}} -\@writefile{lof}{\contentsline {figure}{\numberline {15.7}{\ignorespaces Buck converter solution calculated by envelope-following.\relax }}{355}} -\newlabel{fig:buck_wave}{{15.7}{355}} -\@writefile{lot}{\contentsline {table}{\numberline {15.1}{\ignorespaces CPU and GPU time comparisons (in seconds) for solving Newton update equation with the proposed Gear-2 sensitivity. \relax }}{355}} -\newlabel{table:circuit}{{15.1}{355}} -\@writefile{toc}{\contentsline {section}{\numberline {15.5}Summary}{356}} -\newlabel{sec:summary}{{15.5}{356}} -\@writefile{toc}{\contentsline {section}{\numberline {15.6}Glossary}{356}} -\@writefile{toc}{\contentsline {section}{Bibliography}{356}} +\@writefile{toc}{\contentsline {section}{\numberline {15.1}Introduction}{343}} +\newlabel{fig:ef1}{{15.1(a)}{345}} +\newlabel{sub@fig:ef1}{{(a)}{345}} +\newlabel{fig:ef2}{{15.1(b)}{345}} +\newlabel{sub@fig:ef2}{{(b)}{345}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.1}{\ignorespaces Transient envelope-following analysis. (Both two figures reflect backward-Euler style envelope-following.)\relax }}{345}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Illustration of one envelope skip.}}}{345}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {The envelope changes in a slow time scale.}}}{345}} +\newlabel{fig:ef_intro}{{15.1}{345}} +\@writefile{toc}{\contentsline {section}{\numberline {15.2}The envelope-following method in a nutshell}{346}} +\newlabel{sec:ef}{{15.2}{346}} +\newlabel{eq:dae}{{15.1}{346}} +\newlabel{eq:Newton}{{15.2}{347}} +\newlabel{eq:A}{{15.3}{347}} +\@writefile{toc}{\contentsline {section}{\numberline {15.3}New parallel envelope-following method}{348}} +\newlabel{sec:gmres}{{15.3}{348}} +\@writefile{toc}{\contentsline {subsection}{\numberline {15.3.1}GMRES solver for Newton update equation}{348}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.2}{\ignorespaces The flow of envelope-following method.\relax }}{349}} +\newlabel{fig:ef_flow}{{15.2}{349}} +\@writefile{loa}{\contentsline {algocf}{\numberline {14}{\ignorespaces Standard GMRES algorithm.\relax }}{350}} +\newlabel{alg:GMRES}{{14}{350}} +\newlabel{line:mvp}{{5}{350}} +\newlabel{line:newnorm}{{11}{350}} +\@writefile{toc}{\contentsline {subsection}{\numberline {15.3.2}Parallelization on GPU platforms}{350}} +\newlabel{sec:gpu}{{15.3.2}{350}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.3}{\ignorespaces GPU parallel solver for envelope-following update.\relax }}{351}} +\newlabel{fig:gmres}{{15.3}{351}} +\@writefile{toc}{\contentsline {subsection}{\numberline {15.3.3}Gear-2 based sensitivity calculation}{352}} +\newlabel{sec:gear}{{15.3.3}{352}} +\newlabel{eq:BE}{{15.4}{352}} +\newlabel{eq:sens1}{{15.5}{352}} +\newlabel{eq:Gear_t2}{{15.6}{353}} +\newlabel{eq:sens2}{{15.7}{353}} +\newlabel{eq:Gear_t3}{{15.8}{353}} +\newlabel{eq:sensM}{{15.9}{353}} +\@writefile{loa}{\contentsline {algocf}{\numberline {15}{\ignorespaces The matrix-free method for Krylov subspace construction.\relax }}{354}} +\newlabel{alg:mf_Gear}{{15}{354}} +\newlabel{line:mf_Gear_loop}{{4}{354}} +\newlabel{line:shift}{{8}{354}} +\@writefile{toc}{\contentsline {section}{\numberline {15.4}Numerical examples}{354}} +\newlabel{sec:exp}{{15.4}{354}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.4}{\ignorespaces Diagram of a zero-voltage quasi-resonant flyback converter.\relax }}{355}} +\newlabel{fig:flyback}{{15.4}{355}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.5}{\ignorespaces Illustration of power/ground network model.\relax }}{355}} +\newlabel{fig:pg}{{15.5}{355}} +\newlabel{fig:flybackWhole}{{15.6(a)}{356}} +\newlabel{sub@fig:flybackWhole}{{(a)}{356}} +\newlabel{fig:flybackZoom}{{15.6(b)}{356}} +\newlabel{sub@fig:flybackZoom}{{(b)}{356}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.6}{\ignorespaces Flyback converter solution calculated by envelope-following. The red curve is traditional SPICE simulation result, and the back curve is the envelope-following output with simulation points marked.\relax }}{356}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {The whole plot}}}{356}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Detail of one EF simulation period}}}{356}} +\newlabel{fig:flyback_wave}{{15.6}{356}} +\@writefile{lof}{\contentsline {figure}{\numberline {15.7}{\ignorespaces Buck converter solution calculated by envelope-following.\relax }}{357}} +\newlabel{fig:buck_wave}{{15.7}{357}} +\@writefile{lot}{\contentsline {table}{\numberline {15.1}{\ignorespaces CPU and GPU time comparisons (in seconds) for solving Newton update equation with the proposed Gear-2 sensitivity. \relax }}{357}} +\newlabel{table:circuit}{{15.1}{357}} +\@writefile{toc}{\contentsline {section}{\numberline {15.5}Summary}{358}} +\newlabel{sec:summary}{{15.5}{358}} +\@writefile{toc}{\contentsline {section}{\numberline {15.6}Glossary}{358}} +\@writefile{toc}{\contentsline {section}{Bibliography}{358}} \@setckpt{Chapters/chapter16/ch16}{ -\setcounter{page}{358} +\setcounter{page}{360} \setcounter{equation}{9} \setcounter{enumi}{2} \setcounter{enumii}{0} diff --git a/BookGPU/Chapters/chapter18/ch18.aux b/BookGPU/Chapters/chapter18/ch18.aux index db2a3d1..e8d12a3 100644 --- a/BookGPU/Chapters/chapter18/ch18.aux +++ b/BookGPU/Chapters/chapter18/ch18.aux @@ -2,44 +2,44 @@ \@writefile{toc}{\author{Rapha\IeC {\"e}l Couturier}{}} \@writefile{toc}{\author{Christophe Guyeux}{}} \@writefile{loa}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {chapter}{\numberline {16}Pseudorandom Number Generator on GPU}{359}} +\@writefile{toc}{\contentsline {chapter}{\numberline {16}Pseudorandom Number Generator on GPU}{361}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\newlabel{chapter18}{{16}{359}} -\@writefile{toc}{\contentsline {section}{\numberline {16.1}Introduction}{359}} -\@writefile{toc}{\contentsline {section}{\numberline {16.2}Basic Recalls}{361}} -\newlabel{section:BASIC RECALLS}{{16.2}{361}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.2.1}A Short Presentation of Chaos}{361}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.2.2}On Devaney's Definition of Chaos}{361}} -\newlabel{sec:dev}{{16.2.2}{361}} -\newlabel{Devaney}{{16.1}{361}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.2.3}Chaotic iterations}{362}} -\newlabel{subsection:Chaotic iterations}{{16.2.3}{362}} -\newlabel{Chaotic iterations}{{2}{362}} -\newlabel{eq:generalIC}{{16.4}{363}} -\newlabel{equation Oplus}{{16.5}{363}} -\@writefile{toc}{\contentsline {section}{\numberline {16.3}Toward Efficiency and Improvement for CI PRNG}{363}} -\newlabel{sec:efficient PRNG}{{16.3}{363}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.1}First Efficient Implementation of a PRNG based on Chaotic Iterations}{363}} -\newlabel{algo:seqCIPRNG}{{16.1}{363}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {16.1}C code of the sequential PRNG based on chaotic iterations}{363}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.2}Efficient PRNGs based on Chaotic Iterations on GPU}{364}} -\newlabel{sec:efficient PRNG gpu}{{16.3.2}{364}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.3}Naive Version for GPU}{364}} -\@writefile{loa}{\contentsline {algocf}{\numberline {16}{\ignorespaces Main kernel of the GPU ``naive'' version of the PRNG based on chaotic iterations\relax }}{365}} -\newlabel{algo:gpu_kernel}{{16}{365}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.4}Improved Version for GPU}{365}} -\newlabel{IR}{{17}{366}} -\@writefile{loa}{\contentsline {algocf}{\numberline {17}{\ignorespaces Main kernel for the chaotic iterations based PRNG GPU efficient version\relax }}{366}} -\newlabel{algo:gpu_kernel2}{{17}{366}} -\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.5}Chaos Evaluation of the Improved Version}{366}} -\@writefile{toc}{\contentsline {section}{\numberline {16.4}Experiments}{367}} -\newlabel{sec:experiments}{{16.4}{367}} -\@writefile{lof}{\contentsline {figure}{\numberline {16.1}{\ignorespaces Quantity of pseudorandom numbers generated per second with the xorlike-based PRNG\relax }}{368}} -\newlabel{fig:time_xorlike_gpu}{{16.1}{368}} -\@writefile{toc}{\contentsline {section}{Bibliography}{369}} +\newlabel{chapter18}{{16}{361}} +\@writefile{toc}{\contentsline {section}{\numberline {16.1}Introduction}{361}} +\@writefile{toc}{\contentsline {section}{\numberline {16.2}Basic Recalls}{363}} +\newlabel{section:BASIC RECALLS}{{16.2}{363}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.2.1}A Short Presentation of Chaos}{363}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.2.2}On Devaney's Definition of Chaos}{363}} +\newlabel{sec:dev}{{16.2.2}{363}} +\newlabel{Devaney}{{16.1}{363}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.2.3}Chaotic iterations}{364}} +\newlabel{subsection:Chaotic iterations}{{16.2.3}{364}} +\newlabel{Chaotic iterations}{{2}{364}} +\newlabel{eq:generalIC}{{16.4}{365}} +\newlabel{equation Oplus}{{16.5}{365}} +\@writefile{toc}{\contentsline {section}{\numberline {16.3}Toward Efficiency and Improvement for CI PRNG}{365}} +\newlabel{sec:efficient PRNG}{{16.3}{365}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.1}First Efficient Implementation of a PRNG based on Chaotic Iterations}{365}} +\newlabel{algo:seqCIPRNG}{{16.1}{365}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {16.1}C code of the sequential PRNG based on chaotic iterations}{365}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.2}Efficient PRNGs based on Chaotic Iterations on GPU}{366}} +\newlabel{sec:efficient PRNG gpu}{{16.3.2}{366}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.3}Naive Version for GPU}{366}} +\@writefile{loa}{\contentsline {algocf}{\numberline {16}{\ignorespaces Main kernel of the GPU ``naive'' version of the PRNG based on chaotic iterations\relax }}{367}} +\newlabel{algo:gpu_kernel}{{16}{367}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.4}Improved Version for GPU}{367}} +\newlabel{IR}{{17}{368}} +\@writefile{loa}{\contentsline {algocf}{\numberline {17}{\ignorespaces Main kernel for the chaotic iterations based PRNG GPU efficient version\relax }}{368}} +\newlabel{algo:gpu_kernel2}{{17}{368}} +\@writefile{toc}{\contentsline {subsection}{\numberline {16.3.5}Chaos Evaluation of the Improved Version}{368}} +\@writefile{toc}{\contentsline {section}{\numberline {16.4}Experiments}{369}} +\newlabel{sec:experiments}{{16.4}{369}} +\@writefile{lof}{\contentsline {figure}{\numberline {16.1}{\ignorespaces Quantity of pseudorandom numbers generated per second with the xorlike-based PRNG\relax }}{370}} +\newlabel{fig:time_xorlike_gpu}{{16.1}{370}} +\@writefile{toc}{\contentsline {section}{Bibliography}{371}} \@setckpt{Chapters/chapter18/ch18}{ -\setcounter{page}{371} +\setcounter{page}{373} \setcounter{equation}{5} \setcounter{enumi}{2} \setcounter{enumii}{0} diff --git a/BookGPU/Chapters/chapter3/biblio3.bib b/BookGPU/Chapters/chapter3/biblio3.bib index 927c841..56fda5a 100755 --- a/BookGPU/Chapters/chapter3/biblio3.bib +++ b/BookGPU/Chapters/chapter3/biblio3.bib @@ -471,4 +471,39 @@ volume={98}, number={3}, pages={493-504}, doi={10.1016/j.bpj.2009.10.037}, +} +@article{Sanchez-2-2012, +year={2012}, +issn={1939-8018}, +journal={Journal of Signal Processing Systems}, +doi={10.1007/s11265-012-0715-1}, +title={Highly Parallelable Bidimensional Median Filter for Modern Parallel Programming Models}, +url={http://dx.doi.org/10.1007/s11265-012-0715-1}, +publisher={Springer US}, +keywords={Nonlinear filters; Parallel algorithms; Image processing}, +author={Sánchez, RicardoM. and Rodríguez, PaulA.}, +pages={1-15}, +language={English} +} + +@inproceedings{Batcher:1968:SNA:1468075.1468121, + author = {Batcher, K. E.}, + title = {Sorting networks and their applications}, + booktitle = {Proceedings of the April 30--May 2, 1968, spring joint computer conference}, + series = {AFIPS '68 (Spring)}, + year = {1968}, + location = {Atlantic City, New Jersey}, + pages = {307--314}, + numpages = {8}, + url = {http://doi.acm.org/10.1145/1468075.1468121}, + doi = {10.1145/1468075.1468121}, + acmid = {1468121}, + publisher = {ACM}, + address = {New York, NY, USA}, +} +@book{cormen2001introduction, + title={Introduction to algorithms}, + author={Cormen, Thomas H and Leiserson, Charles E and Rivest, Ronald L and Stein, Clifford}, + year={2001}, + publisher={MIT press} } \ No newline at end of file diff --git a/BookGPU/Chapters/chapter3/ch3.aux b/BookGPU/Chapters/chapter3/ch3.aux index 0d1505e..cb0cf96 100644 --- a/BookGPU/Chapters/chapter3/ch3.aux +++ b/BookGPU/Chapters/chapter3/ch3.aux @@ -23,6 +23,7 @@ \@writefile{toc}{\contentsline {chapter}{\numberline {4}Implementing a fast median filter}{29}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} +\@writefile{toc}{\author{Gilles Perrot}{}} \@writefile{toc}{\contentsline {section}{\numberline {4.1}Introduction}{29}} \@writefile{toc}{\contentsline {section}{\numberline {4.2}Median filtering}{30}} \@writefile{toc}{\contentsline {subsection}{\numberline {4.2.1}Basic principles}{30}} @@ -61,56 +62,60 @@ \newlabel{lst:kernelMedian3RegTri9}{{4.2}{36}} \@writefile{lol}{\contentsline {lstlisting}{\numberline {4.2}3$\times $3 median filter kernel using one register per neighborhood pixel and bubble sort}{36}} \@writefile{toc}{\contentsline {subsection}{\numberline {4.4.2}Further optimization}{36}} -\@writefile{lof}{\contentsline {figure}{\numberline {4.4}{\ignorespaces Comparison of pixel throughputs on GPU C2070 and CPU for generic median, in 3$\times $3 median register-only and \textit {libJacket}.\relax }}{37}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.4}{\ignorespaces Comparison of pixel throughputs on GPU C2070 and CPU for generic median, 3$\times $3 median register-only and \textit {libJacket}.\relax }}{37}} \newlabel{fig:compMedians1}{{4.4}{37}} -\@writefile{lof}{\contentsline {figure}{\numberline {4.5}{\ignorespaces Forgetful selection with the minimal element register count. Illustration for 3$\times $3 pixel window represented in a row and supposed sorted.\relax }}{37}} -\newlabel{fig:forgetful_selection}{{4.5}{37}} -\@writefile{toc}{\contentsline {subsubsection}{\numberline {4.4.2.1}Reducing register count}{37}} +\@writefile{toc}{\contentsline {subsubsection}{\numberline {4.4.2.1}Reducing register count }{37}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.5}{\ignorespaces Forgetful selection with the minimal element register count. Illustration for 3$\times $3 pixel window represented in a row and supposed sorted.\relax }}{38}} +\newlabel{fig:forgetful_selection}{{4.5}{38}} \newlabel{lst:medianForget1pix3}{{4.3}{38}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {4.3}3$\times $3 median filter kernel using the minimum register count of 6 to find the median value by forgetful selection method}{38}} -\@writefile{toc}{\contentsline {subsubsection}{\numberline {4.4.2.2}More data output per thread}{39}} -\@writefile{lof}{\contentsline {figure}{\numberline {4.6}{\ignorespaces Illustration of how window overlapping is used to combine 2 pixel selections in a 3$\times $3 median kernel.\relax }}{40}} -\newlabel{fig:median3_overlap}{{4.6}{40}} -\newlabel{lst:medianForget2pix3}{{4.4}{40}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {4.4}3$\times $3 median filter kernel processing 2 output pixel values per thread using combined forgetful selection.}{40}} -\@writefile{lof}{\contentsline {figure}{\numberline {4.7}{\ignorespaces Comparison of pixel throughput on GPU C2070 for the different 3$\times $3 median kernels.\relax }}{41}} -\newlabel{fig:compMedians2}{{4.7}{41}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {4.3}3$\times $3 median filter kernel using the minimum register count of 6 to find the median value by forgetful selection method. The optimal thread block size is 128 on GTX280 and 256 on C2070.}{38}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.6}{\ignorespaces Determination of the Median value by the forgetful selection process, applied to a $3\times 3$ neighborhood window.\relax }}{39}} +\newlabel{fig:forgetful3}{{4.6}{39}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.7}{\ignorespaces Illustration of how window overlapping is used to combine 2 pixel selections in a 3$\times $3 median kernel.\relax }}{40}} +\newlabel{fig:median3_overlap}{{4.7}{40}} +\@writefile{toc}{\contentsline {subsubsection}{\numberline {4.4.2.2}More data output per thread}{40}} +\newlabel{lst:medianForget2pix3}{{4.4}{41}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {4.4}3$\times $3 median filter kernel processing 2 output pixel values per thread using combined forgetful selection.}{41}} \@writefile{toc}{\contentsline {section}{\numberline {4.5}A 5$\times $5 and more median filter }{41}} -\@writefile{lof}{\contentsline {figure}{\numberline {4.8}{\ignorespaces Reducing register count in a 5$\times $5 register-only median kernel outputting 2 pixels simultaneously. The first 7 forgetful selection stages are common to both processed center pixels. Only the last 5 selections have to be done separately.\relax }}{42}} -\newlabel{fig:median5overlap}{{4.8}{42}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.8}{\ignorespaces Comparison of pixel throughput on GPU C2070 for the different 3$\times $3 median kernels.\relax }}{42}} +\newlabel{fig:compMedians2}{{4.8}{42}} \newlabel{sec:median5}{{4.5.1}{42}} \@writefile{toc}{\contentsline {subsection}{\numberline {4.5.1}A register-only 5$\times $5 median filter }{42}} \newlabel{lst:medianForget2pix5}{{4.5}{42}} \@writefile{lol}{\contentsline {lstlisting}{\numberline {4.5}kernel 5$\times $5 median filter processing 2 output pixel values per thread by a combined forgetfull selection.}{42}} -\@writefile{lot}{\contentsline {table}{\numberline {4.2}{\ignorespaces Performance of various 5$\times $5 median kernel implementations, applied on 4096$\times $4096 pixel image with C2070 GPU card.\relax }}{44}} -\newlabel{tab:median5comp}{{4.2}{44}} -\@writefile{toc}{\contentsline {subsection}{\numberline {4.5.2}Fast approximated n$\times $n median filter }{44}} -\@writefile{lot}{\contentsline {table}{\numberline {4.3}{\ignorespaces Measured performance of one generic pseudo-separable median kernel applied to 4096$\times $4096 pixel image with various window sizes.\relax }}{45}} -\newlabel{tab:medianSeparable}{{4.3}{45}} -\newlabel{lst:medianSeparable}{{4.6}{45}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {4.6}generic pseudo median kernel.}{45}} -\newlabel{img:sap_example_ref}{{4.9(a)}{46}} -\newlabel{sub@img:sap_example_ref}{{(a)}{46}} -\newlabel{img:sap_example_sep_med3}{{4.9(b)}{46}} -\newlabel{sub@img:sap_example_sep_med3}{{(b)}{46}} -\newlabel{img:sap_example_sep_med5}{{4.9(c)}{46}} -\newlabel{sub@img:sap_example_sep_med5}{{(c)}{46}} -\newlabel{img:sap_example_sep_med3_it2}{{4.9(d)}{46}} -\newlabel{sub@img:sap_example_sep_med3_it2}{{(d)}{46}} -\@writefile{lof}{\contentsline {figure}{\numberline {4.9}{\ignorespaces Exemple of separable median filtering (smoother), applied to salt \& pepper noise reduction.\relax }}{46}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Airplane image, corrupted with by salt and pepper noise of density 0.25}}}{46}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Image denoised by a $3\times 3$ separable smoother}}}{46}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(c)}{\ignorespaces {Image denoised by a $5\times 5$ separable smoother}}}{46}} -\@writefile{lof}{\contentsline {subfigure}{\numberline{(d)}{\ignorespaces {Image background estimation by a $55\times 55$ separable smoother}}}{46}} -\newlabel{fig:sap_examples2}{{4.9}{46}} -\@writefile{toc}{\contentsline {section}{Bibliography}{47}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.9}{\ignorespaces Reducing register count in a 5$\times $5 register-only median kernel outputting 2 pixels simultaneously. The first 7 forgetful selection stages are common to both processed center pixels. Only the last 5 selections have to be done separately.\relax }}{43}} +\newlabel{fig:median5overlap}{{4.9}{43}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.10}{\ignorespaces First iteration of the $5\times 5$ selection process, with $k_{25}=14$, which shows how Instruction Level Parallelism is maximized by the use of an incomplete sorting network. Arrows represent the result of the swapping function, with the lowest value at the starting point and the highest value at the end point.\relax }}{43}} +\newlabel{fig:median5overlap}{{4.10}{43}} +\@writefile{lot}{\contentsline {table}{\numberline {4.2}{\ignorespaces Performance of various 5$\times $5 median kernel implementations, applied on 4096$\times $4096 pixel image with C2070 GPU card.\relax }}{45}} +\newlabel{tab:median5comp}{{4.2}{45}} +\@writefile{toc}{\contentsline {subsection}{\numberline {4.5.2}Fast approximated n$\times $n median filter }{45}} +\@writefile{lot}{\contentsline {table}{\numberline {4.3}{\ignorespaces Measured performance of one generic pseudo-separable median kernel applied to 4096$\times $4096 pixel image with various window sizes.\relax }}{46}} +\newlabel{tab:medianSeparable}{{4.3}{46}} +\newlabel{lst:medianSeparable}{{4.6}{46}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {4.6}generic pseudo median kernel.}{46}} +\newlabel{img:sap_example_ref}{{4.11(a)}{47}} +\newlabel{sub@img:sap_example_ref}{{(a)}{47}} +\newlabel{img:sap_example_sep_med3}{{4.11(b)}{47}} +\newlabel{sub@img:sap_example_sep_med3}{{(b)}{47}} +\newlabel{img:sap_example_sep_med5}{{4.11(c)}{47}} +\newlabel{sub@img:sap_example_sep_med5}{{(c)}{47}} +\newlabel{img:sap_example_sep_med3_it2}{{4.11(d)}{47}} +\newlabel{sub@img:sap_example_sep_med3_it2}{{(d)}{47}} +\@writefile{lof}{\contentsline {figure}{\numberline {4.11}{\ignorespaces Exemple of separable median filtering (smoother), applied to salt \& pepper noise reduction.\relax }}{47}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(a)}{\ignorespaces {Airplane image, corrupted with by salt and pepper noise of density 0.25}}}{47}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(b)}{\ignorespaces {Image denoised by a $3\times 3$ separable smoother}}}{47}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(c)}{\ignorespaces {Image denoised by a $5\times 5$ separable smoother}}}{47}} +\@writefile{lof}{\contentsline {subfigure}{\numberline{(d)}{\ignorespaces {Image background estimation by a $55\times 55$ separable smoother}}}{47}} +\newlabel{fig:sap_examples2}{{4.11}{47}} +\@writefile{toc}{\contentsline {section}{Bibliography}{49}} \@setckpt{Chapters/chapter3/ch3}{ -\setcounter{page}{49} +\setcounter{page}{50} \setcounter{equation}{0} \setcounter{enumi}{3} \setcounter{enumii}{0} \setcounter{enumiii}{0} -\setcounter{enumiv}{9} +\setcounter{enumiv}{10} \setcounter{footnote}{0} \setcounter{mpfootnote}{0} \setcounter{part}{1} @@ -120,9 +125,9 @@ \setcounter{subsubsection}{0} \setcounter{paragraph}{0} \setcounter{subparagraph}{0} -\setcounter{figure}{9} +\setcounter{figure}{11} \setcounter{table}{3} -\setcounter{numauthors}{0} +\setcounter{numauthors}{1} \setcounter{parentequation}{0} \setcounter{subfigure}{0} \setcounter{lofdepth}{1} diff --git a/BookGPU/Chapters/chapter3/ch3.tex b/BookGPU/Chapters/chapter3/ch3.tex index e04a36c..94c3e97 100755 --- a/BookGPU/Chapters/chapter3/ch3.tex +++ b/BookGPU/Chapters/chapter3/ch3.tex @@ -109,9 +109,9 @@ It remains that global runtime can still be faster than similar processes run on Therefore, to fully optimize global runtimes, it is important to pay attention to how memory transfers are done. This leads us to propose, in the following section, an overall code structure to be used with all our kernel examples. -Obviously, our code originally accepts various image dimensions and can process color images. +Obviously, our code originally accepts various image dimensions and can process color images when an extrapolated definition of the median filter is choosen. However, so as to propose concise and more readable code, we will assume the following limitations: -8 or 16~bit-coded gray-level input images whose dimensions $H\times W$ are multiples of 512 pixels. +16~bit-coded gray-level input images whose dimensions $H\times W$ are multiples of 512 pixels. \begin{algorithm} \SetNlSty{}{}{:} @@ -131,18 +131,18 @@ However, so as to propose concise and more readable code, we will assume the fol \section{Data transfers, memory management.} This section deals with the following issues: \begin{enumerate} -\item Data transfer from CPU memory to GPU global memory: several GPU memory areas are available as destination memory but the 2-D caching mechanism of texture memory, specifically designed for fetching neighboring pixels, is currently the fastest way to fetch gray-level pixel values inside a kernel computation. This has lead us to choose \textbf{texture memory} as primary GPU memory area for images. -\item Data fetching from GPU global memory to kernel local memory: as said above, we use texture memory. Depending on which process is run, texture data is used either by direct fetching in kernel local memory or through a prefetching in shared memory. +\item Data transfer from CPU memory to GPU global memory: several GPU memory areas are available as destination memory but the 2-D caching mechanism of texture memory \index{memory~hierarchy!texture~memory}, specifically designed for fetching neighboring pixels, is currently the fastest way to fetch gray-level pixel values inside a kernel computation. This has led us to choose \textbf{texture memory} as primary GPU memory area for images. +\item Data fetching from GPU global memory to kernel local memory: as said above, we use texture memory \index{memory~hierarchy!texture~memory}. Depending on which process is run, texture data is used either by direct fetching in kernel local memory or through a prefetching \index{prefetching} in shared memory \index{memory~hierarchy!shared~memory}. \item Data outputting from kernels to GPU memory: there is actually no alternative to global memory, as kernels can not directly write into texture memory and as copying from texture to CPU memory would not be faster than from simple global memory. -\item Data transfer from GPU global memory to CPU memory: it can be drastically accelerated by use of \textbf{pinned memory}, keeping in mind it has to be used sparingly. +\item Data transfer from GPU global memory to CPU memory: it can be drastically accelerated by use of \textbf{pinned memory} \index{memory~hierarchy!pinned~memory}, keeping in mind it has to be used sparingly. \end{enumerate} -Algorithm \ref{algo:memcopy} summarizes all the above considerations and describe how data are handled in our examples. For more information on how to handle the different types of GPU memory, we suggest to refer to CUDA programmer's guide. +Algorithm \ref{algo:memcopy} summarizes all the above considerations and describes how data are handled in our examples. For more information on how to handle the different types of GPU memory, we suggest to refer to the CUDA programmer's guide. -At debug stage, for simplicity's sake, we use the \textbf{cutil} library supplied by the NVidia developpement kit (SDK). Thus, in order to easily implement our examples, we suggest readers download and install the latest NVidia-SDK (ours is SDK4.0), create a new directory \textit{SDK-root-dir/C/src/fast\_kernels} and adapt the generic \textit{Makefile} that can be found in each sub-directory of \textit{SDK-root-dir/C/src/}. Then, only two more files will be enough to have a fully operational environnement: \textit{main.cu} and \textit{fast\_kernels.cu}. +At debug stage, for simplicity's sake, we use the \textbf{cutil} \index{Cutil library} library supplied by the NVidia development kit (SDK). Thus, in order to easily implement our examples, we suggest readers to download and to install the latest NVidia-SDK (ours is SDK4.0), create a new directory \textit{SDK-root-dir/C/src/fast\_kernels} and adapt the generic \textit{Makefile} that can be found in each sub-directory of \textit{SDK-root-dir/C/src/}. Then, only two more files will be enough to have a fully operational environnement: \textit{main.cu} and \textit{fast\_kernels.cu}. Listings \ref{lst:main1}, \ref{lst:fkern1} and \ref{lst:mkfile} implement all the above considerations minimally, while remaining functional. The main file of Listing \ref{lst:main1} is a simplified version of our actual main file. -It has to be noticed that cutil functions \texttt{cutLoadPGMi} and \texttt{cutSavePGMi} only operate on unsigned integer data. As data is coded in short integer format for performance reasons, the use of these functions involves casting data after loading and before saving. This may be overcome by use of a different library. Actually, our choice was to modify the above mentioned cutil functions. +It has to be noticed that cutil functions \texttt{cutLoadPGMi} \index{Cutil library!cutLoadPGMi} and \texttt{cutSavePGMi} \index{Cutil library!cutSavePGMi} only operate on unsigned integer data. As data is coded in short integer format for performance reasons, the use of these functions involves casting data after loading and before saving. This may be overcome by use of a different library. Actually, our choice was to modify the above mentioned cutil functions. Listing \ref{lst:fkern1} gives a minimal kernel skeleton that will serve as the basis for all other kernels. Lines 5 and 6 determine the coordinates $(i, j)$ of the pixel to be processed, each pixel being associated to one thread. The instruction in line 8 combines writing the output gray-level value into global memory and fetching the input gray-level value from 2-D texture memory. @@ -156,9 +156,9 @@ The Makefile given in Listing \ref{lst:mkfile} shows how to adapt examples given \section{Performance measurements} -As our goal is to design very fast implementations of basic image processing algorithms, we need to make quite accurate time-measurements, within the order of magnitude of $0.01~ms$. Again, the easiest way of doing so is to use the helper functions of the cutil library. As usual, as the durations we are measuring are short and possibly suject to non neglectable variations, a good practice is to measure multiple executions and issue the mean runtime. All time results given in this chapter have been obtained through 1000 calls to each kernel. +As our goal is to design very fast implementations of basic image processing algorithms, we need to make quite accurate time-measurements, within the order of magnitude of $0.01~ms$. Again, the easiest way of doing so is to use the helper functions of the cutil library. As usual, as the durations we are measuring are short and possibly subject to non neglectable variations, a good practice is to measure multiple executions and issue the mean runtime. All time results given in this chapter have been obtained through 1000 calls to each kernel. -Listing \ref{lst:chronos} shows how to use the dedicated cutil functions. Timer declaration and creation only need to be performed once while reset, start and stop can be used as often as necessary. Synchronization is mandatory before stopping the timer (Line 7), to avoid runtime measurement being biased. +Listing \ref{lst:chronos} shows how to use the dedicated cutil functions \index{Cutil library!Timer usage}. Timer declaration and creation only need to be performed once while reset, start and stop can be used as often as necessary. Synchronization is mandatory before stopping the timer (Line 7), to avoid runtime measurement being biased. \lstinputlisting[label={lst:chronos},caption=Time measurement technique using cutil functions]{Chapters/chapter3/code/exChronos.cu} In an attempt to provide relevant speedup values, we either implemented CPU versions of the algorithms studied, or used the values found in existing literature. Still, the large number and diversity of hardware platforms and GPU cards makes it impossible to benchmark every possible combination and significant differences may occur between the speedups we announce and those obtained with different devices. As a reference, our developing platform details as follows: @@ -181,15 +181,15 @@ Last, like many authors, we chose to use the pixel throughput value of each proc In order to estimate the potential for improvement of each kernel, a reference throughput measurement, involving identity kernel of Listing \ref{lst:fkern1}, was performed. As this kernel only fetches input values from texture memory and outputs them to global memory without doing any computation, it represents the smallest, thus fastest, possible process and is taken as the reference throughput value (100\%). The same measurement was performed on CPU, with a maximum effective pixel throughput of 130~Mpixel per second. On GPU, depending on grid parameters it amounts to 800~MPixels/s on GTX280 and 1300~Mpixels/s on C2070. - \chapter{Implementing a fast median filter} +\chapterauthor{Gilles Perrot}{FEMTO-ST Institute} \section{Introduction} Median filtering is a well-known method used in a wide range of application frameworks as well as a standalone filter especially for \textit{salt and pepper} denoising. It is able to highly reduce power of noise without blurring edges too much. First introduced by Tukey in \cite{tukey77}, it has been widely studied since then, and many researchers have proposed efficient implementations of it, adapted to various hypothesis, architectures and processors. -Originally, its main drawbacks were its compute complexity, its non linearity and its data-dependent runtime. Several researchers have adress these issues and designed, for example, efficient histogram-based median filter with predictible runtime \cite{Huang:1981:TDS:539567, Weiss:2006:FMB:1179352.1141918}. +Originally, its main drawbacks were its compute complexity, its non linearity and its data-dependent runtime. Several researchers have addressed these issues and designed, for example, efficient histogram-based median filter with predictible runtimes \cite{Huang:1981:TDS:539567, Weiss:2006:FMB:1179352.1141918}. -More recently, the advent of GPUs opened new perspectives in terms of image processing performance, and some researchers managed to take advantage of the new graphic capabilities: in that respect, we can cite the Branchless Vectorized Median filter (BVM) \cite{5402362, chen09} which allows very interesting runtimes on CUDA-enabled devices but, as far as we know, the fastest implementation to date is the histogram-based CCDS median filter \cite{6288187}. +More recently, the advent of GPUs opened new perspectives in terms of image processing performance, and some researchers managed to take advantage of the new graphic capabilities: in that respect, we can cite the Branchless Vectorized Median filter (BVM) \cite{5402362, chen09} which allows very interesting runtimes on CUDA-enabled devices but, as far as we know, the fastest implementation to date is the histogram-based PCMF median filter \cite{Sanchez-2-2012}. Some of the following implementations feature very fast runtimes; They are targeted on Nvidia Tesla GPU (Fermi architecture, compute capability 2.x) but may easily be adapted to other models e.g. those of compute capability 1.3. @@ -213,9 +213,9 @@ Figure \ref{fig:sap_examples} shows an example of a $512\times 512$ pixel image, As mentioned earlier, the selection of the median value can be performed by more than one technique, using either histogram-based or sorting methods, each of them having its own benefits and drawbacks as will be discussed further down. \subsection{A naive implementation} -As a reference, Listing \ref{lst:medianGeneric} gives a simple, not to say simplistic implementation of a CUDA kernel (\texttt{kernel\_medianR}) achieving generic $n\times n$ histogram-based median filtering. Its runtime has a very low data dependency, but this implementation does not suit very well GPU architecture. Each pixel loads the whole of its $n\times n$ neighborhood meaning that one pixel is loaded multiple times inside one single thread block, and above all, the use of a local vector (histogram[]) considerably downgrades performance, as the compiler automatically stores such vectors in local memory (slow). +As a reference, Listing \ref{lst:medianGeneric} gives a simple, not to say simplistic implementation of a CUDA kernel (\texttt{kernel\_medianR}) achieving generic $n\times n$ histogram-based median filtering. Its runtime has a very low data dependency, but this implementation does not suit very well GPU architecture. Each pixel loads the whole of its $n\times n$ neighborhood meaning that one pixel is loaded multiple times inside one single thread block, and above all, the use of a local vector (histogram[]) considerably downgrades performance, as the compiler automatically stores such vectors in local memory (slow) \index{memory~hierarchy!local~memory}. -Table \ref{tab:medianHisto1} displays measured runtimes of \texttt{kernel\_medianR} and pixel throughputs for each GPU version and for both CPU and GPU implementations. Usual window sizes of $3\times 3$, $5\times 5$ and $7\times 7$ are shown. Though some specific applications require larger window sizes and dedicated algorithms , such small square window sizes are most widely used in general purpose image processing. GPU runtimes have been obtained with a grid of 64-thread blocks. This block size, is a good compromise in this case. +Table \ref{tab:medianHisto1} displays measured runtimes of \texttt{kernel\_medianR} and pixel throughputs for each GPU version and for both CPU and GPU implementations. Usual window sizes of $3\times 3$, $5\times 5$ and $7\times 7$ are shown. Though some specific applications require larger window sizes and dedicated algorithms, such small square window sizes are most widely used in general purpose image processing. GPU runtimes have been obtained with a grid of 64-thread blocks. This block size, is a good compromise in this case. The first observation to make when analysing results of Table \ref{tab:medianHisto1} is that, on CPU, window size has almost no influence on the effective pixel throughput. Since inner loops that fill the histogram vector contain very few fetching instructions (from 9 to 49, depending on the window size), it is not surprising to note their neglectable impact compared to outer loops that fetch image pixels (from 256k to 16M instructions). @@ -292,35 +292,35 @@ copy data from GPU global memory to CPU memory\label{algoMedianGeneric:memcpyD2H \section{NVidia GPU tuning recipes} When designing GPU code, besides thinking of the actual data computing process, one must choose the memory type into which to store temporary data. Three types of GPU memory are available: \begin{enumerate} -\item \textbf{Global memory, the most versatile:}\\Offers the largest storing space and global scope but is slowest (400 cycles latency). \textbf{Texture memory} is physically included into it, but allows access through an efficient 2-D caching mechanism. -\item \textbf{Registers, the fastest:}\\Allow access wihtout latency, but only 63 registers are available per thread (thread scope), with a maximum of 32K per Symetric Multiprocessor (SM). -\item \textbf{Shared memory, a complex compromise:}\\All threads in one block can access 48~KBytes of shared memory, which is faster than global memory (20 cycles latency) but slower than registers. +\item \textbf{Global memory, the most versatile:} \index{memory~hierarchy!global~memory}\\Offers the largest storing space and global scope but is slowest (400 cycles latency). \textbf{Texture memory} is physically included into it, but allows access through an efficient 2-D caching mechanism. +\item \textbf{Registers, the fastest:} \index{memory~hierarchy!registers}\\Allow access wihtout latency, but only 63 registers are available per thread (thread scope), with a maximum of 32K per Symetric Multiprocessor (SM) \index{register count}w. +\item \textbf{Shared memory, a complex compromise:} \index{memory~hierarchy!shared~memory}\\All threads in one block can access 48~KBytes of shared memory, which is faster than global memory (20 cycles latency) but slower than registers. However, bank conflicts can occur if two threads of a warp try to access data stored in one single memory bank. In such cases, the parallel process is re-serialized which may cause significant performance decrease. One easy way to avoid it is to ensure that two consecutive threads in one block always access 32-bit data at two consecutive adresses. \end{enumerate} \noindent As observed earlier, designing a median filter GPU implementation using only global memory is fairly straightforward, but its performance remains quite low even if it is faster than CPU. -To overcome this, the most frequent choice made in efficient implementations found in literature is to use shared memory. Such option implies prefetching data prior to doing the actual computations, a relevant choice, as each pixel of an image belongs to n$\times$n different neighborhoods. Thus, it can be expected that fetching each gray-level value from global memory only once should be more efficient than doing it each time it is required. One of the most efficient implementations using shared memory is presented in \cite{5402362}. In the case of the generic kernel of Listing \ref{lst:medianGeneric}, using shared memory without further optimization would not bring valuable speedup because that would just move redundancy from texture to shared memory fetching and would generate bank conflicts. For information, we wrote such a version of the generic median kernel and our measurements showed a speedup of around 3\% (as an example: 32ms for 5$\times$5 median on a 1024$^2$ pixel image, i.e. 33~Mpixel/s ). +To overcome this, the most frequent choice made in efficient implementations found in literature is to use shared memory. Such option implies prefetching \index{prefetching}data prior to doing the actual computations, a relevant choice, as each pixel of an image belongs to n$\times$n different neighborhoods. Thus, it can be expected that fetching each gray-level value from global memory only once should be more efficient than doing it each time it is required. One of the most efficient implementations using shared memory is presented in \cite{5402362}. In the case of the generic kernel of Listing \ref{lst:medianGeneric}, using shared memory without further optimization would not bring valuable speedup because that would just move redundancy from texture to shared memory fetching and would generate bank conflicts. For information, we wrote such a version of the generic median kernel and our measurements showed a speedup of around 3\% (as an example: 32ms for 5$\times$5 median on a 1024$^2$ pixel image, i.e. 33~Mpixel/s ). -As for registers, designing a generic median filter that would only use that type of memory seems difficult, due to the above mentioned 63 register-per-thread limitation. +As for registers, designing a generic median filter that would only use that type of memory seems difficult, due to the above mentioned 63 register-per-thread limitation \index{register count}. Yet, nothing forbids us to design fixed-size filters, each of them specific to one of the most popular window sizes. It might be worth the effort as dramatic increase in performance could be expected. -Another track to follow in order to improve performance of GPU implementations consists in hiding latencies generated by arithmetic instruction calls and memory accesses. Both can be partially hidden by introducing Instruction-Level Parallelism (ILP) and by increasing the data count output by each thread. Though such techniques may seem to break the NVidia occupancy paradigm, they can lead to dramatically higher data throughput values. +Another track to follow in order to improve performance of GPU implementations consists in hiding latencies generated by arithmetic instruction calls and memory accesses. Both can be partially hidden by introducing Instruction-Level Parallelism \index{Instruction-Level Parallelism}(ILP) and by increasing the data count output by each thread. Though such techniques may seem to break the NVidia occupancy paradigm, they can lead to dramatically higher data throughput values. The following sections illustrate these ideas and detail the design of the fastest CUDA median filter known to date. \section{A 3$\times$3 median filter: using registers } Designing a median filter dedicated to the smallest possible square window size is a good challenge to start using registers. -One first issue is that the exclusive use of registers forbids us to implement a naive histogram-based method. In a \textit{8-bit gray level pixel per thread} rule, each histogram requires one 256-element vector to store its values, i.e. four times the maximum register count allowed per thread (63). Considering that a 3$\times$3 median filter involves only 9 pixel values per thread, it seem obvious they can be sorted within the 63-register limit. +One first issue is that the exclusive use of registers forbids us to implement a naive histogram-based method. In a \textit{8-bit gray level pixel per thread} rule, each histogram requires one 256-element vector to store its values, i.e. four times the maximum register count allowed per thread (63)\index{register count}. Considering that a 3$\times$3 median filter involves only 9 pixel values per thread, it seem obvious they can be sorted within the 63-register limit. \subsection{The simplest way} -In the case of a 3$\times$3 median filter, the simplest solution consists in associating one register to each gray-level value, then sorting those 9 values and selecting the fifth one, i.e. the median value. For such a small amount of data to sort, a simple selection method is well indicated. As shown in Listing \ref{lst:kernelMedian3RegTri9} (\texttt{kernelMedian3RegTri9()}), the constraint of only using registers leads to adopt an unusual manner of coding. However, results are persuasive: runtimes are divided by around 120 on GTX280 and 80 on C2070, while only reduced by a 3.5 factor on CPU. -The diagram of Figure \ref{fig:compMedians1} summarizes these first results. Only C2070 throughputs are shown and compared to CPU results. We included the maximum effective pixel throughput in order to see the improvement potential of the different implementations. We also introduced throughputd achieved by \textit{libJacket}, a commercial implementation, as it was the fastest known implementation of a 3$\times$3 median filter to date, as illustrated in \cite{chen09}. One of the authors of libJacket kindly posted the CUDA code of its 3$\times$3 median filter, that we inserted into our own coding structure. The algorithm itself is quite similar to ours, but running it in our own environement produced higher throughput values than those published in \cite{chen09}, not due to different hardware capabilities between our GTX280 and the GTX260 used in the paper, but to the way we perform memory transfers and to our register-only method of storing temporary data. +In the case of a 3$\times$3 median filter, the simplest solution consists in associating one register to each gray-level value, then sorting those 9 values and selecting the fifth one, i.e. the median value. For such a small amount of data to sort, a simple selection method is well indicated. As shown in Listing \ref{lst:kernelMedian3RegTri9} (\texttt{kernelMedian3RegSort9()}), the constraint of only using registers leads to adopt an unusual manner of coding. However, results are persuasive: runtimes are divided by around 120 on GTX280 and 80 on C2070, while only reduced by a 3.5 factor on CPU. +The diagram of Figure \ref{fig:compMedians1} summarizes these first results, obtained with a block size of 128 threads on GTX280 and 256 on C2070. Only C2070 throughputs are shown and compared to CPU results. We included the maximum effective pixel throughput in order to see the improvement potential of the different implementations. We also introduced throughputd achieved by \textit{libJacket}, a commercial implementation, as it was the fastest known implementation of a 3$\times$3 median filter to date, as illustrated in \cite{chen09}. One of the authors of libJacket kindly posted the CUDA code of its 3$\times$3 median filter, that we inserted into our own coding structure. The algorithm itself is quite similar to ours, but running it in our own environement produced higher throughput values than those published in \cite{chen09}, not due to different hardware capabilities between our GTX280 and the GTX260 used in the paper, but to the way we perform memory transfers and to our register-only method of storing temporary data. \lstinputlisting[label={lst:kernelMedian3RegTri9},caption= 3$\times$3 median filter kernel using one register per neighborhood pixel and bubble sort]{Chapters/chapter3/code/kernMedianRegTri9.cu} \begin{figure} \centering - \includegraphics[width=11cm]{Chapters/chapter3/img/debitPlot1.pdf} - \caption{Comparison of pixel throughputs on GPU C2070 and CPU for generic median, in 3$\times$3 median register-only and \textit{libJacket}.} + \includegraphics[width=15cm]{Chapters/chapter3/img/debitPlot1.pdf} + \caption{Comparison of pixel throughputs on GPU C2070 and CPU for generic median, 3$\times$3 median register-only and \textit{libJacket}.} \label{fig:compMedians1} \end{figure} @@ -332,8 +332,8 @@ Running the above register-only 3$\times$3 median filter through the NVidia CUDA \end{itemize} -\subsubsection{Reducing register count} -Our current kernel (\texttt{kernelMedian3RegTri9}) uses one register per gray-level value, which amounts to 9 registers for the entire 3$\times$3 window. +\subsubsection{Reducing register count \index{register count}} +Our current kernel (\texttt{kernelMedian3RegSort9}) uses one register per gray-level value, which amounts to 9 registers for the entire 3$\times$3 window. This count can be reduced by use of an iterative sorting process called \textit{forgetful selection}, where both \textit{extrema} are eliminated at each sorting stage, until only 3 elements remain. The question is to find out the minimal register count $k_{n^2}$ that allows the selection of the median amoung $n^2$ values. The answer can be evaluated considering that, when eliminating the maximum and the minimum values, one has to make sure not to eliminate the global median value. Such a situation is illustrated in Figure \ref{fig:forgetful_selection} for a 3$\times$3 median filter. For better comprehension, the 9 elements of the 3$\times$3 pixel window have been represented in a row. \begin{figure} \centering @@ -353,11 +353,19 @@ which leads to: $$k_{n^2}=\lceil \frac{n^2}{2}\rceil+1 $$ This rule can be applied to the first eliminating stage and remains true with the next ones as each stage suppresses exactly two values, one above and one below the median value. In our 3$\times$3 pixel window example, the minimum register count becomes $k_9=\lceil 9/2\rceil+1 = 6$. +This iterative process is illustrated in Figure \ref{fig:forgetful3}, where it achieves one entire $3\times 3$ median selection, beginning with $k_9=6$ elements. +\begin{figure} + \centering + \includegraphics[width=5cm]{Chapters/chapter3/img/forgetful_selectionb.png} + \caption{Determination of the Median value by the forgetful selection process, applied to a $3\times 3$ neighborhood window.} + \label{fig:forgetful3} +\end{figure} + The \textit{forgetful selection} method, used in \cite{mcguire2008median} does not imply full sorting of values, but only selecting minimum and maximum values, which, at the price of a few iteration steps ($n^2-k$), reduces arithmetic complexity. -Listing \ref{lst:medianForget1pix3} details this process where forgetful selection is achieved by use of simple 2-value sorting function ($s()$, lines 1 to 5) that swaps input values if necessary. Moreover, whenever possible, in order to increase the Instruction-Level Parallelism, successive calls to $s()$ are done with independant elements as arguments. This is illustrated by the macro definitions of lines 7 to 14. +Listing \ref{lst:medianForget1pix3} details this process where forgetful selection is achieved by use of simple 2-value swapping function ($s()$, lines 1 to 5) that swaps input values if necessary, so as to achieve the first steps of a \textit{bitonic}-like sorting network (\cite{Batcher:1968:SNA:1468075.1468121}). Moreover, whenever possible, in order to increase the Instruction-Level Parallelism \index{Instruction-Level Parallelism}, successive calls to $s()$ are done with independant elements as arguments. This is illustrated by the macro definitions of lines 7 to 14 and by Figure \ref{fig:bitonic} which details the first iteration of the $5\times 5$ selection, starting with $k_{25}=14$ elements. -\lstinputlisting[label={lst:medianForget1pix3},caption= 3$\times$3 median filter kernel using the minimum register count of 6 to find the median value by forgetful selection method]{Chapters/chapter3/code/kernMedianForget1pix3.cu} +\lstinputlisting[label={lst:medianForget1pix3},caption= 3$\times$3 median filter kernel using the minimum register count of 6 to find the median value by forgetful selection method. The optimal thread block size is 128 on GTX280 and 256 on C2070. ]{Chapters/chapter3/code/kernMedianForget1pix3.cu} Our such modified kernel provides significantly improved runtimes: a speedup of around 16\% is obtained, and pixel throughput reaches around 1000~MPixel/s on C2070. @@ -379,7 +387,7 @@ Running this $3\times 3$ kernel saves another 10\% runtime, as shown in Figure \ \begin{figure} \centering - \includegraphics[width=11cm]{Chapters/chapter3/img/debitPlot2.pdf} + \includegraphics[width=15cm]{Chapters/chapter3/img/debitPlot2.pdf} \caption{Comparison of pixel throughput on GPU C2070 for the different 3$\times$3 median kernels.} \label{fig:compMedians2} \end{figure} @@ -392,11 +400,18 @@ The next two sections will first detail the particular case of the 5$\times$5 me The minimum register count required to apply the forgetful selection method to a 5$\times$5 median filter is $k_{25}=\lceil 25/2\rceil+1 = 14$. Moreover, two adjacent overlapping windows share 20 pixels ($n^2-one\_column$) so that, when processing 2 pixels simultaneously, a count of 7 common selection stages can be carried out from the first selection stage with 14 common values to the processing of the last common value. That allows to limit register count to 22 per thread. Figure \ref{fig:median5overlap} describes the distribution of overlapping pixels, implemented in Listing \ref{lst:medianForget2pix5}: common selection stages take place from line 25 to line 37, while the remaining separate selection stages occur between lines 45 and 62 after the separation of line 40. \begin{figure} \centering - \includegraphics[width=6cm]{Chapters/chapter3/img/median5_overlap.png} + \includegraphics[width=6cm]{Chapters/chapter3/img/median5_overlap4.png} \caption{Reducing register count in a 5$\times$5 register-only median kernel outputting 2 pixels simultaneously. The first 7 forgetful selection stages are common to both processed center pixels. Only the last 5 selections have to be done separately.} \label{fig:median5overlap} \end{figure} +\begin{figure} + \centering + \includegraphics[width=6cm]{Chapters/chapter3/img/forgetful_selection4.png} + \caption{First iteration of the $5\times 5$ selection process, with $k_{25}=14$, which shows how Instruction Level Parallelism is maximized by the use of an incomplete sorting network. Arrows represent the result of the swapping function, with the lowest value at the starting point and the highest value at the end point.} + \label{fig:median5overlap} +\end{figure} + \lstinputlisting[label={lst:medianForget2pix5},caption=kernel 5$\times$5 median filter processing 2 output pixel values per thread by a combined forgetfull selection.]{Chapters/chapter3/code/kernMedian2pix5.cu} Timing results follow the same variations with image size as in previously presented kernels. That is why Table \ref{tab:median5comp} shows only throughput values obtained for C2070 card and 4096$\times$4096 pixel image. diff --git a/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu b/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu index 363b181..29922a3 100755 --- a/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu +++ b/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu @@ -1,4 +1,4 @@ -__global__ void kernel_Median3RegTri9( short *output, +__global__ void kernel_Median3RegSort9( short *output, int i_dim, int j_dim) { int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ; diff --git a/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu~ b/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu~ index cd173eb..363b181 100755 --- a/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu~ +++ b/BookGPU/Chapters/chapter3/code/kernMedianRegTri9.cu~ @@ -1,19 +1,11 @@ -__global__ void kernel_Median3RegTri9( short *output, int i_dim, int j_dim) +__global__ void kernel_Median3RegTri9( short *output, + int i_dim, int j_dim) { - - // coordonnees absolues du point int j = __mul24(blockIdx.x,blockDim.x) + threadIdx.x ; int i = __mul24(blockIdx.y,blockDim.y) + threadIdx.y ; - - /************************************************************************** - * tri(s) - **************************************************************************/ - int a0, a1, a2, a3, a4, a5, a6, a7, a8 ; + int a0, a1, a2, a3, a4, a5, a6, a7, a8 ; // 1 register per pixel - /******************************************************************************** - * affectation des valeurs - ********************************************************************************/ - a0 = tex2D(tex_img_ins, j-1, i-1) ; + a0 = tex2D(tex_img_ins, j-1, i-1) ; // fetching values a1 = tex2D(tex_img_ins, j , i-1) ; a2 = tex2D(tex_img_ins, j+1, i-1) ; a3 = tex2D(tex_img_ins, j-1, i) ; @@ -23,10 +15,7 @@ __global__ void kernel_Median3RegTri9( short *output, int i_dim, int j_dim) a7 = tex2D(tex_img_ins, j , i+1) ; a8 = tex2D(tex_img_ins, j+1, i+1) ; - //tri selection - bubTriReg9(&a0, &a1, &a2, &a3, &a4, &a5, &a6, &a7, &a8); + bubReg9(&a0,&a1,&a2,&a3,&a4,&a5,&a6,&a7,&a8); // bubble sort - //median au milieu ! - output[ __mul24(i, j_dim) +j ] = a4 ; - -} + output[ __mul24(i, j_dim) +j ] = a4 ; // median at the middle + } diff --git a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu index f4ad2c6..bdf7023 100755 --- a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu +++ b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu @@ -1,7 +1,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) { - int idc, val, min, max, inf, egal, sup, mxinf, minsup, estim ; + int idc, val, min, max, inf, equal, sup, mxinf, minsup, estim ; //coordinates in the block int ib = threadIdx.y ; @@ -40,7 +40,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) while (1) { estim = (min+max)/2 ; - inf = sup = egal = 0 ; + inf = sup = equal = 0 ; mxinf = min ; minsup= max ; for (idc =0; idc< 2*r+1 ; idc++) @@ -54,7 +54,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) { sup++; if( val < minsup) minsup = val ; - } else egal++ ; + } else equal++ ; } if ( (inf <= (r+1))&&(sup <=(r+1)) ) break ; else if (inf>sup) max = mxinf ; @@ -62,7 +62,7 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) } if ( inf >= r+1 ) val = mxinf ; - else if (inf+egal >= r+1) val = estim ; + else if (inf+equal >= r+1) val = estim ; else val = minsup ; output[ __mul24(j, i_dim) +i ] = val ; diff --git a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu~ b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu~ index 5c79c82..f4ad2c6 100755 --- a/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu~ +++ b/BookGPU/Chapters/chapter3/code/kernMedianSeparable.cu~ @@ -6,17 +6,14 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) //coordinates in the block int ib = threadIdx.y ; int jb = threadIdx.x ; - int idx_h = __mul24(ib+r,blockDim.x) + jb ; // index pixel deans shmem (bloc+halo) + int idx_h = __mul24(ib+r,blockDim.x) +jb; // base pixel index int offset = __mul24(blockDim.x,r) ; - // coordonnees absolues du point int j = __mul24(blockIdx.x,blockDim.x) + jb ; int i = __mul24(blockIdx.y,blockDim.y) + ib ; - extern __shared__ int buff[] ; - /*********************************************************************************** - * CHARGEMENT DATA EN SHARED MEM - ***********************************************************************************/ + // DATA PREFETCHING INTO SHARED MEM + extern __shared__ int buff[] ; buff[ idx_h ] = tex2D(tex_img_ins, j, i) ; if (ib < r) @@ -29,10 +26,8 @@ __global__ void kernel_medianV_sh( short *output, int i_dim, int j_dim, int r) } __syncthreads() ; - /********************************************************************************************** - * TRI VERTICAL par algo TORBEN MOGENSEN - * (a little bit slow but saves memory => faster !) - **********************************************************************************************/ + + // TORBEN MOGENSEN SORTING min = max = buff[ ib*blockDim.x +jb] ; for (idc= 0 ; idc< 2*r+1 ; idc++ ) diff --git a/BookGPU/Chapters/chapter3/img/debitPlot1.pdf b/BookGPU/Chapters/chapter3/img/debitPlot1.pdf index 204c5620c1f08c475b009c5804d456281cecc197..9cd73d06e4b68ca59437c127ed0598312b22e1ee 100644 GIT binary patch delta 12892 zcmb8VWl&w+(ykrcH9&B8hsDArxVu9L8r*{W#GN1ucL~9QI|O%!5G=U6yMH`;zq`J3 z_WSG1s!`qJ9zD9JRE;0k>lc$y* z8Wjp4TCgW>KKE_U&-uVUwxLLlXyn-jb%EI#=PwBBiTI1JDh}`K4Mj&z%ly$kmIW-8< zcNZXouHNnu=Krps*0*_@X9|~SR?Kw8-5rcsdXD6|^nunIoS|Yc_Sg+e79`8iSHq^z zD-CfTb`&?J-HGCLU1>$4gk6dW5$v%;O*VFM#B@r_&G_Yofe`KA!pXN|6SNzpN0}t{4-rUdr$1@hsJ=ME+a>}jAtdnNhS^tJwk{?GWbc14nS}P zp$3H_b!wY7a)(!8k=!aE8}GE2GkP6s)iGyK(Pm;8M|jpZwd`FZfbMQ>7z0gHY+ zZ*}sVcRT)3h;c5_Dr+AabUGD{8+P9| z$DTLGKD9&egJa|h2qXObhaJ-!P2Edqb2C4bqAy-|6pZ@)b5G*uiSYeQEt5mgH=Dn_ z(JkP2TUM|+TnXztOA(I5;_He_wU}$XXR;Z;yDlG-Dy41LPK<9~5V$%cyWkjK(n()9 z0_w~zHlZ|it6tBlAEIUc4D`fKpRzl3pB9DXxaE+BKN6uR1L-1L;U38Nih^M0d5RS! ziUPFCB$rXzf%xxJPiRT+ljAI#7_9~p;88*oF+HnP1l8I9pcsoHDKL;if)wrf=s4lP z=q)n}vNd{Onx9(l4IXWW(}PA6ZPC1p7YY&^Fz^l^6NeLnXADj7__u>byU03Ble_sIqm^6Eg3A=&31}4@b`&&~>HBQ3m~fJF9%j6f_hU`+ zrp_`{+H)I8n#t-4jO6y05;*6=C!wQRH%(xzgg*NJOy+Rd#YGVBf#cz%KZ~gmW(fXb z*+z|RKx9*iSZ1N5NVs=&vlPXa*UKELa@6}fDa~Yk8n9q?T07qk9(W(UJfGc_Y{G~T z5E9Y+KGB|aF}#w$G^YCZqo#~<<`^9n$H8md;fU<7JATy`rR_%=2|#_AhK)ME@m z5tU5gD-;wH*INy!zEpVZvpZ;Z;s+b$xK9sQUP{6PQpN??v?(@bOrxqs*193szF7>g z?vcApc0k^p^#`1RxdXAUKE4$=#S~!Vo~(qs3Vr&L?m+CI42F2`87ZXpN-4)8o%RLH z$ZaCjLZka>MA-2=4U1%?%=79|6s%$ok*|h|u*`V)*$h|3Q}n14iI%2Rqm3JECA&9k z4L-Vy;UO`RP#>|pPuG*}HfLr;8s=1yC9U82mH`NxRo%&diLUJ#r}-uSCMBH9bmc6t z$1okqAM6>YKoySh6&J9rX(SjM&QD`rpuJ>=Ltdw!7)MjT+HJGuAPbC5dc}PZjm*S_oTy2H%T`?QyXo6s zrB)zg<|MxT%CvgURuny`!hCJssCOs}M1Rnn(}KHlt}N_yn0xZX^p)O` zc%LF{vSHGvH1E7lc8_;lj|?FD^?6JXdePp6;O~EBsmNx{F0tCKXw}#}I3VaoW4I6G zBl}x8*p>B=bwfOKr%56>uxTC@CKRSGawMZklI^J#-_p(t=`M9k=83hIzCD|8EC=eH zL&R75KgHe>j8E|d!onZbKKFp=X6&iGA+yfHwrX=ZH+&9351EbYj5(VapY9hi^r8({ zhcIx}ti9u3e=Nrf>)(%&RDPBA-SFBG{983^ZNmDPf@wNO*i!X5&M1Rq<0YwI@+9^1 z;65}EETQxvvXy&kD1?SD>yvgu83hn_;1}o(qx1nL(}llnv2?%F=h_^}b}A~1=j5XU zf6nqsfLhJLM_7Zjdh9cV$0n7FrxN`kv|%V7|Ikfs1X4Q8@i8r7vKzFp>tAB?s(YWh z?>#d=wwz`-jzvq~kAanU8>p1@rGIsWtqX7RnMMw>ki$u7(;_$#IkRAo-Q)s-A*VzL zlnu_^@IGF7YbskUctPUl+ofy%@>`5N6v?ZOWCYrB8sWtb-68U0MFuu|sN~}+lQMIa z1aFUc0^+`>`V}MIH(8ZsCg-S><1#+v)$>hZjOV^bg(ERw`n12sZ^1oFz;9%1P=|jU zY*%U6{Vh21s|P{Yh_M0JPA71W(tRNx!qMsX-5^Z2=|TAC@T(#`%(^qH2zvp4i|+AQ{|+A8X8#;rFu~g&E#$Drl;ZM%jo9?_d(3QV!XnvK_|RJ!Xy`k zCWA$EghMphY=1|N#Ja&MmrpHEepVe${t+sdL9W`YFCJIFfoSpT2E7}#bTko<{eYs0 za{Bk^0~aX>t(|!SZ8=@k@p)aA$8Q(f&F|{EfnKYxWH|IANH0ET?Z@rj<87zyD3sBd zjA9s60CHX$oc^xq?5j7NSGNDH2Uj(&r`ts^$<`{R2JW}&E_QF*CkITmaj_f$coG`i zzxZ_`Xm;whEZzr&4h-mGkD8VuIWi%E(!kd}nqUn-OXW&9LGWUL2_qiD3}|qLQttsMYCI1r`=9eH0QaG9rVvW>JxR zVD^tjSyHJ}%|w>_p4WLeT9K6WZA7fVpC|(z2n=y{ovk zrTAKl)%)w_T8oM2hpbu)$qpu_Iik$B!>2LIczCUco%QxAalmRaB_+asZ8Br7vLnr# z#XyrkT(m(aB^85qw^7i%nxx6FZOB)p!BF{2uSBiBP9a(vpghes zOdV+Y(NNO2eev$r#xh20T3Ur+-a^G6q5*HTS`|vx94eCZ*jv9Op3ieO(!9;Ki@x3g zE3td}OWvkWxqxTIg3;_WDz&h4#qq_0oA9r}KZ#Oc{? z&2{#bJXQ5gdrnqddgY64O{ty$aLDyyoD1aUU>m1JPqi9cwxcr@FtJ^FN-m4EJJH1N zonRZ}aS=A!y#A)+LT(;fn&Sbr0$zSAyAF13UT^?fO45w=pY((6g@txuNJ)D|+w%6_ zFD?!x$}?24Iz&p9h`25mBhDt&QbA@x1MAucUuWm(#-fs_m*$Yp;2s*F^``&O@bPc0 zGMVpvUs_ab+r`Y`oO^hMqS4+XX3R5U1YAxFZ1%;U;=-La(dMue!8DC+XNz%SL9;{0 z)>vYp+Uu==YSol!qc~S>aH`%S3G3X@_I8(XfB2!Ruh0mBvhTy0Q)_L0_0%+2IInRr zR_N?*bM#>$+lWSp;y`GVqM49siwKPU8&7L4Jv0t z%>;l)KXT~uBjbYoutNrxK=2>xS@1m(rK)c_-6(9=udH z`z@3$B{?OT)KDiC{}vr%Uf_s?21}R;F7?^y!U~2U% z@*!PKPqwC8F9|qNHfrRPvv{RF&JQl}(#A`qO4QkA15&oz(G_iNBOniRuaaImBbVR#g+B#9y=^@o#GPB7KAXVn=9kUd`ph@YfS75*}`79ugGcWtB{$dHB}EBxFNGSby$pg zOTr7KiQB96QK$R(0>k-Tr`c0fX{ryxSg0}4>0{A;xxhVObaOVJ8jt+Ps-8b<8jP3v zs>G{C&45RT_Ilf{ruwBCw%=sxvdA<(;yjr)ngS)r;6qn>SnsLmkIKOJW*WTAXiq!A zh6#vg!{g&$!;pm0PcrP=S0*Qy?YQuxmzUAdU>5@fI`V`On$?hQF|4<$Fj4Q{Q>tof zAM?vQ@zlow1J0-SPTLSaVM0mHoDnOH`o+*iTQ_;o4K=;8dcA?z7Db>EF4hhCSHQA)!MpqK) z>BXmeF7~%uTNlTOEG*daR{MMawrNUI)UEZP0t>X#0_ugbxwYTxh8I$_^e6I}Olc-A+LRU!N_E9=9p{*cAB@`Qb z%X(LmX{@<+plQJvUQ?C5;?t=fGSk*U$_z5x&7-jVeU!fqv3ZA6@1dQQI-_Q(S1$;m zZ-`t?ynwyZNg8^O>a)Wfd-P{*eRXpcflrq}T=SNIFt-ARTA~g#si!U@WJqpbf44+G z-5Pf<_)_v2t~d8txY=9#bzs{@bqCUwJ2!!^KzllQzPa$xXms)K&OWJ+)mJ2@GCjRt z;sZY(9u@16v&Mh_`TKc#Nx^<=+A*1yPG4`$e7apjQwv_Ek}U>~Sjo;nGHpZ7>T7i! z5T1q%2Zc6p<+MY(Vg9r^XJLT~pW^Le zyL57R*P! z9rWz$abpxPT1&ewNDZAtJ^Fv-P+vB z$$0$K#->kkY? zEFo2>z5lSChg^XJ-b`jQ>^rtYLBX*1JCWQ-+v8hxjt}A6z>L0*UL-*bXDgaM&E|9L zJf8Zq^dhty)dBzEsqR~M+P3hp(cyZI?L1@s$0f6<1E2fzO?M&hcx90G|Sq;=hyMG+Bb4+I~Kf7>t7Wb`@g zHisQ|v|&2m=3bhbc9foTa(8*Tpiy2ghr4wV25`JED)9%TVt-(e5bjxA;q#5Gq_EOhAJ3J@brgK=x2-g9p5*tIHs3>%cdcunIfESA^ z_#f>`1+356pH+1w)s`-t#Ohogi&hOk;dDKVMfComGK*z$_lt7&?|{DBaq-3xYtY!d z{hX27DInRGl;j2Ywk9A)`v$bz-C-_VwUiETsB!$v4nsW(K@~*B(^{pmxJ_P#?BB2A z%y?%DGPCBj#qFOPDbF=SVVIFBvV4;^bf~S@C%nE!E*L));#^!rS=TSJvpPsSI<&D5 zpz&b2S5_VypCw4^oHun8x|hz#$V~rBs4i|Hy*AY6$&3ZaM}6;0uyK;43C^D=E-%;8 z)GR%yccrbJ=DR-nlPE>C$f`Z(wEG~kr~9#B*JM3Av$cpfNc!NHcbC^Y0ii`p(i<}2 zbQxlyjq40TB1OMbtl#`cWZXO}{y2}lM@vgb?^AYmB!->i_wFR#`_T3m#5-g3<-gw? z!;5Aa+k8I-%!^kJY~l}WRsx{*$aCvY;lXlW;hDu#pLSYr#af97Nq+_BjFk_WfeC9% zoI6HJIA>zvzWt2(R^OnZwL#KyZ>IR^6=uy z?bc0zSg1{?ExRIlAm@IR>dyQGdu34b^WBlNGW6kG?h3|7*iyb(0h9yiq{vrao0}#y zvT3#pwwvo2eDLw?x^$Mk58sD{F5A?H&F+s5<<$*i{8z;zNvBCoW?yLL%nAzwBBrG) zR{nk`10JV_{`j`Q;;q2oS$9!l(IgO#C58{vT|8A{VO*~qSyMhSwx*}vThvb(S=H%8 zqbZupn+D}2+mAonoB&R*&$IeX$V`qz-ke$d1Thm60|j*I`w-?~ZHWJp%**PA>Qi%| zllQ>#+}v$7T%m+j({n><-IX5=jiPKXHiPRSfIZ~VjB>@LVba z)>9mHcNRY)PjbUw;BIMN(KukMHt5hfbU?=ILhtM1GTO8mRrO)~XWS}%O#ZM#tug9{ z)~lQI2fgx&m2G08tX0ac)um4bs=LX(u^)|%8(^WQ@Qo$V!PA^(z_IiMnc-hIr~4c| zK&!+cyJsXe_U5XaO#$N6;%Y{H@pIK3=9S!SEq89!GC4_uF;P>~&;%6?0&!5O+`dnS zS>1Wmr1pn6L*4RTNSc_w4u;Py&MiK1Cx$-51O*8R2;8)jVo$|5ED9R5wfVy2M|_K) zYHq&otM2UG8kBKwqzUUJX#8~3z7p^P*ehq$0!Btgj_=}M>ALlun;&)u7$06fzBJ!W zDWCi_L$>Xesab@xDLxhu*3d2z-5Uqwe()-Wkbz=Ug8g8RJBhj+pvGBnU_8-QF3+;_ zUc;ej?3t!N~XB4S?rf(zhZgCwELry59}ENd2_wH zQzKR8L_5N4HF;%!rY*BWn-dX)l*{6wm{)#SrM_|%c@rJQ2iY4DZ;VEa_guce5?+a& zSThWNI{QWW08{uQJI_yMaZmh8yASW2u5UsVZ4Z(e|Mc-9uTIPB^8=fthN|IIfh3@P zC0gBhXLdg4$=V4#Qhb9m7rrO~`e0h0H`OA@VRC-6isDg}a4`+ySjP#I54v;_b0j1M z!|JL{vPp{5NQttF;r=ENWo3<%mzevGf#Sqa$|8$Y6AfVq-Z!xn>N!?2L!~kn#mjTL zk{A7v+wP+nw;#x7@;%;1KaK|OS?#LoyL4Ps9J<@u_6BsbgkYpF=9bzPkoY7K_a*a4 zoT2Vpg9>4vDo&pY9!K$e*G@KO0U2Q&3E9;6;3>>Oq`izlK4UYuh{5kUA?K(*VR;rX zTkL~YakaUzAsvrX$a6U%4)8Hf4>jIeVNITHlYgO5<9utDnBk9uv@n@RPW5@y$u_0E zglnb+En@-`{dIe@(I;2ofYm6Vzd|*&1HOUA|3j$0yg6O2g=8F6_Yx}IF30;Prls`Q z-54y#^3ErG9_;l8O0tnT{-ZQ8IwNXWq2SwJ&+!IPLnj2|K@zvLS4C5d7TyVqppo)# z$JKP`_;aEr>ue+YNZmvqj-s6B!;`*RO5b`UkAB*o-8x@cwFqf22aH$vWw7wq+s01g z)@?^~-B<%ahgFbnclZV! z%REql;Z3p+5noTf1Jp)J9v|ETFTZ?N4y@%KHjYJh=60K)JuQ1$LV!DBuUSo+nb$CBCLI!=V7leIn{{#DD)sg1&3Q7T+=<%QW@ z36Y3-UcvS+V^!?vx39A_u+hvG=aFl48+}hlzmC5elbnm!G`vC86mPm@|@^KnS=;ZK-=~(}((X%wLD4I9t1-elm_#A{(51yP3^+x+ zGEYTZ;2ZsEdt4tA(55x7!;WRF7BAL1J6-bLmfq$|xLD6~9*NFVW|=xY7x6XZY?oyir*=G`rXHCx~7$JJG zb{rYi9`yZY{*se9F>tl0kK`FE8&@1@I@U0WKdTfR;n=^Zf6uUt?a^UOSf)m&)5p@a17I*8 z`9Fh`^B)WTm;WD)o8#Z}f7W=xT;yP0E^^{Qu*QQW0Zuyj5Z0OwvjyKopThL z{3j1ljwTs9BY%WHN^$U3$vE!rBjD;r)U zlsYaw+;i1=>ErSkI5P1Q`xy|1icIzWAsfCu4v(U8%EYCU?x%H(E%Yr2N2}5AhL%`# z4$EX!wIduEE4Sxo^UpeovT7aC`PChs8Edd!!QAxD?E=Y_Ir{oZ$Ngc-=V4%J1Jp=q zQAktPr#Jq3c89~t;>_jZQf$~8(L1Zfm{eVoqR4NsBZ7}RE|H!w;W{fYy3J2geGhxC zVG3nS8_bMX&K1y;_a?u#`?S7q*$nLUMcu9EMt$9gZowW&5 zipJk3^qNAz5~S;x);wyfssiwH>D@Q@zrUhYK$*U0YZ7q2H(~#ib?(3D?sisx6^CiZr=Szr&V$}3zwO$$1tz)lCo4ci@t*uA# zt$qrt0jz_PUE6?CjkGGvfyCe+MxrZd;`fs8V3CY6bh{*;s2df#tj7S8EIVsL0zO}_ zmuH7h-S}l^JHaPsy@e%I(!bDfuVd^Tq@@SyL*~z_`i7TO_tADP_{&ypYLia6kJLbY ztDmD|P#5g(hG#|xlDwI>EuM}gLw1zW;NY{A$`)GPeFTYL2JI`$D!6es^EXNft1tq@ z5B147RAXYz#LCGp2<(7f%f@9 z-Qi&rWZlWShU)-pHI{t2@W3%khI&kz#R`k%n{s0wC9~|xacHaSB*}S+`oVT79=S}{b zKTV&;E3-0hJjV*tTIbC%c>c_2TIt;Ggt8iHTP)6rfFScv!qcwhjR>ZSpI7?XaGy(& z+@RBpZWCgu-m$? zg+8q3-Sf$!j-!z|@^Ur`@KcIOTUdl)?(Z3uvLSZgdBA+y+zfPFx1!+Z{;L!F)pv7G zN5kXc0IR9O?-2i!Q-uz~l@gp-!zT%;^g8YYu7VI(Vt#kvByCq|?7`=Y`H>qh9$!W| zVXK|FdPKQ>72w9*mAUk^UH{_u)8tR7@P@;V;GsE2oDVg=U2a#{a&d;J%-<^4o9({9Zyg(!GADk0 zzmUb-KQ8V$cXHMp9m|mB={1MrgQI=>4r;_6DmNlP-N)4{^3UMTGj=i2kq`<|;OA>JHMbTp#VvBhj^Yr3czH5m4$?%lY~cP>0XJClv2Hf^OGKIexd z?_FBj)Jb3&cpR8`HgSg0UO$MLSpD^M5!4u-?2Kq-UtF1MuVV(2b={NjMaQYA2<)C| zAFBrZXc<4g@YHV`H?N%Ay{?@Thc=7zYi>>Z%2vY(+GWuuV?U^};m8XT+6Q7GmWn~FG%I>u0fPiKY zbR3=V^D~{h-&Jh}k0EfKRJPW)HS1kkHmM32s5E~M)g#f3yWwsXbhq*$uHQLARX_P{ zn+M_S5fL%3=N3?oadDgB5$jr~OAx^rAHDN=Sy>TZ?I%>^80FV420D7QG{Pq;1U+8L zoLyIf2EGZhfHsoV|I@un@+R#L?n@+9l}DKg7b4**Ed3oB9n8E z`2dguR02*;Cjux@K?paF3Wf9%=t{g9C^1O@JtvL=32+lKG~|VfsXoL%+7P`I#u&WNofMMs=5(KBAmi`;XL;7D{0GK zIV6Gm%(v)YzuUH6o>O568&+IO5w$l|o|eGdlsC6D_j%^rFG}T9)|Ryv*|D^+b6MBw zH?3r+K{kYA2cm?cuMZ-uewzOoT`>`dT4{1%$FJkB|7%P}6d}bb78DaR3kV>iez}9; zxho56Cs2Rrub^ImRhKz@ciBW3OXhR2GGd`*>hmxR0Z$nFH5T<~hvC`Fz3s592rGKNW9P>%g?awAR?sV9^ZQsgHsR46FbH;nm1{sIR zLsxbtYimM1MkcZ$rGUJ7I7S|aI8-!;XRT`}4MQZtWPzfwl*9vjI3tedgobn$-?vih)1d@R+&u1%M1wJO7`RaQ-n8; z;|7Sm2>P54XU;A8BlpPEjs=2B5A|V+OOYpSpEGvgXA-?Ns}Msto~Bg1oGkCeT94ooKFxh&;0n??=E2bmxv_*GkLCLudxmIH1-&qY z%;x&|%H`Et0y!%Mlb*#YL@JgB=;anKYONKM_m+(>d?DN_dT@ZrnlijsfF2!JXjYsU zDEFErL@UpCPcy_^m7t|y(psa!Q#}t{tZ=CwGMkS=*dex?M_5HZ-9oHaNh(IhA19km8Ui=v0DfeMoO4 z@6e&>&|DDra3-<$HrH=?(~EqN+WEL7u4M(35K1zM>RKkr>#0;3&`E+d-^G17-Z?S9 z-O3tg(>XvW!I#|5PhVUIUB4O;LdVGaGr@DL`zxvL<2&x)^8Xx+OdNlZ|2aAhG9(2l zom-ItTLmt33ZCn1F`C%6w3EA9D3k@3PCa#U&DlSo!dFVihdl z(zn=}Z|ShAIVAA4q!_&cvA`mW_gsEfgdjdKN=GEh<`0?0@AmiUh+yKg6MXMX30`aD za4b`LKhoA6yGQSEmeEN~Nfn=BNz%40r#MyA)L7uIzSr1LMYW!W(0lbaL4MoXE&LbZwo71r=RUA095%N*L21%4I&KWS*x_v&%75R(X z25za5XQTIa_7r2LbQU;L_Q9D3B$Ah7TfE!OPn1#HFYNt=x!G)F!jP5>Ad+*WZcreN zL;D6~lVb)URaH%&FOegk8!?kmHkfiSzDMb$+W!d4Yo+>)c0C*JFrqGX=AAwRZ-U%+ z*+e5UG58_ujm*SvJ#umRk;UR$391yynD2-77#oRk167rR=>z>32;+EPm~-?BA9y*^ zWr7ksvr$Hal+^;OBHW;E0BmEcH0sRMLD3Qqc}OO{LG zjs$7}32B{goi*>WCmPU;<#_`U z|3HejmNbqxNWt~yzX1#`j(_09|A2|Nb~W%Di~xZ_RU#jUCJl`d=E)+spc|I1uN*G~iAC{|ke-{&yfR-~UedMqD`lJ0UOM z8~XXbi*T@WaQ^qhaj=6x|Hln-{%Z%v+n4)4#Q)KFIJo{hAs2|36EeypfdU4DIFV^+ JB$cI*|3Ba9cJ=@O delta 11129 zcmaiaWmFvNwrwCGxCZy&ZrwDEYw*S;xVyUr8VwMf4(=A*-9vCExVr~;0)flk=iKw& z+xL#~YK*G2zFBjvvL97**4OWw<``=Z#}#Ms1(}nRgM*iwmy4_VHys?(%Yc&-#0jXS zhev~h5)Vhya38H2C|l$gL(=DEXo7ptTBS*2H2^q>J#fbtSjO0>h<+4HUTQ6>s#CTe z<43O^h%1kqNIO-F2@F>Kg#TRaU2X2(+#W#gy5@A(J3Hmq{Q3u=jYF|1Tee$R70Z6!y+(D~|rXP8| z!L^TV54*5~#!&J)EUjUAZdAP3Nwl9}o@OMDJlI~xl2_WMt*&$U`DgQ2jy>Z&NWf6s z_AFf_KZUqRJU7A3NWuw@)vY5v`GE7rr6F~3i*Lefj}1O?dK0X9+~s5!F2_;Hi@Oo-40PSj+cO1)41ygF)M1hlGA%zfNMBYn@%M#2gkR=?P~TU zFn^9Cy)jw~_qjtg4-7RdeTT;ueK}TeI1$LMK%<;$>z4bY*J)L>1cUYNNWv)*v2G?D zg`nTdLe~tfp^zf&N@aR8+^jrQ=j3J`8 z<~jSvW`LhLbMifeJ1IDmq`Y{@wyJ6!*Opkc@ytvXvKifB1S*mIw(VG#UMtq_m z(fd;t?cxwa%tSmk*-me!~zy^rufJ69|lHhb3s4YfNgSC8`ahSDM;r&Xgch+Wppw=rOk*M zIA&LNMm%QE=cpBh2v|au7mHxYogXJ<>gxtSK=>+RbA1$m^w8r}rZNF;?0oG-q{XLv zv5P(x48N{lH~5+O>5~`mGVh02V|cP&$DjWc4GJklh^%YNUg2$QOWqUfmD(dAO+5Kj zK}0uC1Yv4DGEW|;MH%pK+LN$QZ80}8eRz7Hll zmMKQ@R)ENIb!gGq)iGvd@8Xl`T>Vb9*)qIN`tO^5q>^Jrc@-UU_6tO zhR=x7)a|E2cU_dtdn3jP0Ve7Z;iNvGq_(xQqwAHE-crhD(8?dMS(oO0|4M7d4h%u+ z{c~!C|I*CL(t+wz%pcp92Qv%NF+Bu?FJi5Ry5^gQvrr{JiJaW|addGl9Jb`_WNB17 zd{HCOf`Eg|T(jG{Urs9!(O$@bL8LBojpT>^(YBG{$ThVqbxDL}m9Er4XB~4w0h5Ra zH2Y%1%a$Y$oD%}H8CFU2|;JX#Mllcc(OQj+)HLqREt6cSO1Y+Gds5_IJHH|74Z(FA+QoxJ%*dRjVf^Q=&{ z7I~Kn_I0pfR%y<%(&7pC3iU@AX7{@+`OH#I&7-TU8$r2xJHQk^K^EBYf^N`g$lG24 zWq449@zVraLhhN^`7m(^!Iw75@~@9IbxM?nDJqO8?6~AOB)uk`Z*`|cHbQ5)ZdkA~2q zIbOt!(^rr;1o;3m%6V?x-IUqhJP*606I3$~iri;4*nIk2OEcgGz{KSGV@e3gM9uhg zq{KjEk9#%2fWjlY-BqrFTY;O!p(gv88A05V*tAIn>5~p@L$WI@D%Y`K+K8c{OB!0? zrT#jtek_ZtXf{MFE=@n@g+}g`$$V+`+&cxsbWLkiumvc50+C|gbI_h{qvmd@`(@`= zL8;}0uajGRSLhpKXD!>}#f&d8+k*jMgvE!%!{TrHY(aEm7+2kG^>h2AyPEhLIZ3lX zUH0o_&o_#G{KqYmDTurCg3Ow_P|_3cqSCJW<0DUD0LC7;&EhmupH5MpMA0A)aPd{f zQ?8pYd`a9_A;Xyh&1!;Q57|nDk1>CzMJHwo>tmTdrd%>Xkf?I}KLgvP z_x5%h2UKUzOupw$37DJC%e#tOLIU>fDN3+c^JVuDR_jKsBDi;O#hZtkB6#ujGWY%{ z3JT&O);VSBq`(feggv2%$nUPtf6n@|uuW(U2ID4c`9`iOEw+g`y$tLGf4Njec^kV+Jn(5+j)c5GaSaPpZufnE*kF%{BrDpNk!B> zcs4X=JUYIM;{9IS1>bLPHG=(~1$JVlS_AxVP3z32`(E2_k=MALK$XFMD~qD%*m0ze zi0{?n=qU)UAD=ZZoJKl08~PapjfDhUZUnE(r!->QGWAcN-fqmF;J^aVNM3Vs0%2Kb z0@&h4_O3?GX12DhDrOdLwnomd4KxBA052ycD;E!#l9PiA^s?aw!;sNQ*#Al>QP?HS zTujXDU0)br(92NV$WhwN%EHq1Z{iD}>S|`ELCFn*DWe14fOxpTurPEk=%}Wb+E>j) z=~y^=a^8vK`R+ScqqOE?5v+z2%|oZO~hI%#QOg-lxAv4k<< z)-0X1aUtN?)Run5f-d6(HVDql-@0;2Cxq zCl{@^)FMd8l>A6?(pF~BR3lj{p=dzNR(o?V_r#;M?Gf&z2Qu<`1K$w@@~lx86{|%0 z&V^aw%_@wPipCZ@nxNj@vq53rtVy^Gsb5M@FX|@ONe3tTnXKQ8?52Fhlyi_?A~o7X zy|BIed}DldDQb&JZBz9J=ZC7KvtSONWM(zgwRb|4-iGjmpU;Y*w13E`NB)>?&y#OH z@9gZU$MJ+6!u9w12-5^dgc7C;$**-GjxbAv5MX@>fMpyB>l@5>nluOb z{+?r%2RsbAQj}(PgfpLyESvkKzuEBPu9T1M8{Cq2D!YK34i>~72#%Y3v0rR$J&^mI zP3DLX%yBY>k*0sKvdFvFyxcyaOM&kCk*{A9OTQ<0V@(LmkYrXSM8}`T zdaEvGnwb0DuS^?%$qvEc)IPZuS0ov+DP*zu z1U8osOV<#vuvM`jE2?uEO;RJH?xxiZkJ&GLY<1%qfrf4b0ft}inGD2*FD&T-ek*<8 z;Crvf+!7|sqk|?JAqn|I`>u!peLC;lQbhH4`k$|A4i(LA`^OaoBLSk}bK?%2X5-_- ztDToOc7`Q7dIV&J9(X?Ql6hBits7V!f9-_ca*d>pm_@KtfPx%F{3X&F z2OloR$)Rd97397%zecSW)a%Mx`(7C?uMe`5TGyxy8_haCOgO`1OiUZTv0n>%Hg>V_ zuHX#4KhF|C+M97*X|1I=r+suWMaEz$;>r}S z)dztH4$IvfTgvIqA#2VBZco!|CZ?x*I%z7HR6`r)@%2{gX!&ja9O9fTHUVWpm};|| zy(Dc}ik=1!ZhD9!qXW(i(GK(XTP)6k99TRYB+HSoG!)K0UjZkez@a z8z8{681cwVs~DnG>eccjL$={ZCq@%=ahK* zQtonAks1ayP1Er~I~Ilgauav+Iee&g9t&@Gz(x)nzyjY&mMSQumHXnm(t>Q~^ZeAl zFz*M0X$E`5me++yNIqNLb_t@Ts-k0?>AktRvb16C=j7JqM~2(1>d2e8G@4? zg0$b7$^>$LIae$azFVYLTv}DutJ$NZ`DWS%R0FyKi@_BoY&rRzZB9u!f~?+CoBbF; z9l|VJ0$3WReFW;Q8jhl~NT?ju3(uF@-8t#jh&>H~YMGB0A_pdrj|k?!^YSL%5@5BQ zWUs!r4Dx(#=FE$vL!Ua^K52q>4c>I9H-rh?Waimw4}Zpc-9wSIC3JfPttxL~M{yA` z!Zz~v9?Lz?$C|m!>T6sTZEHR151ue6p}g^!>X(JH zJ_i;8$q`mO?}<6caAk|Thg(!+N5;+J@baHLsrOcHatg9C3IZbM>wrg>2>}&kEHmAV z2D(~+NZra6wM0ju-2Pfw)txT8o%EP?YF%Sa};UnY}P*wq*q|aX`_jT`-g|SNuvouGb%9{eG6VbDn4hR z;Q&_;bchzZ28+9lChl&Fdt0Q#zB1a8ZP^(3*3JA%38~mck zP-Tc&ZM2p5@Jo|Zu|lThDY44H5YDe~?UJX`wn8v7X-X9J=_Xyym)`ppt}GZ%&(5nt z9)}PKibZ{azci1c|UdMJ#nKN) zP593&ufCf?PrqOPJZBltM>@sxNd-`Slj~SIoW^b&eRMu~6|3jYqQP zzv_3TD|j{Dug2YHv6M?-hBo{ToiqKo%fGpWe3umoN}ei=*-V=CCydPeS-NVT&a%!q zO!4<1`cjP;9YRT zVc(z7Mjv8sgg!!bjg%7+<1TTHiS4q3(8r*ZSFOf&5!dyp`(?pR_e1`%y@S(Kp|{C4%jh+);!;3ILI{@hcJTu<&Us(m zUprP76t+n=s*#~6ZMOqCIpI(*7rP1UR3#2INQf4adDbthQ-r^wt4++N`ngn}R`K+V zk%~l>=cE0^_*j1nIL(;LS>Si_CWhC)%?!r4`RwsZvQ2cA!O^FG1OJn#o7jH(<}EC= z2P$iGel6o%5;Y>7=B&zxz(7}MsQT5*@xZ~x_~yh+fum~Z0fjHJNe!H?ffIZ3F4`nuxgM2@L7`HDhT%F6C0YM4NRTx>>1q}aNvDj?{8ShS8ojHr(X~zeTI%M+oIox1)N0{4lNrD;uP##X`9djA0j2_`>n^X5}Pp_cLNZ3rl zQo(sYDuZunDF+jPvY&gVY-Gs2jmJvXo<0Pc+mnyT7iPW18PLC(LW`?9y|YTBac1;O zrB*pH5*~$N*-t8>f}vB2h|qW}p?Sq|`i)N=Sr~MM!2mzLU;!GhA;n>t?Cnz}6h}h# zNV%`lR+Jk@$ikzsST`kE<}yUaAu+qs0>68uE=}LxHf7;EKGkGl^xQxAc?)UvuwPAe zRYe>8DioJ)$sc=Z)U-{jBGMwwLqsXqBcmIKL1op0q5A=C9?E+dhIbyI;@uYS zt! z4X&h1?-CDFrG`E%2CbQk=F$r=(8KaKb|j0Z#aF<4lB_F!&@}fAThtEr-aqp`1g>tF zS704&!IGXnngr3p+HZue-o7=~b{4R_{^M#75HG)u8ptAaYRgiE7PRm{I_%x~G`8cD zusA;x_m%X@&d)m|YmLdr$04X>a?8=3&s#0rZY&zZXR9U=qjieb9)gZiV(r$V9-@~) zu5PATTe!9g-BQ{J9$whrcg?d^E*Y9SR8q%ZQY%=AR#-Dma}cd5Ax?01cPqIn{;fs< zc-9~KR$aOjE>BEaQSoGoSE4|F%;u3av>M7cPoj&V-Z?c@Qd5d@Ii&8cXs%S+Kb4V^ zxNCpp$%hH z{wSR|eS0`9-_|({xn?EraHErFxl9@~6ZvfEFQzt>GohcX2P~W-nfcw;d zbV2rZo1e3?8!5LD&x*^!L9c{ursTH}HX~F+Vkn6!@nIuOI|s41>D>8#Dv+h$Q5EmX zwRwe2)RpL7BXVk_DgvOu+2<=s>({wJazq|81@V=9vhJe3t9&CA2xahpq=kX**kkLv zo-G!ivZ{6PD1p2$V-WDA21v=x!9z$10P(&w1OKbpOA{dQFA^_In}!A+z)>5= zOAiOn1%egRGQn}fMroVhaBzS?urGAA|I0-^@iW3~=`}I_k54A%U~Bp>lkDYqe_LRj zoZLKrd0&A4Hpu}0ACv5_{r@q^gc$&DxB&nV%$`)$K+WN*XP` zl0r91BX6F{m&rGlZ$-@GH2E0}8u)W_5a5)qV8Jyd%s1sL(7xu^`mnNg638KDNz{VK ziNxIqC6=F?|8eVIgSp!JR(I=E|cF z#JDROVgrV5ysP*#_x_S@rqHi1bAV~24?AMQ4)ZND5gOT4`%Bs5&E?%UIf_s(!r%M> zP*mf=_go=VSX4gisM^7ucT`-)HIZ0#Xm$C=P-F=qRw>0JQL)mZGRxOg9vQT!p%v6U z-?kD&{KBF)oFBILelGbZ8$6zKprug^Fa;=)bK(pr1pmWR9 zs9{q9V|H2ij3Pb+6+GA*X9(~Oeq2~pGk!DDc>L8(2i+3c5TW`LQOayvJZt-Ea;Fh$ z^p^;NtmrTB3VkFL0Gs(zqzetbY`h>&S@q<+AGxt5`AjKVOdtCh>a^_=83j2fb&J3B zHUX}t7V@NItXBd!A(Zc}W$Rd=@sMFq!TJ;cEO_-KKcK?B3X9ni9byX$!l6n;8G2o7 zO7wr^!;sK2j2#e&6! z=|bPxT3&vcV$+3Vb~5s?%_)T}l!2aV-y zJNuoaL3<_9mK$Q$nz?)3o>BUgC36w6{RkKEmP^y|&j~;$X#BKDyY}{I#ia1G;trs} z6A4rUAJ)xJ>{gT=R4PL!lI=Usc{C83nL?v++o*cnx*iamYrp<9Q=DCG)+a-bgXsBq zw(SS=Y|*&*=_eKaC>>&+k{*35u9s?4n5e0zk~ zQURycYgdyp)O})ANg#|wK$H!UXjBrnOKpIV^pO#I2;4nh|*ahfCAH|&4+_P_*es(wLZKu`(;=H@es`~XFM}`3w6HFLb61$3nG8M06S#WP! z_4p)tFB4fqM!O0vZjVZKUG=bS9II|lj{U(Fx;j~qtE|KjHJ2Y8p;&a^#k_=>d7ZWf zDZg9QBctA71aqbis9^;t6onA&zAJM3I>3+IreGsRtBusmo38cgr^n`KjKYfgrDto4 zc58>nifK}{p=QVNj7eXSM-@t} zsVE5IB;|NSYEuYX3w_kQQgg&xG)mu}#Qo|#VJ7wKX*sCg+!?pB>(qv6sisn&oEzGn z2v4RB=_;E^k9{blGjwyB8Tx27*pNfD=N%-fBjY{;J*i3Ue9ep@is0!Of@5uMRfv?4%^cu#rDF4bR;rnC`CP@P|Nq$~} zfFwq1T^hP}L0VVRtS(!#V_V!=2(Stpm-PJt` zYFcDyo4-uWk`uG~&XyXEiCdcI7dv8@C#ioMQ;g$AmuletPHtYZ>B4AA857QQuZIs3 z7BiJWMvl<9L`FlLCJFb-%-oD8cgTE9UtU&+W@YHDnI%Rr;T`X7;6APobmv6m{5q}B zR4ht0XujJS%g|(pRBFDyRgEpqs0x(|t?x7UKKxXd)>qLy;gjQ*1G&PosK2D% zUcdIBx*2cpBJ6(n)nA4&OE+1fU&cB>lmH|}Z_Dqw8!de2AaGS1g?BqQs4(J=;LCmj zwV_&E&N2EF_?dL{vLKdByUJsGD6v8KIdxrs!${H)-|d^p$j?w9&-z) z*uc1peId{Q7>vYAWlOd3RY96^A>T7%iN;6Ietx+qh+nLGS9hUqdB358&P>XYm~L+g zPy8_PK+fQJ`F^KYGgbZFq1&n9%R9*z8l-r}gR{z0xLZV$9rKeiT`|dsJvtaW1%7y4 ztv68=2k?w(-v6j0(U+c;f~hbPyYlYQjpb43_cSWM_oi$|cGCFL(j%=hR#}&Ho%D|@ z4T6?8w>@mZ_PY28UJ7zLOc6>@3H($lD&WxDxabU#jBU8OBtnsHWmD6bJiSOnD5Q!& ziT=UmE;IPl>aIk&FgfIuX7M|5z>Nt|COQw<37)T40WM~XArbBMs+e7fZ!0?QPx32J z%jb%&52f!M`q?*Pq-mQ&)6J<;Fm~U##6_97l|}fb3-)Dha6>+X#L}XY-+AjqYqC!` z2|1w>=1K4J!oBO^szgl|wRSJ?hic&t-=Kiu&qU!OUMV7jan7XRqs0!xtl^~cg2Nu( zkOoDEci=d(XtO=R{9#JlHPW$RT4C%?=aY8JntoM^`7)q((0xT!e7|hV$n^vEdfI(k5@Z^R=a+1^#?zCY$4YFeIu!%Q15B zI^(V9_8WA2*_j|UWYgl;`%u4+s;I6!?2ULz!r7E%=WtRA$)eY%K^-LCMy(EuLtUQ` zg&Xl^0-)e&w3~@$IbNRC9|79cqGw-WS=|A4-cNcZ3+2uPbq8C5ySTuEkV;er(M+59 zU#Ng`j?GZc=osaBOpT8lo6QLKVc=@@$Ono`5n3iK3w`vRht6QaF2c1AL32QP$}c=cqo?iWR%shex>5rCnD{$EYzQ zi|_VZ8VhA6Hy1w(Oe)S|JH`R&3Doy|!X_2pD_8k|i^KJAj}#hL`pb!3F)(!=h7doO zD-d+5vB%Z%ru$XQZb2~|%QK%uCJt*|L>ze^j_id8b_<;0LNjd2ZMIX}-90>=O+3Xg z8sWY=doQOeJ4}bBAlope3B}AI%X=TQDRP9qo?-lXc;N&7MYwNx?6%OH4Om|<)3y76 zYJrX`LmJ@UIAC~od!N4QZ*YnT5$u2ycVWFIp=k1tb!YQbN`TIgy~@yg#oRT@hQwTQ z)wJ&MzSJ7wbm}91mk<9tdnYT;&j?J6(CXB6Nh#I}1lA*6>_&#zJY#>upF64!X&o}{ zY3V=9e@)nhXP{b%t%-fN>?&^lt?T>#7Lx9^_Fao$SC`e6ds>Y`{-S#eM}=HnoXJUD zC}!<6&@sDmhrL{rrbP|&$@%iZ-&f{lwW2pL=8qDxbXfW@ulXEilVG>IkYmK6GW~B_ zgCyx+BSWg`pL&DMSkG#1v|qy&W^sua=)mzm+1w77sdFOas<98NN%1NA@?!?0rZ#R z1HW+oQg|=wk((1p3FhJWo6pVl5_w+KAAs`(dr^=tD)O%!@S-TWfn2pVyy|eUem-h= zE)LiZA4n7g5TX=#5quK=3Y>ojo_~WM{vL6G{|>-^Q-J>hdH)6S!gTqQ{;zAe=YM_! z;CYdgj{@>gz+Y#_@z+9WPl@t(!U5(5ar_hG;Q5~zfCC8rAH2VG?LYc)@Vwlye_$^% znD+&P{yzr>aDf5;B=YdS23E=o2#Vi1mpzg<^>_3;zHFf>=0MSlDasX`tSN**HnrfE;z7 zp+7=F2u7j_n=bhDCe2_hYr^eF0iTAIz7j%zXcIsYL+Qq^F*wFY^FBW-CZ>F3&&XUk z&^&?vMkeo4qrp(SdOZ8swesY3cYnL_JedB9;C;-~aZoU0=V$%L?+ZxK*^^@6j4_Ya!Ry;@g$hQ0hK8!sV)ski)I=i2bF(yZqU;Vakwq(lCb-y~lh*YQ<&s7!wg zIyG^`m&h9MLt+-G84-j#hC7KqQU*s5pEf~dl_D6Isiy2*M6_&WuCbo_(%H;v69Y2M z;BwrSz0BFXFFA|&Gj<`PhkFkG_zqlBmXjd{&W8GIj9{(iUc~%S`QLoxoWP+ei=H!7mFKJ za;3%hXTu8LveG}}H4AX*df3)yoH7g~?Kysxf`HTLJu5#WeX*7@Zl*&qvPM*Gr&Hu% z(Ohaz+Bi%C&+k@*GfNqtA(3&QB5(dY3}e~b&Hcv{gER@MGH&+N12%dTZr-Zog(WT} z$Qu;)g>E|(NHrOJWP3C3i!9{?K z2s%p|P|ef& z1f|N4is_rtv|nIER$icaZ!lkBRbDbog+vfCWJs6GQXI^%)*-U8GF#NI=X2Hm8a({n zH}CFON^0p2z=&+QP;9?c_rxF|`DO-*%lcx7-cJ3^XduZf)e1^R3kh)(yy9wtvH-tLg*dGAn#_`hhAnxdy!-kJ9FCsT*8e5@uz+ zwjXb6uK}4uGV9dLYb5xcDcfxzD1udY^$W4p&!1OE8~}^S)9!+hM|0*9GrK%Dbp7lS z0!c0_$!2<85KmW7pv;Oz8z{6vcE3?EhmldAiRc%2Os{L@ltWA2qPUfj(wv**q)x3o zs(uZW@Ql?ayO4nY3&%42C^c;?b)2A|Bg^c@!Dtk_R9G(gp=S3_*ix@wcudK;j~}}| z{TQKInk-pEvbj2hsc^fu;Cqpq{?=AQE55CB{`)THA9s=8>8{JZKi9YEHK3>noSE8V zG=&7Xv8Wg+X++6b?Ai~m-=LqouB=~%0ev}6d>`L=spZ1^(b;~V*lvUKgKEGPyXuVo zyetNNP#*nZS*%vlMr57K(1^`NWP@vy5QKnfOEJubfLTcaS?vSm*=vVLc|WV2v>H=a#f51hV7Ikx>baXQhARMPNpneQlJW(bb)M4*Zw<)F zMr#Rio02d>s4@U_6=RC z|Clcse4jxM`SOf=s_lZm4C=)-Wd6*JdWET8#Zh7@o1%~A;8Wo|eR23WFX1~t=;LE{ zpRnP4EQ&)Kc_f{w&|+iQIeRr6ZMy`gQ?Y)dif$s3PhQRqdw+@Xt$2FH;((qRazaP!;_>1yrK8(b=-D~K{oes(2Q(ypGVL{^N9@Wt zzb}TPA0lv8p27I({%f8=QtQ1pmHBp$t=Z?8T|70@>VzDR#V}M{X*RzP#W3zvD89#( zVVo?EL&Dr=lh>V?9kVs|dPn+L<4gKjgP9`DKy&W9N>PT0Gm{)>{a-4z1jgb)ykr<; zSM^aEXMI_*#cgrWHrEfb1^a`UzRm}aU)H2s+55`EnAC4cvddoYyyIT0tK9XTL&&bv z*1yxxo+97LpWEe*_m{qWESq0yy7Of^*bOrCFf;0cKiiGRxE5@x^!MTO^Z6~ESEY;o z+^UBnAbV3I*0Sr3}IAIBy2VT|C#n#y+L7x%_ zjRnL(%E-bACS_&g1dxJ3zyw7U0!A<^kQ9+w#Kg(S#Mb$*j>s%*U@u`}VP@`3`d3zV zHnCA7ZyF7@*)nhci5lU8W6-W{Y;iuIX?$jg-r3(lp3z-%p)%QYM0Zq(UxvP6WR>zE2 z*Jq-%!SwbV%Nc}xqJPZxsv=yFU0-ccf4!dozuuFq-M=~o0fb6E@CHFX2NvxfpG?fZ z+;Q1w1?`Ej)!?!4pEQ5GvitPC9o;Cv;O1_Aja0=2-#rT#AVpU{nEMQ0FIzKne8AjE z`ASs}fj^m`RLWiNe*H2^S`aDdWytRCiOik{XAVJHlmRyXOdB9%GVi^c;toY;+{B5+ zC7ZkzQ4hs_Vtzqo3p#~BH>t=dqoYyftqJEnPKScS1L*1};|766%41Bq&Ap?G z4`3v&8az_L#X_fr*i{tAODA_#v81M@+uWFPO>qI+9o;fJOJ*UzRmzgSJJyar-|l-} zFsk?+lM0?*b1u!yAhxI&CTTFvn>55O&Hs>Ap`Dr8c5U<;*2)^X|8w@dIs9iY`>Ee+ zZz951Tdk6#xg8^cn4>qbyB}aXu^l*P@+&zx>~#ItQj=X*rWe(1(2||e(~Zrjshn6b z%DfHe(XVuWc+?ZZ0|HwFSUCc*-;Xe@K(ZegHD-}br>ei}C)N1P4IeTp^rfr+rl?1z z3jh#z1+Fl+f33;FJBJfaURgzHTe1F@PEbU{I{4s4P?6+VKJx%$Rct#PPN`IrQY6~c zcouHXB6^x~e7um&G$!NbivLlxzjn$|bJ8kYtK$35-}+4uFpZaII(}JMZwTtd#Kahr zkLH=Pn{(?*4zGUL1M=Mr>LPt%e2+HrPt9efrjers7j0K=ZA&aJ36YYsnIpoJ4xf$z zH816aM&gL9E;r%6eVvnlXP*<-^x*;}gD-`~9~l{=i9>oVD=O%{WDHoi4Sw7)%N4%e z5DGueNN|rK?29LI$rTkFBZWJKg*3an1B2cHKG7g??ew~iwqg@KA03a>FS$zUR>sLm zT$znxns`w#N`|+u5lp!-Ly?P6v>rPMT8VUpzhh&9TS>)CQQQBx4yL6+zfgpHdtpz< zw$V2E^JncNGBRF6QIK${nda=?G;3bAQMTR6+^IDHqDD)V_j7NHCL}lt`R;PMW$l|P z?T#L}UV*l%UAI(;og*ejl>AIDCB477sIg|@wOeUh%CC?;3w*F?@&*?iy+TR{0{ z1jqTHll<~fCPqh)-Uy>F*bVlTUBh|f<*{DviTwCVz)#d8M-px>LA%iV9p9)OT$tyC zsXYulgxT4(e=GZO`Sz!Lh;TxJZPV=5pD5e$iMAE|D7B-D(Frr(4iETOdrWf4uxgos zecn?|NP5&{Qu4h|-Z+U~+cybk*U(Bu{HQ@6`wfIBA=}!kD;ZgmjEv0d_ctePzQ_=w za2M1Zi38kHt4KpX!U0i%(z#h#&h&kFZ0hLg2=E*v!HjE?W_b=7KThp4g5*AU-#pHD?482Z>ym)fm< zPw_Ud>?F6dmj0Z`;`O=Z3nl^lOlbAvv<=$9 zDH1$`(H|E^@N}%KLYiKWw~a^mTAm*JAyn)Qk6IyeSGVdT5rytZ$(UAt8&zf>xiZC+ z+eMRy%IF%dD-BHwg%ov7+3w+SVzQQB(n_vtzIM%sx!THBOil9x>LWsE!}I#@R>73V zol;>0d59kPXT!r8WmvkR-Lp%>Ont1_UW*BLuxF^c3Y$cNzI zU`A}QwYfHTPA_NUK8dJViw3UD5ndqIzs+DENSyy1HDy zdDs1ys**1M(Y(OJu+a?0zF?j#2<747_e}khQs~tVbLTyJLB4=+W8bEc63YhE#(ARI zA}z*&&Tf5A%kLGkh5i5!E_nFIm{b1sgI_H3#>YpQ>EQ1s&Mq!es)~{p6|<%U+sI2(W=NHrPZXc8O>h%>;Eo!CPG-Fum zRg}Zd%^m3`(wAX6<{TXoj{o@OZ0Xplm^;1n_Q=E-EGoMBZ6$!00eRX-?PI3EIDBXJ z!L=lcyAuS$jav6I1Q2pav8JgSO3S5A?rhm`wnm2JXEVj0`+AXKXh@x!l zs&|DBQag6iJWq(4nsIkT$_;X^^it$4PB$RH1AiK`aL0WW>l!;^Hm=ofBvbp6~U=rE#eM8d%`ICJpQRztu;U!-bVCmXu zq4?XPtyLgn=R-y`C`?z#X`=(dU(fo&SV>Srpym0J+9rs>K=b2_bl!)DzUGqE)0F$3 z;|#GCbC))*8`P_G`JYz$dw|yNmz^znv*ZAZURkxz!?y#k5IgoDEXQjr%OL3YPuaQ9 zqJ1TR$0Q>@J*NnzQl%M;W2m1802M0EF~d%VwgdZdH<~>Ry^d=dnCDYQGQHm17Lw9j zMGd`X1obL z{;t*27f_pfbCY=0M2q<2`&Q9fAn%sz1pKjUM`t8?y6E(QA?HS^y}{rG&guG=1MGMK6H4IUTEV zn;K-E6W6^K^rrIi?cnlyXXmHUnxk9LzL_t?neO@%_B?V!@&n~>nn7VXLGBMezlX&W z?==AXNC0@lt1yg($w{=ePHa1~*W_KC>e}xu3mo>|c};UFWxo!}qaL-a($a<&Za)5A zcz#Sz--s-ER#Mf(J1$v5qR5v^Zmi6caj2+o*3c4>j{TG*LKu*_XaJfgveBFH|MeZG;BPsJ-w`g9SR-tAjoZoM03-~=H z+mgNGoBVj@w%V<#;io;XeDJ%}+}jldA?8hz$~IK&WP*l{+-{pIyOS=w*!@;qT)Y=o zsF=rQ-Bm`%=%e8fXy;rbBDK=?(-ApZ*IHIsc&~^0&vn28!+=M-7 zRaCB?;NxYjP$F`$&zS#ifx8g>3~9qY+BattGjut5v9m`pRE6RW2+<(+wS4L(*xGV0 zD`~D(sJrb;B1*}4V^1J z48?WoipA6{0C_OA_`GP*FKh(iGQIET6&ci17<5)xm0|cOT$)kU3Z1p zKtRg|HQL0@HGB0MQ#9Kag39`lm0 z>|pVs;a8GP}xX}bIGf#QXbt9rr_#d$7-}UoY-6 z|oH4Az2wR-z%H|69(7O3Z6zJJMS<7BCo1H{2DJtZ@*2OJC z1R`Qp(y+`C*(FFY%v<#%afhXZpc!e`r6&c#6mudJf}EvtA|Ocn=#$6?%FsIDmeNYl zj9cc22KWGfzWxe1^N+j+e%QH=S5AIK>)zMF0@H9teG(M+(E`$lcS=`@1=2YYu_$Wm z2b<8XB0Q0NnEOiG@lPT-CWI&AwDrLtMf`q6%zg!ImJ*@{!A#e8JK}rnK7 z@%2X>L2@scJB#|BgRJe@QO7!?wyUs!C@X=Q^ASf!|MNb$=gkZokrB@N{5EoFu}XWV z%Nsr()1X#UgJpaLwEd0zv8vb&=g}NB6hYV^wBJ1y-!`lbe_?VkBgDHS^x1x!cFBwo zEBYa)yGsX|)7^94G{W5~A3|gaw|_AZKqba=kZbcccGs)4pDo84(ha%+h9G`RPv3lM z8o!UcRe*06q44FOSHPp&$ZAIgE;FnK^W$8Ild+7SB)*;2v-!+*!_?%2_4+8sM)4q) zBW|yR7**fnE&mf*+df^Q*&ui0K%!Xm%~yTCm|95S`Yq+0@W8o`gYB|yshZp(K28+u z=(zoAsDqY`{qiRc8VdcmIsTCoj7?;0XKF4R?0s>uLCzC&|0ssyi`q{#=>1Ro?WYa~ z#9@NiYU5|;?O3R{JEtlAaIkR=_Y7*GH!EJCm=R1GW~%kV!gtyP%M93=yT7_8hEI<# zg+4;8?Ntm1m4GxOWK7w9?hy!6#5g>CXYYMMF*oJaUL?WEE4u3rS>kupO@F2cTiv#( z?1p=RdaGR)OeBLXceP(4<~WckdcF`NI(pF8Z5{EeL$G1JcUjf=+6FlvXLE?>H&)ke z2}a#xbJg4~G3+u|vbC0E2J>qerdN0E0WT1JU{Yj=KY}SaK_knk2u~k1t96Q|eDDHW zLR8IY_s+GkYbKr?Rh_q6s0m?XqW~aImMRsdTR(AhC%phJs`9Bo6V#gZp(qukExz1_AOAg5n z)$g&UCYA*flwKiX$ONY{5QJYiZkAagxF*S(1}tz6MVJW~^Wg17jlrwWQ))K-0P0%0 zIz>gnc3Y>!B2cqd;DLx@>H$Z1pkA4Znqd#zA0qSmuDr;EwbYSckg4Ya5UK&#@}}K+ z5_DsMKkzP&safFbzYG`6Hcyt7?m*oAO@vBk8m<`>KVwMMeR{90DXEDtNs1xWjo60F z>}<$`CsiIYych<{L5(IOHW#3{UT)@Yhkm+M?P_aU6gINf1K<6X|CNJ3E*Z2n3V={{I_;NkJ^^|Hg2#vyrlM zvXKJWSxDJgI7!(#0RP5;fNy?|w*v$qW#wf3n+xFhSLb+(VFR&{vVy^+Y~Z&X5a_MO zTm66PaIk#(cfbaIbN*A1m9=h%R1FH7g^!d601yQHYgqrsuqP;z$H0mS0RRcxt4)*^LUZ(#C!pi~#BurDX!LYFb5)i502(LPl zHNjwNmL`*ASRR97qR6nPJ8gKuXK73l^6PKblCw*V)+4p_BMo;N@$`z6y*gTI%Y8R9 zmgtXP*d6@5NSo=kZV{IYCBHgFaI-mVAcYHu|At37hibqm5KNOsk*GRoXpzPUYlWC5 z1CuyAbQta@M)_Y~pC=YXamiG^LYT<0ghFB(JSs=D2U(HgkmlJ+#V7}Iw>7nKnHgfpIo(B)np1;N$b%3P1*#~UT z6l6KcDbybY*hTjw|g{7FA??W52900PVDQ2_CDbQ1k*GEQ6ni> zk>$ARQyqI$v$8}>@GZ%53}+ypkt&>aeu=UROb23bS0>z1iq2?ED5yGC5G_yzQsGBZ zAUb3PAUHl3Q8}5EO?k!?{e6<%38noDIdHydW+ERQb{}B%z5Mf>Yace|)m9X|Z6~sB*E!&lf;XcM@I6bp3}jb$L;$xK8al zyrvrPs4W&$H#s}|XGPXkS=E4S@*cNF#R%S7EO7MgG4z3p5!p4ps-ySpI)(1y_{%#`;_w2lmZyCJh+Oo{b%GCTxZ0D(Z~TcbP#8}LPCn|xp4Dn(^xj41Q@+BJm5Yj4PC zm|pc;@Bn;7p_a)U)&Wb|q;-4cV#$h5lqy5nBsL~p3&a7+(FoNNre%kl^@o{qyfU_4 zsqOm?1qZC~wRt49BfrU@^%f_jZkT)!w<}vzG%OA!OFh8zhwVh?>I@cEYinCupt(AZ zvYX;3Y8oPYSvka>ewv{9K6fXK0ok!F43>>yEAjQ#n=vE%;Tppe1@4&-nt#yn%ztr7 z{Q77*>kHX`x-ZTGv!DG@O{-qA)!Qtef}vqmtXOOa=#MMH7&Jwf}8 z%1wX*DccqauXkwI@vT>TXQy7KPIu!#4NT2YAqu2qkx8#D;9yCpUljfcs!VC4cRJmz z?$v7}uUjERQKuBj&1GVt%-PfX{?bm(ABmzsi_{V>h^r&Zx$ycJzfG#9=_gpfcqQhy zoe7gOY=w?wMEh?ARySYEaD>;}=O*8~(FqO?d&bJ8yW=Q&W z32;p8ZqWnXVUXbXB$#`&1;NSpwtl;5`UIl(^EZ2|4xe4y+`{4YLwgp8pb>ZEV_P_s z+l`A&?y9+!g^i7kWu;GBXQ>Bs>}+@sIWogRgez8ZlG^twl#rf0=hKV+{=o@JLrymD zlfq=Z4ws{1{^mayXZ&C^ok&D8?BHPR@ecx7ciXpbf9-Y4+pTzB?<7|WCWrXB<HuT~PM znX8`aF39CaaW?q^No~DFIdfCL*>*wcXf+m+|*zHJTlu-Uw~>kCy(Y@VpZ+?PKTHZ9A`YK|+*T;?I%+$hie z>T;V(`4?SpnE9fe zZ@88(3g4Ntz%N%$)6=3sMf{K+yvZppFkZy2+=mjqeT#wiMan^8uga9Mqk=-6RPG(S zs_&z#4NEvkQt@z&R%(z&j6pj!=r zJE@np&-HYqBb_Dok7CFPz3Pgd?OA#2$B+5>c@fn7B@L_7BO-jl?> zj*+T8A_uAE@-8L{)s;wi3Fgj7RCx*IvX@9CtC=${7XN{Kc7~#xgvV?J#LTj+O;^%8 zqb*i=iRE4=s&$Kh^fu1fD{&Pv(myZryRz8#tA{``TTCk(esbGlhvivEe;zpw0cIKqvCMp`REgWikQGAa6LRg zo-=1UksF0oFev(SJ9(kg4LYM>>)wdZ04?aPyGIP7PVUTfS=cG%H%_S)nDlA%Q}*7u zcq1Cd33(K8TnF?Yi+Yh7ni9FN)98~Q6knrgbKLLtN;BT^`1{M}G@>O!N9J?qNpU9c zXJvH6M62_uA>1hmi0X?**Ov@^R$*bilL_%#U@!JR;(!CQ0btv*bEx!QkY$iD{F#NUZ7(Tymiu zq#N{7tKpRd^Tm+m_X(@ek(TzoZuo*nHpAB-7|ucGLsrl`OSCg`g@xv9o7LI{tD%Z2 z^v)oW+fXT{pLbMRG<^g-Z%xE%f6mcPaYb>|t+#16bI{+|2B8~DTh1{BORRuXWsj2@ z&tfvS5%p(05$6#i2L2_tS!*Wd1QlJS;pA`a%UZc@?#|JsA?@Oem6uP18o4UoOG>p6 z@jSfbBZ1m}%IR@B;`k%MRYtS60+qA;OOs-pZJr5}JGzQ?wz1R9`}n$T#Z2Zsy-R|^ z&e+ejgsW`!ffgi)zdEUGA0Om1dG8;zn;xIa;~QG|#9H|d?pZ8d5amKa&FcFnurDD# zRvv9u!=NT)D&;Z@GW-P2n30p+VCXQ2vCs#&&E-T)SI6bVPn#IzB)h;cK=y6{?gSC4 zDY{ZHfh0Ah2e{A4w%CgEQ*u1QsXpl^Dz6;4>;?yP9XSq}W$zbz7a9NxCfvOSDe*>poVP-$vtt~pq+Ik64S6#FHJcfNU`v&% zy9pfzpt;CRVG!4l>%TYQBZH%a+D0G`#Y9(JP@=8wqX}K*sWUE}vJ8)b{b4#!3Xg7B zHdGuWfs~Bq5yzdeiNX5ghv?ncM9JQdlL#3jej?PL-Xj?a2|-7Y53}Kfzn}82mx_c+ zlvV7Q7%{QP?`(6#Unnyy7GK>1tXA~+Z>@XOH zfF%R631&ptbFjl59f?GQTtBnOaPKyS`?R;pCn`SK?IIWhUvR80I>KRQKgF3H6)itf zB^6i<$WJXi-nc+?I-loR}hHrU=!1t-fJuJ{L4ykU+1(f-NjWM}`!pFqh$0mIG;tV1HBfr7Ab@{wx4 z@e1HyP{ID!fWZIK3ywGOO#`w2HNtqwII>I;@{f9zss}zH6W1$ZE*5`D9I!E6_%a#tO9qD>$Oa?@{PQBU zB}M#q004v85)8O8Xy2gF|58|4+5amV@OB~Iq9JeR`@f<&SpUn%{@*@MFz~;sf`J^M z|DDST2C)1u9}58RMtS~OSNi{KfCT`2yMX`qn_*#l>)87*ALzff_!bA`NbuwlMP%dP NK%}A)l@~+&{{TJ5j1~X@ delta 12599 zcmai*1yCH#`tBjYJrE>FaJOYy+#P}iC%C(Nunh!vSllJJI|L6J++Bi0aCf--eL3gc z|NYLrb*HMQ-+B7!cY3y`YHNSp8{_lybBq}*FxET_i4A%H0Gz<;J6c#IHUK*Z87GLX zmi84IELdAMnyL}K{id?0)^8v83J(kJRACm@38^X>dU8VQ``nq#R&z!U7V?(Kw^EyP zX(%J}=Xee!y707!kIc#%VRY10Pg2Ji20 zpPJq~-Fn|fycBlmV+5I!9)qK&J81ZxFkT+OmiL=kF?W=xAr7OKAL!*+gH_XG#zgN; z&?c$IV<}E=x=Gv9e0^=>o8P22q&ClpGD+={Z~eU6idwjj8ON~iB#{Mi$~VM&yU?8O zkq3LD(I;^)Y;8^?6Og=1i{KNLVIvc|HFju@dF*a*bWN1lODfF5BsF<%v2*OTvgTg^ zhi?)*A|9zJIQD-LTly(F`r*65V9yEtkH^OV`}e~W4@kuynyRKd$86Dlcco+1X;*f6 z)e?NoEbg~~mkmj%4T9Wwh-UddV$iPPr(EBA4m%>mSzYw0@sY&cqNyPY36bc~yT09N zij%MO%W7dOLUdLxUNJ*&5MPV#8dD4&1*0Sww4?=1-HgwnHhYLD`Uq_IWG{*;xUSj^ zei_RdWfBrH#MIx_qqfns`nAi6aAgoBzi68Sd+3DbUGP2wm#ItLjWF>fr!pyS=q+4M zTc*>%NEtJ=b(#?x4n`)g3%LZn*;)i;D8Z2ESN_Jo?yoa#hl-AuXsns~BM_Gf+@(HB zsIV%gZ=YpiTuU!CFxvz%F=i&VJ)>rZpugRWo6LdJxB?buFkp-B*m}t|zty!pIi#S!44G zAhx)CuT2iE!WkP!&o~DDkk2@F3Y(dTtN^y8&oqA+1GYJnQ0ZqJlr_%e z{M9H%VxmG4=HJX)Jt=RdPs=-ugvFd!9j%idZ~P%<@F$R5YnZNU7Y9tcxQB(tExR71 zKxIy@m5M`xn9Pl#K{t}~WA{^>^tY;t2?lZ&rB4KTwQ|lBK02b+Gof;&n>G;zysCE_ z;tC5nTffB%W*q0~)q@$pinRV?QUL0#iYtnF<39Lhnfv%Ci~*MAM?dX*SN9e zX~5&2D7J|HXf#9IQ#9Z+zL7_!D6>0$K9ZFY7XUk{&{3q@_$8A5Wq9l4Os0gV)RzlY z-gANebS?sl$UfpA$N{0y{Q&5vM(Slx{5u0`)pkbm5tlHv=u?S~fw}4u0iU3*PsnJa z;%lULt#~1kqqm4{mV|3(-}SZ_iiFitEdTVJ!u} zItCH=psuXk6w9%&p<%^1B~8JZL*OSgEsXR-qAZ%q36!N2)_V!P{*r4QY?u1gr6{=w z&Q1bZ?m?dWnQTEm?wNonQ68Lj5_zD?tG8Ae^hqo%Bpo(AhyDHC!K6h9KejSI``q5n zA%6oU3PrObN{4{I@C{#1ZAMR3enj#W5O{Wh89Vw>^Fx%CGtn4tZf5X|Vv9G2Ow?{i z{0&3KIf!9w5QV)I3H7i*ZH=8Ve#DnCs3Q4FseV^pwx6pNABT4eoJZ3nkE6c&**mQ= z{Yq(qw*II7lN%$#V|MD~4l$afpAu_lVl=E|e}uX&LdgUej$~}i&40?bin=~`5c6pk zi%;P!EQ7I@>`74I9AhSdzSB(Sdet!Hv#%g43Kj>0t$7r#8l&4qt-;rt< zA->2a9h@;D%$OP5i(_L2Uuh} zIgCF&TsqZ)D^KB-{9WfK(wvnxzGM#WU#{~W!ddz5Om7+XQH@uF>_i1!vrmZWIu&&5 z=!NAGLFkZ~F1wBBh3wGRtn>yYW!lLR;Z)7L9KHh0EBjOHzeq-axs^ko{!22}c-M^}-u{wR#k5*)R-h0h2|BQF0U+Da;=CpS0WT3%4wZ zgS0hgk>Lr7CfnMiwdW}|-a}F{!%2Bs>#vW%?WkEl`xKkcRE}>~Rr#Q82dVtAk@j(= z-V}@-nt3+40SDWeST!WmjAB0;T{2m{X&6aMvl+%rwueHws<)XQ*5%uDSjAnrvfO|M zS(gM3)Fi_*j_ymg@Y=5mzbVTCxEt^A`pWkz`gBVJG^~1MKMd!98rLa1UyAD{e1Asa z=_rx)%+M@8muAx8&fJrF9oxcFHm9{0kB7Ch1ojjPAaQ@+*tR2)0Nsei0K6w`Q#%*yjDTdDr$ z;HE3j;A;IIBacr4U`N{>mhcr;D?jm=crDs|&l7}3B>9`?pQG1rxTbJkwUwvEl`wB>LQMc?&u-c|S~ z7=E`WJ)H1Dcgw9nyGD}ys%tl_62h2Yz9t?*8!%>M3~;ID&0d|==3visc>vkd1d;cY z`I{d&H(&=*wrj-jdXKJxsW9v@*UE$vk>zw>>jyBbgyG(rJ%1wbYZ!=z)Lo^0VSR$z z7EqU-&3aE*M(!wACN|oNbiJw9^T9~>P7;01Sa*KqeJRZ7GzDUMGs}WewfSx549(;) zVj^B)dmPUs*?5y*T+0B83}ICb`DAckFR(q%<+6w{+1l(Ac`5oDzxYR(C&D|h5JO+U zQ@fJm(fegQfS_U)<_p5PDEv25I}>LYCsRXP6li2`WQD>8AOn#76#%&a00=cYKd+df zor|H9sf`V@vZ=YNjiD0?tE7#gxicAv3>w+nxrmC|dyr`}v-7Z#F$1}{$=CouPUwk) z8xoC9#QIkNHHw=$8=Km>Kpoti&`U8x2WeAFa|@ThjZi_w#ne`vjDr(0ht3Yi$M=tA?B&ztdlqxT($rz&3Iod8DV%Qa%6Xjx6R_m{uC6jqS#3I5E!89S)a1F zR=K)uzZ-F3etf+;U&BoC=5W-~75C9$!_>mm(f3TZ#oGR>X)8YS#Ul#jhgN-~lUgr@ zu}3!j+kp$cthEOn-!Y&EZD2cXJ;fXoNkSR%Bgb7 z&f5ARdgCZy_@q~`7KiVLy{zN$BfSl)quX5rNz#y6^%Ft^W{a1WEn_CV>Kp8t2wUqU zTWwYun$jvClqJ!|@}cLhv9*=2w&|qT^y|q++`=BB$v$W^--T|`PjCbPhG=3Q;xbua zoPg%qxb%4`QN^wjf30(CiskL5o=(I7hfu8O6{&rK*0EMOxJMX6X+H8poI_d}G2YfTKg&+NUzlYu-_Q4tL1XpkhpZ(`vRVeF6%ypmL~qTG7OZ1j(+dT_!{H_xRizAqVI|o(z*!Xu%2z`Pnz_4L|Gh`%=!JFN7D%!uvEL=Ztd04 zhI%O(CEm%+IYV;QxGcZXztbg`_5v1$H$d9D{2w&T)vIUUo7`_SP897WJT#;)g)>J}JEY6TGobSSx`5N7-*<7_=+)3hL%Ad14 zhuK;LyRF)7rHdmNb%6~WyWCnJ$X+m(lvM;@747}2d9)CXXdC5=yo77`2FhT#)6I;| zt=Lqx_e~;=u8Xy^sg?QqTK&l@)cWmT%eL40^V;Gv*jRiF8a;Gs2;nlh-~lqZMIoZ+1<#<9Le`@>R3nz0ezxG~ltJ`h@iM zVn(<96~8q`v~>IZhPujetN*fCi`_Cz!RWqkudk=O4B6-UO1;L`J`L49?JR{HyFg6+ z`w7^fPf2Y~`+<38F~o_&H5}u39{aoQ)?5jZPZU@Kj)yFI3)=QD05IpRIxc#nDLKFY z-4IESMU27gXhg|h7^_b>X(u;2hEeA@Ld`@YrTeY!JAxZScZCB$8%->rYYjosBs>Fk zp|VyTHisv<2SNDW>(`mVSxI)wlRKYueI3lD&V%flTHk*w6l`v?u{v*F`CLxDh<&h8 z(oo*aLo_tl%!pINeFyd-Gf7~wMe-JW=W!L;$MWTVV&-Gqe9b|Bq*$c}_13pl%t<<2 z36MAv7!;F$f=xJ6tR70^q)niCKseB==CeL?f@zE{E z-m0BnYScO_F4J|ul)qviJOsgT+g)>lB=M}hHdw+h>)U^u7=deq< z^(WVxcOsiqKHH+TA81m04V2!QTBS%vetC;c7(T+W#1I$9fMDH5;^=9tu)iMLyHxPd z&?d+W`zf=76x0}1SxM+@0?TyskZY9pZIdYMw%`k;&lsh1+pbiY(r!$0CA!EpWHU!I zZ~_G4;ON!MQmX~e89rQsMy=9Yy4ZIN%Qo;Y34>%0F!<B5Lx3c zZ8}_?nqub#U#s7<>TMho3t`ZkMzv0#w~Bv>%b~q3%=W26ARE<8Y-efVd;Il#F?JoL z40FMFo6_sW*4AD=(OX2JoU&u1>t1ybVVA9A)X-NVJY8a>LP-UiEtMJi-OSD^n~lqW zq@KH@)mb*iKYjFavQL=?KkrIE?tIH5^jb${F(}Lf95V8m+%O`1-eEkw-llIz_0_E+ zJ2R#Zk--*LMb(U{iRrO!r&1f=<4G-PO=E0*Et@Fr2Q*lYdKf`GF@T6XuwIps1~pO% zE(e%gj~xae7)Br*K1BJ2pUyMeOv%yyt%jj^SG0YRfOK2w+sR%D32a%lk;7T0`|Q{2xF^H7Xiaz|Va(stpzh&LSaF6Ey@I~k?t!+O-u-&Dr*=dt zoF*t2zBF#)o1xbG1qcANvY>AjNpE1k#Asd*KJ>@_(gQ)M`|0fLUKN)i-lzGljFh;q zC4N?qWkB;CVQiN;M4CF8cVH}+-bKV$9iCnbjRu{MgnRem6h!UcC}wC#`e5s50!G2a1N2EriO=EVB(%6#{tp3<8_}mG*^A6#| zYu1Z5Grfy>NII@{_VUIxP^Lehv6V zKd?RquI#r!HP%p}Z346UuOG|60Bw~IU zcQmM}n-Uuvjrc>5^ya17sVY^E8yQ?6;FaQod9pn)M>ew|MCo;>`fIZ)q%*_0>WcVb z^$?#1hln%HAsYTpatpCK3_ed0uS^$#nurrw?DZ_|sGOqE?A%%(+MQG~L{TsRw&o4X z2>%F8{H39eeP&B<@2@s#g*uAe;EFe0-Q%(a%MQ_pbuPlVl&evL4Hezm$!_44Szzhh z2kO$srklgB@Gy)BW0#&|80BADZCR@n23(^OG>jbw-409Zj-7W#qTWV7*Q{OqrZZFD z-;+>$>7A&Lch?xt!_z34{VQQ2SOCNhR?38M%6}z319@bUzR#Tlrvwfxw zu#)WFZhzW{DS(!7G%AI8^w|I$W9v8E;4@Rd8O21x{=PUSaE!m_+dBgn;(^VR3WiQ1 zrS;)gY*F@X3}yZ2A2}LAy5#Z}BJ>hEkuzkd{R%g&wQ-!w!;?&+=;F}N(yxxQn=6oy zwF!9mPXv2}{MCH$+(T((9Z0R%RZi9Y-oGgivw3rlR!;8$GV|?i7>v&ct1U4&eV4n0 zneFz%$5gmM*CR|rCHT}tEzqy(<=qinTN0`Y*zjbmkx%)YG5u7ghel3D4;Pe+dED!90 z!K0-I^GAdZPi6zZDesG6FmgK&;4R&?zIg;d;*CM1kv-*Z&=g-by-RPe+?vqCl4Zdp z7~vCs_(&I}OU133_gl4gYK9y3`^}4%=og=7h8m`EKGQwe`_icA22`(rnC7wD(hUWD za9ovD{l&!A+jiaCQNVsOWMT^?Nde&wwuWSz-uvd%c#DLW=97F7xM{NoRWX)4qFIt9 z5N-YTo%fitt4T7Se3{KFp5P-(1qt;DN3$S;h47vsmYx$uX@;Xs8fhFZ1w9jlkH-|0 zhisESWCpl*zqyH+*^miQR>o=g33E<&f^>|&-G1suV~+-M%}Y1<3H*2t-Q@No+RWU1 zPDxu=xg~7>x;EB94-SprqsoYT9j6kv7DB5)0l!xmBhh}*%LgGQC!h86dY z!gLMeCzH&MvKUaxWI0+CsJP=iMgGbEX->?EJ7r8NXzMdosy=QUx6oMS)^FwMk+?S%zwCc+)7VTrj3TP!Z4ZL97Pufi+;u z;jjZIZa0?tF*k-7ndtTkn(Ai^XOIbeXriDZiVWsH_17LTogR_&p5Tt!I2p!OXE>wI zhpZ?Q^4BUd(OKK`EtxJ5PNRNYy>wd~ZI;KAPBz95y0`ZZt@6vo=;oKk*3|nC;RnlW znh8ORUR4?X;6;jOeQ!_ij?l`tH^<_(%ye@hUdqg>thx*2zkmPWDG&Kc{((jOG+`Dr zO(T}iElTh#!hdKDnIx*Q#JuKp zXni%eDPLDr=sv#b#O)*MWF2NOzF(t^&fDhOT-Dw(a(s4gvtnCe7i{3_+U$MJN!~m% zYp=g>1HNK&cc3e;FJv_R+|zh7bs+kXu*mzM6+oD^GGlv#^pm8bf<^n_gH+PUHeOPJ zg!nYCu3k=2;S%liW!(dU&HI)wBz795ZA;)3&|KiH5u+U3eph7Q;Yb0FSi$CfzjKYd*!m<71eDj$YitUPE&!J>B;R7*&lbvI_ z>Jo{M9TCdSk5alVgx}LP-W@Gzr{x0v7jns(JFJkQFGRVpVOa6H19Cw-~G>@jHn@J$NrftLv*(>#bw( zq<>kT0}hISfLvM5H_Ynz9~abiWu{Vx#p4hp;|BoP0e>pCzbg>{B!xN}4!WYr31OmH`M*39ChZrl|NW$* z_BJN}FDe8b%-_t1jg5orFWAHWZz{z8|Di&Ewf~0-q0)g~a{xgQQ92;_NYm@1j#{eF zT3chfKwYdrz(6*)S{CAffSj|T37jG$aj&uZFYn48cY zI9mZxma1iw?AtYPvvOULmM=~VEp=oj=}qCTHIn;Pv)1erD!#H~|vO#(Ml{Ktlc|c3#mD?cVM2@a~kgDog%z6Td*ikEF1#RWB&ljEJ^;K_O9FIN0w( zd{%9V>M$Sg77(jCQe(0uf3QiV&{I$b@Pu>>biJUGV&ee((8OMWi&^Fsuy@8Jwzfnl z>Nnyx@QAnU9ftdVpo=!aN%yP=?-A=5j$*$jB}3b%3=rtjldyh`#8!&?+ageFqTg(+ zV`1#FT4T4dHuY|(2_xSL-&86j%mO)7JsG!pXPy{b`}!;CLW-Oi`Z`X**BQ8<85mL zK+Os&oW)nO{eTLey#s0JW#tk0og#-8U{MSw55z5QFnKVg-HcI=+9=jjXADU#axEhD zdRJ!alVaP*^=&A{vd#C&M89i5)_v})rT5NRMQo~SJs`0G+zx<;V=o&#NDhoS%?>DY z{fUDPVWz}A``CJN)SZ^ezl=g|NG`$z(VTCdt*WL#_oLE&({@pVD8ug-PniGCQAv9m zPd5rTu(ZrGLqSXp{&>2W=M+V%Q!sV)+Xs31;=UleYX((P57++raQAri_ntmu*B4%4FC=vw-5kA9 z3<3n&%hcOG+g4U0t4pR`=I8a>?uzDaB z&ADv&u=?7@pU3s$btcD?HsK6P_xHZT;p@Jo4V=-@e7uy)GJa)^gbs%wdfAF;ItYyC zBK{~Fc&5C+EgwdyRGESjBr}m6B8#NWRcXp1Nhf9+auL+~>6*~0ko;4TehojFu&D$) zebWcibKrbALr2L3#gCPLD}a58UU3uxiHUtW+|8r#M?(aSu9@3zWzit>(8fBZ&J zBSl#+`Th7dEZ~OZAj-JXD$buZpijb&vjfBi)={dmCL3EY0sWRue!_LX#ABjUrwXWQ znlF=$j#as&VY^Vz##Fz(V37(ycg+I@$_5+AkCq=Z$*`2p!Rdd!{hpT>sEAp-AZxJ| zL4m`PjZ2sI-?!(?=n-7a*Zv(QrK~+S_XXF4jn-% zuPUnCL!07z05|6@yPRDxW%p!G&RLlPm{V5+;BBIMr>MW=Gj^bNLAnvaduP`~DxgES>7!7kDXYwYjwJKkuP6ex=F58+Tnh2} zIwm!~@X-?#aHi~=vfOG>q6@Di4c*pqS?yBNR0f}4GhLv8#b1Q)ZASb7V%qZM;Qnu{ z@{x^YZp6ncrlgZfgL#mzjCszJJXl6OeU$Qij2SBZgInyYMu%O z!UqJyzbZR$GApxX9R0Pq*{jq0eSkXW8DufdREof!M1_>)L+||xv7#mkPdUYXx4y~l zW@9%UVXoLA>_tKlCkKhRZ6r*s7mt}r#FlNQ)bxOmws!%*Hg$}^*vh$|T zLhN1RC!9A2%lPTAH}e8I?&f+BR* z_U;Ubzy~L#hUD8lb2T&py{dE-(^oyUlD@KNnolza1#?--_T!iWwktlQc`J{57X{!^FI!+mnHgffZN`jPkyMj~*} zkU@f~boSt?_!IJk$PPAGo?E!(v}ZXKM^0yEWy!U@D2D#C0TzB%V*RVpF=7AHvj_uv zl9NF{>G<>E4SZnqA$MEtmvjjGU~-YB0yQr03%V#8lXX;l-(hiK;210SVWsG>$_iC8Dt`ASKDv>o6dQf$ldtLIP_wCdq~?VM(N+X1+kN}Ql8?`ohrc@M z^GXkNsk<#z>uGDqb-O*1U~Y%=gv9;G@p$RRa$CX6%h`HYSmsGY3br%#a-dOkO65yq zqP45(wrg52rjP{K*=BcRw+6iil5CtJYhL9N!VD{H-I|UZ5{W%^GV4Whv7g%3tp=mk z?KQh~MAClLupX*&!D+rmPWt8K(d(%#`pSKJ`y9sG&PVEnlh(a&FO%29W-$#pX1~SZ zv9noZ(?<2f@zIaNv=fx1tTLm{xDLbH_T6Ql4qpEG_%u;)CH^H1LwI*c&pmFrsxgz+ z{k$Li>!Or{8cO|@z3mE;ZZ$-Fb5q(kN}i~yMG02DaYW9E7?Q>B4)o2|8A~TBZ*MZmsl8?^)UoX58KKSbvukUytSKnHKpVXvI>0eit2P zQ~e1PmYnY~CA+PzD$h?J@Z8_eSs@||Qz2Ex%~9P@D@~e=H}ZoY@Fvt+%MMIIGodxJ z*icQ*BU%s4GwXN$QMu!usM%-}sqg`f1B6$!(6Vy`Hk9FzFOS4sKvqqAB5?m%Bv0&? zsE^mz+fdTdx1s@kF$d(=hO1G$g&&q0w3Guk;+}b9iwbm#|vj_^Nya9!_L>W>| z8C6v%m9tyYGFxE_LEKR;^%P#>_Ss(IJ9*sco+AX|EMA>_C60SCXH0){!o0EIX_xsv z4?K@pgKocD>wyAy#~y@7P-SEM1MhiowfYnz&xCb^DMEf2fJoCbCiv1qdpIHWodTu3 zfV5f2232Fr*aCf=DV}h=QfeK=H~Jc@=9cuFEq3a>U>#{6k{?HJ=U-P__*4BD`0P0# z4kJa=Bva$}zz6(kDZ^(huDAOWj*>Pn9UNNcjA!&Kv&&U!iIczwQ9&jldT?RLXUoM)9GTcB$NL zX+;?qtVwYE9#(Ebf{Y7YXxs|pgyELVT_ZM$f;z%uns|Rqb|+ap5^*2GCU(9Gq4e7a z+0cTQ&!^wTYrm!XNaga>gfT)a-hn5(v6-drl|nH@@3W|R15`J8jtj+PQ0@m5gGIs* z@uBHzUMtABpoOB`rp9AqzN6SQ3E-9ya>OECuP(TZQGFcwp~8S;Z-NjDT!%#EmG{Vd zhLX5@IM&M_@|@WVjm}VTgj8bgdPyj>H5u_;hn=pztQ|S8)G)rz_o~VY15f(DNIl8j zu9oO?oWB>gT#@Dq=P!QCkaWCEGf3FOAyA~!AZu%Y!vC{S{hdH1`9Eps%YU zLsmjciuJ%CC}r|a-M!MM=b*OPE92O=eU@kMa>46vrmf-VI(K5*nFai!Dj97K;1IZd zc$7=Q?yy?wwBTAhtmRv|e&c(^x%(@-ys!v7zT*D1;sCklvN z%Mg3wXx?v^@tt!|k0r3%bl2=r&4sZ-*4M@}+}_C6ktF=X)4QPa2ub^~5a)s2k<;9Q z%vfJzG5B72mvluGK77G%(>vEWqbyXiSZ4U&6XHE|Zr$li7^p6bSuakX$Z?)`)L&5e z;iX4MYz9TuEga#k40kVX(_VdMKtN;)A`Ww#jV`_+F?Tu;JpT&r9PI(**t1NUAWsiR z3`giZ_j42Z2Gn+qHOnz@Oi#bj+3M}54&qr?YpOK@gzC72;{9Dy{ik4YAGgxpCEel2?dSWY=o!`6LgeO>owRxnbD=P zVPS3oc8{&Jv2bzzcCRI!1HfR1O)s~OoaR&cmk2* z4h3_90Aw7{eKOF6S29ijI~fSNgNB>ykH*0TML2Arzp)U@U;%7!@pup!Prq==LAZ$!i?`FF(y^*}A0;(w|? z_=r>D&lUUM74T1Yu773W`CH`r)5i5D8P}iGpnqin{cQoW|I5SuC!6S>O6*VnKOWE@ z3;Ul@{2O9%{cQpLsW|>M?0;qXBLe?i0sjmJ^sh0A2$JdkHC;w#PKI{o5FWm-82_0j z=D!#h1e2B?a>_RW2jYRUs9t_KVGtMeQw02VklB%;{2c&zpnm}Prv~8Sgl??*KYi>R z>>!Ae01o)SG61-s@56t&q4WDsjSGtW{-JRLp!57s4aCm*UmA2?|7d#uW9UFA9{fjV s9uO4V{ht~e`(J$W9~y`a`ZE1PW9NeI=n@f>Kmh=ODAd#vijpY*3)0qP4gdfE diff --git a/BookGPU/Chapters/chapter6/ch6.aux b/BookGPU/Chapters/chapter6/ch6.aux index ef830ac..2be0894 100644 --- a/BookGPU/Chapters/chapter6/ch6.aux +++ b/BookGPU/Chapters/chapter6/ch6.aux @@ -3,110 +3,110 @@ \@writefile{toc}{\author{Stephane Vialle}{}} \@writefile{toc}{\author{Jens Gustedt}{}} \@writefile{loa}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {chapter}{\numberline {6}Development methodologies for GPU and cluster of GPUs}{81}} +\@writefile{toc}{\contentsline {chapter}{\numberline {6}Development methodologies for GPU and cluster of GPUs}{83}} \@writefile{lof}{\addvspace {10\p@ }} \@writefile{lot}{\addvspace {10\p@ }} -\@writefile{toc}{\contentsline {section}{\numberline {6.1}Introduction}{82}} -\newlabel{ch6:intro}{{6.1}{82}} -\@writefile{toc}{\contentsline {section}{\numberline {6.2}General scheme of synchronous code with computation/communication overlapping in GPU clusters}{82}} -\newlabel{ch6:part1}{{6.2}{82}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.1}Synchronous parallel algorithms on GPU clusters}{82}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.1}{\ignorespaces Native overlap of internode CPU communications with GPU computations.\relax }}{84}} -\newlabel{fig:ch6p1overlapnative}{{6.1}{84}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.2}Native overlap of CPU communications and GPU computations}{84}} -\newlabel{algo:ch6p1overlapnative}{{6.1}{85}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.1}Generic scheme implicitly overlapping MPI communications with CUDA GPU computations}{85}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.2}{\ignorespaces Overlap of internode CPU communications with a sequence of CPU/GPU data transfers and GPU computations.\relax }}{86}} -\newlabel{fig:ch6p1overlapseqsequence}{{6.2}{86}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.3}Overlapping with sequences of transfers and computations}{86}} -\newlabel{algo:ch6p1overlapseqsequence}{{6.2}{87}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.2}Generic scheme explicitly overlapping MPI communications with sequences of CUDA CPU/GPU transfers and CUDA GPU computations}{87}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.3}{\ignorespaces Overlap of internode CPU communications with a streamed sequence of CPU/GPU data transfers and GPU computations.\relax }}{88}} -\newlabel{fig:ch6p1overlapstreamsequence}{{6.3}{88}} -\newlabel{algo:ch6p1overlapstreamsequence}{{6.3}{89}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.3}Generic scheme explicitly overlapping MPI communications with streamed sequences of CUDA CPU/GPU transfers and CUDA GPU computations}{89}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.4}{\ignorespaces Complete overlap of internode CPU communications, CPU/GPU data transfers and GPU computations, interleaving computation-communication iterations\relax }}{91}} -\newlabel{fig:ch6p1overlapinterleaved}{{6.4}{91}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.4}Interleaved communications-transfers-computations overlapping}{91}} -\newlabel{algo:ch6p1overlapinterleaved}{{6.4}{92}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.4}Generic scheme explicitly overlapping MPI communications, CUDA CPU/GPU transfers and CUDA GPU computations, interleaving computation-communication iterations}{92}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.5}Experimental validation}{94}} -\newlabel{ch6:p1expes}{{6.2.5}{94}} -\newlabel{ch6:p1block-cyclic}{{6.2.5}{94}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.5}{\ignorespaces Experimental performances of different synchronous algorithms computing a dense matrix product\relax }}{95}} -\newlabel{fig:ch6p1syncexpematrixprod}{{6.5}{95}} -\@writefile{toc}{\contentsline {section}{\numberline {6.3}General scheme of asynchronous parallel code with computation/communication overlapping}{96}} -\newlabel{ch6:part2}{{6.3}{96}} -\@writefile{loa}{\contentsline {algocf}{\numberline {3}{\ignorespaces Synchronous iterative scheme\relax }}{96}} -\newlabel{algo:ch6p2sync}{{3}{96}} -\@writefile{loa}{\contentsline {algocf}{\numberline {4}{\ignorespaces Asynchronous iterative scheme\relax }}{96}} -\newlabel{algo:ch6p2async}{{4}{96}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.1}A basic asynchronous scheme}{97}} -\newlabel{ch6:p2BasicAsync}{{6.3.1}{97}} -\newlabel{algo:ch6p2BasicAsync}{{6.5}{98}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.5}Initialization of the basic asynchronous scheme}{98}} -\newlabel{algo:ch6p2BasicAsyncComp}{{6.6}{99}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.6}Computing function in the basic asynchronous scheme}{99}} -\newlabel{algo:ch6p2BasicAsyncSendings}{{6.7}{100}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.7}Sending function in the basic asynchronous scheme}{100}} -\newlabel{algo:ch6p2BasicAsyncReceptions}{{6.8}{101}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.8}Reception function in the basic asynchronous scheme}{101}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.2}Synchronization of the asynchronous scheme}{102}} -\newlabel{ch6:p2SsyncOverAsync}{{6.3.2}{102}} -\newlabel{algo:ch6p2Sync}{{6.9}{103}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.9}Initialization of the synchronized scheme}{103}} -\newlabel{algo:ch6p2SyncComp}{{6.10}{104}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.10}Computing function in the synchronized scheme}{104}} -\newlabel{algo:ch6p2SyncReceptions}{{6.11}{105}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.11}Reception function in the synchronized scheme}{105}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.3}Asynchronous scheme using MPI, OpenMP and CUDA}{106}} -\newlabel{ch6:p2GPUAsync}{{6.3.3}{106}} -\newlabel{algo:ch6p2AsyncSyncComp}{{6.12}{107}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.12}Computing function in the final asynchronous scheme}{107}} -\newlabel{algo:ch6p2syncGPU}{{6.13}{109}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.13}Computing function in the final asynchronous scheme}{109}} -\newlabel{algo:ch6p2FullOverAsyncMain}{{6.14}{111}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.14}Initialization of the main process of complete overlap with asynchronism}{111}} -\newlabel{algo:ch6p2FullOverAsyncComp1}{{6.15}{112}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.15}Computing function in the final asynchronous scheme with CPU/GPU overlap}{112}} -\newlabel{algo:ch6p2FullOverAsyncComp2}{{6.16}{113}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.16}Auxiliary computing function in the final asynchronous scheme with CPU/GPU overlap}{113}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.4}Experimental validation}{114}} -\newlabel{sec:ch6p2expes}{{6.3.4}{114}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.6}{\ignorespaces Computation times of the test application in synchronous and asynchronous modes.\relax }}{115}} -\newlabel{fig:ch6p2syncasync}{{6.6}{115}} -\@writefile{lof}{\contentsline {figure}{\numberline {6.7}{\ignorespaces Computation times with or without overlap of Jacobian updatings in asynchronous mode.\relax }}{116}} -\newlabel{fig:ch6p2aux}{{6.7}{116}} -\@writefile{toc}{\contentsline {section}{\numberline {6.4}Perspective: A unifying programming model}{117}} -\newlabel{sec:ch6p3unify}{{6.4}{117}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.1}Resources}{117}} -\newlabel{sec:ch6p3resources}{{6.4.1}{117}} -\newlabel{algo:ch6p3ORWLresources}{{6.17}{118}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.17}Declaration of ORWL resources for a block-cyclic matrix multiplication}{118}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.2}Control}{118}} -\newlabel{sec:ch6p3ORWLcontrol}{{6.4.2}{118}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.3}Example: block-cyclic matrix multiplication (MM)}{119}} -\newlabel{sec:ch6p3ORWLMM}{{6.4.3}{119}} -\newlabel{algo:ch6p3ORWLBCCMM}{{6.18}{119}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.18}Block-cyclic matrix multiplication, high level per task view}{119}} -\newlabel{algo:ch6p3ORWLlcopy}{{6.19}{120}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.19}An iterative local copy operation}{120}} -\newlabel{algo:ch6p3ORWLrcopy}{{6.20}{120}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.20}An iterative remote copy operation as part of a block cyclic matrix multiplication task}{120}} -\newlabel{algo:ch6p3ORWLtrans}{{6.21}{120}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.21}An iterative GPU transfer and compute operation as part of a block cyclic matrix multiplication task}{120}} -\newlabel{algo:ch6p3ORWLdecl}{{6.22}{121}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.22}Dynamic declaration of handles to represent the resources}{121}} -\newlabel{algo:ch6p3ORWLinit}{{6.23}{122}} -\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.23}Dynamic initialization of access mode and priorities}{122}} -\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.4}Tasks and operations}{122}} -\newlabel{sec:ch6p3tasks}{{6.4.4}{122}} -\@writefile{toc}{\contentsline {section}{\numberline {6.5}Conclusion}{123}} -\newlabel{ch6:conclu}{{6.5}{123}} -\@writefile{toc}{\contentsline {section}{\numberline {6.6}Glossary}{123}} -\@writefile{toc}{\contentsline {section}{Bibliography}{124}} +\@writefile{toc}{\contentsline {section}{\numberline {6.1}Introduction}{84}} +\newlabel{ch6:intro}{{6.1}{84}} +\@writefile{toc}{\contentsline {section}{\numberline {6.2}General scheme of synchronous code with computation/communication overlapping in GPU clusters}{84}} +\newlabel{ch6:part1}{{6.2}{84}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.1}Synchronous parallel algorithms on GPU clusters}{84}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.1}{\ignorespaces Native overlap of internode CPU communications with GPU computations.\relax }}{86}} +\newlabel{fig:ch6p1overlapnative}{{6.1}{86}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.2}Native overlap of CPU communications and GPU computations}{86}} +\newlabel{algo:ch6p1overlapnative}{{6.1}{87}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.1}Generic scheme implicitly overlapping MPI communications with CUDA GPU computations}{87}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.2}{\ignorespaces Overlap of internode CPU communications with a sequence of CPU/GPU data transfers and GPU computations.\relax }}{88}} +\newlabel{fig:ch6p1overlapseqsequence}{{6.2}{88}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.3}Overlapping with sequences of transfers and computations}{88}} +\newlabel{algo:ch6p1overlapseqsequence}{{6.2}{89}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.2}Generic scheme explicitly overlapping MPI communications with sequences of CUDA CPU/GPU transfers and CUDA GPU computations}{89}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.3}{\ignorespaces Overlap of internode CPU communications with a streamed sequence of CPU/GPU data transfers and GPU computations.\relax }}{90}} +\newlabel{fig:ch6p1overlapstreamsequence}{{6.3}{90}} +\newlabel{algo:ch6p1overlapstreamsequence}{{6.3}{91}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.3}Generic scheme explicitly overlapping MPI communications with streamed sequences of CUDA CPU/GPU transfers and CUDA GPU computations}{91}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.4}{\ignorespaces Complete overlap of internode CPU communications, CPU/GPU data transfers and GPU computations, interleaving computation-communication iterations\relax }}{93}} +\newlabel{fig:ch6p1overlapinterleaved}{{6.4}{93}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.4}Interleaved communications-transfers-computations overlapping}{93}} +\newlabel{algo:ch6p1overlapinterleaved}{{6.4}{94}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.4}Generic scheme explicitly overlapping MPI communications, CUDA CPU/GPU transfers and CUDA GPU computations, interleaving computation-communication iterations}{94}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.2.5}Experimental validation}{96}} +\newlabel{ch6:p1expes}{{6.2.5}{96}} +\newlabel{ch6:p1block-cyclic}{{6.2.5}{96}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.5}{\ignorespaces Experimental performances of different synchronous algorithms computing a dense matrix product\relax }}{97}} +\newlabel{fig:ch6p1syncexpematrixprod}{{6.5}{97}} +\@writefile{toc}{\contentsline {section}{\numberline {6.3}General scheme of asynchronous parallel code with computation/communication overlapping}{98}} +\newlabel{ch6:part2}{{6.3}{98}} +\@writefile{loa}{\contentsline {algocf}{\numberline {3}{\ignorespaces Synchronous iterative scheme\relax }}{98}} +\newlabel{algo:ch6p2sync}{{3}{98}} +\@writefile{loa}{\contentsline {algocf}{\numberline {4}{\ignorespaces Asynchronous iterative scheme\relax }}{98}} +\newlabel{algo:ch6p2async}{{4}{98}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.1}A basic asynchronous scheme}{99}} +\newlabel{ch6:p2BasicAsync}{{6.3.1}{99}} +\newlabel{algo:ch6p2BasicAsync}{{6.5}{100}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.5}Initialization of the basic asynchronous scheme}{100}} +\newlabel{algo:ch6p2BasicAsyncComp}{{6.6}{101}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.6}Computing function in the basic asynchronous scheme}{101}} +\newlabel{algo:ch6p2BasicAsyncSendings}{{6.7}{102}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.7}Sending function in the basic asynchronous scheme}{102}} +\newlabel{algo:ch6p2BasicAsyncReceptions}{{6.8}{103}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.8}Reception function in the basic asynchronous scheme}{103}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.2}Synchronization of the asynchronous scheme}{104}} +\newlabel{ch6:p2SsyncOverAsync}{{6.3.2}{104}} +\newlabel{algo:ch6p2Sync}{{6.9}{105}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.9}Initialization of the synchronized scheme}{105}} +\newlabel{algo:ch6p2SyncComp}{{6.10}{106}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.10}Computing function in the synchronized scheme}{106}} +\newlabel{algo:ch6p2SyncReceptions}{{6.11}{107}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.11}Reception function in the synchronized scheme}{107}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.3}Asynchronous scheme using MPI, OpenMP and CUDA}{108}} +\newlabel{ch6:p2GPUAsync}{{6.3.3}{108}} +\newlabel{algo:ch6p2AsyncSyncComp}{{6.12}{109}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.12}Computing function in the final asynchronous scheme}{109}} +\newlabel{algo:ch6p2syncGPU}{{6.13}{111}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.13}Computing function in the final asynchronous scheme}{111}} +\newlabel{algo:ch6p2FullOverAsyncMain}{{6.14}{113}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.14}Initialization of the main process of complete overlap with asynchronism}{113}} +\newlabel{algo:ch6p2FullOverAsyncComp1}{{6.15}{114}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.15}Computing function in the final asynchronous scheme with CPU/GPU overlap}{114}} +\newlabel{algo:ch6p2FullOverAsyncComp2}{{6.16}{115}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.16}Auxiliary computing function in the final asynchronous scheme with CPU/GPU overlap}{115}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.3.4}Experimental validation}{116}} +\newlabel{sec:ch6p2expes}{{6.3.4}{116}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.6}{\ignorespaces Computation times of the test application in synchronous and asynchronous modes.\relax }}{117}} +\newlabel{fig:ch6p2syncasync}{{6.6}{117}} +\@writefile{lof}{\contentsline {figure}{\numberline {6.7}{\ignorespaces Computation times with or without overlap of Jacobian updatings in asynchronous mode.\relax }}{118}} +\newlabel{fig:ch6p2aux}{{6.7}{118}} +\@writefile{toc}{\contentsline {section}{\numberline {6.4}Perspective: A unifying programming model}{119}} +\newlabel{sec:ch6p3unify}{{6.4}{119}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.1}Resources}{119}} +\newlabel{sec:ch6p3resources}{{6.4.1}{119}} +\newlabel{algo:ch6p3ORWLresources}{{6.17}{120}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.17}Declaration of ORWL resources for a block-cyclic matrix multiplication}{120}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.2}Control}{120}} +\newlabel{sec:ch6p3ORWLcontrol}{{6.4.2}{120}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.3}Example: block-cyclic matrix multiplication (MM)}{121}} +\newlabel{sec:ch6p3ORWLMM}{{6.4.3}{121}} +\newlabel{algo:ch6p3ORWLBCCMM}{{6.18}{121}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.18}Block-cyclic matrix multiplication, high level per task view}{121}} +\newlabel{algo:ch6p3ORWLlcopy}{{6.19}{122}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.19}An iterative local copy operation}{122}} +\newlabel{algo:ch6p3ORWLrcopy}{{6.20}{122}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.20}An iterative remote copy operation as part of a block cyclic matrix multiplication task}{122}} +\newlabel{algo:ch6p3ORWLtrans}{{6.21}{122}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.21}An iterative GPU transfer and compute operation as part of a block cyclic matrix multiplication task}{122}} +\newlabel{algo:ch6p3ORWLdecl}{{6.22}{123}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.22}Dynamic declaration of handles to represent the resources}{123}} +\newlabel{algo:ch6p3ORWLinit}{{6.23}{124}} +\@writefile{lol}{\contentsline {lstlisting}{\numberline {6.23}Dynamic initialization of access mode and priorities}{124}} +\@writefile{toc}{\contentsline {subsection}{\numberline {6.4.4}Tasks and operations}{124}} +\newlabel{sec:ch6p3tasks}{{6.4.4}{124}} +\@writefile{toc}{\contentsline {section}{\numberline {6.5}Conclusion}{125}} +\newlabel{ch6:conclu}{{6.5}{125}} +\@writefile{toc}{\contentsline {section}{\numberline {6.6}Glossary}{125}} +\@writefile{toc}{\contentsline {section}{Bibliography}{126}} \@setckpt{Chapters/chapter6/ch6}{ -\setcounter{page}{126} +\setcounter{page}{128} \setcounter{equation}{0} \setcounter{enumi}{4} \setcounter{enumii}{0} -- 2.39.5