Representing the target device information in the LLVM IR

[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"

[ABataev]

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.

[efriedma]

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

docs/LangRef.rst

::
 
     ARCHITECTURE-VENDOR-OPERATING_SYSTEM
     ARCHITECTURE-VENDOR-OPERATING_SYSTEM-ENVIRONMENT
 
 This information is passed along to the backend so that it generates
 code for the proper architecture. It's possible to override this on the
 command line with the ``-mtriple`` command line option.
 
+.. _target_devices_triples:
+
+Devices Triples
+---------------
+A module may specify comma-separated list of triples that describes the
+OpenMP offloading targets to be supported. The list of triples is denoted
+as devices triple, whose syntax is simply:
+
+.. code-block:: llvm
+
+    target device_triples = "x86_64-mic,i386-pc-linux-gnu"
+
+The *devices triples* string consists of a series of substrings separated by
+','. Each substring which consists of a series of identifiers joined by the
+minus sign character ('-'), is in the same format of the *target triple*.
+
+This information is passed along to the backend so that it generates
+code for the proper architecture.
+
 .. _pointeraliasing:
 
 Pointer Aliasing Rules
 ----------------------
 
 Any memory access must be done through a pointer value associated with
 an address range of the memory access, otherwise the behavior is
 undefined. Pointer values are associated with address ranges according

include/llvm/Bitcode/LLVMBitCodes.h

enum ModuleCodes {
   // SOURCE_FILENAME: [namechar x N]
   MODULE_CODE_SOURCE_FILENAME = 16,
 
   // HASH: [5*i32]
   MODULE_CODE_HASH = 17,
 
   // IFUNC: [ifunc value type, addrspace, resolver val#, linkage, visibility]
   MODULE_CODE_IFUNC = 18,
+
+  // TRIPLE:      [strchr x N, ..., strchr x N]
+  MODULE_CODE_DEVICES = 19,
 };
 
 /// PARAMATTR blocks have code for defining a parameter attribute set.
 enum AttributeCodes {
   // FIXME: Remove `PARAMATTR_CODE_ENTRY_OLD' in 4.0
   PARAMATTR_CODE_ENTRY_OLD = 1, // ENTRY: [paramidx0, attr0,
                                 //         paramidx1, attr1...]
   PARAMATTR_CODE_ENTRY = 2,     // ENTRY: [attrgrp0, attrgrp1, ...]

include/llvm/IR/Module.h

private:
   std::string ModuleID;           ///< Human readable identifier for the module
   std::string SourceFileName;     ///< Original source file name for module,
                                   ///< recorded in bitcode.
   std::string TargetTriple;       ///< Platform target triple Module compiled on
                                   ///< Format: (arch)(sub)-(vendor)-(sys0-(abi)
   void *NamedMDSymTab;            ///< NamedMDNode names.
   DataLayout DL;                  ///< DataLayout associated with the module
 
+  std::string TargetDevices;      ///< Target devices
+
   friend class Constant;
 
 /// @}
 /// @name Constructors
 /// @{
 public:
   /// The Module constructor. Note that there is no default constructor. You
   /// must provide a name for the module upon construction.
/// @{
 
   /// Get the data layout for the module's target platform.
   const DataLayout &getDataLayout() const;
 
   /// Get the target triple which is a string describing the target host.
   /// @returns a string containing the target triple.
   const std::string &getTargetTriple() const { return TargetTriple; }
 
+  /// Get the target device information which is a comma-separated string
+  /// describing one or more devices.
+  const std::string &getTargetDevices() const { return TargetDevices; }
+
   /// Get the global data context.
   /// @returns LLVMContext - a container for LLVM's global information
   LLVMContext &getContext() const { return Context; }
 
   /// Get any module-scope inline assembly blocks.
   /// @returns a string containing the module-scope inline assembly blocks.
   const std::string &getModuleInlineAsm() const { return GlobalScopeAsm; }
 
/// @{
 
   /// Set the data layout
   void setDataLayout(StringRef Desc);
   void setDataLayout(const DataLayout &Other);
 
   /// Set the target triple.
   void setTargetTriple(StringRef T) { TargetTriple = T; }
 
+  /// set the target device information.
+  void setTargetDevices(StringRef T) { TargetDevices = T; }
+
   /// Set the module-scope inline assembly blocks.
   /// A trailing newline is added if the input doesn't have one.
   void setModuleInlineAsm(StringRef Asm) {
     GlobalScopeAsm = Asm;
     if (!GlobalScopeAsm.empty() && GlobalScopeAsm.back() != '\n')
       GlobalScopeAsm += '\n';
   }
 

lib/AsmParser/LLLexer.cpp

#define KEYWORD(STR)                                                           \
   KEYWORD(caller);
   KEYWORD(within);
   KEYWORD(from);
   KEYWORD(tail);
   KEYWORD(musttail);
   KEYWORD(notail);
   KEYWORD(target);
   KEYWORD(triple);
+  KEYWORD(device_triples);
   KEYWORD(source_filename);
   KEYWORD(unwind);
   KEYWORD(deplibs);             // FIXME: Remove in 4.0.
   KEYWORD(datalayout);
   KEYWORD(volatile);
   KEYWORD(atomic);
   KEYWORD(unordered);
   KEYWORD(monotonic);

lib/AsmParser/LLParser.cpp

bool LLParser::ParseTargetDefinition() {
   default: return TokError("unknown target property");
   case lltok::kw_triple:
     Lex.Lex();
     if (ParseToken(lltok::equal, "expected '=' after target triple") ||
         ParseStringConstant(Str))
       return true;
     M->setTargetTriple(Str);
     return false;
+  case lltok::kw_device_triples:
+    Lex.Lex();
+    if (ParseToken(lltok::equal, "expected '=' after target devices") ||
+        ParseStringConstant(Str))
+      return true;
+    M->setTargetDevices(Str);
+    return false;
   case lltok::kw_datalayout:
     Lex.Lex();
     if (ParseToken(lltok::equal, "expected '=' after target datalayout") ||
         ParseStringConstant(Str))
       return true;
     if (DataLayoutStr.empty())
       M->setDataLayout(Str);
     return false;

lib/AsmParser/LLToken.h

enum Kind {
   kw_caller,
   kw_within,
   kw_from,
   kw_tail,
   kw_musttail,
   kw_notail,
   kw_target,
   kw_triple,
+  kw_device_triples,
   kw_source_filename,
   kw_unwind,
   kw_deplibs, // FIXME: Remove in 4.0
   kw_datalayout,
   kw_volatile,
   kw_atomic,
   kw_unordered,
   kw_monotonic,

lib/Bitcode/Reader/BitcodeReader.cpp

while (true) {
     }
     case bitc::MODULE_CODE_TRIPLE: {  // TRIPLE: [strchr x N]
       std::string S;
       if (convertToString(Record, 0, S))
         return error("Invalid record");
       TheModule->setTargetTriple(S);
       break;
     }
+    case bitc::MODULE_CODE_DEVICES: { // TRIPLE: [strchr x N, ..., strchr x N]
+      std::string S;
+      if (convertToString(Record, 0, S))
+        return error("Invalid record");
+      TheModule->setTargetDevices(S);
+      break;
+    }
     case bitc::MODULE_CODE_DATALAYOUT: {  // DATALAYOUT: [strchr x N]
       std::string S;
       if (convertToString(Record, 0, S))
         return error("Invalid record");
       TheModule->setDataLayout(S);
       break;
     }
     case bitc::MODULE_CODE_ASM: {  // ASM: [strchr x N]

lib/Bitcode/Writer/BitcodeWriter.cpp

 /// Emit top-level description of module, including target triple, inline asm,
 /// descriptors for global variables, and function prototype info.
 /// Returns the bit offset to backpatch with the location of the real VST.
 void ModuleBitcodeWriter::writeModuleInfo() {
   // Emit various pieces of data attached to a module.
   if (!M.getTargetTriple().empty())
     writeStringRecord(Stream, bitc::MODULE_CODE_TRIPLE, M.getTargetTriple(),
                       0 /*TODO*/);
+  if (!M.getTargetDevices().empty())
+    writeStringRecord(Stream, bitc::MODULE_CODE_DEVICES, M.getTargetDevices(),
+                      0 /*TODO*/);
   const std::string &DL = M.getDataLayoutStr();
   if (!DL.empty())
     writeStringRecord(Stream, bitc::MODULE_CODE_DATALAYOUT, DL, 0 /*TODO*/);
   if (!M.getModuleInlineAsm().empty())
     writeStringRecord(Stream, bitc::MODULE_CODE_ASM, M.getModuleInlineAsm(),
                       0 /*TODO*/);
 
   // Emit information about sections and GC, computing how many there are. Also

lib/IR/AsmWriter.cpp

void AssemblyWriter::printModule(const Module *M) {
   }
 
   const std::string &DL = M->getDataLayoutStr();
   if (!DL.empty())
     Out << "target datalayout = \"" << DL << "\"\n";
   if (!M->getTargetTriple().empty())
     Out << "target triple = \"" << M->getTargetTriple() << "\"\n";
 
+  if (!M->getTargetDevices().empty())
+    Out << "target device_triples = \"" << M->getTargetDevices() << "\"\n";
+
   if (!M->getModuleInlineAsm().empty()) {
     Out << '\n';
 
     // Split the string into lines, to make it easier to read the .ll file.
     StringRef Asm = M->getModuleInlineAsm();
     do {
       StringRef Front;
       std::tie(Front, Asm) = Asm.split('\n');

lib/Transforms/Utils/CloneModule.cpp

std::unique_ptr<Module> llvm::CloneModule(
     const Module &M, ValueToValueMapTy &VMap,
     function_ref<bool(const GlobalValue *)> ShouldCloneDefinition) {
   // First off, we need to create the new module.
   std::unique_ptr<Module> New =
       llvm::make_unique<Module>(M.getModuleIdentifier(), M.getContext());
   New->setSourceFileName(M.getSourceFileName());
   New->setDataLayout(M.getDataLayout());
   New->setTargetTriple(M.getTargetTriple());
+  New->setTargetDevices(M.getTargetDevices());
   New->setModuleInlineAsm(M.getModuleInlineAsm());
 
   // Loop over all of the global variables, making corresponding globals in the
   // new module.  Here we add them to the VMap and to the new Module.  We
   // don't worry about attributes or initializers, they will come later.
   //
   for (Module::const_global_iterator I = M.global_begin(), E = M.global_end();
        I != E; ++I) {
if (!ShouldCloneDefinition(&*I)) {
       continue;
     }
     auto *GA = GlobalAlias::create(I->getValueType(),
                                    I->getType()->getPointerAddressSpace(),
                                    I->getLinkage(), I->getName(), New.get());
     GA->copyAttributesFrom(&*I);
     VMap[&*I] = GA;
   }
-  
   // Now that all of the things that global variable initializer can refer to
   // have been created, loop through and copy the global variable referrers
   // over...  We also set the attributes on the global now.
   //
   for (Module::const_global_iterator I = M.global_begin(), E = M.global_end();
        I != E; ++I) {
     if (I->isDeclaration())
       continue;