[llvm-dev] [LLVM][RFC] Representing the target device information in the LLVM IR

[Friedman, Eli via llvm-dev]

On 4/25/2018 3:05 PM, Lin, Jin via llvm-dev wrote:
>
> RFC: Representing the target device information in the LLVM IR
>
> ===========================================================================
>
> Why this RFC change?
>
> =================
>
> The target device information needs to be passed to the LLVM backend 
> when OpenMP backend outlining is enabled. For example, for multiple 
> target devices, the target compilation has to generate a single host 
> to support all the targets. In order to make sure all the target 
> outlined functions have the same interface, the information of all the 
> target architectures is needed during host and target compilation. In 
> the following example, the firstprivate variable ‘d’ is passed by 
> value under x86_64-mic and passed by reference under 
> i386-pc-linux-gnu. In order to avoid this inconsistency, the compiler 
> needs all the target architecture information so it can find a common 
> interface. In this example, it will change the x86_64-mic interface 
> for ‘d’ to pass by reference.
>
> Existing code: 64-bit firstprivate variable
>
> void foo() {
>
>  double d = 1.0;
>
>  #pragma omp target firstprivate(d)
>
>  {}
>
> }
>
> $clang –fopenmp-backend -fopenmp-targets=x86_64-mic, i386-pc-linux-gnu …
>
> x86_64-mic
>
> define void @__omp_offloading…(i64 %d) #0 {
>
> entry:
>
> …
>
> }
>
> i386-pc-linux-gnu
>
> define void @__omp_offloading…(double* dereferenceable(8) %d) #0 {
>
> entry:
>
>  …
>
> }
>
> There is an inconsistency between host and target part(s) of the program!
>

I don't see how this inconsistency is a problem... at least, not on its 
own.  The host code doesn't call either of these functions directly; it 
calls the OpenMP runtime, which should invoke the offloaded function 
correctly.  (If it doesn't, that's a bug in the OpenMP lowering, not the 
LLVM backend.)

-Eli

-- 
Employee of Qualcomm Innovation Center, Inc.
Qualcomm Innovation Center, Inc. is a member of Code Aurora Forum, a Linux Foundation Collaborative Project

-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20180425/e0b0cf33/attachment.html>