<head><!-- BaNnErBlUrFlE-HeAdEr-start -->
<style>
#pfptBanner7311kgh { all: revert !important; display: block !important;
visibility: visible !important; opacity: 1 !important;
background-color: #D0D8DC !important;
max-width: none !important; max-height: none !important }
.pfptPrimaryButton7311kgh:hover, .pfptPrimaryButton7311kgh:focus {
background-color: #b4c1c7 !important; }
.pfptPrimaryButton7311kgh:active {
background-color: #90a4ae !important; }
</style>
<!-- BaNnErBlUrFlE-HeAdEr-end -->
</head><!-- BaNnErBlUrFlE-BoDy-start -->
<!-- Preheader Text : BEGIN -->
<div style="display:none !important;display:none;visibility:hidden;mso-hide:all;font-size:1px;color:#ffffff;line-height:1px;height:0px;max-height:0px;opacity:0;overflow:hidden;">
I don't have a working example that I can easily send you since the MatShell I have involves a lot of other code. Essentially, I am trying to do something like this: bool graphCreated = false; cudaGraph_t graph; cudaGraphExec_t instance;if
</div>
<!-- Preheader Text : END -->
<!-- Email Banner : BEGIN -->
<div style="display:none !important;display:none;visibility:hidden;mso-hide:all;font-size:1px;color:#ffffff;line-height:1px;height:0px;max-height:0px;opacity:0;overflow:hidden;">ZjQcmQRYFpfptBannerStart</div>
<!--[if ((ie)|(mso))]>
<table border="0" cellspacing="0" cellpadding="0" width="100%" style="padding: 16px 0px 16px 0px; direction: ltr" ><tr><td>
<table border="0" cellspacing="0" cellpadding="0" style="padding: 0px 10px 5px 6px; width: 100%; border-radius:4px; border-top:4px solid #90a4ae;background-color:#D0D8DC;"><tr><td valign="top">
<table align="left" border="0" cellspacing="0" cellpadding="0" style="padding: 4px 8px 4px 8px">
<tr><td style="color:#000000; font-family: 'Arial', sans-serif; font-weight:bold; font-size:14px; direction: ltr">
This Message Is From an External Sender
</td></tr>
<tr><td style="color:#000000; font-weight:normal; font-family: 'Arial', sans-serif; font-size:12px; direction: ltr">
This message came from outside your organization.
</td></tr>
</table>
</td></tr></table>
</td></tr></table>
<![endif]-->
<![if !((ie)|(mso))]>
<div dir="ltr" id="pfptBanner7311kgh" style="all: revert !important; display:block !important; text-align: left !important; margin:16px 0px 16px 0px !important; padding:8px 16px 8px 16px !important; border-radius: 4px !important; min-width: 200px !important; background-color: #D0D8DC !important; background-color: #D0D8DC; border-top: 4px solid #90a4ae !important; border-top: 4px solid #90a4ae;">
<div id="pfptBanner7311kgh" style="all: unset !important; float:left !important; display:block !important; margin: 0px 0px 1px 0px !important; max-width: 600px !important;">
<div id="pfptBanner7311kgh" style="all: unset !important; display:block !important; visibility: visible !important; background-color: #D0D8DC !important; color:#000000 !important; color:#000000; font-family: 'Arial', sans-serif !important; font-family: 'Arial', sans-serif; font-weight:bold !important; font-weight:bold; font-size:14px !important; line-height:18px !important; line-height:18px">
This Message Is From an External Sender
</div>
<div id="pfptBanner7311kgh" style="all: unset !important; display:block !important; visibility: visible !important; background-color: #D0D8DC !important; color:#000000 !important; color:#000000; font-weight:normal; font-family: 'Arial', sans-serif !important; font-family: 'Arial', sans-serif; font-size:12px !important; line-height:18px !important; line-height:18px; margin-top:2px !important;">
This message came from outside your organization.
</div>
</div>
<div style="clear: both !important; display: block !important; visibility: hidden !important; line-height: 0 !important; font-size: 0.01px !important; height: 0px"> </div>
</div>
<![endif]>
<div style="display:none !important;display:none;visibility:hidden;mso-hide:all;font-size:1px;color:#ffffff;line-height:1px;height:0px;max-height:0px;opacity:0;overflow:hidden;">ZjQcmQRYFpfptBannerEnd</div>
<!-- Email Banner : END -->
<!-- BaNnErBlUrFlE-BoDy-end -->
<div dir="ltr">I don't have a working example that I can easily send you since the MatShell I have involves a lot of other code. <div><br></div><div>Essentially, I am trying to do something like this:</div><div><br></div><div><font face="monospace">bool graphCreated = false;<br>cudaGraph_t graph;<br>cudaGraphExec_t instance;<br>if (!graphCreated)<br> {<br> gpuErrchk(cudaStreamBeginCapture(s, cudaStreamCaptureModeGlobal));<br> PetscCall(MatComputeOperator(ShellMat, MATDENSECUDA, &DenseMat));<br> gpuErrchk(cudaStreamEndCapture(s, &graph));<br> gpuErrchk(cudaGraphInstantiate(&instance, graph, NULL, NULL, 0));<br> graphCreated = true;<br> }<br>gpuErrchk(cudaGraphLaunch(instance, s));<br>gpuErrchk(cudaStreamSynchronize(s));<br>gpuErrchk(cudaGraphExecDestroy(instance));<br>gpuErrchk(cudaGraphDestroy(graph));</font><br></div><div><font face="monospace"><br></font></div><div><font face="arial, sans-serif">The ShellMat context has the CUDA Stream "s", and that stream is used in the shell computations. If I replace the MatComputeOperator() call with </font></div><div><font face="arial, sans-serif"><br></font></div><font face="monospace">for (int i = 0; i < N; i++) </font><div><font face="monospace"> MatMult(ShellMat, x, y);</font></div><div><font face="monospace"><br></font></div><div><font face="arial, sans-serif">the code works inside the CUDA graph. I tried modifying the MatConvert_Shell() code and wrapping the main for loop that computes each column with the graph, but this still doesn't work.</font></div><div><font face="arial, sans-serif"><br></font></div><div><font face="arial, sans-serif">I think if you have a MatShell that does some trivial computation (like return y = x), that should be a good enough example to start with. </font></div><div><font face="arial, sans-serif"><br></font></div><div><font face="arial, sans-serif">For now, I ended up making my own version of MatComputeOperator that doesn't involve Mats or Vecs in the main loop (working directly with arrays) so I could have control over the stream operations.</font></div><div><font face="arial, sans-serif"><br></font></div><div><font face="arial, sans-serif">Thanks,</font></div><div><font face="arial, sans-serif">Sreeram</font></div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Tue, May 14, 2024 at 10:15 PM Junchao Zhang <<a href="mailto:junchao.zhang@gmail.com">junchao.zhang@gmail.com</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div dir="ltr"><div><br></div><div>I haven't used CUDA graph with PETSc. Do you happen to have a working example so we can debug?</div><br clear="all"><div><div dir="ltr" class="gmail_signature"><div dir="ltr">--Junchao Zhang</div></div></div><br></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Tue, May 14, 2024 at 6:08 PM Sreeram R Venkat <<a href="mailto:srvenkat@utexas.edu" target="_blank">srvenkat@utexas.edu</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><div>
<div style="font-size:1px;color:rgb(255,255,255);line-height:1px;height:0px;max-height:0px;opacity:0;overflow:hidden;display:none">
I have a MatShell object that I want to convert to a MATDENSECUDA. Normally, I use MatComputeOperator for this. However, I would now also like to use a CUDA Graph so that all the calls to MatMult are captured. I can wrap a code like for (int
</div>
<div style="font-size:1px;color:rgb(255,255,255);line-height:1px;height:0px;max-height:0px;opacity:0;overflow:hidden;display:none">ZjQcmQRYFpfptBannerStart</div>
<u></u>
<div dir="ltr" id="m_-4586763279418122487m_-306059168063068977pfptBannerfgin0u9" style="display:block;text-align:left;margin:16px 0px;padding:8px 16px;border-radius:4px;min-width:200px;background-color:rgb(208,216,220);border-top:4px solid rgb(144,164,174)">
<div id="m_-4586763279418122487m_-306059168063068977pfptBannerfgin0u9" style="float:left;display:block;margin:0px 0px 1px;max-width:600px">
<div id="m_-4586763279418122487m_-306059168063068977pfptBannerfgin0u9" style="display:block;background-color:rgb(208,216,220);color:rgb(0,0,0);font-family:Arial,sans-serif;font-weight:bold;font-size:14px;line-height:18px">
This Message Is From an External Sender
</div>
<div id="m_-4586763279418122487m_-306059168063068977pfptBannerfgin0u9" style="font-weight:normal;display:block;background-color:rgb(208,216,220);color:rgb(0,0,0);font-family:Arial,sans-serif;font-size:12px;line-height:18px;margin-top:2px">
This message came from outside your organization.
</div>
</div>
<div style="height:0px;clear:both;display:block;line-height:0;font-size:0.01px"> </div>
</div>
<u></u>
<div style="font-size:1px;color:rgb(255,255,255);line-height:1px;height:0px;max-height:0px;opacity:0;overflow:hidden;display:none">ZjQcmQRYFpfptBannerEnd</div>
<div dir="ltr"><div dir="ltr" class="gmail_signature"><div dir="ltr"><div style="color:rgb(34,34,34)">I have a MatShell object that I want to convert to a MATDENSECUDA. Normally, I use MatComputeOperator for this. However, I would now also like to use a CUDA Graph so that all the calls to MatMult are captured. I can wrap a code like</div><div style="color:rgb(34,34,34)"><br></div><div style="color:rgb(34,34,34)">for (int i = 0; i < N; i++)</div><div style="color:rgb(34,34,34)"> MatMult(A, x,y);</div><div style="color:rgb(34,34,34)"><br></div><div style="color:rgb(34,34,34)">in a CUDA Graph, and it runs fine. If I try to wrap MatComputeOperator in a graph, I get runtime errors like</div><div style="color:rgb(34,34,34)"> cuda error 906 (cudaErrorStreamCaptureImplicit) : operation would make the legacy stream depend on a capturing blocking stream</div><div style="color:rgb(34,34,34)"><br></div><div style="color:rgb(34,34,34)">I tried modifying the MatConvert_Shell routine to only put the graph around the main for loop, but that still gives the same errors. Is there a way to use CUDA Graphs here (either through a modified MatConvert_Shell or otherwise)?</div><div style="color:rgb(34,34,34)"><br></div><div style="color:rgb(34,34,34)">Thanks,</div><div style="color:rgb(34,34,34)">Sreeram</div></div></div></div>
</div></blockquote></div>
</blockquote></div>