<html><head><meta http-equiv="Content-Type" content="text/html; charset=utf-8"></head><body style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><br class=""><div><br class=""><blockquote type="cite" class=""><div class="">On Jan 23, 2022, at 10:47 PM, Jacob Faibussowitsch <<a href="mailto:jacob.fai@gmail.com" class="">jacob.fai@gmail.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><meta http-equiv="Content-Type" content="text/html; charset=utf-8" class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><blockquote type="cite" class=""><div class="" style="font-family: Menlo-Regular;">The outer LogEventBegin/End captures the entire time, including copies, kernel launches etc.</div></blockquote><div class=""><br class=""></div>Not if the GPU call is asynchronous. To time the call the stream must also be synchronized with the host. The only way to truly time only the kernel calls themselves is to wrap the actual call itself:<div class=""><br class=""></div><div class="">```</div><div class="">cublasXaxpy_petsc(…)</div><div class="">{</div><div class=""> PetscLogGpuTimeBegin();</div><div class=""> cublasXaxpy(…);</div><div class=""> PetscLogGpuTimeEnd();</div><div class="">}</div><div class="">```</div></div></div></blockquote><div><br class=""></div> Indeed, they are wrapped as above.</div><div><br class=""><blockquote type="cite" class=""><div class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><div class=""><br class=""></div><div class="">Note that</div><div class=""><br class=""></div><div class="">```</div><div class="">#define cublasXaxpy_petsc(…) <span style="caret-color: rgb(0, 0, 0);" class="">PetscLogGpuTimeBegin();</span><span style="caret-color: rgb(0, 0, 0);" class="">cublasXaxpy(…);</span><span style="caret-color: rgb(0, 0, 0);" class="">PetscLogGpuTimeEnd();</span></div><div class="">```</div><div class=""><br class=""></div><div class="">Is not sufficient, as this would still include transfers if those transfers happen as direct arguments to the function:</div><div class=""><br class=""></div><div class="">```</div><div class="">cublasXaxpy_petsc(RAII_xfer_to_device(),…);</div></div></div></blockquote><div><br class=""></div><div><br class=""></div> I am not sure what you mean here? RAII_xfer_to_device()? Do you mean unified memory transfers down? I don't think we use those.</div><div><br class=""></div><div> The PetscLogGpuTimeBegin()/End was written by Hong so it works with events to get a GPU timing, it is not suppose to include the CPU kernel launch times or the time to move the scalar arguments to the GPU. It may not be perfect but it is the best we can do to capture the time the GPU is actively doing the numerics, which is what we want.</div><div><br class=""></div><div><br class=""><blockquote type="cite" class=""><div class=""><div style="word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><div class="">```</div><div class=""><br class=""><div class=""><div class="">
<div dir="auto" style="caret-color: rgb(0, 0, 0); letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><div dir="auto" style="caret-color: rgb(0, 0, 0); letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; word-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;" class=""><div class="">Best regards,<br class=""><br class="">Jacob Faibussowitsch<br class="">(Jacob Fai - booss - oh - vitch)<br class=""></div></div></div>
</div>
<div class=""><br class=""><blockquote type="cite" class=""><div class="">On Jan 23, 2022, at 21:37, Barry Smith <<a href="mailto:bsmith@petsc.dev" class="">bsmith@petsc.dev</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><meta charset="UTF-8" class=""><div style="caret-color: rgb(0, 0, 0); font-family: Menlo-Regular; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><br class="Apple-interchange-newline"><br class=""><blockquote type="cite" class=""><div class="">On Jan 23, 2022, at 10:01 PM, Junchao Zhang <<a href="mailto:junchao.zhang@gmail.com" class="">junchao.zhang@gmail.com</a>> wrote:</div><br class="Apple-interchange-newline"><div class=""><div dir="ltr" class=""><div dir="ltr" class=""><br class=""></div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Jan 22, 2022 at 9:00 PM Junchao Zhang <<a href="mailto:junchao.zhang@gmail.com" class="">junchao.zhang@gmail.com</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div dir="ltr" class=""><div dir="ltr" class=""><br class=""><br class=""></div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Jan 22, 2022 at 5:00 PM Barry Smith <<a href="mailto:bsmith@petsc.dev" target="_blank" class="">bsmith@petsc.dev</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div class=""><div class=""><br class=""></div> The GPU flop rate (when 100 percent flops on the GPU) should always be higher than the overall flop rate (the previous column). For large problems they should be similar, for small problems the GPU one may be much higher.<div class=""><br class=""></div><div class=""> If the CPU one is higher (when 100 percent flops on the GPU) something must be wrong with the logging. I looked at the code for the two cases and didn't see anything obvious.</div><div class=""><br class=""></div><div class=""> Junchao and Jacob,</div><div class=""> I think some of the timing code in the Kokkos interface is wrong. </div><div class=""><br class=""></div><div class=""> * The PetscLogGpuTimeBegin/End should be inside the viewer access code not outside it. (The GPU time is an attempt to best time the kernels, not other processing around the use of the kernels, that other stuff is captured in the general LogEventBegin/End.</div></div></blockquote></div></div></blockquote><div class="">What about potential host to device memory copy before calling a kernel? Should we count it in the kernel time?</div></div></div></div></blockquote><div class=""><br class=""></div> Nope, absolutely not. The GPU time represents the time the GPU is doing active work. The outer LogEventBegin/End captures the entire time, including copies, kernel launches etc. No reason to put the copy time in the GPU time because then there would be no need for the GPU since it would be the LogEventBegin/End. The LogEventBegin/End minus the GPU time represents any overhead from transfers.</div><div style="caret-color: rgb(0, 0, 0); font-family: Menlo-Regular; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><br class=""></div><div style="caret-color: rgb(0, 0, 0); font-family: Menlo-Regular; font-size: 12px; font-style: normal; font-variant-caps: normal; font-weight: normal; letter-spacing: normal; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none;" class=""><br class=""><blockquote type="cite" class=""><div class=""><div dir="ltr" class=""><div class="gmail_quote"><div class=""><br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div dir="ltr" class=""><div class="gmail_quote"><div class="">Good point </div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div class=""><div class=""> * The use of WaitForKokkos() is confusing and seems inconsistent. </div></div></blockquote><div class="">I need to have a look. Until now, I have not paid much attention to kokkos profiling.</div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div class=""><div class=""> -For example it is used in VecTDot_SeqKokkos() which I would think has a barrier anyways because it puts a scalar result into update? </div><div class=""> -Plus PetscLogGpuTimeBegin/End is suppose to already have suitable system (that Hong added) to ensure the kernel is complete; reading the manual page and looking at Jacobs cupmcontext.hpp it seems to be there so I don't think WaitForKokkos() is needed in most places (or is Kokkos asynchronous and needs this for correctness?) </div><div class="">But these won't explain the strange result of overall flop rate being higher than GPU flop rate.</div><div class=""><br class=""></div><div class=""> Barry</div><div class=""><br class=""></div><div class=""><br class=""></div><div class=""><br class=""></div><div class=""><br class=""><div class=""><br class=""><blockquote type="cite" class=""><div class="">On Jan 22, 2022, at 11:44 AM, Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank" class="">mfadams@lbl.gov</a>> wrote:</div><br class=""><div class=""><div dir="ltr" class="">I am getting some funny timings and I'm trying to figure it out. <div class="">I figure the gPU flop rates are bit higher because the timers are inside of the CPU timers, but<span class="Apple-converted-space"> </span><b class="">some are a lot bigger or inverted</b> <div class=""><font face="monospace" class=""><br class=""></font></div><div class=""><font face="monospace" class="">--- Event Stage 2: KSP Solve only<br class=""><br class="">MatMult 400 1.0 1.0094e+01 1.2 1.07e+11 1.0 3.7e+05 6.1e+04 0.0e+00 2 55 62 54 0 68 91100100 0 671849 857147 0 0.00e+00 0 0.00e+00 100<br class="">MatView 2 1.0 4.5257e-03 2.5 0.00e+00 0.0 0.0e+00 0.0e+00 2.0e+00 0 0 0 0 0 0 0 0 0 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">KSPSolve 2 1.0 1.4591e+01 1.1 1.18e+11 1.0 3.7e+05 6.1e+04 1.2e+03 2 60 62 54 60 100100100100100 512399 804048 0 0.00e+00 0 0.00e+00 100<br class="">SFPack 400 1.0 2.4545e-03 1.4 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 0 0 0 0 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">SFUnpack 400 1.0 9.4637e-05 1.7 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 0 0 0 0 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">VecTDot 802 1.0 3.0577e+00 2.1 3.36e+09 1.0 0.0e+00 0.0e+00 8.0e+02 0 2 0 0 40 13 3 0 0 67<span class="Apple-converted-space"> </span><b class="">69996 488328</b><span class="Apple-converted-space"> </span> 0 0.00e+00 0 0.00e+00 100<br class="">VecNorm 402 1.0 1.9597e+00 3.4 1.69e+09 1.0 0.0e+00 0.0e+00 4.0e+02 0 1 0 0 20 6 1 0 0 33 54744 571507 0 0.00e+00 0 0.00e+00 100<br class="">VecCopy 4 1.0 1.7143e-0228.6 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 0 0 0 0 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">VecSet 4 1.0 3.8051e-0316.9 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 0 0 0 0 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">VecAXPY 800 1.0 8.6160e-0113.6 3.36e+09 1.0 0.0e+00 0.0e+00 0.0e+00 0 2 0 0 0 6 3 0 0 0<span class="Apple-converted-space"> </span><b class="">247787 448304</b><span class="Apple-converted-space"> </span> 0 0.00e+00 0 0.00e+00 100<br class="">VecAYPX 398 1.0 1.6831e+0031.1 1.67e+09 1.0 0.0e+00 0.0e+00 0.0e+00 0 1 0 0 0 5 1 0 0 0 63107 77030 0 0.00e+00 0 0.00e+00 100<br class="">VecPointwiseMult 402 1.0 3.8729e-01 9.3 8.43e+08 1.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 2 1 0 0 0 138502 262413 0 0.00e+00 0 0.00e+00 100<br class="">VecScatterBegin 400 1.0 1.1947e+0035.1 0.00e+00 0.0 3.7e+05 6.1e+04 0.0e+00 0 0 62 54 0 5 0100100 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">VecScatterEnd 400 1.0 6.2969e+00 8.8 0.00e+00 0.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 10 0 0 0 0 0 0 0 0.00e+00 0 0.00e+00 0<br class="">PCApply 402 1.0 3.8758e-01 9.3 8.43e+08 1.0 0.0e+00 0.0e+00 0.0e+00 0 0 0 0 0 2 1 0 0 0 138396 262413 0 0.00e+00 0 0.00e+00 100<br class="">---------------------------------------------------------------------------------------------------------------------------------------------------------------<br class=""></font></div><div class=""><br class=""></div></div></div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Jan 22, 2022 at 11:10 AM Junchao Zhang <<a href="mailto:junchao.zhang@gmail.com" target="_blank" class="">junchao.zhang@gmail.com</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div dir="ltr" class=""><div dir="ltr" class=""><br class=""><br class=""></div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Sat, Jan 22, 2022 at 10:04 AM Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank" class="">mfadams@lbl.gov</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div dir="ltr" class="">Logging GPU flops should be inside of PetscLogGpuTimeBegin()/End() right?</div></blockquote><div class="">No, PetscLogGpuTime() does not know the flops of the caller.</div><div class=""> </div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Fri, Jan 21, 2022 at 9:47 PM Barry Smith <<a href="mailto:bsmith@petsc.dev" target="_blank" class="">bsmith@petsc.dev</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div class=""><div class=""><br class=""></div><div class=""> Mark,</div><div class=""><br class=""></div> Fix the logging before you run more. It will help with seeing the true disparity between the MatMult and the vector ops.<div class=""><br class=""></div><div class=""><br class=""><div class=""><blockquote type="cite" class=""><div class="">On Jan 21, 2022, at 9:37 PM, Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank" class="">mfadams@lbl.gov</a>> wrote:</div><br class=""><div class=""><div dir="ltr" class="">Here is one with 2M / GPU. Getting better.</div><br class=""><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Fri, Jan 21, 2022 at 9:17 PM Barry Smith <<a href="mailto:bsmith@petsc.dev" target="_blank" class="">bsmith@petsc.dev</a>> wrote:<br class=""></div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div class=""><div class=""><br class=""></div> Matt is correct, vectors are way too small.<div class=""><br class=""></div><div class=""> BTW: Now would be a good time to run some of the Report I benchmarks on Crusher to get a feel for the kernel launch times and performance on VecOps.</div><div class=""><br class=""></div><div class=""> Also Report 2.</div><div class=""><br class=""></div><div class=""> Barry</div><div class=""><br class=""><div class=""><br class=""><blockquote type="cite" class=""><div class="">On Jan 21, 2022, at 7:58 PM, Matthew Knepley <<a href="mailto:knepley@gmail.com" target="_blank" class="">knepley@gmail.com</a>> wrote:</div><br class=""><div class=""><div dir="ltr" class=""><div dir="ltr" class="">On Fri, Jan 21, 2022 at 6:41 PM Mark Adams <<a href="mailto:mfadams@lbl.gov" target="_blank" class="">mfadams@lbl.gov</a>> wrote:<br class=""></div><div class="gmail_quote"><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div dir="ltr" class="">I am looking at performance of a CG/Jacobi solve on a 3D Q2 Laplacian (ex13) on one Crusher node (8 GPUs on 4 GPU sockets, MI250X or is it MI200?).<div class="">This is with a 16M equation problem. GPU-aware MPI and non GPU-aware MPI are similar (mat-vec is a little faster w/o, the total is about the same, call it noise)<br class=""><div class=""><br class=""></div><div class="">I found that MatMult was about 3x faster using 8 cores/GPU, that is all 64 cores on the node, then when using 1 core/GPU. With the same size problem of course.</div><div class="">I was thinking MatMult should be faster with just one MPI process. Oh well, worry about that later.</div><div class=""><br class=""></div><div class="">The bigger problem, and I have observed this to some extent with the Landau TS/SNES/GPU-solver on the V/A100s, is that the vector operations are expensive or crazy expensive.</div>You can see (attached) and the times here that the solve is dominated by not-mat-vec:</div><div class=""><br class=""><div class=""><span class="" style="font-family: monospace;">------------------------------------------------------------------------------------------------------------------------</span><br class=""></div><div class=""><font face="monospace" class="">Event Count Time (sec) Flop --- Global --- --- Stage ---- <b class="">Total GPU<span class="Apple-converted-space"> </span></b> - CpuToGpu - - GpuToCpu - GPU<br class=""> Max Ratio Max Ratio Max Ratio Mess AvgLen Reduct %T %F %M %L %R %T %F %M %L %R<span class="Apple-converted-space"> </span><b class="">Mflop/s Mflop/s</b><span class="Apple-converted-space"> </span>Count Size Count Size %F<br class="">---------------------------------------------------------------------------------------------------------------------------------------------------------------<br class=""></font></div><div class=""><font face="monospace" class="">17:15 main= /gpfs/alpine/csc314/scratch/adams/petsc/src/snes/tests/data$ grep "MatMult 400" jac_out_00*5_8_gpuawaremp*<br class="">MatMult 400 1.0<span class="Apple-converted-space"> </span><b class="">1.2507e+00</b><span class="Apple-converted-space"> </span>1.3 1.34e+10 1.1 3.7e+05 1.6e+04 0.0e+00 1 55 62 54 0 27 91100100 0<span class="Apple-converted-space"> </span><b class="">668874 0</b><span class="Apple-converted-space"> </span> 0 0.00e+00 0 0.00e+00 100<br class="">17:15 main= /gpfs/alpine/csc314/scratch/adams/petsc/src/snes/tests/data$ grep "KSPSolve 2" jac_out_001*_5_8_gpuawaremp*<br class="">KSPSolve 2 1.0<span class="Apple-converted-space"> </span><b class="">4.4173e+00</b><span class="Apple-converted-space"> </span>1.0 1.48e+10 1.1 3.7e+05 1.6e+04 1.2e+03 4 60 62 54 61 100100100100100<span class="Apple-converted-space"> </span><b class="">208923 1094405</b><span class="Apple-converted-space"> </span> 0 0.00e+00 0 0.00e+00 100</font><br class=""></div></div><div class=""><font face="monospace" class=""><br class=""></font></div>Notes about flop counters here, <div class="">* that MatMult flops are not logged as GPU flops but something is logged nonetheless.<div class="">* The GPU flop rate is 5x the total flop rate in KSPSolve :\<br class=""><div class="">* I think these nodes have an FP64 peak flop rate of 200 Tflops, so we are at < 1%.</div></div></div></div></blockquote><div class=""><br class=""></div><div class="">This looks complicated, so just a single remark:</div><div class=""><br class=""></div><div class="">My understanding of the benchmarking of vector ops led by Hannah was that you needed to be much</div><div class="">bigger than 16M to hit peak. I need to get the tech report, but on 8 GPUs I would think you would be</div><div class="">at 10% of peak or something right off the bat at these sizes. Barry, is that right?</div><div class=""><br class=""></div><div class=""> Thanks,</div><div class=""><br class=""></div><div class=""> Matt</div><div class=""> </div><blockquote class="gmail_quote" style="margin: 0px 0px 0px 0.8ex; border-left-width: 1px; border-left-style: solid; border-left-color: rgb(204, 204, 204); padding-left: 1ex;"><div dir="ltr" class=""><div class=""><div class=""><div class="">Anway, not sure how to proceed but I thought I would share.</div><div class="">Maybe ask the Kokkos guys if the have looked at Crusher.</div><div class=""><br class=""></div><div class="">Mark</div></div></div></div></blockquote></div>--<span class="Apple-converted-space"> </span><br class=""><div dir="ltr" class=""><div dir="ltr" class=""><div class=""><div dir="ltr" class=""><div class=""><div dir="ltr" class=""><div class="">What most experimenters take for granted before they begin their experiments is infinitely more interesting than any results to which their experiments lead.<br class="">-- Norbert Wiener</div><div class=""><br class=""></div><div class=""><a href="http://www.cse.buffalo.edu/~knepley/" target="_blank" class="">https://www.cse.buffalo.edu/~knepley/</a><br class=""></div></div></div></div></div></div></div></div></div></blockquote></div><br class=""></div></div></blockquote></div><span id="gmail-m_7178493797419230199gmail-m_-5850858973953305955gmail-m_-8561488623817931590gmail-m_-9217502836458641567gmail-m_-1042935854083030742cid:f_kyp816vp0" class=""><jac_out_001_kokkos_Crusher_6_8_gpuawarempi.txt></span></div></blockquote></div><br class=""></div></div></blockquote></div></blockquote></div></div></blockquote></div><span id="gmail-m_7178493797419230199gmail-m_-5850858973953305955cid:f_kyq28fj80" class=""><jac_out_001_kokkos_Crusher_5_8_notpl.txt></span><span id="gmail-m_7178493797419230199gmail-m_-5850858973953305955cid:f_kyq28fji1" class=""><jac_out_001_kokkos_Crusher_6_8_notpl.txt></span></div></blockquote></div></div></div></blockquote></div></div></blockquote></div></div></div></blockquote></div></div></blockquote></div><br class=""></div></div></div></div></blockquote></div><br class=""></body></html>