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

[Lin, Jin via llvm-dev]

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>