[cfe-dev] [RFC] Delayed target-specific diagnostic when compiling for the devices.

[Justin Lebar via cfe-dev]

>> 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?
> Yes, right. OpenMP does not require explicit marking of the device
function. It supports implicit target functions.

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?

(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.)

On Tue, Jan 15, 2019 at 2:26 PM Alexey Bataev <a.bataev at outlook.com> wrote:

> Yes, right. OpenMP does not require explicit marking of the device
> function. It supports implicit target functions.
>
> Best regards,
> Alexey Bataev
>
> 15 янв. 2019 г., в 17:24, Justin Lebar <jlebar at google.com> написал(а):
>
> > 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.
>
> 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?
>
> On Tue, Jan 15, 2019 at 2:20 PM Alexey Bataev <a.bataev at outlook.com>
> wrote:
>
>> This is not only for asm, we need to delay all target-specific
>> diagnostics.
>> I'm not saying that we need to move the host diagnostic, only the
>> diagnostic for the device compilation.
>> 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.
>>
>> Best regards,
>> Alexey Bataev
>>
>> 15 янв. 2019 г., в 17:07, Reid Kleckner <rnk at google.com> написал(а):
>>
>> 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.
>>
>> On Tue, Jan 15, 2019 at 1:58 PM John McCall via cfe-dev <
>> cfe-dev at lists.llvm.org> wrote:
>>
>>>
>>>
>>> On 14 Jan 2019, at 23:23, Finkel, Hal J. via cfe-dev wrote:
>>>
>>> > On 1/14/19 6:22 PM, Alexey Bataev wrote:
>>> >> Yes, it is implemented in the frontend. No, it is not implemented in
>>> >> the Sema analysis, it is implemented in Codegen part of the frontend.
>>> >> And it is almost impossible to implement it in Sema. It requires
>>> >> recursive analysis of the functions used (not only called directly,
>>> >> but also indirectly). The better place to implement it is the CodeGen
>>> >> of the frontend.
>>> >
>>> > Your proposal makes sense to me. And, I agree, just moving the logic
>>> > to
>>> > CodeGen seems like the most-straightforward solution. I don't see any
>>> > conceptual problem with inline assembly being checked a bit later in
>>> > the
>>> > pipeline (it is, after all, quite target specific).
>>>
>>> We generally prefer to diagnose things in Sema rather than IRGen, and
>>> one
>>> reason is that we want these diagnostics to show up in clients like
>>> IDEs.
>>>
>>> Obviously, ASM constraint validation is a pretty minor diagnostic, but I
>>> don't see what naturally limits this to just that.  In fact, hasn't this
>>> come up before?
>>>
>>> If this is really just specific to ASM constraint validation, I think we
>>> can find a way to delay the diagnosis of those to IRGen conditionally.
>>> Otherwise, I think you need to find a way to suppress diagnostics from
>>> functions that you don't care about.
>>>
>>> John.
>>> _______________________________________________
>>> cfe-dev mailing list
>>> cfe-dev at lists.llvm.org
>>> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev
>>>
>>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20190115/c08f91f3/attachment.html>