<div dir="ltr"><div dir="ltr"><div dir="ltr">>> Do you mean that because the OpenMP programmer does not explicitly mark device functions, the general approach taken to these sorts of errors in CUDA is inapplicable?</div><div dir="ltr"></div><div dir="ltr">> Yes, right. OpenMP does not require explicit marking of the device function. It supports implicit target functions.<div><br></div><div>Would you be willing to elaborate on how the fact that OpenMP has implicit device functions means that the general approach taken by CUDA to this problem is inapplicable to OpenMP?</div><div><br></div><div>(For example, I'd have naively thought that CUDA __host__ __device__ functions are very similar to an OpenMP implicit-device function, in basically all respects other than the fact that the CUDA functions have an explicit attribute.)</div></div></div></div><br><div class="gmail_quote"><div dir="ltr">On Tue, Jan 15, 2019 at 2:26 PM Alexey Bataev <<a href="mailto:a.bataev@outlook.com" target="_blank">a.bataev@outlook.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="auto">
Yes, right. OpenMP does not require explicit marking of the device function. It supports implicit target functions.<br>
<br>
<div id="gmail-m_4506947491132299509gmail-m_-5021161699907549984AppleMailSignature" dir="ltr">Best regards,
<div>Alexey Bataev</div>
</div>
<div dir="ltr"><br>
15 янв. 2019 г., в 17:24, Justin Lebar <<a href="mailto:jlebar@google.com" target="_blank">jlebar@google.com</a>> написал(а):<br>
<br>
</div>
<blockquote type="cite">
<div dir="ltr">
<div dir="ltr">
<div dir="ltr">> As for Cuda, it is a little but different. In Cuda the programmer must explicitly mark the device functions, while in OpenMP it must be done implicitly. Thus, we cannot reuse the solution used for Cuda.</div>
<div dir="ltr"><br>
</div>
<div>Do you mean that because the OpenMP programmer does not explicitly mark device functions, the general approach taken to these sorts of errors in CUDA is inapplicable?</div>
</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Tue, Jan 15, 2019 at 2:20 PM Alexey Bataev <<a href="mailto:a.bataev@outlook.com" target="_blank">a.bataev@outlook.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="auto">This is not only for asm, we need to delay all target-specific diagnostics.
<div>I'm not saying that we need to move the host diagnostic, only the diagnostic for the device compilation.</div>
<div>As for Cuda, it is a little but different. In Cuda the programmer must explicitly mark the device functions, while in OpenMP it must be done implicitly. Thus, we cannot reuse the solution used for Cuda.<br>
<br>
<div id="gmail-m_4506947491132299509gmail-m_-5021161699907549984gmail-m_-5303006821952141123AppleMailSignature" dir="ltr">Best regards,
<div>Alexey Bataev</div>
</div>
<div dir="ltr"><br>
15 янв. 2019 г., в 17:07, Reid Kleckner <<a href="mailto:rnk@google.com" target="_blank">rnk@google.com</a>> написал(а):<br>
<br>
</div>
<blockquote type="cite">
<div dir="ltr">
<div dir="ltr">
<div dir="ltr">Didn't we already go through this for CUDA? It has all the same issues. I see some comments about Sema::CUDADeferredDiags. I would look at that and try to generalize it.</div>
</div>
<br>
<div class="gmail_quote">
<div dir="ltr" class="gmail-m_4506947491132299509gmail-m_-5021161699907549984gmail-m_-5303006821952141123gmail_attr">On Tue, Jan 15, 2019 at 1:58 PM John McCall via cfe-dev <<a href="mailto:cfe-dev@lists.llvm.org" target="_blank">cfe-dev@lists.llvm.org</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">
<br>
<br>
On 14 Jan 2019, at 23:23, Finkel, Hal J. via cfe-dev wrote:<br>
<br>
> On 1/14/19 6:22 PM, Alexey Bataev wrote:<br>
>> Yes, it is implemented in the frontend. No, it is not implemented in <br>
>> the Sema analysis, it is implemented in Codegen part of the frontend. <br>
>> And it is almost impossible to implement it in Sema. It requires <br>
>> recursive analysis of the functions used (not only called directly, <br>
>> but also indirectly). The better place to implement it is the CodeGen <br>
>> of the frontend.<br>
><br>
> Your proposal makes sense to me. And, I agree, just moving the logic <br>
> to<br>
> CodeGen seems like the most-straightforward solution. I don't see any<br>
> conceptual problem with inline assembly being checked a bit later in <br>
> the<br>
> pipeline (it is, after all, quite target specific).<br>
<br>
We generally prefer to diagnose things in Sema rather than IRGen, and <br>
one<br>
reason is that we want these diagnostics to show up in clients like <br>
IDEs.<br>
<br>
Obviously, ASM constraint validation is a pretty minor diagnostic, but I<br>
don't see what naturally limits this to just that. In fact, hasn't this<br>
come up before?<br>
<br>
If this is really just specific to ASM constraint validation, I think we<br>
can find a way to delay the diagnosis of those to IRGen conditionally.<br>
Otherwise, I think you need to find a way to suppress diagnostics from<br>
functions that you don't care about.<br>
<br>
John.<br>
_______________________________________________<br>
cfe-dev mailing list<br>
<a href="mailto:cfe-dev@lists.llvm.org" target="_blank">cfe-dev@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev</a><br>
</blockquote>
</div>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote></div>