[AMDGPU] replace hostcall module flag with function attribute

[sameerds]

Summary

The module flag to indicate use of hostcall is insufficient to catch
all cases where hostcall might be in use by a kernel. This is now
replaced by a function attribute that gets propagated to top-level
kernel functions via their respective call-graph.

If the attribute "amdgpu-no-hostcall-ptr" is absent on a kernel, the
default behaviour is to emit kernel metadata indicating that the
kernel uses the hostcall buffer pointer passed as an implicit
argument.

The attribute may be placed explicitly by the user, or inferred by the
AMDGPU attributor by examining the call-graph. The attribute is
inferred only if the function is not being sanitized, and the
implictarg_ptr does not result in a load of any byte in the hostcall
pointer argument.

[jdoerfert]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

You can use AAPointerInfo for the call site return IRPosition. It will (through the iterations) gather all accesses and put them into "bins" based on offset and size. It deals with uses in calls, etc. and if there is stuff missing it is better to add it in one place so we benefit throughout.

588

Always check the return value, if it is false something went wrong and some calls might have been missed. Basically,

if (!A.check...)
  NeedsHostCallPtr = true;

then you also can just return false in the callback w/o setting it.

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

I am not following what you have in mind. "implicitarg_ptr" is a pointer returned by an intrinsic that reads an ABI-defined register. I need to check that for a given call-graph, a particular range of bytes relative to that base pointer are never accessed. The above DFS on the uses conservatively assumes that such a load exists unless it can conclusively trace every use of the base pointer. This may include the pointer being passed to an extern function or being stored into a different memory location (although we don't expect ABI registers being capture this way). I am not seeing how to construct this around AAPointerInfo. As far as I can see, the public interface only talks about uses that are recognized as loads and stores.

[sameerds]

eliminated NeedsHostcall by simply checking the return value of the walk over all calls

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

588

Thanks! This eliminated the variable. Also used the opportunity to get more creative with the callback names!

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

198

The ABI should not be a property of the subtarget, and the global subtarget shouldn't be used

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-v5.ll

56

Can you also add cases where the load is wider and the hostcall pointer is a subset of it?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

520

Can you use checkForAllUses instead of creating your own worklist?

[jdoerfert]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

Not actually tested, replaces the function body. Depends on D119249.

const auto PointerInfoAA = A.getAAFor<AAPointerInfo>(*this, IRPosition::callback_returned(cast<CallBase>(Ptr)), DepClassTy::Required);
if (!PointerInfoAA.getState().isValidState())
  return true; // Abort (which is weird as false is abort in the other CB).
AAPointerInfo::OffsetAndSize OAS(*Position, /* probably look pointer width up in DL */ 8);
return !forallInterferingAccesses(OAS, [](const AAPointerInfo::Access &Acc, bool IsExact) {
   return Acc.getRemoteInst()->isDroppable(); });

580–583

I'd suggest to switch the return value of ArgUsed... so it matches this functions (and other callbacks).

[jdoerfert]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

You don't actually need the state check.
And as I said, this will take care of following pointers passed into callees or through memory to other places, all while ignoring dead code, etc.

[sameerds]

further simplified the callback return value; moved hostcall info out of the subtarget

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

198

This is still depending on the global subtarget. There's no reason isHsaAbiVersion* should depend on the subtarget. All it's doing is checking global properties

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

198

I don't understand what the objection is. Existing functions that check the ABI version clearly do so by accessing the subtarget. I am merely following existing practice.

I have now rearranged the code a bit. Maybe this works? To be honest, I am not very familiar with how ABI information is tracked. I will heartily welcome any advice on how to retrieve the ABI version and then determine the location of the hostcall implicit arg.

520

checkForAllUses does not track sufficient state to catch every load that happens to overlap with the hostcall pointer. The only pattern likely to happen in real life is "GEP to offset 24, typecast to i64*, load 8 bytes". But unless we can guarantee that this is the only pattern, we need something robust.

@jdoerfert pointed me in the correct direction to make full use of existing machinery, but that is now dependent on D119249

565

I see now. forallInterferingAccesses() does check for valid state on entry, which is sufficient to take care of all the opaque uses like a call to an extern function or a complicated phi or a capturing store. Thanks a ton ... this has been very educational!

580–583

Agreed. In fact switching the return value allowed me to merge the two callbacks into one. I am still keeping the positive name on the outermost function "funcRetrievesHostcallPtr" because it matches the sense near the callsite.

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

198

Fundamentally the ABI is/needs to be a module level property. The subtarget is a function level property, and the subtarget read from the global target machine has been deprecated for years. getHsaABIVersion isn't actually reading anything from the subtarget, and is only using it to query the triple. Really getHsaAbiVersion should be passed the module, which could query the triple/module flag/compiler flag, none of which are subtarget properties

[sameerds]

added tests for i128 load. hostcall position is now independent of subtarget.

[jdoerfert]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

Yes, all "forAll" functions will return false if we cannot visit "all". Though, these methods will utilize the smarts, e.g., ignore what is dead, look at loads if the value is stored in a way we can track it through memory, transfer accesses in a callee to the caller "space" if a pointer is passed to the callee,... etc.

[jdoerfert]

@kpyzhov Usually people don't accept without comment, especially if there is clearly ongoing discussion.

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

Yes, all "forAll" functions will return false if we cannot visit "all". Though, these methods will utilize the smarts, e.g., ignore what is dead, look at loads if the value is stored in a way we can track it through memory, transfer accesses in a callee to the caller "space" if a pointer is passed to the callee,... etc.

@jdoerfert, do you see D119249 landing soon? We are kinda on a short runway here (less than a handful of days) and hoping to land this review quickly because it solves an important issue. I would prefer to have the check that you outlined, but the alternative is to let my version through now, and then update it when the new interface becomes available.

[jdoerfert]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

Did you test my proposed code? I'll land my patch tomorrow, if yours works as you expect with the proposed AAPointerInfo use, great. If it doesn't I don't necessarily mind you merging something else with a clear intention to address issues and remove duplication as we go.

I can lift my commit block as we go and @arsenm can also give it a good-to-go if need be.

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

Lifting your commit block would be useful in general. I certainly do not intend to submit something in a hurry.

I applied your change and tested the proposed code. It gives more pessimistic results than my original crude version, i.e., it marks some uses of the implicitarg_ptr as accessing the target range, when I am pretty sure it does not. See test_kernel42 in the lit test hsa-metadata-hostcall-v3.ll in this review. It is intended to access eight bytes from 16 onwards, which clearly does not overlap with offset 24, which is the start of the target range (hostcall pointer).

Strangely, in the debug messages, the AAPointerInfoCallsiteReturned() constructor is called, but I don't see any messages marked [AAPointerInfo].

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

Good news. The PointerInfo works as expected ... it's just that AMDGPUAttributor has a very short Allowed list, and I needed to add PointerInfo there.

So @jdoerfert, D119249 works for me with the suggested changes. If that makes it to mainline, I'll update this change to use it.

@arsenm do you see any issues with enabling AAPointerInfo in the AMDGPUAttributor? Will there be concerns about it being "too heavy" or something?

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

565

No, that's what it's there for

[sameerds]

use AAPointerInfo; add more tests to demonstrate the same

[jdoerfert]

Attributor use looks good. @arsenm will need to look at the rest.

clang/lib/CodeGen/CodeGenModule.cpp

if (LangOpts.Sanitize.has(SanitizerKind::Address)) {
           FT, llvm::GlobalValue::ExternalLinkage,
           "__amdgpu_device_library_preserve_asan_functions", &getModule());
       auto *Var = new llvm::GlobalVariable(
           getModule(), FT->getPointerTo(),
           /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage, F,
           "__amdgpu_device_library_preserve_asan_functions_ptr", nullptr,
           llvm::GlobalVariable::NotThreadLocal);
       addCompilerUsedGlobal(Var);
-      if (!getModule().getModuleFlag("amdgpu_hostcall")) {
-        getModule().addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
-      }
     }
     // Emit amdgpu_code_object_version module flag, which is code object version
     // times 100.
     // ToDo: Enable module flag for all code object version when ROCm device
     // library is ready.
     if (getTarget().getTargetOpts().CodeObjectVersion == TargetOptions::COV_5) {
       getModule().addModuleFlag(llvm::Module::Error,
                                 "amdgpu_code_object_version",

clang/test/CodeGenCUDA/amdgpu-asan-printf.cu

-// RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
-// RUN:   -fcuda-is-device -target-cpu gfx906 -fsanitize=address \
-// RUN:   -O3 -x hip | FileCheck -check-prefixes=MFCHECK %s
-
-// MFCHECK: !{{.*}} = !{i32 4, !"amdgpu_hostcall", i32 1}
-
-// Test to check hostcall module flag metadata is generated correctly
-// when a program has printf call and compiled with -fsanitize=address.
-#include "Inputs/cuda.h"
-__device__ void non_kernel() {
-  printf("sanitized device function");
-}
-
-__global__ void kernel() {
-  non_kernel();
-}
-

clang/test/CodeGenCUDA/amdgpu-asan.cu

 // Create a sample address sanitizer bitcode library.
 
 // RUN: %clang_cc1 -x ir -fcuda-is-device -triple amdgcn-amd-amdhsa -emit-llvm-bc \
 // RUN:   -disable-llvm-passes -o %t.asanrtl.bc %S/Inputs/amdgpu-asanrtl.ll
 
 // Check sanitizer runtime library functions survive
 // optimizations without being removed or parameters altered.
 
 // RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx906 -fsanitize=address \
 // RUN:   -mlink-bitcode-file %t.asanrtl.bc -x hip \
-// RUN:   | FileCheck -check-prefixes=ASAN,MFCHECK %s
+// RUN:   | FileCheck -check-prefixes=ASAN %s
 
 // RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx906 -fsanitize=address \
 // RUN:   -O3 -mlink-bitcode-file %t.asanrtl.bc -x hip \
-// RUN:   | FileCheck -check-prefixes=ASAN,MFCHECK %s
+// RUN:   | FileCheck -check-prefixes=ASAN %s
 
 // RUN: %clang_cc1 %s -emit-llvm -o - -triple=amdgcn-amd-amdhsa \
 // RUN:   -fcuda-is-device -target-cpu gfx906 -x hip \
 // RUN:   | FileCheck %s
 
 // REQUIRES: amdgpu-registered-target
 
 // ASAN-DAG: define weak void @__amdgpu_device_library_preserve_asan_functions()
 // ASAN-DAG: @__amdgpu_device_library_preserve_asan_functions_ptr = weak addrspace(1) constant void ()* @__amdgpu_device_library_preserve_asan_functions
 // ASAN-DAG: @llvm.compiler.used = {{.*}}@__amdgpu_device_library_preserve_asan_functions_ptr
 // ASAN-DAG: define weak void @__asan_report_load1(i64 %{{.*}})
 
-// MFCHECK: !{{.*}} = !{i32 4, !"amdgpu_hostcall", i32 1}
-
 // CHECK-NOT: @__amdgpu_device_library_preserve_asan_functions
 // CHECK-NOT: @__asan_report_load1

llvm/lib/Target/AMDGPU/AMDGPUAttributes.def

 //===----------------------------------------------------------------------===//
 
 // NOTE: NO INCLUDE GUARD DESIRED!
 
 AMDGPU_ATTRIBUTE(DISPATCH_PTR, "amdgpu-no-dispatch-ptr")
 AMDGPU_ATTRIBUTE(QUEUE_PTR, "amdgpu-no-queue-ptr")
 AMDGPU_ATTRIBUTE(DISPATCH_ID, "amdgpu-no-dispatch-id")
 AMDGPU_ATTRIBUTE(IMPLICIT_ARG_PTR, "amdgpu-no-implicitarg-ptr")
+AMDGPU_ATTRIBUTE(HOSTCALL_PTR, "amdgpu-no-hostcall-ptr")
 AMDGPU_ATTRIBUTE(WORKGROUP_ID_X, "amdgpu-no-workgroup-id-x")
 AMDGPU_ATTRIBUTE(WORKGROUP_ID_Y, "amdgpu-no-workgroup-id-y")
 AMDGPU_ATTRIBUTE(WORKGROUP_ID_Z, "amdgpu-no-workgroup-id-z")
 AMDGPU_ATTRIBUTE(WORKITEM_ID_X, "amdgpu-no-workitem-id-x")
 AMDGPU_ATTRIBUTE(WORKITEM_ID_Y, "amdgpu-no-workitem-id-y")
 AMDGPU_ATTRIBUTE(WORKITEM_ID_Z, "amdgpu-no-workitem-id-z")
 
 #undef AMDGPU_ATTRIBUTE

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

 //===- AMDGPUAttributor.cpp -----------------------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 /// \file This pass uses Attributor framework to deduce AMDGPU attributes.
 //
 //===----------------------------------------------------------------------===//
 
 #include "AMDGPU.h"
 #include "GCNSubtarget.h"
+#include "Utils/AMDGPUBaseInfo.h"
 #include "llvm/CodeGen/TargetPassConfig.h"
 #include "llvm/IR/IntrinsicsAMDGPU.h"
 #include "llvm/IR/IntrinsicsR600.h"
 #include "llvm/Target/TargetMachine.h"
 #include "llvm/Transforms/IPO/Attributor.h"
 
 #define DEBUG_TYPE "amdgpu-attributor"
 
static bool isDSAddress(const Constant *C) {
   if (!GV)
     return false;
   unsigned AS = GV->getAddressSpace();
   return AS == AMDGPUAS::LOCAL_ADDRESS || AS == AMDGPUAS::REGION_ADDRESS;
 }
 
 /// Returns true if the function requires the implicit argument be passed
 /// regardless of the function contents.
-static bool funcRequiresImplicitArgPtr(const Function &F) {
+static bool funcRequiresHostcallPtr(const Function &F) {
   // Sanitizers require the hostcall buffer passed in the implicit arguments.
   return F.hasFnAttribute(Attribute::SanitizeAddress) ||
          F.hasFnAttribute(Attribute::SanitizeThread) ||
          F.hasFnAttribute(Attribute::SanitizeMemory) ||
          F.hasFnAttribute(Attribute::SanitizeHWAddress) ||
          F.hasFnAttribute(Attribute::SanitizeMemTag);
 }
 
bool needsQueuePtr(const Constant *C, Function &Fn) {
     if (IsNonEntryFunc && (Access & DS_GLOBAL))
       return true;
 
     return !HasAperture && (Access & ADDR_SPACE_CAST);
   }
 
 private:
   /// Used to determine if the Constant needs the queue pointer.
   DenseMap<const Constant *, uint8_t> ConstantStatus;
 };
 
 struct AAAMDAttributes : public StateWrapper<
   BitIntegerState<uint16_t, ALL_ARGUMENT_MASK, 0>, AbstractAttribute> {
   using Base = StateWrapper<BitIntegerState<uint16_t, ALL_ARGUMENT_MASK, 0>,
                             AbstractAttribute>;
 
   AAAMDAttributes(const IRPosition &IRP, Attributor &A) : Base(IRP) {}
struct AAAMDAttributesFunction : public AAAMDAttributes {
   AAAMDAttributesFunction(const IRPosition &IRP, Attributor &A)
       : AAAMDAttributes(IRP, A) {}
 
   void initialize(Attributor &A) override {
     Function *F = getAssociatedFunction();
 
     // If the function requires the implicit arg pointer due to sanitizers,
     // assume it's needed even if explicitly marked as not requiring it.
-    const bool NeedsImplicit = funcRequiresImplicitArgPtr(*F);
-    if (NeedsImplicit)
+    const bool NeedsHostcall = funcRequiresHostcallPtr(*F);
+    if (NeedsHostcall) {
       removeAssumedBits(IMPLICIT_ARG_PTR);
+      removeAssumedBits(HOSTCALL_PTR);
+    }
 
     for (auto Attr : ImplicitAttrs) {
-      if (NeedsImplicit && Attr.first == IMPLICIT_ARG_PTR)
+      if (NeedsHostcall &&
+          (Attr.first == IMPLICIT_ARG_PTR || Attr.first == HOSTCALL_PTR))
         continue;
 
       if (F->hasFnAttribute(Attr.second))
         addKnownBits(Attr.first);
     }
 
     if (F->isDeclaration())
       return;
ChangeStatus updateImpl(Attributor &A) override {
     if (!NeedsQueuePtr) {
       NeedsQueuePtr = checkForQueuePtr(A);
     }
 
     if (NeedsQueuePtr) {
       removeAssumedBits(QUEUE_PTR);
     }
 
+    if (funcRetrievesHostcallPtr(A)) {
+      removeAssumedBits(IMPLICIT_ARG_PTR);
+      removeAssumedBits(HOSTCALL_PTR);
+    }
+
     return getAssumed() != OrigAssumed ? ChangeStatus::CHANGED
                                        : ChangeStatus::UNCHANGED;
   }
 
   ChangeStatus manifest(Attributor &A) override {
     SmallVector<Attribute, 8> AttrList;
     LLVMContext &Ctx = getAssociatedFunction()->getContext();
 
for (BasicBlock &BB : *F) {
               return true;
           }
         }
       }
     }
 
     return false;
   }
+
+  bool funcRetrievesHostcallPtr(Attributor &A) {
+    auto Pos = llvm::AMDGPU::getHostcallImplicitArgPosition();
+
+    // Check if this is a call to the implicitarg_ptr builtin and it
+    // is used to retrieve the hostcall pointer. The implicit arg for
+    // hostcall is not used only if every use of the implicitarg_ptr
+    // is a load that clearly does not retrieve any byte of the
+    // hostcall pointer. We check this by tracing all the uses of the
+    // initial call to the implicitarg_ptr intrinsic.
+    auto DoesNotLeadToHostcallPtr = [&](Instruction &I) {
+      auto &Call = cast<CallBase>(I);
+      if (Call.getIntrinsicID() != Intrinsic::amdgcn_implicitarg_ptr)
+        return true;
+
+      const auto &PointerInfoAA = A.getAAFor<AAPointerInfo>(
+          *this, IRPosition::callsite_returned(Call), DepClassTy::REQUIRED);
+
+      AAPointerInfo::OffsetAndSize OAS(Pos, 8);
+      return PointerInfoAA.forallInterferingAccesses(
+          OAS, [](const AAPointerInfo::Access &Acc, bool IsExact) {
+            return Acc.getRemoteInst()->isDroppable();
+          });
+    };
+
+    bool UsedAssumedInformation = false;
+    return !A.checkForAllCallLikeInstructions(DoesNotLeadToHostcallPtr, *this,
+                                              UsedAssumedInformation);
+  }
 };
 
 AAAMDAttributes &AAAMDAttributes::createForPosition(const IRPosition &IRP,
                                                     Attributor &A) {
   if (IRP.getPositionKind() == IRPosition::IRP_FUNCTION)
     return *new (A.Allocator) AAAMDAttributesFunction(IRP, A);
   llvm_unreachable("AAAMDAttributes is only valid for function position");
 }
struct AAAMDFlatWorkGroupSize
   }
 
   ChangeStatus updateImpl(Attributor &A) override {
     ChangeStatus Change = ChangeStatus::UNCHANGED;
 
     auto CheckCallSite = [&](AbstractCallSite CS) {
       Function *Caller = CS.getInstruction()->getFunction();
       LLVM_DEBUG(dbgs() << "[AAAMDFlatWorkGroupSize] Call " << Caller->getName()
                         << "->" << getAssociatedFunction()->getName() << '\n');
 
       const auto &CallerInfo = A.getAAFor<AAAMDFlatWorkGroupSize>(
           *this, IRPosition::function(*Caller), DepClassTy::REQUIRED);
 
       Change |=
           clampStateAndIndicateChange(this->getState(), CallerInfo.getState());
 
       return true;
     };
 
     bool AllCallSitesKnown = true;
     if (!A.checkForAllCallSites(CheckCallSite, *this, true, AllCallSitesKnown))
       return indicatePessimisticFixpoint();
 
     return Change;
   }
 
   ChangeStatus manifest(Attributor &A) override {
     SmallVector<Attribute, 8> AttrList;
     Function *F = getAssociatedFunction();
     LLVMContext &Ctx = F->getContext();
 
     auto &InfoCache = static_cast<AMDGPUInformationCache &>(A.getInfoCache());
     unsigned Min, Max;
     std::tie(Min, Max) = InfoCache.getMaximumFlatWorkGroupRange(*F);
 
     // Don't add the attribute if it's the implied default.
     if (getAssumed().getLower() == Min && getAssumed().getUpper() - 1 == Max)
       return ChangeStatus::UNCHANGED;
 
     SmallString<10> Buffer;
for (Function &F : M) {
         Functions.insert(&F);
     }
 
     CallGraphUpdater CGUpdater;
     BumpPtrAllocator Allocator;
     AMDGPUInformationCache InfoCache(M, AG, Allocator, nullptr, *TM);
     DenseSet<const char *> Allowed(
         {&AAAMDAttributes::ID, &AAUniformWorkGroupSize::ID,
-         &AAAMDFlatWorkGroupSize::ID, &AACallEdges::ID});
+         &AAAMDFlatWorkGroupSize::ID, &AACallEdges::ID, &AAPointerInfo::ID});
 
     Attributor A(Functions, InfoCache, CGUpdater, &Allowed);
 
     for (Function &F : M) {
       if (!F.isIntrinsic()) {
         A.getOrCreateAAFor<AAAMDAttributes>(IRPosition::function(F));
         A.getOrCreateAAFor<AAUniformWorkGroupSize>(IRPosition::function(F));
         if (!AMDGPU::isEntryFunctionCC(F.getCallingConv())) {

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

void MetadataStreamerV2::emitHiddenKernelArgs(const Function &Func,
   auto Int8PtrTy = Type::getInt8PtrTy(Func.getContext(),
                                       AMDGPUAS::GLOBAL_ADDRESS);
 
   // Emit "printf buffer" argument if printf is used, otherwise emit dummy
   // "none" argument.
   if (HiddenArgNumBytes >= 32) {
     if (Func.getParent()->getNamedMetadata("llvm.printf.fmts"))
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenPrintfBuffer);
-    else if (Func.getParent()->getFunction("__ockl_hostcall_internal")) {
+    else if (!Func.hasFnAttribute("amdgpu-no-hostcall-ptr")) {
       // The printf runtime binding pass should have ensured that hostcall and
       // printf are not used in the same module.
       assert(!Func.getParent()->getNamedMetadata("llvm.printf.fmts"));
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenHostcallBuffer);
     } else
       emitKernelArg(DL, Int8PtrTy, Align(8), ValueKind::HiddenNone);
   }
 
void MetadataStreamerV3::emitKernelArg(
   Args.push_back(Arg);
 }
 
 void MetadataStreamerV3::emitHiddenKernelArgs(const MachineFunction &MF,
                                               unsigned &Offset,
                                               msgpack::ArrayDocNode Args) {
   auto &Func = MF.getFunction();
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
+  const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
 
   unsigned HiddenArgNumBytes = ST.getImplicitArgNumBytes(Func);
   if (!HiddenArgNumBytes)
     return;
 
   const Module *M = Func.getParent();
   auto &DL = M->getDataLayout();
   auto Int64Ty = Type::getInt64Ty(Func.getContext());
auto Int8PtrTy =
       Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
 
   // Emit "printf buffer" argument if printf is used, emit "hostcall buffer"
   // if "hostcall" module flag is set, otherwise emit dummy "none" argument.
   if (HiddenArgNumBytes >= 32) {
     if (M->getNamedMetadata("llvm.printf.fmts"))
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
                     Args);
-    else if (M->getModuleFlag("amdgpu_hostcall")) {
+    else if (MFI.hasHostcallPtr()) {
       // The printf runtime binding pass should have ensured that hostcall and
       // printf are not used in the same module.
       assert(!M->getNamedMetadata("llvm.printf.fmts"));
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
                     Args);
     } else
       emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_none", Offset, Args);
   }

 
 void MetadataStreamerV5::emitHiddenKernelArgs(const MachineFunction &MF,
                                               unsigned &Offset,
                                               msgpack::ArrayDocNode Args) {
   auto &Func = MF.getFunction();
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
   const Module *M = Func.getParent();
   auto &DL = M->getDataLayout();
+  const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
 
   auto Int64Ty = Type::getInt64Ty(Func.getContext());
   auto Int32Ty = Type::getInt32Ty(Func.getContext());
   auto Int16Ty = Type::getInt16Ty(Func.getContext());
 
   emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_x", Offset, Args);
   emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_y", Offset, Args);
   emitKernelArg(DL, Int32Ty, Align(4), "hidden_block_count_z", Offset, Args);
auto Int8PtrTy =
       Type::getInt8PtrTy(Func.getContext(), AMDGPUAS::GLOBAL_ADDRESS);
 
   if (M->getNamedMetadata("llvm.printf.fmts")) {
     emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_printf_buffer", Offset,
                   Args);
   } else
     Offset += 8; // Skipped.
 
-  if (M->getModuleFlag("amdgpu_hostcall")) {
+  if (MFI.hasHostcallPtr()) {
     emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_hostcall_buffer", Offset,
                   Args);
   } else
     Offset += 8; // Skipped.
 
   emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_multigrid_sync_arg", Offset,
                 Args);
 
void MetadataStreamerV5::emitHiddenKernelArgs(const MachineFunction &MF,
 
   // hidden_private_base and hidden_shared_base are only used by GFX8.
   if (ST.getGeneration() == AMDGPUSubtarget::VOLCANIC_ISLANDS) {
     emitKernelArg(DL, Int32Ty, Align(4), "hidden_private_base", Offset, Args);
     emitKernelArg(DL, Int32Ty, Align(4), "hidden_shared_base", Offset, Args);
   } else
     Offset += 8; // Skipped.
 
-  const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
   if (MFI.hasQueuePtr())
     emitKernelArg(DL, Int8PtrTy, Align(8), "hidden_queue_ptr", Offset, Args);
 }
 
 } // end namespace HSAMD
 } // end namespace AMDGPU
 } // end namespace llvm

llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.h

private:
   // Private memory buffer
   // Compute directly in sgpr[0:1]
   // Other shaders indirect 64-bits at sgpr[0:1]
   bool ImplicitBufferPtr : 1;
 
   // Pointer to where the ABI inserts special kernel arguments separate from the
   // user arguments. This is an offset from the KernargSegmentPtr.
   bool ImplicitArgPtr : 1;
+  bool HostcallPtr : 1;
 
   bool MayNeedAGPRs : 1;
 
   // The hard-wired high half of the address of the global information table
   // for AMDPAL OS type. 0xffffffff represents no hard-wired high half, since
   // current hardware only allows a 16 bit value.
   unsigned GITPtrHigh;
 
public:
   bool hasWorkItemIDZ() const {
     return WorkItemIDZ;
   }
 
   bool hasImplicitArgPtr() const {
     return ImplicitArgPtr;
   }
 
+  bool hasHostcallPtr() const {
+    return HostcallPtr;
+  }
+
   bool hasImplicitBufferPtr() const {
     return ImplicitBufferPtr;
   }
 
   AMDGPUFunctionArgInfo &getArgInfo() {
     return ArgInfo;
   }
 

llvm/lib/Target/AMDGPU/SIMachineFunctionInfo.cpp

: AMDGPUMachineFunction(MF),
     WorkGroupIDZ(false),
     WorkGroupInfo(false),
     PrivateSegmentWaveByteOffset(false),
     WorkItemIDX(false),
     WorkItemIDY(false),
     WorkItemIDZ(false),
     ImplicitBufferPtr(false),
     ImplicitArgPtr(false),
+    HostcallPtr(false),
     GITPtrHigh(0xffffffff),
     HighBitsOf32BitAddress(0),
     GDSSize(0) {
   const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
   const Function &F = MF.getFunction();
   FlatWorkGroupSizes = ST.getFlatWorkGroupSizes(F);
   WavesPerEU = ST.getWavesPerEU(F);
 
if (!AMDGPU::isGraphics(CC)) {
     if (!F.hasFnAttribute("amdgpu-no-dispatch-ptr"))
       DispatchPtr = true;
 
     if (!F.hasFnAttribute("amdgpu-no-queue-ptr"))
       QueuePtr = true;
 
     if (!F.hasFnAttribute("amdgpu-no-dispatch-id"))
       DispatchID = true;
+
+    if (!F.hasFnAttribute("amdgpu-no-hostcall-ptr"))
+      HostcallPtr = true;
   }
 
   // FIXME: This attribute is a hack, we just need an analysis on the function
   // to look for allocas.
   bool HasStackObjects = F.hasFnAttribute("amdgpu-stack-objects");
 
   // TODO: This could be refined a lot. The attribute is a poor way of
   // detecting calls or stack objects that may require it before argument

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.h

 bool isHsaAbiVersion4(const MCSubtargetInfo *STI);
 /// \returns True if HSA OS ABI Version identification is 5,
 /// false otherwise.
 bool isHsaAbiVersion5(const MCSubtargetInfo *STI);
 /// \returns True if HSA OS ABI Version identification is 3 or 4,
 /// false otherwise.
 bool isHsaAbiVersion3AndAbove(const MCSubtargetInfo *STI);
 
+/// \returns The offset of the hostcall pointer argument from implicitarg_ptr
+unsigned getHostcallImplicitArgPosition();
+
 struct GcnBufferFormatInfo {
   unsigned Format;
   unsigned BitsPerComp;
   unsigned NumComponents;
   unsigned NumFormat;
   unsigned DataFormat;
 };
 

llvm/lib/Target/AMDGPU/Utils/AMDGPUBaseInfo.cpp

bool isHsaAbiVersion5(const MCSubtargetInfo *STI) {
   return false;
 }
 
 bool isHsaAbiVersion3AndAbove(const MCSubtargetInfo *STI) {
   return isHsaAbiVersion3(STI) || isHsaAbiVersion4(STI) ||
          isHsaAbiVersion5(STI);
 }
 
+// FIXME: All such magic numbers about the ABI should be in a
+// central TD file.
+unsigned getHostcallImplicitArgPosition() {
+  switch (AmdhsaCodeObjectVersion) {
+  case 2:
+  case 3:
+  case 4:
+    return 24;
+  case 5:
+    return 80;
+  default:
+    llvm_unreachable("Unexpected code object version");
+    return 0;
+  }
+}
+
 #define GET_MIMGBaseOpcodesTable_IMPL
 #define GET_MIMGDimInfoTable_IMPL
 #define GET_MIMGInfoTable_IMPL
 #define GET_MIMGLZMappingTable_IMPL
 #define GET_MIMGMIPMappingTable_IMPL
 #define GET_MIMGBiasMappingTable_IMPL
 #define GET_MIMGOffsetMappingTable_IMPL
 #define GET_MIMGG16MappingTable_IMPL

llvm/lib/Transforms/Utils/AMDGPUEmitPrintf.cpp

static Value *fitArgInto64Bits(IRBuilder<> &Builder, Value *Arg) {
 
   llvm_unreachable("unexpected type");
 }
 
 static Value *callPrintfBegin(IRBuilder<> &Builder, Value *Version) {
   auto Int64Ty = Builder.getInt64Ty();
   auto M = Builder.GetInsertBlock()->getModule();
   auto Fn = M->getOrInsertFunction("__ockl_printf_begin", Int64Ty, Int64Ty);
-  if (!M->getModuleFlag("amdgpu_hostcall")) {
-    M->addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
-  }
   return Builder.CreateCall(Fn, Version);
 }
 
 static Value *callAppendArgs(IRBuilder<> &Builder, Value *Desc, int NumArgs,
                              Value *Arg0, Value *Arg1, Value *Arg2, Value *Arg3,
                              Value *Arg4, Value *Arg5, Value *Arg6,
                              bool IsLast) {
   auto Int64Ty = Builder.getInt64Ty();

llvm/test/CodeGen/AMDGPU/addrspacecast-constantexpr.ll

 
 attributes #0 = { argmemonly nounwind }
 attributes #1 = { nounwind }
 ;.
 ; AKF_HSA: attributes #[[ATTR0:[0-9]+]] = { argmemonly nofree nounwind willreturn }
 ; AKF_HSA: attributes #[[ATTR1]] = { nounwind }
 ;.
 ; ATTRIBUTOR_HSA: attributes #[[ATTR0:[0-9]+]] = { argmemonly nofree nounwind willreturn }
-; ATTRIBUTOR_HSA: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa-call.ll

 ; ATTRIBUTOR_HSA-NEXT:    [[FADD:%.*]] = fadd float [[F]], 1.000000e+00
 ; ATTRIBUTOR_HSA-NEXT:    ret float [[FADD]]
 ;
   %f = call float @llvm.amdgcn.rcp.f32(float %arg)
   %fadd = fadd float %f, 1.0
   ret float %fadd
 }
 
-; Implicit arguments need to be enabled for sanitizers
+; Hostcall needs to be enabled for sanitizers
 define amdgpu_kernel void @kern_sanitize_address() #4 {
 ; AKF_HSA-LABEL: define {{[^@]+}}@kern_sanitize_address
 ; AKF_HSA-SAME: () #[[ATTR5:[0-9]+]] {
 ; AKF_HSA-NEXT:    store volatile i32 0, i32 addrspace(1)* null, align 4
 ; AKF_HSA-NEXT:    ret void
 ;
 ; ATTRIBUTOR_HSA-LABEL: define {{[^@]+}}@kern_sanitize_address
 ; ATTRIBUTOR_HSA-SAME: () #[[ATTR17:[0-9]+]] {
 ; ATTRIBUTOR_HSA-NEXT:    store volatile i32 0, i32 addrspace(1)* null, align 4
 ; ATTRIBUTOR_HSA-NEXT:    ret void
 ;
   store volatile i32 0, i32 addrspace(1)* null
   ret void
 }
 
-; Implicit arguments need to be enabled for sanitizers
+; Hostcall needs to be enabled for sanitizers
 define void @func_sanitize_address() #4 {
 ; AKF_HSA-LABEL: define {{[^@]+}}@func_sanitize_address
 ; AKF_HSA-SAME: () #[[ATTR5]] {
 ; AKF_HSA-NEXT:    store volatile i32 0, i32 addrspace(1)* null, align 4
 ; AKF_HSA-NEXT:    ret void
 ;
 ; ATTRIBUTOR_HSA-LABEL: define {{[^@]+}}@func_sanitize_address
 ; ATTRIBUTOR_HSA-SAME: () #[[ATTR17]] {
 ; ATTRIBUTOR_HSA-NEXT:    store volatile i32 0, i32 addrspace(1)* null, align 4
 ; ATTRIBUTOR_HSA-NEXT:    ret void
 ;
   store volatile i32 0, i32 addrspace(1)* null
   ret void
 }
 
-; Implicit arguments need to be enabled for sanitizers
+; Hostcall needs to be enabled for sanitizers
 define void @func_indirect_sanitize_address() #3 {
 ; AKF_HSA-LABEL: define {{[^@]+}}@func_indirect_sanitize_address
 ; AKF_HSA-SAME: () #[[ATTR3]] {
 ; AKF_HSA-NEXT:    call void @func_sanitize_address()
 ; AKF_HSA-NEXT:    ret void
 ;
 ; ATTRIBUTOR_HSA-LABEL: define {{[^@]+}}@func_indirect_sanitize_address
 ; ATTRIBUTOR_HSA-SAME: () #[[ATTR18:[0-9]+]] {
 ; ATTRIBUTOR_HSA-NEXT:    call void @func_sanitize_address()
 ; ATTRIBUTOR_HSA-NEXT:    ret void
 ;
   call void @func_sanitize_address()
   ret void
 }
 
-; Implicit arguments need to be enabled for sanitizers
+; Hostcall needs to be enabled for sanitizers
 define amdgpu_kernel void @kern_indirect_sanitize_address() #3 {
 ; AKF_HSA-LABEL: define {{[^@]+}}@kern_indirect_sanitize_address
 ; AKF_HSA-SAME: () #[[ATTR4]] {
 ; AKF_HSA-NEXT:    call void @func_sanitize_address()
 ; AKF_HSA-NEXT:    ret void
 ;
 ; ATTRIBUTOR_HSA-LABEL: define {{[^@]+}}@kern_indirect_sanitize_address
 ; ATTRIBUTOR_HSA-SAME: () #[[ATTR18]] {

 ; AKF_HSA: attributes #[[ATTR1]] = { nounwind "target-cpu"="fiji" }
 ; AKF_HSA: attributes #[[ATTR2]] = { nounwind "target-cpu"="gfx900" }
 ; AKF_HSA: attributes #[[ATTR3]] = { nounwind }
 ; AKF_HSA: attributes #[[ATTR4]] = { nounwind "amdgpu-calls" }
 ; AKF_HSA: attributes #[[ATTR5]] = { nounwind sanitize_address }
 ; AKF_HSA: attributes #[[ATTR6:[0-9]+]] = { nounwind sanitize_address "amdgpu-no-implicitarg-ptr" }
 ;.
 ; ATTRIBUTOR_HSA: attributes #[[ATTR0:[0-9]+]] = { nounwind readnone speculatable willreturn }
-; ATTRIBUTOR_HSA: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR3]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR4]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR5]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR6]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR7]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR8]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR9]] = { nounwind "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR10]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR11]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR12]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="gfx900" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR13]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="gfx900" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR14]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR3]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR4]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR5]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR6]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR7]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR8]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR9]] = { nounwind "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR10]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR11]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR12]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="gfx900" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR13]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="gfx900" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR14]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "target-cpu"="fiji" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR15]] = { nounwind "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR16]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR16]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR17]] = { nounwind sanitize_address "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR18]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR19:[0-9]+]] = { nounwind sanitize_address "amdgpu-no-implicitarg-ptr" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_HSA: attributes #[[ATTR20]] = { nounwind }
 ;.

llvm/test/CodeGen/AMDGPU/annotate-kernel-features-hsa.ll

 attributes #1 = { nounwind }
 
 ;.
 ; AKF_HSA: attributes #[[ATTR0:[0-9]+]] = { nounwind readnone speculatable willreturn }
 ; AKF_HSA: attributes #[[ATTR1]] = { nounwind }
 ; AKF_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-stack-objects" }
 ;.
 ; ATTRIBUTOR_HSA: attributes #[[ATTR0:[0-9]+]] = { nounwind readnone speculatable willreturn }
-; ATTRIBUTOR_HSA: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR3]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR4]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR5]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR6]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR7]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR8]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR9]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR10]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_HSA: attributes #[[ATTR11]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR3]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR4]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR5]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR6]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR7]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR8]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR9]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR10]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_HSA: attributes #[[ATTR11]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/annotate-kernel-features.ll

 ; NOHSA: attributes #[[ATTR7]] = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-item-id-y" "uniform-work-group-size"="false" }
 ; NOHSA: attributes #[[ATTR8]] = { nounwind "amdgpu-work-item-id-y" "amdgpu-work-item-id-z" "uniform-work-group-size"="false" }
 ; NOHSA: attributes #[[ATTR9]] = { nounwind "amdgpu-work-group-id-y" "amdgpu-work-group-id-z" "amdgpu-work-item-id-y" "amdgpu-work-item-id-z" "uniform-work-group-size"="false" }
 ;.
 ; AKF_CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind readnone speculatable willreturn }
 ; AKF_CHECK: attributes #[[ATTR1]] = { nounwind }
 ;.
 ; ATTRIBUTOR_CHECK: attributes #[[ATTR0:[0-9]+]] = { nounwind readnone speculatable willreturn }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR3]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR4]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR5]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR6]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR7]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR8]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
-; ATTRIBUTOR_CHECK: attributes #[[ATTR9]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR1]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR2]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR3]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR4]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR5]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR6]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR7]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR8]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_CHECK: attributes #[[ATTR9]] = { nounwind "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workitem-id-x" "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/direct-indirect-call.ll

 ; CHECK-SAME: () #[[ATTR1]] {
 ; CHECK-NEXT:    call void @direct()
 ; CHECK-NEXT:    ret void
 ;
   call void @direct()
   ret void
 }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; CHECK: attributes #[[ATTR1]] = { "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/duplicate-attribute-indirect.ll

;
   ret void
 }
 
 attributes #0 = { "amdgpu-no-dispatch-id" }
 
 ;.
 ; AKF_GCN: attributes #[[ATTR0]] = { "amdgpu-calls" "amdgpu-no-dispatch-id" "amdgpu-stack-objects" }
 ;.
-; ATTRIBUTOR_GCN: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_GCN: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_GCN: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/hsa-metadata-enqueue-kernel-v3.ll

 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
 ; CHECK-NEXT:       - .offset:         24
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         32
 ; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_kind:     hidden_hostcall_buffer
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_default_queue
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         48
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_completion_action

llvm/test/CodeGen/AMDGPU/hsa-metadata-enqueue-kernel.ll

 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
+; CHECK-NEXT:     - Size:          8
+; CHECK-NEXT:       Align:         8
+; CHECK-NEXT:       ValueKind:     HiddenHostcallBuffer
 ; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
 ; CHECK-NOT:        ValueKind:     HiddenCompletionAction
 define amdgpu_kernel void @test_non_enqueue_kernel_caller(i8 %a) #0
     !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
     !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
   ret void
 }
 

 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenNone
+; CHECK-NEXT:       ValueKind:     HiddenHostcallBuffer
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenDefaultQueue
 ; CHECK-NEXT:       AddrSpaceQual: Global
 ; CHECK-NEXT:     - Size:          8
 ; CHECK-NEXT:       Align:         8
 ; CHECK-NEXT:       ValueKind:     HiddenCompletionAction

llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll

 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
 ; CHECK-NEXT:       - .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         48
 ; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_kind:     hidden_hostcall_buffer
 ; CHECK:          .name:           test32
 ; CHECK:          .symbol:         test32.kd
 define amdgpu_kernel void @test32(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #3 {
 entry:
   %a.val = load half, half addrspace(1)* %a

 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
 ; CHECK-NEXT:       - .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         48
 ; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_kind:     hidden_hostcall_buffer
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         56
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none

 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
 ; CHECK-NEXT:       - .offset:         40
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         48
 ; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_none
+; CHECK-NEXT:         .value_kind:     hidden_hostcall_buffer
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         56
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none
 ; CHECK-NEXT:       - .address_space:  global
 ; CHECK-NEXT:         .offset:         64
 ; CHECK-NEXT:         .size:           8
 ; CHECK-NEXT:         .value_kind:     hidden_none

llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v5.ll

 entry:
   %a.val = load half, half addrspace(1)* %a
   %b.val = load half, half addrspace(1)* %b
   %r.val = fadd half %a.val, %b.val
   store half %r.val, half addrspace(1)* %r
   ret void
 }
 
-!llvm.module.flags = !{!0}
 !llvm.printf.fmts = !{!1, !2}
 
-!0 = !{i32 1, !"amdgpu_hostcall", i32 1}
 !1 = !{!"1:1:4:%d\5Cn"}
 !2 = !{!"2:1:8:%g\5Cn"}
 
 attributes #0 = { optnone noinline "calls-enqueue-kernel" }
 

llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args.ll

 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
-; CHECK-NEXT:       ValueKind:       HiddenNone
+; CHECK-NEXT:       ValueKind:       HiddenHostcallBuffer
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:   CodeProps:
 define amdgpu_kernel void @test32(
     half addrspace(1)* %r,
     half addrspace(1)* %a,
     half addrspace(1)* %b) #3 {
 entry:
   %a.val = load half, half addrspace(1)* %a

 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
-; CHECK-NEXT:       ValueKind:       HiddenNone
+; CHECK-NEXT:       ValueKind:       HiddenHostcallBuffer
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone

 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetY
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenGlobalOffsetZ
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
-; CHECK-NEXT:       ValueKind:       HiddenNone
+; CHECK-NEXT:       ValueKind:       HiddenHostcallBuffer
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone
 ; CHECK-NEXT:       AddrSpaceQual:   Global
 ; CHECK-NEXT:     - Size:            8
 ; CHECK-NEXT:       Align:           8
 ; CHECK-NEXT:       ValueKind:       HiddenNone

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-absent-v3.ll

-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
-
-; CHECK:              ---
-; CHECK:      amdhsa.kernels:
-; CHECK:        - .args:
-; CHECK-NEXT:       - .name:           a
-; CHECK-NEXT:         .offset:         0
-; CHECK-NEXT:         .size:           1
-; CHECK-NEXT:         .type_name:      char
-; CHECK-NEXT:         .value_kind:     by_value
-; CHECK-NEXT:       - .offset:         8
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       - .offset:         16
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       - .offset:         24
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
-
-; CHECK-NOT:          .value_kind:     hidden_hostcall_buffer
-
-; CHECK:          .language:       OpenCL C
-; CHECK-NEXT:     .language_version:
-; CHECK-NEXT:       - 2
-; CHECK-NEXT:       - 0
-; CHECK:          .name:           test_kernel
-; CHECK:          .symbol:         test_kernel.kd
-
-define amdgpu_kernel void @test_kernel(i8 %a) #0
-    !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
-    !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
-  ret void
-}
-
-; CHECK:  amdhsa.version:
-; CHECK-NEXT: - 1
-; CHECK-NEXT: - 0
-
-attributes #0 = { optnone noinline "amdgpu-implicitarg-num-bytes"="48" }
-
-!1 = !{i32 0}
-!2 = !{!"none"}
-!3 = !{!"char"}
-!4 = !{!""}
-
-!opencl.ocl.version = !{!90}
-!90 = !{i32 2, i32 0}
-
-; PARSER: AMDGPU HSA Metadata Parser Test: PASS

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-absent.ll

-; RUN: llc -mtriple=amdgcn-amd-amdhsa --amdhsa-code-object-version=2 -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa --amdhsa-code-object-version=2 -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
-
-; CHECK: ---
-; CHECK:  Version: [ 1, 0 ]
-; CHECK:  Kernels:
-
-; CHECK:      - Name:            test_kernel
-; CHECK-NEXT:   SymbolName:      'test_kernel@kd'
-; CHECK-NEXT:   Language:        OpenCL C
-; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
-; CHECK-NEXT:   Args:
-; CHECK-NEXT:     - Name:          a
-; CHECK-NEXT:       TypeName:      char
-; CHECK-NEXT:       Size:          1
-; CHECK-NEXT:       Align:         1
-; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       AccQual:       Default
-; CHECK-NEXT:     - Size:          8
-; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:     - Size:          8
-; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:     - Size:          8
-; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NOT:        ValueKind:     HiddenHostcallBuffer
-; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
-; CHECK-NOT:        ValueKind:     HiddenCompletionAction
-
-define amdgpu_kernel void @test_kernel(i8 %a) #0
-    !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
-    !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
-  ret void
-}
-
-attributes #0 = { optnone noinline "amdgpu-implicitarg-num-bytes"="48" }
-
-!1 = !{i32 0}
-!2 = !{!"none"}
-!3 = !{!"char"}
-!4 = !{!""}
-
-!opencl.ocl.version = !{!90}
-!90 = !{i32 2, i32 0}
-
-; PARSER: AMDGPU HSA Metadata Parser Test: PASS

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present-v3-asan.ll

 !1 = !{i32 0}
 !2 = !{!"none"}
 !3 = !{!"char"}
 !4 = !{!""}
 
 !opencl.ocl.version = !{!90}
 !90 = !{i32 2, i32 0}
 
-!llvm.module.flags = !{!0}
-!0 = !{i32 4, !"amdgpu_hostcall", i32 1}
-
 ; CHECK: AMDGPU HSA Metadata Parser Test: PASS

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present-v3.ll

-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
-
-; CHECK:              ---
-; CHECK:      amdhsa.kernels:
-; CHECK:        - .args:
-; CHECK-NEXT:       - .name:           a
-; CHECK-NEXT:         .offset:         0
-; CHECK-NEXT:         .size:           1
-; CHECK-NEXT:         .type_name:      char
-; CHECK-NEXT:         .value_kind:     by_value
-; CHECK-NEXT:       - .offset:         8
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_global_offset_x
-; CHECK-NEXT:       - .offset:         16
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_global_offset_y
-; CHECK-NEXT:       - .offset:         24
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_global_offset_z
-; CHECK-NEXT:       - .address_space:  global
-; CHECK-NEXT:         .offset:         32
-; CHECK-NEXT:         .size:           8
-; CHECK-NEXT:         .value_kind:     hidden_hostcall_buffer
-; CHECK:          .language:       OpenCL C
-; CHECK-NEXT:     .language_version:
-; CHECK-NEXT:       - 2
-; CHECK-NEXT:       - 0
-; CHECK:          .name:           test_kernel
-; CHECK:          .symbol:         test_kernel.kd
-
-define amdgpu_kernel void @test_kernel(i8 %a) #0
-    !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
-    !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
-  ret void
-}
-
-; CHECK:  amdhsa.version:
-; CHECK-NEXT: - 1
-; CHECK-NEXT: - 0
-
-attributes #0 = { optnone noinline "amdgpu-implicitarg-num-bytes"="48" }
-
-!1 = !{i32 0}
-!2 = !{!"none"}
-!3 = !{!"char"}
-!4 = !{!""}
-
-!opencl.ocl.version = !{!90}
-!90 = !{i32 2, i32 0}
-
-!llvm.module.flags = !{!0}
-!0 = !{i32 1, !"amdgpu_hostcall", i32 1}
-
-; PARSER: AMDGPU HSA Metadata Parser Test: PASS

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present.ll

-; RUN: llc -mtriple=amdgcn-amd-amdhsa --amdhsa-code-object-version=2 -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck %s
-; RUN: llc -mtriple=amdgcn-amd-amdhsa --amdhsa-code-object-version=2 -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s
-
-; CHECK: ---
-; CHECK:  Version: [ 1, 0 ]
-; CHECK:  Kernels:
-
-; CHECK:      - Name:            test_kernel
-; CHECK-NEXT:   SymbolName:      'test_kernel@kd'
-; CHECK-NEXT:   Language:        OpenCL C
-; CHECK-NEXT:   LanguageVersion: [ 2, 0 ]
-; CHECK-NEXT:   Args:
-; CHECK-NEXT:     - Name:          a
-; CHECK-NEXT:       TypeName:      char
-; CHECK-NEXT:       Size:          1
-; CHECK-NEXT:       Align:         1
-; CHECK-NEXT:       ValueKind:     ByValue
-; CHECK-NEXT:       AccQual:       Default
-; CHECK-NEXT:     - Size:          8
-; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetX
-; CHECK-NEXT:     - Size:          8
-; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetY
-; CHECK-NEXT:     - Size:          8
-; CHECK-NEXT:       Align:         8
-; CHECK-NEXT:       ValueKind:     HiddenGlobalOffsetZ
-; CHECK-NEXT:     - Size:            8
-; CHECK-NEXT:       Align:           8
-; CHECK-NEXT:       ValueKind:       HiddenHostcallBuffer
-; CHECK-NEXT:       AddrSpaceQual:   Global
-; CHECK-NOT:        ValueKind:     HiddenDefaultQueue
-; CHECK-NOT:        ValueKind:     HiddenCompletionAction
-
-declare <2 x i64> @__ockl_hostcall_internal(i8*, i32, i64, i64, i64, i64, i64, i64, i64, i64)
-
-define amdgpu_kernel void @test_kernel(i8 %a) #0
-    !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
-    !kernel_arg_base_type !3 !kernel_arg_type_qual !4 {
-  ret void
-}
-
-attributes #0 = { optnone noinline "amdgpu-implicitarg-num-bytes"="48" }
-
-!1 = !{i32 0}
-!2 = !{!"none"}
-!3 = !{!"char"}
-!4 = !{!""}
-
-!opencl.ocl.version = !{!90}
-!90 = !{i32 2, i32 0}
-
-; PARSER: AMDGPU HSA Metadata Parser Test: PASS

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-v3.ll

+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=3 < %s | FileCheck --check-prefix=CHECK %s
+
+declare void @function1()
+
+declare void @function2() #0
+
+; Function Attrs: noinline
+define void @function3(i8 addrspace(4)* %argptr, i8 addrspace(4)* addrspace(1)* %sink) #4 {
+  store i8 addrspace(4)* %argptr, i8 addrspace(4)* addrspace(1)* %sink, align 8
+  ret void
+}
+
+; Function Attrs: noinline
+define void @function4(i64 %arg, i64* %a) #4 {
+  store i64 %arg, i64* %a
+  ret void
+}
+
+; Function Attrs: noinline
+define void @function5(i8 addrspace(4)* %ptr, i64* %sink) #4 {
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 8
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %sink
+  ret void
+}
+
+; Function Attrs: nounwind readnone speculatable willreturn
+declare align 4 i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr() #1
+
+; CHECK: amdhsa.kernels:
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel10
+define amdgpu_kernel void @test_kernel10(i8* %a) #2 {
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Call to an extern function
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel20
+define amdgpu_kernel void @test_kernel20(i8* %a) #2 {
+  call void @function1()
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Explicit attribute on kernel
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel21
+define amdgpu_kernel void @test_kernel21(i8* %a) #3 {
+  call void @function1()
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Explicit attribute on extern callee
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel22
+define amdgpu_kernel void @test_kernel22(i8* %a) #2 {
+  call void @function2()
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Access more bytes than the pointer size
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel30
+define amdgpu_kernel void @test_kernel30(i128* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  %cast = bitcast i8 addrspace(4)* %gep to i128 addrspace(4)*
+  %x = load i128, i128 addrspace(4)* %cast
+  store i128 %x, i128* %a
+  ret void
+}
+
+; Typical load of hostcall buffer pointer
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel40
+define amdgpu_kernel void @test_kernel40(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 24
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Typical usage, overriden by explicit attribute on kernel
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel41
+define amdgpu_kernel void @test_kernel41(i64* %a) #3 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 24
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Access to implicit arg before the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel42
+define amdgpu_kernel void @test_kernel42(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Access to implicit arg after the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel43
+define amdgpu_kernel void @test_kernel43(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 32
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Accessing a byte just before the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel44
+define amdgpu_kernel void @test_kernel44(i8* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 23
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Accessing a byte inside the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel45
+define amdgpu_kernel void @test_kernel45(i8* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 24
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Accessing a byte inside the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel46
+define amdgpu_kernel void @test_kernel46(i8* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 31
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Accessing a byte just after the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel47
+define amdgpu_kernel void @test_kernel47(i8* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 32
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Access with an unknown offset
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel50
+define amdgpu_kernel void @test_kernel50(i8* %a, i32 %b) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 %b
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Multiple geps reaching the hostcall pointer argument.
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel51
+define amdgpu_kernel void @test_kernel51(i8* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep1 = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  %gep2 = getelementptr inbounds i8, i8 addrspace(4)* %gep1, i64 8
+  %x = load i8, i8 addrspace(4)* %gep2, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Multiple geps not reaching the hostcall pointer argument.
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel52
+define amdgpu_kernel void @test_kernel52(i8* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep1 = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  %gep2 = getelementptr inbounds i8, i8 addrspace(4)* %gep1, i64 16
+  %x = load i8, i8 addrspace(4)* %gep2, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Hostcall pointer used inside a function call
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel60
+define amdgpu_kernel void @test_kernel60(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 24
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  call void @function4(i64 %x, i64* %a)
+  ret void
+}
+
+; Hostcall pointer retrieved inside a function call; chain of geps
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel61
+define amdgpu_kernel void @test_kernel61(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  call void @function5(i8 addrspace(4)* %gep, i64* %a)
+  ret void
+}
+
+; Pointer captured
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel70
+define amdgpu_kernel void @test_kernel70(i8 addrspace(4)* addrspace(1)* %sink) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 42
+  store i8 addrspace(4)* %gep, i8 addrspace(4)* addrspace(1)* %sink, align 8
+  ret void
+}
+
+; Pointer captured inside function call
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel71
+define amdgpu_kernel void @test_kernel71(i8 addrspace(4)* addrspace(1)* %sink) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 42
+  call void @function3(i8 addrspace(4)* %gep, i8 addrspace(4)* addrspace(1)* %sink)
+  ret void
+}
+
+; Ineffective pointer capture
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel72
+define amdgpu_kernel void @test_kernel72() #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 42
+  store i8 addrspace(4)* %gep, i8 addrspace(4)* addrspace(1)* undef, align 8
+  ret void
+}
+
+attributes #0 = { "amdgpu-no-hostcall-ptr" }
+attributes #1 = { nounwind readnone speculatable willreturn }
+attributes #2 = { "amdgpu-implicitarg-num-bytes"="48" }
+attributes #3 = { "amdgpu-implicitarg-num-bytes"="48" "amdgpu-no-hostcall-ptr" }
+attributes #4 = { noinline }

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-v5.ll

+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=5 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 --amdhsa-code-object-version=5 < %s | FileCheck --check-prefix=CHECK %s
+
+declare void @function1()
+
+declare void @function2() #0
+
+; Function Attrs: noinline
+define void @function3(i8 addrspace(4)* %argptr, i8 addrspace(4)* addrspace(1)* %sink) #2 {
+  store i8 addrspace(4)* %argptr, i8 addrspace(4)* addrspace(1)* %sink, align 8
+  ret void
+}
+
+; Function Attrs: noinline
+define void @function4(i64 %arg, i64* %a) #2 {
+  store i64 %arg, i64* %a
+  ret void
+}
+
+; Function Attrs: noinline
+define void @function5(i8 addrspace(4)* %ptr, i64* %sink) #2 {
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 64
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %sink
+  ret void
+}
+
+; Function Attrs: nounwind readnone speculatable willreturn
+declare align 4 i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr() #1
+
+; CHECK: amdhsa.kernels:
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel10
+define amdgpu_kernel void @test_kernel10(i8* %a) {
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Call to an extern function
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel20
+define amdgpu_kernel void @test_kernel20(i8* %a) {
+  call void @function1()
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Explicit attribute on kernel
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel21
+define amdgpu_kernel void @test_kernel21(i8* %a) #0 {
+  call void @function1()
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Explicit attribute on extern callee
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel22
+define amdgpu_kernel void @test_kernel22(i8* %a) {
+  call void @function2()
+  store i8 3, i8* %a, align 1
+  ret void
+}
+
+; Access more bytes than the pointer size
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel30
+define amdgpu_kernel void @test_kernel30(i128* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 72
+  %cast = bitcast i8 addrspace(4)* %gep to i128 addrspace(4)*
+  %x = load i128, i128 addrspace(4)* %cast
+  store i128 %x, i128* %a
+  ret void
+}
+
+; Typical load of hostcall buffer pointer
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel40
+define amdgpu_kernel void @test_kernel40(i64* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 80
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Typical usage, overriden by explicit attribute on kernel
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel41
+define amdgpu_kernel void @test_kernel41(i64* %a) #0 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 80
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Access to implicit arg before the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel42
+define amdgpu_kernel void @test_kernel42(i64* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 72
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Access to implicit arg after the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel43
+define amdgpu_kernel void @test_kernel43(i64* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 88
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  store i64 %x, i64* %a
+  ret void
+}
+
+; Accessing a byte just before the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel44
+define amdgpu_kernel void @test_kernel44(i8* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 79
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Accessing a byte inside the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel45
+define amdgpu_kernel void @test_kernel45(i8* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 80
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Accessing a byte inside the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel46
+define amdgpu_kernel void @test_kernel46(i8* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 87
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Accessing a byte just after the hostcall pointer
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel47
+define amdgpu_kernel void @test_kernel47(i8* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 88
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Access with an unknown offset
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel50
+define amdgpu_kernel void @test_kernel50(i8* %a, i32 %b) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 %b
+  %x = load i8, i8 addrspace(4)* %gep, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Multiple geps reaching the hostcall pointer argument.
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel51
+define amdgpu_kernel void @test_kernel51(i8* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep1 = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  %gep2 = getelementptr inbounds i8, i8 addrspace(4)* %gep1, i64 64
+  %x = load i8, i8 addrspace(4)* %gep2, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Multiple geps not reaching the hostcall pointer argument.
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel52
+define amdgpu_kernel void @test_kernel52(i8* %a) {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep1 = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  %gep2 = getelementptr inbounds i8, i8 addrspace(4)* %gep1, i64 16
+  %x = load i8, i8 addrspace(4)* %gep2, align 1
+  store i8 %x, i8* %a, align 1
+  ret void
+}
+
+; Hostcall pointer used inside a function call
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel60
+define amdgpu_kernel void @test_kernel60(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 80
+  %cast = bitcast i8 addrspace(4)* %gep to i64 addrspace(4)*
+  %x = load i64, i64 addrspace(4)* %cast
+  call void @function4(i64 %x, i64* %a)
+  ret void
+}
+
+; Hostcall pointer retrieved inside a function call; chain of geps
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel61
+define amdgpu_kernel void @test_kernel61(i64* %a) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i64 16
+  call void @function5(i8 addrspace(4)* %gep, i64* %a)
+  ret void
+}
+
+; Pointer captured
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel70
+define amdgpu_kernel void @test_kernel70(i8 addrspace(4)* addrspace(1)* %sink) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 42
+  store i8 addrspace(4)* %gep, i8 addrspace(4)* addrspace(1)* %sink, align 8
+  ret void
+}
+
+; Pointer captured inside function call
+
+; CHECK:  - .args:
+; CHECK: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel71
+define amdgpu_kernel void @test_kernel71(i8 addrspace(4)* addrspace(1)* %sink) #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 42
+  call void @function3(i8 addrspace(4)* %gep, i8 addrspace(4)* addrspace(1)* %sink)
+  ret void
+}
+
+; Ineffective pointer capture
+
+; CHECK:  - .args:
+; CHECK-NOT: hidden_hostcall_buffer
+; CHECK-LABEL:    .name:           test_kernel72
+define amdgpu_kernel void @test_kernel72() #2 {
+  %ptr = tail call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %gep = getelementptr inbounds i8, i8 addrspace(4)* %ptr, i32 42
+  store i8 addrspace(4)* %gep, i8 addrspace(4)* addrspace(1)* undef, align 8
+  ret void
+}
+
+attributes #0 = { "amdgpu-no-hostcall-ptr" }
+attributes #1 = { nounwind readnone speculatable willreturn }
+attributes #2 = { noinline }

llvm/test/CodeGen/AMDGPU/propagate-flat-work-group-size.ll

 attributes #1 = { "amdgpu-flat-work-group-size"="64,128" }
 attributes #2 = { "amdgpu-flat-work-group-size"="64,64" }
 attributes #3 = { "amdgpu-flat-work-group-size"="128,256" }
 attributes #4 = { "amdgpu-flat-work-group-size"="512,1024" }
 attributes #5 = { "amdgpu-flat-work-group-size"="128,512" }
 attributes #6 = { "amdgpu-flat-work-group-size"="512,512" }
 attributes #7 = { "amdgpu-flat-work-group-size"="64,256" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-flat-work-group-size"="64,128" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR2]] = { "amdgpu-flat-work-group-size"="128,512" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR3]] = { "amdgpu-flat-work-group-size"="64,64" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR4]] = { "amdgpu-flat-work-group-size"="128,128" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR5]] = { "amdgpu-flat-work-group-size"="512,512" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR6]] = { "amdgpu-flat-work-group-size"="64,256" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR7]] = { "amdgpu-flat-work-group-size"="128,256" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR8]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-flat-work-group-size"="1,256" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-flat-work-group-size"="64,128" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-flat-work-group-size"="128,512" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR3]] = { "amdgpu-flat-work-group-size"="64,64" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR4]] = { "amdgpu-flat-work-group-size"="128,128" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR5]] = { "amdgpu-flat-work-group-size"="512,512" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR6]] = { "amdgpu-flat-work-group-size"="64,256" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR7]] = { "amdgpu-flat-work-group-size"="128,256" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR8]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/simple-indirect-call.ll

; GFX9-NEXT:    s_endpgm
   %fp = load void()*, void()** %fptr.cast
   call void %fp()
   ret void
 }
 
 ;.
 ; AKF_GCN: attributes #[[ATTR0]] = { "amdgpu-calls" "amdgpu-stack-objects" }
 ;.
-; ATTRIBUTOR_GCN: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; ATTRIBUTOR_GCN: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ; ATTRIBUTOR_GCN: attributes #[[ATTR1]] = { "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/uniform-work-group-attribute-missing.ll

 ; CHECK-NEXT:    ret void
 ;
   call void @foo()
   ret void
 }
 
 attributes #0 = { "uniform-work-group-size"="true" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ;.

llvm/test/CodeGen/AMDGPU/uniform-work-group-multistep.ll

 ; CHECK-NEXT:    ret void
 ;
   call void @internal2()
   ret void
 }
 
 attributes #0 = { "uniform-work-group-size"="true" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
 ;.

llvm/test/CodeGen/AMDGPU/uniform-work-group-nested-function-calls.ll

 ; CHECK-NEXT:    ret void
 ;
   call void @func2()
   ret void
 }
 
 attributes #2 = { "uniform-work-group-size"="true" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
 ;.

llvm/test/CodeGen/AMDGPU/uniform-work-group-prevent-attribute-propagation.ll

 ; CHECK-NEXT:    ret void
 ;
   call void @func()
   ret void
 }
 
 attributes #1 = { "uniform-work-group-size"="true" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR1]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
 ;.

llvm/test/CodeGen/AMDGPU/uniform-work-group-recursion-test.ll

;
   store i32 %r2, i32 addrspace(1)* %m
   ret void
 }
 
 ; nounwind and readnone are added to match attributor results.
 attributes #0 = { nounwind readnone }
 attributes #1 = { "uniform-work-group-size"="true" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { nounwind readnone "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
-; CHECK: attributes #[[ATTR1]] = { nounwind readnone "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
-; CHECK: attributes #[[ATTR2]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
+; CHECK: attributes #[[ATTR0]] = { nounwind readnone "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR1]] = { nounwind readnone "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
+; CHECK: attributes #[[ATTR2]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
 ;.

llvm/test/CodeGen/AMDGPU/uniform-work-group-test.ll

 ;
   call void @func2()
   call void @func3()
   ret void
 }
 
 attributes #0 = { "uniform-work-group-size"="false" }
 ;.
-; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
+; CHECK: attributes #[[ATTR0]] = { "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="false" }
 ;.

mlir/lib/Dialect/GPU/Transforms/SerializeToHsaco.cpp

for (std::unique_ptr<llvm::Module> &libModule : mbModules.getValue()) {
     if (err) {
       getOperation().emitError(
           "Unrecoverable failure during device library linking.");
       // We have no guaranties about the state of `ret`, so bail
       return nullptr;
     }
   }
 
-  // Set amdgpu_hostcall if host calls have been linked, as needed by newer LLVM
-  // FIXME: Is there a way to set this during printf() lowering that makes sense
-  if (ret->getFunction("__ockl_hostcall_internal"))
-    if (!ret->getModuleFlag("amdgpu_hostcall"))
-      ret->addModuleFlag(llvm::Module::Override, "amdgpu_hostcall", 1);
   return ret;
 }
 
 LogicalResult
 SerializeToHsacoPass::optimizeLlvm(llvm::Module &llvmModule,
                                    llvm::TargetMachine &targetMachine) {
   int optLevel = this->optLevel.getValue();
   if (optLevel < 0 || optLevel > 3)