<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body text="#000000" bgcolor="#FFFFFF">
<p><br>
</p>
<div class="moz-cite-prefix">On 3/20/19 12:25 PM, Alexey Bataev wrote:<br>
</div>
<blockquote type="cite" cite="mid:AM5PR0902MB19238A30BF5E0741760FC64196410@AM5PR0902MB1923.eurprd09.prod.outlook.com">
<p>Hi Hal, it is hard to tell. I can try to add the option that will lead to definition of this macro to clang, if you want to try it.<br>
</p>
</blockquote>
<p><br>
</p>
<p>That would be great, thanks!</p>
<p> -Hal<br>
</p>
<p><br>
</p>
<blockquote type="cite" cite="mid:AM5PR0902MB19238A30BF5E0741760FC64196410@AM5PR0902MB1923.eurprd09.prod.outlook.com">
<p></p>
<pre class="moz-signature" cols="72">-------------
Best regards,
Alexey Bataev</pre>
<div class="moz-cite-prefix">20.03.2019 13:12, Finkel, Hal J. пишет:<br>
</div>
<blockquote type="cite" cite="mid:ea79364c-681d-5ea9-8096-b016819df4e9@anl.gov">
<p>Thanks, Ye. I suppose that I thought it always worked that way :-)</p>
<p>Alexey, Doru, do you know if there's any semantic problem or other concerns with enabling this option and/or making it the default?</p>
<p> -Hal<br>
</p>
<div class="moz-cite-prefix">On 3/20/19 11:32 AM, Ye Luo via Openmp-dev wrote:<br>
</div>
<blockquote type="cite" cite="mid:CACiEoHngRT+ANa47qSc6s0xWhAdFcVft9yYZ9ib2ugz4tZzL3Q@mail.gmail.com">
<div dir="ltr">
<div dir="ltr">
<div dir="ltr">
<div dir="ltr">
<div dir="ltr">
<div>Hi all,</div>
<div>After going through the source, I didn't find CUDA stream support.</div>
<div>Luckily, I only need to add</div>
<div>#define CUDA_API_PER_THREAD_DEFAULT_STREAM</div>
<div>before</div>
<div>#include <cuda.h></div>
<div>in libomptarget/plugins/cuda/src/rtl.cpp</div>
<div>Then the multiple target goes to different streams and may execute concurrently.<br>
</div>
<div>
<div>#pragma omp parallel</div>
<div>{</div>
<div> #pragma omp target</div>
<div> {</div>
<div> //offload computation<br>
</div>
<div> }<br>
</div>
<div>}</div>
<div>This is exactly I want.</div>
<div><br>
</div>
<div>I know the XL compiler uses streams in a different way but achieves similar effects.<br>
</div>
</div>
<div>Is there anyone working on using streams with openmp target in llvm?</div>
<div>Will clang-ykt get something similar to XL and upstream to the mainline?</div>
<div><br>
</div>
<div>If we just add #define CUDA_API_PER_THREAD_DEFAULT_STREAM in the cuda rtl, will it be a trouble?</div>
<div>As a compiler user, I'd like to have a better solution rather than having a patch just for myself.</div>
<div><br>
</div>
<div>Best,<br>
</div>
<div>Ye<br>
</div>
<div dir="ltr">
<div>
<div dir="ltr" class="gmail_signature">
<div dir="ltr">
<div>
<div dir="ltr">===================<br>
Ye Luo, Ph.D.<br>
Computational Science Division & Leadership Computing Facility<br>
Argonne National Laboratory</div>
</div>
</div>
</div>
</div>
<br>
</div>
<br>
<div class="gmail_quote">
<div dir="ltr" class="gmail_attr">Ye Luo <<a href="mailto:xw111luoye@gmail.com" moz-do-not-send="true">xw111luoye@gmail.com</a>> 于2019年3月17日周日 下午2:26写道：<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>Hi,</div>
<div>How to turn on streams when using OpenMP offload?</div>
<div>When different host threads individually start target regions (even not using nowait). The offloaded computation goes to different CUDA streams and may execute concurrently. This is currently available in XL.</div>
<div>With Clang, nvprof shows only the run only uses the default stream.</div>
<div>Is there a way to do that with Clang?</div>
<div>On the other hand,<br>
</div>
<div>nvcc has option<span class="gmail-m_-5563546046592982142gmail-st"> --<i>default</i>-<i>stream per</i>-<i>thread</i></span></div>
<div>I'm not familar with clang CUDA, is there a similar option?</div>
<div>Best,<br>
</div>
<div>Ye<br>
</div>
<div>
<div>
<div dir="ltr" class="gmail-m_-5563546046592982142gmail_signature">
<div dir="ltr">
<div>
<div dir="ltr">===================<br>
Ye Luo, Ph.D.<br>
Computational Science Division & Leadership Computing Facility<br>
Argonne National Laboratory</div>
</div>
</div>
</div>
</div>
</div>
</div>
</blockquote>
</div>
</div>
</div>
</div>
</div>
</div>
<br>
<fieldset class="mimeAttachmentHeader"></fieldset>
<pre class="moz-quote-pre" wrap="">_______________________________________________
Openmp-dev mailing list
<a class="moz-txt-link-abbreviated" href="mailto:Openmp-dev@lists.llvm.org" moz-do-not-send="true">Openmp-dev@lists.llvm.org</a>
<a class="moz-txt-link-freetext" href="https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev" moz-do-not-send="true">https://lists.llvm.org/cgi-bin/mailman/listinfo/openmp-dev</a>
</pre>
</blockquote>
<pre class="moz-signature" cols="72">--
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory</pre>
</blockquote>
</blockquote>
<pre class="moz-signature" cols="72">--
Hal Finkel
Lead, Compiler Technology and Programming Languages
Leadership Computing Facility
Argonne National Laboratory</pre>
</body>
</html>