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

Given a global variable @gg, the compiler has to generate code on the host to
specify whether it is passed by value or passed by reference. In the following
example, if the compiler generates the code for passing by value, the outlined
function on the target i386-pc-linux-gnu cannot get the correct value since it
assumes the variable @gg is passed by reference.

Here is the corresponding IR on the host side.
  %0 = load double, double* @gg, align 8, !tbaa !3
  %1 = bitcast double %0 to i64
   ...
  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0,
i32 2
  %13 = bitcast i8** %12 to i64*
  store i64 %1, i64* %13, align 8
  ...
  %20 = call i32 @__tgt_target(i64 -1, i8* @.omp_offload.region_id, i32 4, i8**
%4, i8** %6, i64* getelementptr inbounds ([4 x i64], [4 x i64]* @.offload_sizes,
i32 0, i32 0), i64* getelementptr inbounds ([4 x i64], [4 x i64]*
@.offload_maptypes, i32 0, i32 0))

Thanks,

Jin

From: Friedman, Eli [mailto:efriedma at codeaurora.org]
Sent: Wednesday, April 25, 2018 3:18 PM
To: Lin, Jin <jin.lin at intel.com>; 'llvm-dev at lists.llvm.org'
<llvm-dev at lists.llvm.org>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information
in the LLVM IR

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/43369acf/attachment.html>

On 4/25/2018 3:48 PM, Lin, Jin wrote:
>
> Given a global variable @gg, the compiler has to generate code on the 
> host to specify whether it is passed by value or passed by reference. 
> In the following example, if the compiler generates the code for 
> passing by value, the outlined function on the target 
> i386-pc-linux-gnucannot get the correct value since it assumes the 
> variable @gg is passed by reference.
>
> Here is the corresponding IR on the host side.
>
>   %0 = load double, double* @gg, align 8, !tbaa !3
>
>   %1 = bitcast double %0 to i64
>
>    …
>
>   %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* 
> %.offload_baseptrs, i32 0, i32 2
>
>   %13 = bitcast i8** %12 to i64*
>
>   store i64 %1, i64* %13, align 8
>
Could you describe the overall process of calling an offloaded function 
in a bit more detail?  How do you describe the ABI of the called 
function to the OpenMP runtime?

I suspect you shouldn't be trying to store things which aren't pointers 
into offload_baseptrs.

-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/18567e1b/attachment.html>

For the firstprivate clause, the compiler generates code to pass it  by value or
by reference to the outlined function. The reason the first private scalars is
generally passed by value is for the performance reason.
For this particular case, the compiler cannot generate code to pass the double
@gg by value under i386-pc-linux-gnu since the value is 64 bit while the
architecture is 32bit.
For the host compilation, the compiler generates the code to pass the data as
well as the outlined function name to the OMP runtime.
For the target compilation, the compiler generates the outlined function so that
it can be called by the OMP runtime.
So, the compiler is required to generate a single call on the host to support
all the targets. All the target versions must have the same interface. So the
common interface of the outline function should be used. For this particular
example, the variable @gcc should be passed by reference under x86_64-mic.
Please let me know if you have more questions.
Jin

From: Friedman, Eli [mailto:efriedma at codeaurora.org]
Sent: Wednesday, April 25, 2018 4:14 PM
To: Lin, Jin <jin.lin at intel.com>; 'llvm-dev at lists.llvm.org'
<llvm-dev at lists.llvm.org>
Subject: Re: [llvm-dev] [LLVM][RFC] Representing the target device information
in the LLVM IR

On 4/25/2018 3:48 PM, Lin, Jin wrote:
Given a global variable @gg, the compiler has to generate code on the host to
specify whether it is passed by value or passed by reference. In the following
example, if the compiler generates the code for passing by value, the outlined
function on the target i386-pc-linux-gnu cannot get the correct value since it
assumes the variable @gg is passed by reference.

Here is the corresponding IR on the host side.
  %0 = load double, double* @gg, align 8, !tbaa !3
  %1 = bitcast double %0 to i64
   ...
  %12 = getelementptr inbounds [4 x i8*], [4 x i8*]* %.offload_baseptrs, i32 0,
i32 2
  %13 = bitcast i8** %12 to i64*
  store i64 %1, i64* %13, align 8

Could you describe the overall process of calling an offloaded function in a bit
more detail?  How do you describe the ABI of the called function to the OpenMP
runtime?

I suspect you shouldn't be trying to store things which aren't pointers
into offload_baseptrs.

-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/20180426/01291d18/attachment.html>