[CUDA] Allow parsing of host and device code simultaneously.

[tra]

Summary

The patch:

• adds -aux-triple option to specify target triple
• propagates aux target info to AST context and Preprocessor
• pulls in target specific preprocessor macros.
• pulls in target-specific builtins from aux target.
• sets appropriate host or device attribute on builtins.

Rationale:

In order to compile CUDA source with mixed host/device code without physically separating host and device code, we need to be able to parse code that may contain target-specific constructs from both host and device targets.

During device-side compilation we need to be able to include any host includes we may encounter. Similarly, during host compilation, we do need to include device-side includes that may be required for device-specific code in the file to parse. In both cases, we need to fake target environment well enough for the headers to work (I.e. x86 host's headers want to see amd64 or i386 defined, CUDA includes are looking for NVPTX-specific macros).

We also need to be able to parse target-specific builtins from both host and device targets in the same TU.

Generally speaking it's not possible to achieve this in all cases. Fortunately, CUDA's case is simpler and proposed patch works pretty well in practice:

• clang already implements attribute-based function overloading which allows avoiding name clashes between host and device functions.
• basic host and device types are expected to match, so a lot of type-related predefined macros from host and device targets have the same value.
• host includes (x86 on linux) do not use predefined NVPTX-specific macros, so including them is not a problem.
• builtins from the aux target are only used for parsing and AST construction. CUDA never generates IR for them. If a builtin is used from a wrong context, it will violate calling restriction and produce an error.
• this change includes *all* builtins from the aux target which provides us with a superset of builtins that would be available on the opposite side of the compilation which will allow creating different ASTs on host and device sides. IMO it's similar to already-existing issue of diverging host/device code caused by common (ab)use of CUDA_ARCH macro.

[jpienaar]

Nice, so this will allow parsing/AST construction with builtins from 2 architectures but will fail to compile if a builtin for the host/device is called from device/host.

You mention this is not generally possible. Can you give some examples?

include/clang/Driver/CC1Options.td

329

You use aux target in all the errors to the user so perhaps for consistency "Triple for aux target". It could be more self-documenting too ("Triple for aux target used during CUDA compilation."?) as I don't know if a lot of people would be able to guess what the auxiliary triple is or where it is used.

include/clang/Frontend/CompilerInstance.h

354

Nit: period at the end for uniformity.

[tra]

cosmetic fixes.

[tra]

In D12917#247868, @jpienaar wrote:

Nice, so this will allow parsing/AST construction with builtins from 2 architectures but will fail to compile if a builtin for the host/device is called from device/host.

It CUDA target call checks are in effect, we'll get a usual error that H->D or D->H call is illegal. If we're compiling with -fcuda-disable-target-call-checks, then we'll fail during codegen phase complaining that builtin has not been implemented.

You mention this is not generally possible. Can you give some examples?

It's mostly related to mixing include files and preprocessor macros from multiple targets.
If an include is written to conditionally provide mutually exclusive constructs for targets A and B, then compiling for A with -aux-triple B will produce something uncompilable. For instance following include will fail due to redefinition of foo() if compiled with -triple x86_64 -aux-triple nvptx64

#ifdef __amd64__
int foo() {return 1;}
#endif
#ifdef __PTX__
int foo() {return 2;}
#endif

include/clang/Driver/CC1Options.td

329

Changed to "Auxiliary target triple." While CUDA uses the aux target triple for preprocessor macros and builtins, nothing stops anyone else using this information for other purposes, so I kept the name and the comment generic.

[echristo]

One inline request and one inline comment, otherwise looks pretty good!

Thanks :)

-eric

include/clang/Frontend/CompilerInstance.h

354

Go ahead and just commit this separately please. :)

lib/Frontend/CompilerInstance.cpp

82–87

Seems like an icky place to do this, perhaps where we create the Target?

[tra]

Updated to address Eric's comments.

[tra]

lib/Frontend/CompilerInstance.cpp

82–87

Added separate setter method for AuxTarget and moved TargetInfo creation to CompilerInstance::ExecuteAction as you've suggested.

[echristo]

Works for me. Thanks!

-eric

include/clang/AST/ASTContext.h

private:
   bool AddrSpaceMapMangling;
 
   friend class ASTDeclReader;
   friend class ASTReader;
   friend class ASTWriter;
   friend class CXXRecordDecl;
 
   const TargetInfo *Target;
+  const TargetInfo *AuxTarget;
   clang::PrintingPolicy PrintingPolicy;
   
 public:
   IdentifierTable &Idents;
   SelectorTable &Selectors;
   Builtin::Context &BuiltinInfo;
   mutable DeclarationNameTable DeclarationNames;
   IntrusiveRefCntPtr<ExternalASTSource> ExternalSource;
public:
   /// Return the total memory used for various side tables.
   size_t getSideTableAllocatedMemory() const;
   
   PartialDiagnostic::StorageAllocator &getDiagAllocator() {
     return DiagAllocator;
   }
 
   const TargetInfo &getTargetInfo() const { return *Target; }
-  
+  const TargetInfo *getAuxTargetInfo() const { return AuxTarget; }
   /// getIntTypeForBitwidth -
   /// sets integer QualTy according to specified details:
   /// bitwidth, signed/unsigned.
   /// Returns empty type if there is no appropriate target types.
   QualType getIntTypeForBitwidth(unsigned DestWidth,
                                  unsigned Signed) const;
   /// getRealTypeForBitwidth -
   /// sets floating point QualTy according to specified bitwidth.
private:
   void operator=(const ASTContext &) = delete;
 
 public:
   /// \brief Initialize built-in types.
   ///
   /// This routine may only be invoked once for a given ASTContext object.
   /// It is normally invoked after ASTContext construction.
   ///
-  /// \param Target The target 
-  void InitBuiltinTypes(const TargetInfo &Target);
-  
+  /// \param Target The target
+  void InitBuiltinTypes(const TargetInfo &Target,
+                        const TargetInfo *AuxTarget = nullptr);
 private:
   void InitBuiltinType(CanQualType &R, BuiltinType::Kind K);
 
   // Return the Objective-C type encoding for a given type.
   void getObjCEncodingForTypeImpl(QualType t, std::string &S,
                                   bool ExpandPointedToStructures,
                                   bool ExpandStructures,
                                   const FieldDecl *Field,

include/clang/Basic/Builtins.h

 struct Info {
   const char *Name, *Type, *Attributes, *HeaderName;
   LanguageID Langs;
   const char *Features;
 };
 
 /// \brief Holds information about both target-independent and
 /// target-specific builtins, allowing easy queries by clients.
+///
+/// Builtins from an optional auxiliary target are stored in
+/// AuxTSRecords. Their IDs are shifted up by NumTSRecords and need to
+/// be translated back with getAuxBuiltinID() before use.
 class Context {
   const Info *TSRecords;
+  const Info *AuxTSRecords;
   unsigned NumTSRecords;
+  unsigned NumAuxTSRecords;
+
 public:
   Context();
 
   /// \brief Perform target-specific initialization
-  void initializeTarget(const TargetInfo &Target);
-  
+  /// \param AuxTarget Target info to incorporate builtins from. May be nullptr.
+  void InitializeTarget(const TargetInfo &Target, const TargetInfo *AuxTarget);
   /// \brief Mark the identifiers for all the builtins with their
   /// appropriate builtin ID # and mark any non-portable builtin identifiers as
   /// such.
   void initializeBuiltins(IdentifierTable &Table, const LangOptions& LangOpts);
 
   /// \brief Return the identifier name for the specified builtin,
   /// e.g. "__builtin_abs".
   const char *getName(unsigned ID) const {
public:
   bool isConstWithoutErrno(unsigned ID) const {
     return strchr(getRecord(ID).Attributes, 'e') != nullptr;
   }
 
   const char *getRequiredFeatures(unsigned ID) const {
     return getRecord(ID).Features;
   }
 
+  /// \brief Return true if builtin ID belongs to AuxTarget.
+  bool isAuxBuiltinID(unsigned ID) const {
+    return ID >= (Builtin::FirstTSBuiltin + NumTSRecords);
+  }
+
+  /// Return real buitin ID (i.e. ID it would have furing compilation
+  /// for AuxTarget).
+  unsigned getAuxBuiltinID(unsigned ID) const { return ID - NumTSRecords; }
+
 private:
   const Info &getRecord(unsigned ID) const;
 
   /// \brief Is this builtin supported according to the given language options?
   bool builtinIsSupported(const Builtin::Info &BuiltinInfo,
                           const LangOptions &LangOpts);
 
   /// \brief Helper function for isPrintfLike and isScanfLike.
   bool isLike(unsigned ID, unsigned &FormatIdx, bool &HasVAListArg,
               const char *Fmt) const;
 };
 
 }
 } // end namespace clang
 #endif

include/clang/Driver/CC1Options.td

 // This isn't normally used, it is just here so we can parse a
 // CompilerInvocation out of a driver-derived argument vector.
 def cc1 : Flag<["-"], "cc1">;
 def cc1as : Flag<["-"], "cc1as">;
 
 def ast_merge : Separate<["-"], "ast-merge">,
   MetaVarName<"<ast file>">,
   HelpText<"Merge the given AST file into the translation unit being compiled.">;
+def aux_triple : Separate<["-"], "aux-triple">,
+  HelpText<"Auxiliary target triple.">;
 def code_completion_at : Separate<["-"], "code-completion-at">,
   MetaVarName<"<file>:<line>:<column>">,
   HelpText<"Dump code-completion information at a location">;
 def remap_file : Separate<["-"], "remap-file">,
   MetaVarName<"<from>;<to>">,
   HelpText<"Replace the contents of the <from> file with the contents of the <to> file">;
 def code_completion_at_EQ : Joined<["-"], "code-completion-at=">,
   Alias<code_completion_at>;

include/clang/Frontend/CompilerInstance.h

class CompilerInstance : public ModuleLoader {
   IntrusiveRefCntPtr<CompilerInvocation> Invocation;
 
   /// The diagnostics engine instance.
   IntrusiveRefCntPtr<DiagnosticsEngine> Diagnostics;
 
   /// The target being compiled for.
   IntrusiveRefCntPtr<TargetInfo> Target;
 
+  /// Auxiliary Target info.
+  IntrusiveRefCntPtr<TargetInfo> AuxTarget;
+
   /// The virtual file system.
   IntrusiveRefCntPtr<vfs::FileSystem> VirtualFileSystem;
 
   /// The file manager.
   IntrusiveRefCntPtr<FileManager> FileMgr;
 
   /// The source manager.
   IntrusiveRefCntPtr<SourceManager> SourceMgr;
public:
 
   bool hasTarget() const { return Target != nullptr; }
 
   TargetInfo &getTarget() const {
     assert(Target && "Compiler instance has no target!");
     return *Target;
   }
 
   /// Replace the current Target.
   void setTarget(TargetInfo *Value);
 
   /// }
+  /// @name AuxTarget Info
+  /// {
+
+  TargetInfo *getAuxTarget() const { return AuxTarget.get(); }
+
+  /// Replace the current AuxTarget.
+  void setAuxTarget(TargetInfo *Value);
+
+  /// }
   /// @name Virtual File System
   /// {
 
   bool hasVirtualFileSystem() const { return VirtualFileSystem != nullptr; }
 
   vfs::FileSystem &getVirtualFileSystem() const {
     assert(hasVirtualFileSystem() &&
            "Compiler instance has no virtual file system");

include/clang/Frontend/FrontendOptions.h

public:
 
   /// \brief A list of arguments to forward to LLVM's option processing; this
   /// should only be used for debugging and experimental features.
   std::vector<std::string> LLVMArgs;
 
   /// \brief File name of the file that will provide record layouts
   /// (in the format produced by -fdump-record-layouts).
   std::string OverrideRecordLayoutsFile;
-  
+  /// \brief Auxiliary triple for CUDA compilation.
+  std::string AuxTriple;
+
 public:
   FrontendOptions() :
     DisableFree(false), RelocatablePCH(false), ShowHelp(false),
     ShowStats(false), ShowTimers(false), ShowVersion(false),
     FixWhatYouCan(false), FixOnlyWarnings(false), FixAndRecompile(false),
     FixToTemporaries(false), ARCMTMigrateEmitARCErrors(false),
     SkipFunctionBodies(false), UseGlobalModuleIndex(true),
     GenerateGlobalModuleIndex(true), ASTDumpDecls(false), ASTDumpLookups(false),

include/clang/Lex/Preprocessor.h

 /// Lexers know only about tokens within a single source file, and don't
 /// know anything about preprocessor-level issues like the \#include stack,
 /// token expansion, etc.
 class Preprocessor : public RefCountedBase<Preprocessor> {
   IntrusiveRefCntPtr<PreprocessorOptions> PPOpts;
   DiagnosticsEngine        *Diags;
   LangOptions       &LangOpts;
   const TargetInfo  *Target;
+  const TargetInfo  *AuxTarget;
   FileManager       &FileMgr;
   SourceManager     &SourceMgr;
   std::unique_ptr<ScratchBuffer> ScratchBuf;
   HeaderSearch      &HeaderInfo;
   ModuleLoader      &TheModuleLoader;
 
   /// \brief External source of macros.
   ExternalPreprocessorSource *ExternalSource;
Preprocessor(IntrusiveRefCntPtr<PreprocessorOptions> PPOpts,
                TranslationUnitKind TUKind = TU_Complete);
 
   ~Preprocessor();
 
   /// \brief Initialize the preprocessor using information about the target.
   ///
   /// \param Target is owned by the caller and must remain valid for the
   /// lifetime of the preprocessor.
-  void Initialize(const TargetInfo &Target);
+  /// \param AuxTarget is owned by the caller and must remain valid for
+  /// the lifetime of the preprocessor.
+  void Initialize(const TargetInfo &Target,
+                  const TargetInfo *AuxTarget = nullptr);
 
   /// \brief Initialize the preprocessor to parse a model file
   ///
   /// To parse model files the preprocessor of the original source is reused to
   /// preserver the identifier table. However to avoid some duplicate
   /// information in the preprocessor some cleanup is needed before it is used
   /// to parse model files. This method does that cleanup.
   void InitializeForModelFile();
 
   /// \brief Cleanup after model file parsing
   void FinalizeForModelFile();
 
   /// \brief Retrieve the preprocessor options used to initialize this
   /// preprocessor.
   PreprocessorOptions &getPreprocessorOpts() const { return *PPOpts; }
   
   DiagnosticsEngine &getDiagnostics() const { return *Diags; }
   void setDiagnostics(DiagnosticsEngine &D) { Diags = &D; }
 
   const LangOptions &getLangOpts() const { return LangOpts; }
   const TargetInfo &getTargetInfo() const { return *Target; }
+  const TargetInfo *getAuxTargetInfo() const { return AuxTarget; }
   FileManager &getFileManager() const { return FileMgr; }
   SourceManager &getSourceManager() const { return SourceMgr; }
   HeaderSearch &getHeaderSearchInfo() const { return HeaderInfo; }
 
   IdentifierTable &getIdentifierTable() { return Identifiers; }
   const IdentifierTable &getIdentifierTable() const { return Identifiers; }
   SelectorTable &getSelectorTable() { return Selectors; }
   Builtin::Context &getBuiltinInfo() { return BuiltinInfo; }

lib/AST/ASTContext.cpp

: FunctionProtoTypes(this_()), TemplateSpecializationTypes(this_()),
       ObjCClassDecl(nullptr), ObjCProtocolClassDecl(nullptr), BOOLDecl(nullptr),
       CFConstantStringTypeDecl(nullptr), ObjCInstanceTypeDecl(nullptr),
       FILEDecl(nullptr), jmp_bufDecl(nullptr), sigjmp_bufDecl(nullptr),
       ucontext_tDecl(nullptr), BlockDescriptorType(nullptr),
       BlockDescriptorExtendedType(nullptr), cudaConfigureCallDecl(nullptr),
       FirstLocalImport(), LastLocalImport(), ExternCContext(nullptr),
       SourceMgr(SM), LangOpts(LOpts),
       SanitizerBL(new SanitizerBlacklist(LangOpts.SanitizerBlacklistFiles, SM)),
-      AddrSpaceMap(nullptr), Target(nullptr), PrintingPolicy(LOpts),
-      Idents(idents), Selectors(sels), BuiltinInfo(builtins),
-      DeclarationNames(*this), ExternalSource(nullptr), Listener(nullptr),
-      Comments(SM), CommentsLoaded(false),
+      AddrSpaceMap(nullptr), Target(nullptr), AuxTarget(nullptr),
+      PrintingPolicy(LOpts), Idents(idents), Selectors(sels),
+      BuiltinInfo(builtins), DeclarationNames(*this), ExternalSource(nullptr),
+      Listener(nullptr), Comments(SM), CommentsLoaded(false),
       CommentCommandTraits(BumpAlloc, LOpts.CommentOpts), LastSDM(nullptr, 0) {
   TUDecl = TranslationUnitDecl::Create(*this);
 }
 
 ASTContext::~ASTContext() {
   ReleaseParentMapEntries();
 
   // Release the DenseMaps associated with DeclContext objects.

 }
 
 void ASTContext::InitBuiltinType(CanQualType &R, BuiltinType::Kind K) {
   BuiltinType *Ty = new (*this, TypeAlignment) BuiltinType(K);
   R = CanQualType::CreateUnsafe(QualType(Ty, 0));
   Types.push_back(Ty);
 }
 
-void ASTContext::InitBuiltinTypes(const TargetInfo &Target) {
+void ASTContext::InitBuiltinTypes(const TargetInfo &Target,
+                                  const TargetInfo *AuxTarget) {
   assert((!this->Target || this->Target == &Target) &&
          "Incorrect target reinitialization");
   assert(VoidTy.isNull() && "Context reinitialized?");
 
   this->Target = &Target;
-  
+  this->AuxTarget = AuxTarget;
   ABI.reset(createCXXABI(Target));
   AddrSpaceMap = getAddressSpaceMap(Target, LangOpts);
   AddrSpaceMapMangling = isAddrSpaceMapManglingEnabled(Target, LangOpts);
   
   // C99 6.2.5p19.
   InitBuiltinType(VoidTy,              BuiltinType::Void);
 
   // C99 6.2.5p2.

lib/Basic/Builtins.cpp

 #define LIBBUILTIN(ID, TYPE, ATTRS, HEADER, LANGS)                             \
   { #ID, TYPE, ATTRS, HEADER, LANGS, nullptr },
 #include "clang/Basic/Builtins.def"
 };
 
 const Builtin::Info &Builtin::Context::getRecord(unsigned ID) const {
   if (ID < Builtin::FirstTSBuiltin)
     return BuiltinInfo[ID];
-  assert(ID - Builtin::FirstTSBuiltin < NumTSRecords && "Invalid builtin ID!");
+  assert(ID - Builtin::FirstTSBuiltin < (NumTSRecords + NumAuxTSRecords) &&
+         "Invalid builtin ID!");
+  if (isAuxBuiltinID(ID))
+    return AuxTSRecords[getAuxBuiltinID(ID) - Builtin::FirstTSBuiltin];
   return TSRecords[ID - Builtin::FirstTSBuiltin];
 }
 
 Builtin::Context::Context() {
   // Get the target specific builtins from the target.
   TSRecords = nullptr;
+  AuxTSRecords = nullptr;
   NumTSRecords = 0;
+  NumAuxTSRecords = 0;
 }
 
-void Builtin::Context::initializeTarget(const TargetInfo &Target) {
+void Builtin::Context::InitializeTarget(const TargetInfo &Target,
+                                        const TargetInfo *AuxTarget) {
   assert(NumTSRecords == 0 && "Already initialized target?");
   Target.getTargetBuiltins(TSRecords, NumTSRecords);
+  if (AuxTarget)
+    AuxTarget->getTargetBuiltins(AuxTSRecords, NumAuxTSRecords);
 }
 
 bool Builtin::Context::builtinIsSupported(const Builtin::Info &BuiltinInfo,
                                           const LangOptions &LangOpts) {
   bool BuiltinsUnsupported = LangOpts.NoBuiltin &&
                              strchr(BuiltinInfo.Attributes, 'f');
   bool MathBuiltinsUnsupported =
     LangOpts.NoMathBuiltin && BuiltinInfo.HeaderName &&
void Builtin::Context::initializeBuiltins(IdentifierTable &Table,
   for (unsigned i = Builtin::NotBuiltin+1; i != Builtin::FirstTSBuiltin; ++i)
     if (builtinIsSupported(BuiltinInfo[i], LangOpts)) {
       Table.get(BuiltinInfo[i].Name).setBuiltinID(i);
     }
 
   // Step #2: Register target-specific builtins.
   for (unsigned i = 0, e = NumTSRecords; i != e; ++i)
     if (builtinIsSupported(TSRecords[i], LangOpts))
-      Table.get(TSRecords[i].Name).setBuiltinID(i+Builtin::FirstTSBuiltin);
+      Table.get(TSRecords[i].Name).setBuiltinID(i + Builtin::FirstTSBuiltin);
+
+  // Step #3: Register target-specific builtins for AuxTarget.
+  for (unsigned i = 0, e = NumAuxTSRecords; i != e; ++i)
+    Table.get(AuxTSRecords[i].Name)
+        .setBuiltinID(i + Builtin::FirstTSBuiltin + NumTSRecords);
 }
 
 void Builtin::Context::forgetBuiltin(unsigned ID, IdentifierTable &Table) {
   Table.get(getRecord(ID).Name).setBuiltinID(0);
 }
 
 bool Builtin::Context::isLike(unsigned ID, unsigned &FormatIdx,
                               bool &HasVAListArg, const char *Fmt) const {

lib/CodeGen/CGBuiltin.cpp

if (Value *V = EmitTargetBuiltinExpr(BuiltinID, E))
     return RValue::get(V);
 
   ErrorUnsupported(E, "builtin function");
 
   // Unknown builtin, for now just dump it out and return undef.
   return GetUndefRValue(E->getType());
 }
 
-Value *CodeGenFunction::EmitTargetBuiltinExpr(unsigned BuiltinID,
-                                              const CallExpr *E) {
-  switch (getTarget().getTriple().getArch()) {
+static Value *EmitTargetArchBuiltinExpr(CodeGenFunction *CGF,
+                                        unsigned BuiltinID, const CallExpr *E,
+                                        llvm::Triple::ArchType Arch) {
+  switch (Arch) {
   case llvm::Triple::arm:
   case llvm::Triple::armeb:
   case llvm::Triple::thumb:
   case llvm::Triple::thumbeb:
-    return EmitARMBuiltinExpr(BuiltinID, E);
+    return CGF->EmitARMBuiltinExpr(BuiltinID, E);
   case llvm::Triple::aarch64:
   case llvm::Triple::aarch64_be:
-    return EmitAArch64BuiltinExpr(BuiltinID, E);
+    return CGF->EmitAArch64BuiltinExpr(BuiltinID, E);
   case llvm::Triple::x86:
   case llvm::Triple::x86_64:
-    return EmitX86BuiltinExpr(BuiltinID, E);
+    return CGF->EmitX86BuiltinExpr(BuiltinID, E);
   case llvm::Triple::ppc:
   case llvm::Triple::ppc64:
   case llvm::Triple::ppc64le:
-    return EmitPPCBuiltinExpr(BuiltinID, E);
+    return CGF->EmitPPCBuiltinExpr(BuiltinID, E);
   case llvm::Triple::r600:
   case llvm::Triple::amdgcn:
-    return EmitAMDGPUBuiltinExpr(BuiltinID, E);
+    return CGF->EmitAMDGPUBuiltinExpr(BuiltinID, E);
   case llvm::Triple::systemz:
-    return EmitSystemZBuiltinExpr(BuiltinID, E);
+    return CGF->EmitSystemZBuiltinExpr(BuiltinID, E);
   case llvm::Triple::nvptx:
   case llvm::Triple::nvptx64:
-    return EmitNVPTXBuiltinExpr(BuiltinID, E);
+    return CGF->EmitNVPTXBuiltinExpr(BuiltinID, E);
   case llvm::Triple::wasm32:
   case llvm::Triple::wasm64:
-    return EmitWebAssemblyBuiltinExpr(BuiltinID, E);
+    return CGF->EmitWebAssemblyBuiltinExpr(BuiltinID, E);
   default:
     return nullptr;
   }
 }
 
+Value *CodeGenFunction::EmitTargetBuiltinExpr(unsigned BuiltinID,
+                                              const CallExpr *E) {
+  if (getContext().BuiltinInfo.isAuxBuiltinID(BuiltinID)) {
+    assert(getContext().getAuxTargetInfo() && "Missing aux target info");
+    return EmitTargetArchBuiltinExpr(
+        this, getContext().BuiltinInfo.getAuxBuiltinID(BuiltinID), E,
+        getContext().getAuxTargetInfo()->getTriple().getArch());
+  }
+
+  return EmitTargetArchBuiltinExpr(this, BuiltinID, E,
+                                   getTarget().getTriple().getArch());
+}
+
 static llvm::VectorType *GetNeonType(CodeGenFunction *CGF,
                                      NeonTypeFlags TypeFlags,
                                      bool V1Ty=false) {
   int IsQuad = TypeFlags.isQuad();
   switch (TypeFlags.getEltType()) {
   case NeonTypeFlags::Int8:
   case NeonTypeFlags::Poly8:
     return llvm::VectorType::get(CGF->Int8Ty, V1Ty ? 1 : (8 << IsQuad));

lib/Frontend/CompilerInstance.cpp

return (BuildGlobalModuleIndex ||
            getFrontendOpts().GenerateGlobalModuleIndex)) &&
          !ModuleBuildFailed;
 }
 
 void CompilerInstance::setDiagnostics(DiagnosticsEngine *Value) {
   Diagnostics = Value;
 }
 
-void CompilerInstance::setTarget(TargetInfo *Value) {
-  Target = Value;
-}
+void CompilerInstance::setTarget(TargetInfo *Value) { Target = Value; }
+void CompilerInstance::setAuxTarget(TargetInfo *Value) { AuxTarget = Value; }
 
 void CompilerInstance::setFileManager(FileManager *Value) {
   FileMgr = Value;
   if (Value)
     VirtualFileSystem = Value->getVirtualFileSystem();
   else
     VirtualFileSystem.reset();
 }
 
 void CompilerInstance::setSourceManager(SourceManager *Value) {
   SourceMgr = Value;
 }
 
void CompilerInstance::createPreprocessor(TranslationUnitKind TUKind) {
   HeaderSearch *HeaderInfo = new HeaderSearch(&getHeaderSearchOpts(),
                                               getSourceManager(),
                                               getDiagnostics(),
                                               getLangOpts(),
                                               &getTarget());
   PP = new Preprocessor(&getPreprocessorOpts(), getDiagnostics(), getLangOpts(),
                         getSourceManager(), *HeaderInfo, *this, PTHMgr,
                         /*OwnsHeaderSearch=*/true, TUKind);
-  PP->Initialize(getTarget());
+  PP->Initialize(getTarget(), getAuxTarget());
 
   // Note that this is different then passing PTHMgr to Preprocessor's ctor.
   // That argument is used as the IdentifierInfoLookup argument to
   // IdentifierTable's ctor.
   if (PTHMgr) {
     PTHMgr->setPreprocessor(&*PP);
     PP->setPTHManager(PTHMgr);
   }

 
 // ASTContext
 
 void CompilerInstance::createASTContext() {
   Preprocessor &PP = getPreprocessor();
   auto *Context = new ASTContext(getLangOpts(), PP.getSourceManager(),
                                  PP.getIdentifierTable(), PP.getSelectorTable(),
                                  PP.getBuiltinInfo());
-  Context->InitBuiltinTypes(getTarget());
+  Context->InitBuiltinTypes(getTarget(), getAuxTarget());
   setASTContext(Context);
 }
 
 // ExternalASTSource
 
 void CompilerInstance::createPCHExternalASTSource(
     StringRef Path, bool DisablePCHValidation, bool AllowPCHWithCompilerErrors,
     void *DeserializationListener, bool OwnDeserializationListener) {
bool CompilerInstance::ExecuteAction(FrontendAction &Act) {
   raw_ostream &OS = llvm::errs();
 
   // Create the target instance.
   setTarget(TargetInfo::CreateTargetInfo(getDiagnostics(),
                                          getInvocation().TargetOpts));
   if (!hasTarget())
     return false;
 
+  // Create TargetInfo for the other side of CUDA compilation.
+  if (getLangOpts().CUDA && !getFrontendOpts().AuxTriple.empty()) {
+    std::shared_ptr<TargetOptions> TO(new TargetOptions);
+    TO->Triple = getFrontendOpts().AuxTriple;
+    setAuxTarget(TargetInfo::CreateTargetInfo(getDiagnostics(), TO));
+  }
+
   // Inform the target of the language options.
   //
   // FIXME: We shouldn't need to do this, the target should be immutable once
   // created. This complexity should be lifted elsewhere.
   getTarget().adjust(getLangOpts());
 
   // rewriter project will change target built-in bool type from its default. 
   if (getFrontendOpts().ProgramAction == frontend::RewriteObjC)

lib/Frontend/CompilerInvocation.cpp

Opts.CodeCompleteOpts.IncludeCodePatterns
     = Args.hasArg(OPT_code_completion_patterns);
   Opts.CodeCompleteOpts.IncludeGlobals
     = !Args.hasArg(OPT_no_code_completion_globals);
   Opts.CodeCompleteOpts.IncludeBriefComments
     = Args.hasArg(OPT_code_completion_brief_comments);
 
   Opts.OverrideRecordLayoutsFile
     = Args.getLastArgValue(OPT_foverride_record_layout_EQ);
+  Opts.AuxTriple =
+      llvm::Triple::normalize(Args.getLastArgValue(OPT_aux_triple));
+
   if (const Arg *A = Args.getLastArg(OPT_arcmt_check,
                                      OPT_arcmt_modify,
                                      OPT_arcmt_migrate)) {
     switch (A->getOption().getID()) {
     default:
       llvm_unreachable("missed a case");
     case OPT_arcmt_check:
       Opts.ARCMTAction = FrontendOptions::ARCMT_Check;

lib/Frontend/InitPreprocessor.cpp

void clang::InitializePreprocessor(
   // Emit line markers for various builtin sections of the file.  We don't do
   // this in asm preprocessor mode, because "# 4" is not a line marker directive
   // in this mode.
   if (!PP.getLangOpts().AsmPreprocessor)
     Builder.append("# 1 \"<built-in>\" 3");
 
   // Install things like __POWERPC__, __GNUC__, etc into the macro table.
   if (InitOpts.UsePredefines) {
+    if (LangOpts.CUDA && PP.getAuxTargetInfo())
+      InitializePredefinedMacros(*PP.getAuxTargetInfo(), LangOpts, FEOpts,
+                                 Builder);
+
     InitializePredefinedMacros(PP.getTargetInfo(), LangOpts, FEOpts, Builder);
 
     // Install definitions to make Objective-C++ ARC work well with various
     // C++ Standard Library implementations.
     if (LangOpts.ObjC1 && LangOpts.CPlusPlus && LangOpts.ObjCAutoRefCount) {
       switch (InitOpts.ObjCXXARCStandardLibrary) {
       case ARCXX_nolib:
         case ARCXX_libcxx:

lib/Lex/Preprocessor.cpp

 
 Preprocessor::Preprocessor(IntrusiveRefCntPtr<PreprocessorOptions> PPOpts,
                            DiagnosticsEngine &diags, LangOptions &opts,
                            SourceManager &SM, HeaderSearch &Headers,
                            ModuleLoader &TheModuleLoader,
                            IdentifierInfoLookup *IILookup, bool OwnsHeaders,
                            TranslationUnitKind TUKind)
     : PPOpts(PPOpts), Diags(&diags), LangOpts(opts), Target(nullptr),
-      FileMgr(Headers.getFileMgr()), SourceMgr(SM),
-      ScratchBuf(new ScratchBuffer(SourceMgr)),HeaderInfo(Headers),
+      AuxTarget(nullptr), FileMgr(Headers.getFileMgr()), SourceMgr(SM),
+      ScratchBuf(new ScratchBuffer(SourceMgr)), HeaderInfo(Headers),
       TheModuleLoader(TheModuleLoader), ExternalSource(nullptr),
       Identifiers(opts, IILookup),
       PragmaHandlers(new PragmaNamespace(StringRef())),
-      IncrementalProcessing(false), TUKind(TUKind),
-      CodeComplete(nullptr), CodeCompletionFile(nullptr),
-      CodeCompletionOffset(0), LastTokenWasAt(false),
-      ModuleImportExpectsIdentifier(false), CodeCompletionReached(0),
-      MainFileDir(nullptr), SkipMainFilePreamble(0, true), CurPPLexer(nullptr),
-      CurDirLookup(nullptr), CurLexerKind(CLK_Lexer), CurSubmodule(nullptr),
-      Callbacks(nullptr), CurSubmoduleState(&NullSubmoduleState),
-      MacroArgCache(nullptr), Record(nullptr),
-      MIChainHead(nullptr), DeserialMIChainHead(nullptr) {
+      IncrementalProcessing(false), TUKind(TUKind), CodeComplete(nullptr),
+      CodeCompletionFile(nullptr), CodeCompletionOffset(0),
+      LastTokenWasAt(false), ModuleImportExpectsIdentifier(false),
+      CodeCompletionReached(0), MainFileDir(nullptr),
+      SkipMainFilePreamble(0, true), CurPPLexer(nullptr), CurDirLookup(nullptr),
+      CurLexerKind(CLK_Lexer), CurSubmodule(nullptr), Callbacks(nullptr),
+      CurSubmoduleState(&NullSubmoduleState), MacroArgCache(nullptr),
+      Record(nullptr), MIChainHead(nullptr), DeserialMIChainHead(nullptr) {
   OwnsHeaderSearch = OwnsHeaders;
   
   CounterValue = 0; // __COUNTER__ starts at 0.
   
   // Clear stats.
   NumDirectives = NumDefined = NumUndefined = NumPragma = 0;
   NumIf = NumElse = NumEndif = 0;
   NumEnteredSourceFiles = 0;
Preprocessor::~Preprocessor() {
   for (MacroArgs *ArgList = MacroArgCache; ArgList;)
     ArgList = ArgList->deallocate();
 
   // Delete the header search info, if we own it.
   if (OwnsHeaderSearch)
     delete &HeaderInfo;
 }
 
-void Preprocessor::Initialize(const TargetInfo &Target) {
+void Preprocessor::Initialize(const TargetInfo &Target,
+                              const TargetInfo *AuxTarget) {
   assert((!this->Target || this->Target == &Target) &&
          "Invalid override of target information");
   this->Target = &Target;
-  
+  assert((!this->AuxTarget || this->AuxTarget == AuxTarget) &&
+         "Invalid override of aux target information.");
+  this->AuxTarget = AuxTarget;
+
   // Initialize information about built-ins.
-  BuiltinInfo.initializeTarget(Target);
+  BuiltinInfo.InitializeTarget(Target, AuxTarget);
   HeaderInfo.setTarget(Target);
 }
 
 void Preprocessor::InitializeForModelFile() {
   NumEnteredSourceFiles = 0;
 
   // Reset pragmas
   PragmaHandlersBackup = std::move(PragmaHandlers);

lib/Sema/SemaDecl.cpp

if (Context.BuiltinInfo.isReturnsTwice(BuiltinID) &&
                                          FD->getLocation()));
     if (Context.BuiltinInfo.isNoThrow(BuiltinID) && !FD->hasAttr<NoThrowAttr>())
       FD->addAttr(NoThrowAttr::CreateImplicit(Context, FD->getLocation()));
     if (Context.BuiltinInfo.isConst(BuiltinID) && !FD->hasAttr<ConstAttr>())
       FD->addAttr(ConstAttr::CreateImplicit(Context, FD->getLocation()));
     if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads &&
         Context.BuiltinInfo.isTSBuiltin(BuiltinID) &&
         !FD->hasAttr<CUDADeviceAttr>() && !FD->hasAttr<CUDAHostAttr>()) {
-      // Target-specific builtins are assumed to be intended for use
-      // in this particular CUDA compilation mode and should have
-      // appropriate attribute set so we can enforce CUDA function
-      // call restrictions.
-      if (getLangOpts().CUDAIsDevice)
+      // Assign appropriate attribute depending on CUDA compilation
+      // mode and the target builtin belongs to. E.g. during host
+      // compilation, aux builtins are __device__, the rest are __host__.
+      if (getLangOpts().CUDAIsDevice !=
+          Context.BuiltinInfo.isAuxBuiltinID(BuiltinID))
         FD->addAttr(CUDADeviceAttr::CreateImplicit(Context, FD->getLocation()));
       else
         FD->addAttr(CUDAHostAttr::CreateImplicit(Context, FD->getLocation()));
     }
   }
 
   IdentifierInfo *Name = FD->getIdentifier();
   if (!Name)

test/SemaCUDA/builtins.cu

-// Tests that target-specific builtins have appropriate host/device
-// attributes and that CUDA call restrictions are enforced. Also
-// verify that non-target builtins can be used from both host and
-// device functions.
+// Tests that host and target builtins can be used in the same TU,
+// have appropriate host/device attributes and that CUDA call
+// restrictions are enforced. Also verify that non-target builtins can
+// be used from both host and device functions.
 //
 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -triple x86_64-unknown-unknown \
+// RUN:     -aux-triple nvptx64-unknown-cuda \
 // RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
 // RUN: %clang_cc1 -triple nvptx64-unknown-cuda -fcuda-is-device \
+// RUN:     -aux-triple x86_64-unknown-unknown \
 // RUN:     -fcuda-target-overloads -fsyntax-only -verify %s
 
+#if !(defined(__amd64__) && defined(__PTX__))
+#error "Expected to see preprocessor macros from both sides of compilation."
+#endif
 
-#ifdef __CUDA_ARCH__
-// Device-side builtins are not allowed to be called from host functions.
 void hf() {
-  int x = __builtin_ptx_read_tid_x(); // expected-note  {{'__builtin_ptx_read_tid_x' declared here}}
+  int x = __builtin_ia32_rdtsc();
+  int y = __builtin_ptx_read_tid_x(); // expected-note  {{'__builtin_ptx_read_tid_x' declared here}}
   // expected-error@-1 {{reference to __device__ function '__builtin_ptx_read_tid_x' in __host__ function}}
   x = __builtin_abs(1);
 }
+
 __attribute__((device)) void df() {
   int x = __builtin_ptx_read_tid_x();
+  int y = __builtin_ia32_rdtsc(); // expected-error {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
+                                  // expected-note@20 {{'__builtin_ia32_rdtsc' declared here}}
   x = __builtin_abs(1);
 }
-#else
-// Host-side builtins are not allowed to be called from device functions.
-__attribute__((device)) void df() {
-  int x = __builtin_ia32_rdtsc();   // expected-note {{'__builtin_ia32_rdtsc' declared here}}
-  // expected-error@-1 {{reference to __host__ function '__builtin_ia32_rdtsc' in __device__ function}}
-  x = __builtin_abs(1);
-}
-void hf() {
-  int x = __builtin_ia32_rdtsc();
-  x = __builtin_abs(1);
-}
-#endif