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

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!

Change Made
=========We proposed new module level attribute to represent target device
information.

/// Get the target device information, which is a comma-separated string
/// describing one or more devices.
const std::string &getTargetDevices() const { return TargetDevices; }

/// set the target device information.
void setTargetDevices(StringRef T) { TargetDevices = T; }

IR Dump (the extension indicated in red font)
target triple = "x86_64-unknown-linux-gnu"
target device_triples = "x86_64-mic,i386-pc-linux-gnu"

=======Summary:
=======We propose a new module-level attribute to represent the target device
information embedded in the LLVM IR. The change is simple and straightforward.
The patch is uploaded with this RFC for code review.
https://reviews.llvm.org/D46071
https://reviews.llvm.org/D46074


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

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>

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>