Page MenuHomePhabricator

Representing the target device information in the LLVM IR
Needs ReviewPublic

Authored by jinlin on Apr 25 2018, 10:39 AM.



The target device information needs to be passed to the LLVM backend when the OMP backend outlining is enabled. For example, for multiple target devices, the target compilation has to generate one single host to support all the targets. In order to make sure all the target outline function have the same interface, the information of all the target architecture are needed during host and target compilation. In the following example, the firstprivate variable d is represented as passing by value under x86_64 mic and passing by reference under i386-pc-linux-gnu. In order to fix this inconsistency issue, the compiler can change the form of firstprivate d from passing by value under x86_64 mic to passing by reference after the compiler has all the target architecture information.

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 …


define void @__omp_offloading…(i64 %d) #0 {



define void @__omp_offloading…(double* dereferenceable(8) %d) #0 {


There is an inconsistency between host and target part(s) of the program!

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

/ Get the target device information which is a string separated by the
/ comma to describe one or more than one device.
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"

Diff Detail

Event Timeline

jinlin created this revision.Apr 25 2018, 10:39 AM
jinlin updated this revision to Diff 143964.Apr 25 2018, 10:48 AM
jinlin added a reviewer: efriedma.
jinlin updated this revision to Diff 143965.Apr 25 2018, 10:50 AM
jinlin updated this revision to Diff 143966.Apr 25 2018, 10:56 AM
jinlin updated this revision to Diff 143967.
jinlin set the repository for this revision to rOMP OpenMP.Apr 25 2018, 11:00 AM

Hmm, I don't think this is correct. I think it is against the OpenMP standard.
According to OpenMP 4.5:
When an original variable is mapped to a device data environment and the associated 19 corresponding variable is not present in the device data environment, a new corresponding variable 20 (of the same type and size as the original variable) is created in the device data environment.

If you try to offload the code from x86_64 to i386 with long int type, for example, you won't be able to create the variable of the same size. I think, that this kind of offloading should be just not allowed.

Also, it means, that the code generated by the backend will be different rather than the code emitted by the frontend. You're going to break the ABI using some LLVM transformations.

That's very strange that you started with the patches. This is a very important thing that requires a discussion. I think you'd better start with the RFC. Currently, all this stuff looks like some kind of a hack.

I agree this needs an RFC; I don't understand how you plan to use this information.