<html aria-label="message body"><head><meta http-equiv="content-type" content="text/html; charset=utf-8"></head><body style="overflow-wrap: break-word; -webkit-nbsp-mode: space; line-break: after-white-space;"><div><br></div> I have updated the MR with what I think is now correct code for computing the diagonal on the GPU, could you please try it again and let me know if it works and how much time it saves (I think it is should be significant).<div><br></div><div> Thankts for your patients,</div><div><br></div><div> Barry</div><div><br id="lineBreakAtBeginningOfMessage"><div><br><blockquote type="cite"><div>On Oct 2, 2025, at 1:16 AM, LEDAC Pierre <Pierre.LEDAC@cea.fr> wrote:</div><br class="Apple-interchange-newline"><div><meta charset="UTF-8"><div id="divtagdefaultwrapper" dir="ltr" style="font-style: normal; font-variant-caps: normal; font-weight: 400; 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; font-size: 12pt; font-family: Calibri, Helvetica, sans-serif;"><div style="margin-top: 0px; margin-bottom: 0px;">Yes, probably the reason I saw also a crash in my test case after a quick fix of the integer conversion. </div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div id="Signature"><div id="divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;"><div style="font-family: Tahoma; font-size: 13px;"><div class="BodyFragment"><font size="2"><span style="font-size: 10pt;"><div class="PlainText">Pierre LEDAC<br>Commissariat à l’énergie atomique et aux énergies alternatives<br>Centre de SACLAY<br>DES/ISAS/DM2S/SGLS/LCAN<br>Bâtiment 451 – point courrier n°41<br>F-91191 Gif-sur-Yvette<br>+33 1 69 08 04 03<br>+33 6 83 42 05 79</div></span></font></div></div></div></div></div><hr tabindex="-1" style="font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; letter-spacing: normal; orphans: auto; text-align: start; text-indent: 0px; text-transform: none; white-space: normal; widows: auto; word-spacing: 0px; -webkit-text-stroke-width: 0px; text-decoration: none; display: inline-block; width: 1043.6875px;"><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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; float: none; display: inline !important;"></span><div id="divRplyFwdMsg" dir="ltr" style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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;"><font face="Calibri, sans-serif" style="font-size: 11pt;"><b>De :</b><span class="Apple-converted-space"> </span>Barry Smith <<a href="mailto:bsmith@petsc.dev">bsmith@petsc.dev</a>><br><b>Envoyé :</b><span class="Apple-converted-space"> </span>jeudi 2 octobre 2025 02:16:40<br><b>À :</b><span class="Apple-converted-space"> </span>LEDAC Pierre<br><b>Cc :</b><span class="Apple-converted-space"> </span>Junchao Zhang;<span class="Apple-converted-space"> </span><a href="mailto:petsc-users@mcs.anl.gov">petsc-users@mcs.anl.gov</a><br><b>Objet :</b><span class="Apple-converted-space"> </span>Re: [petsc-users] [GPU] Jacobi preconditioner</font><div> </div></div><div style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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;"><div><br></div> Sorry about that. The current code is buggy anyways; I will let you know when I have tested it extensively so you can try again.<div><br></div><div> Barry</div><div><br id="lineBreakAtBeginningOfMessage"><div><br><blockquote type="cite"><div>On Oct 1, 2025, at 3:47 PM, LEDAC Pierre <<a href="mailto:Pierre.LEDAC@cea.fr">Pierre.LEDAC@cea.fr</a>> wrote:</div><br class="Apple-interchange-newline"><div><div id="divtagdefaultwrapper" dir="ltr" style="font-style: normal; font-variant-caps: normal; font-weight: 400; 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; font-size: 12pt; font-family: Calibri, Helvetica, sans-serif;"><div style="margin-top: 0px; margin-bottom: 0px;">Sorry the correct error is:</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><p style="margin-top: 0px; margin-bottom: 0px;"></p><div style="font-family: Calibri, Helvetica, sans-serif, serif, EmojiFont; font-size: 16px;">/export/home/catA/pl254994/trust/petsc/lib/src/LIBPETSC/build/petsc-barry-2025-09-30-add-matgetdiagonal-cuda/src/mat/impls/aij/seq/seqcusparse/aijcusparse.cu(3765): error: argument of type "<span style="color: rgb(255, 0, 0);">int*</span>" is incompatible with parameter of type "const PetscInt *"</div><div style="font-family: Calibri, Helvetica, sans-serif, serif, EmojiFont; font-size: 16px;"> GetDiagonal_CSR<<<(int)((n + 255) / 256), 256, 0, PetscDefaultCudaStream>>>(cusparsestruct->rowoffsets_gpu->data().get(), matstruct->cprowIndices->data().get(), cusparsestruct->workVector->data().get(), n, darray);</div><br><p style="margin-top: 0px; margin-bottom: 0px;"></p><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div id="Signature"><div id="divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;"><div style="font-family: Tahoma; font-size: 13px;"><div class="BodyFragment"><font size="2"><span style="font-size: 10pt;"><div class="PlainText">Pierre LEDAC<br>Commissariat à l’énergie atomique et aux énergies alternatives<br>Centre de SACLAY<br>DES/ISAS/DM2S/SGLS/LCAN<br>Bâtiment 451 – point courrier n°41<br>F-91191 Gif-sur-Yvette<br>+33 1 69 08 04 03<br>+33 6 83 42 05 79</div></span></font></div></div></div></div></div><hr tabindex="-1" style="font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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; display: inline-block; width: 862.390625px;"><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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; float: none; display: inline !important;"></span><div id="divRplyFwdMsg" dir="ltr" style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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;"><font face="Calibri, sans-serif" style="font-size: 11pt;"><b>De :</b><span class="Apple-converted-space"> </span>LEDAC Pierre<br><b>Envoyé :</b><span class="Apple-converted-space"> </span>mercredi 1 octobre 2025 21:46:00<br><b>À :</b><span class="Apple-converted-space"> </span>Barry Smith<br><b>Cc :</b><span class="Apple-converted-space"> </span>Junchao Zhang;<span class="Apple-converted-space"> </span><a href="mailto:petsc-users@mcs.anl.gov">petsc-users@mcs.anl.gov</a><br><b>Objet :</b><span class="Apple-converted-space"> </span>RE: [petsc-users] [GPU] Jacobi preconditioner</font><div> </div></div><div style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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;"><div id="divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif;"><div style="margin-top: 0px; margin-bottom: 0px;">Hi all,</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;">Thanks for the MR, there is a build issue cause we use <span>--with-64-bit-indices:</span></div><p style="margin-top: 0px; margin-bottom: 0px;"><span></span></p><div><br></div><div>/export/home/catA/pl254994/trust/petsc/lib/src/LIBPETSC/build/petsc-barry-2025-09-30-add-matgetdiagonal-cuda/src/mat/impls/aij/seq/seqcusparse/aijcusparse.cu(3765): error: argument of type "PetscInt" is incompatible with parameter of type "const PetscInt *"</div><div> GetDiagonal_CSR<<<(int)((n + 255) / 256), 256, 0, PetscDefaultCudaStream>>>(cusparsestruct->rowoffsets_gpu->data().get(), matstruct->cprowIndices->data().get(), cusparsestruct->workVector->data().get(), n, darray);</div><div><br></div><div></div><div>Thanks,</div><p style="margin-top: 0px; margin-bottom: 0px;"></p><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div id="Signature"><div id="divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;"><div style="font-family: Tahoma; font-size: 13px;"><div class="BodyFragment"><font size="2"><span style="font-size: 10pt;"><div class="PlainText">Pierre LEDAC<br>Commissariat à l’énergie atomique et aux énergies alternatives<br>Centre de SACLAY<br>DES/ISAS/DM2S/SGLS/LCAN<br>Bâtiment 451 – point courrier n°41<br>F-91191 Gif-sur-Yvette<br>+33 1 69 08 04 03<br>+33 6 83 42 05 79</div></span></font></div></div></div></div></div><hr tabindex="-1" style="display: inline-block; width: 862.390625px;"><div id="divRplyFwdMsg" dir="ltr"><font face="Calibri, sans-serif" style="font-size: 11pt;"><b>De :</b><span class="Apple-converted-space"> </span>Barry Smith <<a href="mailto:bsmith@petsc.dev">bsmith@petsc.dev</a>><br><b>Envoyé :</b><span class="Apple-converted-space"> </span>mercredi 1 octobre 2025 18:48:37<br><b>À :</b><span class="Apple-converted-space"> </span>LEDAC Pierre<br><b>Cc :</b><span class="Apple-converted-space"> </span>Junchao Zhang;<span class="Apple-converted-space"> </span><a href="mailto:petsc-users@mcs.anl.gov">petsc-users@mcs.anl.gov</a><br><b>Objet :</b><span class="Apple-converted-space"> </span>Re: [petsc-users] [GPU] Jacobi preconditioner</font><div> </div></div><div><div><br></div> I have finally created an MR that moves the Jacobi accessing of the diagonal to the GPU, which should improve the GPU performance of your code. <a href="https://urldefense.us/v3/__https://gitlab.com/petsc/petsc/-/merge_requests/8756__;!!G_uCfscf7eWS!dKYXeFfmzkN4Tc2oZZQPw_GL8AwkmzFXhYH0TzM4Dtrf72oWmFpj8bmmF-L_Fl784As9NbsIqk1pM9g-xSi5XZo$">https://gitlab.com/petsc/petsc/-/merge_requests/8756</a><div><br></div><div> Please give it a try and let us know if it causes any difficulties or, hopefully, improves your code's performance significantly.</div><div><br></div><div> Sorry for the long delay, NVIDIA is hiring too many PETSc developers away from us.</div><div><br></div><div> Barry<br id="lineBreakAtBeginningOfMessage"><div><br><blockquote type="cite"><div>On Jul 31, 2025, at 6:46 AM, LEDAC Pierre <<a href="mailto:Pierre.LEDAC@cea.fr">Pierre.LEDAC@cea.fr</a>> wrote:</div><br class="Apple-interchange-newline"><div><div id="divtagdefaultwrapper" dir="ltr" style="font-style: normal; font-variant-caps: normal; font-weight: 400; 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; font-size: 12pt; font-family: Calibri, Helvetica, sans-serif;"><div style="margin-top: 0px; margin-bottom: 0px;">Thanks Barry, I agree but didn't dare asking for that.</div><br><div id="Signature"><div id="divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;"><div style="font-family: Tahoma; font-size: 13px;"><div class="BodyFragment"><font size="2"><span style="font-size: 10pt;"><div class="PlainText">Pierre LEDAC<br>Commissariat à l’énergie atomique et aux énergies alternatives<br>Centre de SACLAY<br>DES/ISAS/DM2S/SGLS/LCAN<br>Bâtiment 451 – point courrier n°41<br>F-91191 Gif-sur-Yvette<br>+33 1 69 08 04 03<br>+33 6 83 42 05 79</div></span></font></div></div></div></div></div><hr tabindex="-1" style="font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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; display: inline-block; width: 1043.6875px;"><span class="Apple-converted-space"> </span><span class="Apple-converted-space"> </span><span style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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; float: none; display: inline !important;"></span><div id="divRplyFwdMsg" dir="ltr" style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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;"><font face="Calibri, sans-serif" style="font-size: 11pt;"><b>De :</b><span class="Apple-converted-space"> </span>Barry Smith <<a href="mailto:bsmith@petsc.dev">bsmith@petsc.dev</a>><br><b>Envoyé :</b><span class="Apple-converted-space"> </span>mercredi 30 juillet 2025 20:34:26<br><b>À :</b><span class="Apple-converted-space"> </span>Junchao Zhang<br><b>Cc :</b><span class="Apple-converted-space"> </span>LEDAC Pierre;<span class="Apple-converted-space"> </span><a href="mailto:petsc-users@mcs.anl.gov">petsc-users@mcs.anl.gov</a><br><b>Objet :</b><span class="Apple-converted-space"> </span>Re: [petsc-users] [GPU] Jacobi preconditioner</font><div> </div></div><div style="caret-color: rgb(0, 0, 0); font-family: Helvetica; font-size: 18px; font-style: normal; font-variant-caps: normal; font-weight: 400; 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;"><div><br></div> We absolutely should have a <font face="Calibri, Helvetica, sans-serif, EmojiFont, Apple Color Emoji, Segoe UI Emoji, NotoColorEmoji, Segoe UI Symbol, Android Emoji, EmojiSymbols" size="3">MatGetDiagonal_SeqAIJCUSPARSE(). It's somewhat embarrassing that we don't provide this.</font><div><br></div><div> I have found some potential code at <a href="https://urldefense.us/v3/__https://stackoverflow.com/questions/60311408/how-to-get-the-diagonal-of-a-sparse-matrix-in-cusparse__;!!G_uCfscf7eWS!dKYXeFfmzkN4Tc2oZZQPw_GL8AwkmzFXhYH0TzM4Dtrf72oWmFpj8bmmF-L_Fl784As9NbsIqk1pM9g-gN1IJTI$">https://stackoverflow.com/questions/60311408/how-to-get-the-diagonal-of-a-sparse-matrix-in-cusparse</a></div><div><br></div><div> Barry</div><div><br></div><div><br></div><div><br id="lineBreakAtBeginningOfMessage"><div><br><blockquote type="cite"><div>On Jul 28, 2025, at 11:43 AM, Junchao Zhang <<a href="mailto:junchao.zhang@gmail.com">junchao.zhang@gmail.com</a>> wrote:</div><br class="Apple-interchange-newline"><div><div dir="ltr">Yes, MatGetDiagonal_SeqAIJCUSPARSE hasn't been implemented. petsc/cuda and petsc/kokkos backends are separate code. <br>If petsc/kokkos meet your needs, then just use them. For petsc users, we hope it will be just a difference of extra --download-kokkos --download-kokkos-kernels in configuration. <div><br></div><div>--Junchao Zhang</div><br></div><br><div class="gmail_quote gmail_quote_container"><div dir="ltr" class="gmail_attr">On Mon, Jul 28, 2025 at 2:51 AM LEDAC Pierre <<a href="mailto:Pierre.LEDAC@cea.fr">Pierre.LEDAC@cea.fr</a>> wrote:<br></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="msg3805772366042246901"><div dir="ltr"><div id="m_-58953184740190291divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif;"><div style="margin-top: 0px; margin-bottom: 0px;">Hello all,</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;">We are solving with PETSc a linear system updated every time step (constant stencil but coefficients changing).</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;">The matrix is preallocated once with <span>MatSetPreallocationCOO() then filled each time step with <span>MatSetValuesCOO() and we use device pointers for coo_i, coo_j, and coefficients values.</span></span></div><div style="margin-top: 0px; margin-bottom: 0px;"><span><span><br></span></span></div><div style="margin-top: 0px; margin-bottom: 0px;"><span><span>It is working fine with a GMRES Ksp solver and PC Jacobi but we are surprised to see that every time step, during PCSetUp, MatGetDiagonal_SeqAIJ is called whereas the matrix is on the device. Looking at the API, it seems there is no <span style="font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols; font-size: 16px;">MatGetDiagonal_SeqAIJCUSPARSE() but a <span style="font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols; font-size: 16px;">MatGetDiagonal_SeqAIJKOKKOS().</span></span></span></span></div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;">Does it mean we should use Kokkos backend in PETSc to have Jacobi preconditioner built directly on device ? Or I am doing something wrong ?</div><div style="margin-top: 0px; margin-bottom: 0px;">NB: Gmres is running well on device.</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;">I could use -ksp_reuse_preconditioner to avoid Jacobi being recreated each solve on host but it increases significantly the number of iterations.</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;">Thanks,</div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;"><span id="cid:ii_19851ad63faf456b1e51"><pastedImage.png></span><br></div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div style="margin-top: 0px; margin-bottom: 0px;"><br></div><div id="m_-58953184740190291Signature"><div id="m_-58953184740190291divtagdefaultwrapper" dir="ltr" style="font-size: 12pt; font-family: Calibri, Helvetica, sans-serif, EmojiFont, "Apple Color Emoji", "Segoe UI Emoji", NotoColorEmoji, "Segoe UI Symbol", "Android Emoji", EmojiSymbols;"><div style="font-family: Tahoma; font-size: 13px;"><div><font size="2"><span style="font-size: 10pt;"><div>Pierre LEDAC<br>Commissariat à l’énergie atomique et aux énergies alternatives<br>Centre de SACLAY<br>DES/ISAS/DM2S/SGLS/LCAN<br>Bâtiment 451 – point courrier n°41<br>F-91191 Gif-sur-Yvette<br>+33 1 69 08 04 03<br>+33 6 83 42 05 79</div></span></font></div></div></div></div></div></div></div></blockquote></div></div></blockquote></div></div></div></div></blockquote></div></div></div></div></div></blockquote></div></div></div></div></blockquote></div><br></div></body></html>