[HIP][Clang][CodeGen][RFC] Add codegen support for C++ Parallel Algorithm Offload

[AlexVlx]

Summary

This patch adds the CodeGen changes needed by the standard algorithm offload feature being proposed here: https://discourse.llvm.org/t/rfc-adding-c-parallel-algorithm-offload-support-to-clang-llvm/72159/1, which will only be available for the HIP language on AMD targets. The verbose documentation is included in the head of the patch series. This change concludes the set of additions needed in Clang, and essentially relaxes restrictions on what gets emitted on the device path, when compiling in hipstdpar mode (after the previous patch relaxed restrictions on what is semantically correct):

1. Unless a function is explicitly marked __host__, it will get emitted, whereas before only __device__ and __global__ functions would be emitted;
2. Unsupported builtins are ignored as opposed to being marked as an error, as the decision on their validity is deferred to the hipstdpar specific code selection pass we are adding, which will be the topic of the final patch in this series;
3. We add the stdpar specific passes to the opt pipeline, independent of optimisation level:
   • When compiling for the accelerator / offload device, we add a code selection pass;
   • When compiling for the host, iff the user requested it via the --hipstdpar-interpose-alloc flag, we add a pass which replaces canonical allocation / deallocation functions with accelerator aware equivalents.

A test to validate that unannotated functions get correctly emitted is added as well. Please note that __device__, __global__ and __host__ are used to match existing nomenclature, they would not be present in user code.

[efriedma]

clang/lib/CodeGen/CGBuiltin.cpp

5785

This doesn't make sense; we can't just ignore bits of the source code. I guess this is related to "the decision on their validity is deferred", but I don't see how you expect this to work.

[AlexVlx]

clang/lib/CodeGen/CGBuiltin.cpp

5785

This is one of the weirder parts, so let's consider the following example:

cpp
void foo() { __builtin_ia32_pause(); }
void bar() { __builtin_trap(); }

void baz(const vector<int>& v) {
    return for_each(par_unseq, cbegin(v), cend(v), [](auto&& x) { if (x == 42) bar(); });
}

In the case above, what we'd offload to the accelerator, and ask the target BE to lower, is the implementation of for_each, and bar, because it is reachable from the latter. foo is not reachable by any execution path on the accelerator side, however it includes a builtin that is unsupported by the accelerator (unless said accelerator is x86, which is not impossible, but not something we're dealing with at the moment). If we were to actually error out early, in the FE, in these cases, there's almost no appeal to what is being proposed, because standard headers, as well as other libraries, are littered with various target specific builtins that are not going to be supported. This all builds on the core invariant of this model / extension / thingamabob, which is that the algorithms, and only the algorithms, are targets for offload. It thus follows that as long as code that is reachable from an algorithm's implementation is safe, all is fine, but we cannot know this in the FE / on an AST level, because we need the actual CFG. This part is handled in LLVM in the SelectAcceleratorCodePass that's in the last patch in this series.

Now, you will quite correctly observe that there's nothing preventing an user from calling foo in the callable they pass to an algorithm; they might read the docs / appreciate that this won't work, but even there they are not safe, because there via some opaque call chain they might end up touching some unsupported builtin. My intuition here, which is reflected above in letting builtins just flow through, is that such cases are better served with a compile time error, which is what will obtain once the target BE chokes trying to lower an unsupported builtin. It's not going to be a beautiful error, and we could probably prettify it somewhat if we were to check after we've done the accelerator code selection pass, but it will happen at compile time. Another solution would be to emit these as traps (poison + trap for value returning ones), but I am concerned that it would lead to really fascinating debug journeys.

Having said this, if there's a better way to deal with these scenarios, it would be rather nice. Similarly, if the above doesn't make sense, please let me know.

[efriedma]

clang/lib/CodeGen/CGBuiltin.cpp

5785

Oh, I see; you "optimistically" compile everything assuming it might run on the accelerator, then run LLVM IR optimizations, then determine late which bits of code will actually run on the accelerator, which then prunes the code which shouldn't run.

I'm not sure I really like this... would it be possible to infer which functions need to be run on the accelerator based on the AST? I mean, if your API takes a lambda expression that runs on the accelerator, you can mark the lambda's body as "must be emitted for GPU", then recursively mark all the functions referred to by the lambda.

Emiting errors lazily from the backend means you get different diagnostics depending on the optimization level.

If you do go with this codegen-based approach, it's not clear to me how you detect that a forbidden builtin was called; if you skip the error handling, you just get a literal "undef".

[AlexVlx]

clang/lib/CodeGen/CGBuiltin.cpp

5785

I'm not sure I really like this... - actually, I am not a big fan either, however I think it's about the best one can do, given the constraints (consume standard C++, no annotations on the user side etc.). Having tried a few times in the past (and at least once in a different compiler), I don't quite think this can be done on an AST level. It would add some fairly awkward checking during template instantiation (no way to know earlier that a CallableFoo was passed to an offloadable algorithm), and it's a bit unwieldy to basically compute the CFG on the AST and mark reachable Callees at that point. Ignoring those, the main reason for which we cannot do this is that the interface is not constrained to only take lambdas, but callables in general, and that includes pointers to function as well. We don't deal with those today, but plan to, and there's a natural solution when operating on IR, assuming closed / internalised Modules (which is the case at least for AMDGPU at the moment). The final challenge pertains to the AST being per TU, with no cross-TU visibility, whereas with IR you can either pre-link the BC (implicitly or LTO) and then operate on the entire compilation. This is a problem with cases where foo defined in TU0 is reachable from algorithm_bar_offloaded_impl in TU1. So TL;DR, I think it would be more complex to do this on the AST and would end up more brittle / less future proof.

In what regards how to do deferred diagnostics, it think it can be done like this (I crossed streams in my prior reply when discussing this part, so it's actually nonsense): instead of emitting undef here, we can emit a builtin with the same signature, but with the name suffixed with e.g. (__stdpar_unsupported) or something similar. Then, when doing the reachability computation later, if we stumble upon a node in the CFG that contains a builtin suffixed with __stdpar_unsupported we error out, and can provide nice diagnostics since we'd have the call-chain handy. Thoughts?

[yaxunl]

Since we need to support -O0, we need to be prepared that we may not be able to remove all the calls of unsupported functions even though they may never be called at run time.

We could simply replace them with traps in the middle end. This should work if such functions are not called at run time. The only issue is that if they are called at run time, how do we tell users that they used unsupported functions and where. A trap does not help since it only crashes the program without useful information.

We could emit calls of any unsupported functions as calls of __clang_unsupported(file_name, line_number, function_name).

In the middle-end pass where we eliminate functions not referenced by kernels, we could emit reports about calls of __clang_unsupported under a certain -R option. We could turn on that option for -stdpar in clang driver.

We can emit printf of file_name, line_number and function_name for the first active lane then emit trap for a call of __clang_unsupported(file_name, line_number, function_name) under an option in the middle-end pass to facilitate users debugging their code.

[AlexVlx]

@yaxunl interesting point - are you worried about cases where due to missing inlining / const prop an indirect call site that can be replaced with a direct one would remain indirect? I think the problem in that case would actually be different, in that possibly reachable functions would not be identified as such and would be erroneously removed. I'm not sure there's any case where we'd fail to remove a meant to be unreachable function. We can definitely go with the __clang_unsupported approach, but I think I'd prefer these to be compile time errors rather than remarks + runtime printf, not in the least because printf adds some overhead. A way to ensure we don't "miss a spot" might be to check after removal for any remaining unsupported builtins, instead of doing it during reachability computation (this is coupled with the special naming from the prior post).

[efriedma]

In what regards how to do deferred diagnostics, it think it can be done like this (I crossed streams in my prior reply when discussing this part, so it's actually nonsense): instead of emitting undef here, we can emit a builtin with the same signature, but with the name suffixed with e.g. (__stdpar_unsupported) or something similar. Then, when doing the reachability computation later, if we stumble upon a node in the CFG that contains a builtin suffixed with __stdpar_unsupported we error out, and can provide nice diagnostics since we'd have the call-chain handy. Thoughts?

Sure, something like that. If you stick a SourceLocation on it, you can even recover the original clang source location.

We can definitely go with the __clang_unsupported approach, but I think I'd prefer these to be compile time errors rather than remarks + runtime printf, not in the least because printf adds some overhead.

The overhead should be pretty minimal if the code doesn't actually run.

So TL;DR, I think it would be more complex to do this on the AST and would end up more brittle / less future proof.

Since we need to support -O0

The biggest downside of working in the backend is that it becomes very hard for users to predict what will compile, and will not compile. Particularly if you want to support -O0. (I was sort of assuming you just wouldn't support -O0.) If you work on the AST, fewer constructs will be accepted, but you can actually define rules about which constructs will/will not be accepted.

[yaxunl]

In D155850#4523051, @AlexVlx wrote:

@yaxunl interesting point - are you worried about cases where due to missing inlining / const prop an indirect call site that can be replaced with a direct one would remain indirect? I think the problem in that case would actually be different, in that possibly reachable functions would not be identified as such and would be erroneously removed. I'm not sure there's any case where we'd fail to remove a meant to be unreachable function. We can definitely go with the __clang_unsupported approach, but I think I'd prefer these to be compile time errors rather than remarks + runtime printf, not in the least because printf adds some overhead. A way to ensure we don't "miss a spot" might be to check after removal for any remaining unsupported builtins, instead of doing it during reachability computation (this is coupled with the special naming from the prior post).

For programs having multiple TUs we cannot decide whether an unsupported function is used by a kernel during the compilation of a single TU. We can only decide that when we have the IR for the whole program. Currently, the HIP toolchain uses LTO of lld for multiple TUs, I am not sure whether we can emit clang diagnostics from lld. If not, then we need to use remarks. If we are confident to remove most unreachable unsupported functions at -O0, we may not need to use printf at run time. Remarks at LTO should be sufficient.

[AlexVlx]

This adds more ecumenical handling of unsupported builtins, as per the review discussion (a suffixed equivalent stub is emitted instead); it's paired with an associated change in accelerator code selection pass, where the actual check for these stubs occurs. I've also adjusted where the latter pass gets added to the opt pipeline, for the AMDGCN target; for the latter it's better, for the moment, to run it later because we essentially do LTCG, and therefore can unambiguously determine reachability by operating on the full module.

[efriedma]

clang/lib/CodeGen/CGBuiltin.cpp

5786

Else-after-return.

clang/lib/CodeGen/CodeGenModule.cpp

5339

You can't just pretend a thread-local variable isn't thread-local. If the intent here is that thread-local variables are illegal in device code, you need to figure out some way to produce a diagnostic. (Maybe by generating a call to __stdpar_unsupported_threadlocal or something like that if code tries to refer to such a variable.)

[AlexVlx]

clang/lib/CodeGen/CodeGenModule.cpp

5339

Oh, this is actually an error that slipped through, I botched the diff it appears, I'll correct it, apologies.

[AlexVlx]

Remove noise, correct style.

[AlexVlx]

Extend handling of unsupported builtins to include dealing with the target attribute.

[efriedma]

LGTM (but please don't merge until we reach consensus on the overall feature)

[AlexVlx]

In D155850#4570336, @efriedma wrote:

LGTM (but please don't merge until we reach consensus on the overall feature)

Of course, and thank you for the review. Please, do stick around if you don't mind, because this'll still get at least one update.

[yaxunl]

LGTM from HIP side. Thanks.

[arsenm]

clang/lib/CodeGen/BackendUtil.cpp

1105–1106

Formatting

[keryell]

clang/lib/CodeGen/CGBuiltin.cpp

5785

There is a lot of interesting design information in this discussion thread which will be lost forever after this is merged.
Is there a way to keep a summary as a comment somewhere to help the future readers/maintainers/historians?

[AlexVlx]

Add support for handling certain cases of unambiguously accelerator unsupported ASM i.e. cases where constraints are clearly mismatched. When that happens, we instead emit an ASM__stdpar_unsupported stub which takes as its single argument the constexpr string value of the ASM block. Later, in the AcceleratorCodeSelection pass, if such a stub is reachable from an accelerator callable, we error out and print the offending ASM alongside the location.

[AlexVlx]

Fix typo.

[yaxunl]

clang/lib/CodeGen/CGStmt.cpp

2424

maybe prefix with __ to avoid potential name collision with users' code

[AlexVlx]

Switch to __ASM prefix.

[AlexVlx]

Updating to reflect the outcome of the RFC, which is that this will be added as a HIP extension exclusively.

[AlexVlx]

Rebase.

[AlexVlx]

Use unmangled names in test.

[AlexVlx]

Rebase.

[hnrklssn]

clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp is failing: https://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/38041/testReport/junit/Clang/CodeGenHipStdPar/unannotated_functions_get_emitted_cpp/

project/clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp:15:22: error: NO-HIPSTDPAR-DEV: expected string not found in input
// NO-HIPSTDPAR-DEV: define {{.*}} void @bar({{.*}})
                     ^
<stdin>:1:1: note: scanning from here
; ModuleID = '/Users/buildslave/jenkins/workspace/clang-stage1-cmake-RA-incremental/llvm-project/clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp'
^
<stdin>:7:1: note: possible intended match here
define void @bar(ptr noundef %a, float noundef %b) #0 {
^

It looks like it may be due to the matcher having whitespace on both sides of {{.*}}, while the output only has a single space between define and void, but I'm not too well versed in FileCheck edge cases to know for sure.

[AlexVlx]

In D155850#4654146, @hnrklssn wrote:

clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp is failing: https://green.lab.llvm.org/green/job/clang-stage1-cmake-RA-incremental/38041/testReport/junit/Clang/CodeGenHipStdPar/unannotated_functions_get_emitted_cpp/

project/clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp:15:22: error: NO-HIPSTDPAR-DEV: expected string not found in input
// NO-HIPSTDPAR-DEV: define {{.*}} void @bar({{.*}})
                     ^
<stdin>:1:1: note: scanning from here
; ModuleID = '/Users/buildslave/jenkins/workspace/clang-stage1-cmake-RA-incremental/llvm-project/clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp'
^
<stdin>:7:1: note: possible intended match here
define void @bar(ptr noundef %a, float noundef %b) #0 {
^

It looks like it may be due to the matcher having whitespace on both sides of {{.*}}, while the output only has a single space between define and void, but I'm not too well versed in FileCheck edge cases to know for sure.

Thank you for the ping... this is pretty confusing since it's not tripping any of the buildbots, or flaring locally, let me look into it.

[AlexVlx]

Simplify test.

clang/lib/CodeGen/BackendUtil.cpp

 #include "llvm/Transforms/Instrumentation/PGOInstrumentation.h"
 #include "llvm/Transforms/Instrumentation/SanitizerBinaryMetadata.h"
 #include "llvm/Transforms/Instrumentation/SanitizerCoverage.h"
 #include "llvm/Transforms/Instrumentation/ThreadSanitizer.h"
 #include "llvm/Transforms/ObjCARC.h"
 #include "llvm/Transforms/Scalar/EarlyCSE.h"
 #include "llvm/Transforms/Scalar/GVN.h"
 #include "llvm/Transforms/Scalar/JumpThreading.h"
+#include "llvm/Transforms/HipStdPar/HipStdPar.h"
 #include "llvm/Transforms/Utils/Debugify.h"
 #include "llvm/Transforms/Utils/EntryExitInstrumenter.h"
 #include "llvm/Transforms/Utils/ModuleUtils.h"
 #include <memory>
 #include <optional>
 using namespace clang;
 using namespace llvm;
 
if (CodeGenOpts.FatLTO) {
     if (CodeGenOpts.UnifiedLTO && !TheModule->getModuleFlag("UnifiedLTO"))
       TheModule->addModuleFlag(Module::Error, "UnifiedLTO", uint32_t(1));
   }
 
   // Print a textual, '-passes=' compatible, representation of pipeline if
   // requested.
   if (PrintPipelinePasses) {
     MPM.printPipeline(outs(), [&PIC](StringRef ClassName) {
       auto PassName = PIC.getPassNameForClassName(ClassName);
       return PassName.empty() ? ClassName : PassName;
     });
     outs() << "\n";
     return;
   }
 
+  if (LangOpts.HIPStdPar && !LangOpts.CUDAIsDevice &&
+      LangOpts.HIPStdParInterposeAlloc)
+    MPM.addPass(HipStdParAllocationInterpositionPass());
+
   // Now that we have all of the passes ready, run them.
   {
     PrettyStackTraceString CrashInfo("Optimizer");
     llvm::TimeTraceScope TimeScope("Optimizer");
     MPM.run(*TheModule, MAM);
   }
 }
 

clang/lib/CodeGen/CGBuiltin.cpp

if (CGF.Builder.getIsFPConstrained() &&
       CGF.Builder.getDefaultConstrainedExcept() != fp::ebIgnore) {
     if (Value *Result =
             CGF.getTargetHooks().testFPKind(V, BuiltinID, CGF.Builder, CGF.CGM))
       return Result;
   }
   return nullptr;
 }
 
+static RValue EmitHipStdParUnsupportedBuiltin(CodeGenFunction *CGF,
+                                              const FunctionDecl *FD) {
+  auto Name = FD->getNameAsString() + "__hipstdpar_unsupported";
+  auto FnTy = CGF->CGM.getTypes().GetFunctionType(FD);
+  auto UBF = CGF->CGM.getModule().getOrInsertFunction(Name, FnTy);
+
+  SmallVector<Value *, 16> Args;
+  for (auto &&FormalTy : FnTy->params())
+    Args.push_back(llvm::PoisonValue::get(FormalTy));
+
+  return RValue::get(CGF->Builder.CreateCall(UBF, Args));
+}
+
 RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID,
                                         const CallExpr *E,
                                         ReturnValueSlot ReturnValue) {
   const FunctionDecl *FD = GD.getDecl()->getAsFunction();
   // See if we can constant fold this builtin.  If so, don't emit it at all.
   // TODO: Extend this handling to all builtin calls that we can constant-fold.
   Expr::EvalResult Result;
   if (E->isPRValue() && E->EvaluateAsRValue(Result, CGM.getContext()) &&
case TEK_Aggregate:
       return RValue::getAggregate(ReturnValue.getValue(),
                                   ReturnValue.isVolatile());
     case TEK_Complex:
       llvm_unreachable("No current target builtin returns complex");
     }
     llvm_unreachable("Bad evaluation kind in EmitBuiltinExpr");
   }
 
+  if (getLangOpts().HIPStdPar && getLangOpts().CUDAIsDevice)
+    return EmitHipStdParUnsupportedBuiltin(this, FD);
+
   ErrorUnsupported(E, "builtin function");
 
   // Unknown builtin, for now just dump it out and return undef.
   return GetUndefRValue(E->getType());
 }
 
 static Value *EmitTargetArchBuiltinExpr(CodeGenFunction *CGF,
                                         unsigned BuiltinID, const CallExpr *E,
                                         ReturnValueSlot ReturnValue,
                                         llvm::Triple::ArchType Arch) {
+  // When compiling in HipStdPar mode we have to be conservative in rejecting
+  // target specific features in the FE, and defer the possible error to the
+  // AcceleratorCodeSelection pass, wherein iff an unsupported target builtin is
+  // referenced by an accelerator executable function, we emit an error.
+  // Returning nullptr here leads to the builtin being handled in
+  // EmitStdParUnsupportedBuiltin.
+  if (CGF->getLangOpts().HIPStdPar && CGF->getLangOpts().CUDAIsDevice &&
+      Arch != CGF->getTarget().getTriple().getArch())
+    return nullptr;
+
   switch (Arch) {
   case llvm::Triple::arm:
   case llvm::Triple::armeb:
   case llvm::Triple::thumb:
   case llvm::Triple::thumbeb:
     return CGF->EmitARMBuiltinExpr(BuiltinID, E, ReturnValue, Arch);
   case llvm::Triple::aarch64:
   case llvm::Triple::aarch64_32:

clang/lib/CodeGen/CGStmt.cpp

if ((i < ResultTypeRequiresCast.size()) && ResultTypeRequiresCast[i]) {
         return;
       }
       Dest = CGF.MakeAddrLValue(A, Ty);
     }
     CGF.EmitStoreThroughLValue(RValue::get(Tmp), Dest);
   }
 }
 
+static void EmitHipStdParUnsupportedAsm(CodeGenFunction *CGF,
+                                        const AsmStmt &S) {
+  constexpr auto Name = "__ASM__hipstdpar_unsupported";
+
+  StringRef Asm;
+  if (auto GCCAsm = dyn_cast<GCCAsmStmt>(&S))
+    Asm = GCCAsm->getAsmString()->getString();
+
+  auto &Ctx = CGF->CGM.getLLVMContext();
+
+  auto StrTy = llvm::ConstantDataArray::getString(Ctx, Asm);
+  auto FnTy = llvm::FunctionType::get(llvm::Type::getVoidTy(Ctx),
+                                      {StrTy->getType()}, false);
+  auto UBF = CGF->CGM.getModule().getOrInsertFunction(Name, FnTy);
+
+  CGF->Builder.CreateCall(UBF, {StrTy});
+}
+
 void CodeGenFunction::EmitAsmStmt(const AsmStmt &S) {
   // Pop all cleanup blocks at the end of the asm statement.
   CodeGenFunction::RunCleanupsScope Cleanups(*this);
 
   // Assemble the final asm string.
   std::string AsmString = S.generateAsmString(getContext());
 
   // Get all the output and input constraints together.
   SmallVector<TargetInfo::ConstraintInfo, 4> OutputConstraintInfos;
   SmallVector<TargetInfo::ConstraintInfo, 4> InputConstraintInfos;
 
-  for (unsigned i = 0, e = S.getNumOutputs(); i != e; i++) {
+  bool IsHipStdPar = getLangOpts().HIPStdPar && getLangOpts().CUDAIsDevice;
+  bool IsValidTargetAsm = true;
+  for (unsigned i = 0, e = S.getNumOutputs(); i != e && IsValidTargetAsm; i++) {
     StringRef Name;
     if (const GCCAsmStmt *GAS = dyn_cast<GCCAsmStmt>(&S))
       Name = GAS->getOutputName(i);
     TargetInfo::ConstraintInfo Info(S.getOutputConstraint(i), Name);
     bool IsValid = getTarget().validateOutputConstraint(Info); (void)IsValid;
+    if (IsHipStdPar && !IsValid)
+      IsValidTargetAsm = false;
+    else
       assert(IsValid && "Failed to parse output constraint");
     OutputConstraintInfos.push_back(Info);
   }
 
-  for (unsigned i = 0, e = S.getNumInputs(); i != e; i++) {
+  for (unsigned i = 0, e = S.getNumInputs(); i != e && IsValidTargetAsm; i++) {
     StringRef Name;
     if (const GCCAsmStmt *GAS = dyn_cast<GCCAsmStmt>(&S))
       Name = GAS->getInputName(i);
     TargetInfo::ConstraintInfo Info(S.getInputConstraint(i), Name);
     bool IsValid =
       getTarget().validateInputConstraint(OutputConstraintInfos, Info);
-    assert(IsValid && "Failed to parse input constraint"); (void)IsValid;
+    if (IsHipStdPar && !IsValid)
+      IsValidTargetAsm = false;
+    else
+      assert(IsValid && "Failed to parse input constraint");
     InputConstraintInfos.push_back(Info);
   }
 
+  if (!IsValidTargetAsm)
+    return EmitHipStdParUnsupportedAsm(this, S);
+
   std::string Constraints;
 
   std::vector<LValue> ResultRegDests;
   std::vector<QualType> ResultRegQualTys;
   std::vector<llvm::Type *> ResultRegTypes;
   std::vector<llvm::Type *> ResultTruncRegTypes;
   std::vector<llvm::Type *> ArgTypes;
   std::vector<llvm::Type *> ArgElemTypes;

clang/lib/CodeGen/CMakeLists.txt

 set(LLVM_LINK_COMPONENTS
   AggressiveInstCombine
   Analysis
   BitReader
   BitWriter
   CodeGenTypes
   Core
   Coroutines
   Coverage
   Demangle
   Extensions
   FrontendHLSL
   FrontendOpenMP
+  HIPStdPar
   IPO
   IRPrinter
   IRReader
   InstCombine
   Instrumentation
   LTO
   Linker
   MC

clang/lib/CodeGen/CodeGenFunction.cpp

void CodeGenFunction::checkTargetFeatures(SourceLocation Loc,
 
   // Grab the required features for the call. For a builtin this is listed in
   // the td file with the default cpu, for an always_inline function this is any
   // listed cpu and any listed features.
   unsigned BuiltinID = TargetDecl->getBuiltinID();
   std::string MissingFeature;
   llvm::StringMap<bool> CallerFeatureMap;
   CGM.getContext().getFunctionFeatureMap(CallerFeatureMap, FD);
+  // When compiling in HipStdPar mode we have to be conservative in rejecting
+  // target specific features in the FE, and defer the possible error to the
+  // AcceleratorCodeSelection pass, wherein iff an unsupported target builtin is
+  // referenced by an accelerator executable function, we emit an error.
+  bool IsHipStdPar = getLangOpts().HIPStdPar && getLangOpts().CUDAIsDevice;
   if (BuiltinID) {
     StringRef FeatureList(CGM.getContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
     if (!Builtin::evaluateRequiredTargetFeatures(
-        FeatureList, CallerFeatureMap)) {
+        FeatureList, CallerFeatureMap) && !IsHipStdPar) {
       CGM.getDiags().Report(Loc, diag::err_builtin_needs_feature)
           << TargetDecl->getDeclName()
           << FeatureList;
     }
   } else if (!TargetDecl->isMultiVersion() &&
              TargetDecl->hasAttr<TargetAttr>()) {
     // Get the required features for the callee.
 
for (const auto &F : CalleeFeatureMap) {
         ReqFeatures.push_back(F.getKey());
     }
     if (!llvm::all_of(ReqFeatures, [&](StringRef Feature) {
       if (!CallerFeatureMap.lookup(Feature)) {
         MissingFeature = Feature.str();
         return false;
       }
       return true;
-    }))
+    }) && !IsHipStdPar)
       CGM.getDiags().Report(Loc, diag::err_function_needs_feature)
           << FD->getDeclName() << TargetDecl->getDeclName() << MissingFeature;
   } else if (!FD->isMultiVersion() && FD->hasAttr<TargetAttr>()) {
     llvm::StringMap<bool> CalleeFeatureMap;
     CGM.getContext().getFunctionFeatureMap(CalleeFeatureMap, TargetDecl);
 
     for (const auto &F : CalleeFeatureMap) {
       if (F.getValue() && (!CallerFeatureMap.lookup(F.getKey()) ||
-                           !CallerFeatureMap.find(F.getKey())->getValue()))
+                           !CallerFeatureMap.find(F.getKey())->getValue()) &&
+          !IsHipStdPar)
         CGM.getDiags().Report(Loc, diag::err_function_needs_feature)
             << FD->getDeclName() << TargetDecl->getDeclName() << F.getKey();
     }
   }
 }
 
 void CodeGenFunction::EmitSanitizerStatReport(llvm::SanitizerStatKind SSK) {
   if (!CGM.getCodeGenOpts().SanitizeStats)

clang/lib/CodeGen/CodeGenModule.cpp

llvm::GlobalValue::LinkageTypes Linkage =
           : llvm::GlobalValue::InternalLinkage;
   auto *GV = new llvm::GlobalVariable(getModule(), Init->getType(),
                                       /*isConstant=*/true, Linkage, Init, Name);
   setGVProperties(GV, TPO);
   if (supportsCOMDAT())
     GV->setComdat(TheModule.getOrInsertComdat(GV->getName()));
   Emitter.finalize(GV);
 
     return ConstantAddress(GV, GV->getValueType(), Alignment);
 }
 
 ConstantAddress CodeGenModule::GetWeakRefReference(const ValueDecl *VD) {
   const AliasAttr *AA = VD->getAttr<AliasAttr>();
   assert(AA && "No alias?");
 
   CharUnits Alignment = getContext().getDeclAlign(VD);
   llvm::Type *DeclTy = getTypes().ConvertTypeForMem(VD->getType());
void CodeGenModule::EmitGlobal(GlobalDecl GD) {
   // If this is CUDA, be selective about which declarations we emit.
   if (LangOpts.CUDA) {
     if (LangOpts.CUDAIsDevice) {
       if (!Global->hasAttr<CUDADeviceAttr>() &&
           !Global->hasAttr<CUDAGlobalAttr>() &&
           !Global->hasAttr<CUDAConstantAttr>() &&
           !Global->hasAttr<CUDASharedAttr>() &&
           !Global->getType()->isCUDADeviceBuiltinSurfaceType() &&
-          !Global->getType()->isCUDADeviceBuiltinTextureType())
+          !Global->getType()->isCUDADeviceBuiltinTextureType() &&
+          !(LangOpts.HIPStdPar &&
+            isa<FunctionDecl>(Global) &&
+            !Global->hasAttr<CUDAHostAttr>()))
         return;
     } else {
       // We need to emit host-side 'shadows' for all global
       // device-side variables because the CUDA runtime needs their
       // size and host-side address in order to provide access to
       // their device-side incarnations.
 
       // So device-only functions are the only things we skip.
if (Linkage == llvm::GlobalVariable::CommonLinkage) {
     if (!GV->getInitializer()->isNullValue())
       GV->setLinkage(llvm::GlobalVariable::WeakAnyLinkage);
   }
 
   setNonAliasAttributes(D, GV);
 
   if (D->getTLSKind() && !GV->isThreadLocal()) {
     if (D->getTLSKind() == VarDecl::TLS_Dynamic)
       CXXThreadLocals.push_back(D);
     setTLSMode(GV, *D);
   }
 
   maybeSetTrivialComdat(*D, *GV);
 
   // Emit the initializer function if necessary.
   if (NeedsGlobalCtor || NeedsGlobalDtor)
     EmitCXXGlobalVarDeclInitFunc(D, GV, NeedsGlobalCtor);

clang/test/CodeGenHipStdPar/unannotated-functions-get-emitted.cpp

+// RUN: %clang_cc1 -x hip -emit-llvm -fcuda-is-device \
+// RUN:   -o - %s | FileCheck --check-prefix=NO-HIPSTDPAR-DEV %s
+
+// RUN: %clang_cc1 --hipstdpar -emit-llvm -fcuda-is-device \
+// RUN:   -o - %s | FileCheck --check-prefix=HIPSTDPAR-DEV %s
+
+#define __device__ __attribute__((device))
+
+// NO-HIPSTDPAR-DEV-NOT: define {{.*}} void @foo({{.*}})
+// HIPSTDPAR-DEV: define {{.*}} void @foo({{.*}})
+extern "C" void foo(float *a, float b) {
+  *a = b;
+}
+
+// NO-HIPSTDPAR-DEV: define {{.*}} void @bar({{.*}})
+// HIPSTDPAR-DEV: define {{.*}} void @bar({{.*}})
+extern "C" __device__ void bar(float *a, float b) {
+  *a = b;
+}

clang/test/CodeGenHipStdPar/unsupported-ASM.cpp

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu \
+// RUN:   --hipstdpar -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+
+#define __global__ __attribute__((global))
+
+__global__ void foo(int i) {
+    asm ("addl %2, %1; seto %b0" : "=q" (i), "+g" (i) : "r" (i));
+}
+
+// CHECK: declare void @__ASM__hipstdpar_unsupported([{{.*}}])

clang/test/CodeGenHipStdPar/unsupported-builtins.cpp

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -aux-triple x86_64-unknown-linux-gnu \
+// RUN:   --hipstdpar -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+
+#define __global__ __attribute__((global))
+
+__global__ void foo() { return __builtin_ia32_pause(); }
+
+// CHECK: declare void @__builtin_ia32_pause__hipstdpar_unsupported()