jinlin created this revision.
jinlin added a reviewer: hfinkel.
Herald added a subscriber: cfe-commits.
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"
Repository:
rC Clang
https://reviews.llvm.org/D46074
Files:
lib/CodeGen/ModuleBuilder.cpp
Index: lib/CodeGen/ModuleBuilder.cpp
===================================================================
--- lib/CodeGen/ModuleBuilder.cpp
+++ lib/CodeGen/ModuleBuilder.cpp
@@ -131,6 +131,18 @@
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,
Index: lib/CodeGen/ModuleBuilder.cpp
===================================================================
--- lib/CodeGen/ModuleBuilder.cpp
+++ lib/CodeGen/ModuleBuilder.cpp
@@ -131,6 +131,18 @@
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,
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
