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

[Alexey Bataev via cfe-dev]

Ping!

-------------
Best regards,
Alexey Bataev

09.01.2019 10:02, Alexey.Bataev пишет:
>
> Hi, many of the OpenMP users experience troubles when they try to
> compile real-world OpenMP applications, which use offloading constructs.
>
> Problem Description
> ===============
> For example, let’s look at the abstract code:
> ```
> void target_unused() {
>   int a;
>   __asm__("constraints"
>           : "=a"(a)
>           :         // None
>   );
> }
>
> void target_used() {
>   int a;
>   __asm__("constraints"
>           : "=a"(a)
>           :         // None
>   );
> }
>
> void foo() {
>   target_unused();
> #pragma omp target
>   target_used();
> }
> ```
> Assume, we going to compile this code on X86_64 host to run on the
> NVidia NVPTX64 device. When we compile this code for the host,
> everything is good. But when we compile the same code for the NVPTX64
> target, we get the next error messages:
> ```
> 11:13: error: invalid output constraint '=a' in asm
> 20:13: error: invalid output constraint '=a' in asm
> ```
> But, actually, we should see only one error message, the second one,
> for the function `target_used()`, which is actually used in the target
> region. The second function, `target_unused()` is used only on the
> host and we should no produce error message for this function when we
> compile the code for the device.
>
> The main problem with those functions is that they are not marked
> explicitly as the device functions, just like it is required in CUDA.
> In OpenMP, it is not required to mark them explicitly as the
> device-only or both device-host function. They can be marked
> implicitly, depending of the fact that they are used in target-based
> constructs (probably, indirectly, through chain of calls) or not.
>
> That’s why we need to postpone some of the target-related diagnostics
> till the CodeGen phase.
>
> Possible solution.
> ==============
> 1. Move target-specific checks to the CodeGen.
>
> The best solutions of all, does not require significant redesign of
> the existing code base, just requires some copy-paste of the
> diagnostics and the associated logic.
>
> 2. Add special delayed diagnostics class, associate it with the
> declarations somehow and check it for each declaration during the CodeGen.
>
> Requires redesign of the diagnostic subsystem + redesign of the
> codegen subsystem, at least.
>
> 3. Introduce special expression, statement and declaration nodes.
>
> This nodes serve only as the containers for the target-specific error
> messages, generated in Sema. During the codegen phase, if the codegen
> run into one of such constructs, it just emits the diagnostic, stored
> in these nodes.
>
> Requires additional expression/statement/declarations nodes + looks
> like a hack.
>
> -- 
> -------------
> Best regards,
> Alexey Bataev

-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 833 bytes
Desc: OpenPGP digital signature
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20190114/bb0621d9/attachment.sig>