Representing the target device information in the LLVM IR ( CLANG changes)

[jinlin]

Summary

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 …

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!

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"

lib/CodeGen/ModuleBuilder.cpp

llvm::Module *StartModule(llvm::StringRef ModuleName,
       Initialize(*Ctx);
       return M.get();
     }
 
     void Initialize(ASTContext &Context) override {
       Ctx = &Context;
 
       M->setTargetTriple(Ctx->getTargetInfo().getTriple().getTriple());
+
+      // The target device information is represented as module level
+      // attribute.
+      SmallString<128> Res;
+      for (auto &Device : Ctx->getLangOpts().OMPTargetTriples) {
+        if (!Res.empty())
+          Res += ",";
+        Res += Device.getTriple();
+      }
+      if (!Res.empty())
+        M->setTargetDevices(Res);
+
       M->setDataLayout(Ctx->getTargetInfo().getDataLayout());
       Builder.reset(new CodeGen::CodeGenModule(Context, HeaderSearchOpts,
                                                PreprocessorOpts, CodeGenOpts,
                                                *M, Diags, CoverageInfo));
 
       for (auto &&Lib : CodeGenOpts.DependentLibraries)
         Builder->AddDependentLib(Lib);
       for (auto &&Opt : CodeGenOpts.LinkerOptions)