<html>
<head>
<meta http-equiv="content-type" content="text/html; charset=UTF-8">
</head>
<body text="#000000" bgcolor="#FFFFFF">
<p>Hi, many of the OpenMP users experience troubles when they try to
compile real-world OpenMP applications, which use offloading
constructs.<br>
<br>
Problem Description<br>
===============<br>
For example, let’s look at the abstract code:<br>
```<br>
void target_unused() {<br>
int a;<br>
__asm__("constraints"<br>
: "=a"(a)<br>
: // None<br>
);<br>
}<br>
<br>
void target_used() {<br>
int a;<br>
__asm__("constraints"<br>
: "=a"(a)<br>
: // None<br>
);<br>
}<br>
<br>
void foo() {<br>
target_unused();<br>
#pragma omp target<br>
target_used();<br>
}<br>
```<br>
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:<br>
```<br>
11:13: error: invalid output constraint '=a' in asm<br>
20:13: error: invalid output constraint '=a' in asm<br>
```<br>
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.<br>
<br>
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.<br>
<br>
That’s why we need to postpone some of the target-related
diagnostics till the CodeGen phase.<br>
<br>
Possible solution.<br>
==============<br>
1. Move target-specific checks to the CodeGen.<br>
<br>
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.<br>
<br>
2. Add special delayed diagnostics class, associate it with the
declarations somehow and check it for each declaration during the
CodeGen.</p>
<p>Requires redesign of the diagnostic subsystem + redesign of the
codegen subsystem, at least.<br>
</p>
<p>3. Introduce special expression, statement and declaration nodes.<br>
<br>
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.<br>
<br>
Requires additional expression/statement/declarations nodes +
looks like a hack.<br>
</p>
<pre class="moz-signature" cols="72">--
-------------
Best regards,
Alexey Bataev</pre>
</body>
</html>