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

[Lin, Jin via llvm-dev]

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>