AMDGPU: Use the implicit kernargs for code object version 5

[cfang]

Summary

Specifically, for trap handling, for targets that do not support getDoorbellID,
we load the queue_ptr from the implicit kernarg, and move queue_ptr to s[0:1].
To get aperture bases when targets do not have aperture registers, we load
private_base or shared_base directly from the implicit kernarg. In clang, we use
implicitarg_ptr + offsets to implement __builtin_amdgcn_workgroup_size_{xyz}.

[yaxunl]

clang/lib/CodeGen/CGBuiltin.cpp

16253–16254

Is the logic reversed?

[cfang]

clang/lib/CodeGen/CGBuiltin.cpp

16253–16254

You are right. Will correct it. Thanks

[cfang]

Correct the logic ordering error in selecting the intrinsic based on code object version.
Also update the code based on clang-format check and Lint suggestions.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16252–16254

Given that it's an offset from a different base, I think it would be cleaner to just branch around the two cases

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

410–412

This isn't covered by any test changes

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

326–328

You shouldn't merge these into the same enum. This enum should be renamed, this is for a different clover ABI

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

4876

You're repeating this long sequence to get the queue pointer in two places, should common these into a function to get the queue pointer. Alternatively, emit the intrinsic and move this expansion into a lowering of the queue pointer intrinsic

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

155

The code object version should probably come from the IR, not a global opt

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16243

Isn't the total size 256?

16244

If I remember correctly we require 8 byte alignment (not that it makes much difference)

[sameerds]

Please fix the commit description so that the first line is self-contained and separated from the rest by a blank line. This matters a lot when looking at the output of "git log ---oneline". In particular, the start with "in this work we implement" is pretty much unnecessary.

[cfang]

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

155

Can you be explicit how to get code object version from the IR?

[cfang]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

4876

We are loading different implicit kernel arguments in these two place, one is for queue_ptr, and another is for private_base/shared_base. I can try to figure out whether we can factor out some common part.

[scott.linder]

clang/lib/CodeGen/CGBuiltin.cpp

16255

I agree with Matt above; this comment seems like it should also be updated for the v5 case?

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

4388

typo

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

1817

These parens are redundant

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

155

IIUC the global opt var is the best we have right now, and any improvement to that situation is orthogonal to this change. I would vote that this not block the patch under review

[cfang]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

410–412

Are the tests of @llvm_amdgcn_is_private, @llvm_amdgcn_is_shared and @addrspacecast under v5 covering this? For example, for @llvm_amdgcn_is_shared, "s_load_dword s0, s[6:7], 0xcc" is generated to load
the shared_base and s[6:7] holds the argument ptr.

[cfang]

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

326–328

I am not clear what is the clover ABI. But GRID_DIM and GRID_OFFSET have never been used. Can we simply remove them?

[cfang]

Update based on review feedback

1. Update EmitAMDGPUImplicitArgPtr: DereferenceableBytes=256 and Align=8
2. Branch around v5 to emit ImplicitArgPtr or DispatchPtr,
3. Remove unused entries (GROD_DIM and GRID_OFFSET) in ImplicitParameter enum
4. Fix typos and remove unnecessary "(" and ")".

[cfang]

Rebase and update LIT tests.

[sameerds]

Please remove the JIRA ticket number from the commit description.

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

4876

There is still a lot common between this hunk and the previous hunk. It should be refactored into a common function.

[cfang]

1. Introduce a common function, SITargetLowering::loadImplicitKernelArgument, which is used

in both getSegmentAperture and lowerTrapHsaQueuePtr.

2. Define enum ImplicitKernargOffset to get implicit kernel argument offset directly.
   • This eliminates the functions to get the offset.

[arsenm]

clang/lib/CodeGen/CGBuiltin.cpp

16259

Extra space after (

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

1839

Use buildPtrAdd (not sure why we still have materializePtrAdd)

4850

ST is available in the class

llvm/lib/Target/AMDGPU/SIDefines.h

784

Add a COV5 suffix? Probably should also wrap in a namespace

[cfang]

Ping!

[cfang]

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

1839

Use buildPtrAdd ! Thanks.

4850

Remove it. Thanks.

llvm/lib/Target/AMDGPU/SIDefines.h

784

Add COV5 suffix, and wrap in a namespace of ImplicitArg. Rename the type to
Offset. So it is of AMDGPU::ImplicitArg::Offset type.

[cfang]

Update based on Matt's comments:

1. Use buildPtrAdd
2. Remove a space
3. Add suffix for the enum definition and also wrap with a namespace
4. Remove the redundant def of ST (SubTarget)
5. Updated according to clang-format

[arsenm]

llvm/lib/Target/AMDGPU/SIDefines.h

784

I meant suffix on the enum itself, not on each individual field

[sameerds]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

546

I think we should keep using the original getHeapPtrImplicitArgPosition(). Hardcoding the enum here doesn't necessarily make the code more readable. And later if we have a different value in COV6, we will end up reintroducing a check for the code-object-version anyway. That check can be encapsulated within the get...ArgPosition() family of functions.

llvm/lib/Target/AMDGPU/SIDefines.h

786

This should be "UPTO_COV4". Or if we really want to say COV5, then "BEFORE_COV5" or "PRE_COV5". But to me, "UPTO_COV4" is the clearest.

[cfang]

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

546

I personally do not have any preference here to use the offset enum or a function. Similarly I also could not understand why the offset of an existing argument changes its value across code object versions. However, I do think this should not block the current work of code object 5.

llvm/lib/Target/AMDGPU/SIDefines.h

784

OK, will change to the suffix of the enum itself.

786

This issue does not exist if the _COV5 suffix is for the type of this enum because we only consider COV5 in this definition.

[cfang]

A minor change: add suffix to the enum itself instead of the individual field.
Also remove the "Fixes" field in the summary (commit message).

[arsenm]

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

323–325

This isn't FIRST_IMPLICIT unless you reuse the same value. There's also no use of FIRST_IMPLICIT so you might as well remove it

[cfang]

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

323–325

FIRST_IMPLICIT is used in multiple places, for example, getImplicitArgPtr, which is used to lower Intrinsic::amdgcn_implicitarg_ptr. WE need amdgcn_implicitarg_ptr + offset to reference the implicit kernel argument.

clang/lib/CodeGen/CGBuiltin.cpp

   if (!E)
     return Call;
   QualType BuiltinRetType = E->getType();
   auto *RetTy = cast<llvm::PointerType>(CGF.ConvertType(BuiltinRetType));
   if (RetTy == Call->getType())
     return Call;
   return CGF.Builder.CreateAddrSpaceCast(Call, RetTy);
 }
 
+Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) {
+  auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_implicitarg_ptr);
+  auto *Call = CGF.Builder.CreateCall(F);
+  Call->addRetAttr(
+      Attribute::getWithDereferenceableBytes(Call->getContext(), 256));
+  Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(8)));
+  return Call;
+}
+
 // \p Index is 0, 1, and 2 for x, y, and z dimension, respectively.
 Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
-  const unsigned XOffset = 4;
-  auto *DP = EmitAMDGPUDispatchPtr(CGF);
-  // Indexing the HSA kernel_dispatch_packet struct.
-  auto *Offset = llvm::ConstantInt::get(CGF.Int32Ty, XOffset + Index * 2);
+  bool IsCOV_5 = CGF.getTarget().getTargetOpts().CodeObjectVersion ==
+                 clang::TargetOptions::COV_5;
+  Constant *Offset;
+  Value *DP;
+  if ( IsCOV_5) {
+    // Indexing the implicit kernarg segment.
+    Offset = llvm::ConstantInt::get(CGF.Int32Ty, 12 + Index * 2);
+    DP = EmitAMDGPUImplicitArgPtr(CGF);
+  } else {
+    // Indexing the HSA kernel_dispatch_packet struct.
+    Offset = llvm::ConstantInt::get(CGF.Int32Ty, 4 + Index * 2);
+    DP = EmitAMDGPUDispatchPtr(CGF);
+  }
+
   auto *GEP = CGF.Builder.CreateGEP(CGF.Int8Ty, DP, Offset);
   auto *DstTy =
       CGF.Int16Ty->getPointerTo(GEP->getType()->getPointerAddressSpace());
   auto *Cast = CGF.Builder.CreateBitCast(GEP, DstTy);
   auto *LD = CGF.Builder.CreateLoad(
       Address(Cast, CGF.Int16Ty, CharUnits::fromQuantity(2)));
   llvm::MDBuilder MDHelper(CGF.getLLVMContext());
   llvm::MDNode *RNode = MDHelper.createRange(APInt(16, 1),

clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
 // RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
-// RUN:     | FileCheck %s
+// RUN:     | FileCheck -check-prefix=PRECOV5 %s
+
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
+// RUN:     -fcuda-is-device -mcode-object-version=5 -emit-llvm -o - -x hip %s \
+// RUN:     | FileCheck -check-prefix=COV5 %s
 
 #include "Inputs/cuda.h"
 
-// CHECK-LABEL: test_get_workgroup_size
-// CHECK: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
-// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 4
-// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
-// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 6
-// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
-// CHECK: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 8
-// CHECK: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
+// PRECOV5-LABEL: test_get_workgroup_size
+// PRECOV5: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+// PRECOV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 4
+// PRECOV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
+// PRECOV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 6
+// PRECOV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
+// PRECOV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 8
+// PRECOV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
+
+// COV5-LABEL: test_get_workgroup_size
+// COV5: call align 8 dereferenceable(256) i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+// COV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 12
+// COV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
+// COV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 14
+// COV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
+// COV5: getelementptr i8, i8 addrspace(4)* %{{.*}}, i32 16
+// COV5: load i16, i16 addrspace(4)* %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load
 __device__ void test_get_workgroup_size(int d, int *out)
 {
   switch (d) {
   case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;
   case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
   case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
   default: *out = 0;
   }
 }
 
 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}

llvm/lib/Target/AMDGPU/AMDGPUAttributor.cpp

     auto Pos = llvm::AMDGPU::getHostcallImplicitArgPosition();
     AAPointerInfo::OffsetAndSize OAS(Pos, 8);
     return funcRetrievesImplicitKernelArg(A, OAS);
   }
 
   bool funcRetrievesHeapPtr(Attributor &A) {
     if (AMDGPU::getAmdhsaCodeObjectVersion() != 5)
       return false;
-    auto Pos = llvm::AMDGPU::getHeapPtrImplicitArgPosition();
-    AAPointerInfo::OffsetAndSize OAS(Pos, 8);
+    AAPointerInfo::OffsetAndSize OAS(AMDGPU::HEAP_PTR_OFFSET, 8);
     return funcRetrievesImplicitKernelArg(A, OAS);
   }
 
   bool funcRetrievesQueuePtr(Attributor &A) {
     if (AMDGPU::getAmdhsaCodeObjectVersion() != 5)
       return false;
-    auto Pos = llvm::AMDGPU::getQueuePtrImplicitArgPosition();
-    AAPointerInfo::OffsetAndSize OAS(Pos, 8);
+    AAPointerInfo::OffsetAndSize OAS(AMDGPU::QUEUE_PTR_OFFSET, 8);
     return funcRetrievesImplicitKernelArg(A, OAS);
   }
 
   bool funcRetrievesImplicitKernelArg(Attributor &A,
                                       AAPointerInfo::OffsetAndSize OAS) {
     // 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

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.h

 
   SDValue loadInputValue(SelectionDAG &DAG,
                          const TargetRegisterClass *RC,
                          EVT VT, const SDLoc &SL,
                          const ArgDescriptor &Arg) const;
 
   enum ImplicitParameter {
     FIRST_IMPLICIT,
-    GRID_DIM = FIRST_IMPLICIT,
-    GRID_OFFSET,
+    PRIVATE_BASE,
+    SHARED_BASE,
+    QUEUE_PTR,
   };
 
   /// Helper function that returns the byte offset of the given
   /// type of implicit parameter.
   uint32_t getImplicitParameterOffset(const MachineFunction &MF,
                                       const ImplicitParameter Param) const;
 
   MVT getFenceOperandTy(const DataLayout &DL) const override {
     return MVT::i32;
   }
 

llvm/lib/Target/AMDGPU/AMDGPUISelLowering.cpp

   const AMDGPUMachineFunction *MFI = MF.getInfo<AMDGPUMachineFunction>();
   const AMDGPUSubtarget &ST =
       AMDGPUSubtarget::get(getTargetMachine(), MF.getFunction());
   unsigned ExplicitArgOffset = ST.getExplicitKernelArgOffset(MF.getFunction());
   const Align Alignment = ST.getAlignmentForImplicitArgPtr();
   uint64_t ArgOffset = alignTo(MFI->getExplicitKernArgSize(), Alignment) +
                        ExplicitArgOffset;
   switch (Param) {
-  case GRID_DIM:
+  case FIRST_IMPLICIT:
     return ArgOffset;
-  case GRID_OFFSET:
-    return ArgOffset + 4;
+  case PRIVATE_BASE:
+    return ArgOffset + AMDGPU::PRIVATE_BASE_OFFSET;
+  case SHARED_BASE:
+    return ArgOffset + AMDGPU::SHARED_BASE_OFFSET;
+  case QUEUE_PTR:
+    return ArgOffset + AMDGPU::QUEUE_PTR_OFFSET;
   }
   llvm_unreachable("unexpected implicit parameter type");
 }
 
 #define NODE_NAME_CASE(node) case AMDGPUISD::node: return #node;
 
 const char* AMDGPUTargetLowering::getTargetNodeName(unsigned Opcode) const {
   switch ((AMDGPUISD::NodeType)Opcode) {

llvm/lib/Target/AMDGPU/AMDGPULegalizerInfo.cpp

       .addDef(GetReg)
       .addImm(Encoding);
     MRI.setType(GetReg, S32);
 
     auto ShiftAmt = B.buildConstant(S32, WidthM1 + 1);
     return B.buildShl(S32, GetReg, ShiftAmt).getReg(0);
   }
 
+  // TODO: can we be smarter about machine pointer info?
+  MachinePointerInfo PtrInfo(AMDGPUAS::CONSTANT_ADDRESS);
+  // For code object version 5, private_base and shared_base are passed through
+  // implicit kernargs.
+  if (AMDGPU::getAmdhsaCodeObjectVersion() == 5) {
+    AMDGPUTargetLowering::ImplicitParameter Param =
+        AS == AMDGPUAS::LOCAL_ADDRESS ? AMDGPUTargetLowering::SHARED_BASE
+                                      : AMDGPUTargetLowering::PRIVATE_BASE;
+    uint64_t Offset =
+        ST.getTargetLowering()->getImplicitParameterOffset(B.getMF(), Param);
+
+    Register KernargPtrReg = MRI.createGenericVirtualRegister(
+        LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+
+    if (!loadInputValue(KernargPtrReg, B,
+                        AMDGPUFunctionArgInfo::KERNARG_SEGMENT_PTR))
+      return Register();
+
+    MachineMemOperand *MMO = MF.getMachineMemOperand(
+        PtrInfo,
+        MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
+            MachineMemOperand::MOInvariant,
+        LLT::scalar(32), commonAlignment(Align(64), Offset));
+
+    Register LoadAddr;
+    // Pointer address
+    B.materializePtrAdd(LoadAddr, KernargPtrReg, LLT::scalar(64), Offset);
+    // Load address
+    return B.buildLoad(S32, LoadAddr, *MMO).getReg(0);
+  }
+
   Register QueuePtr = MRI.createGenericVirtualRegister(
     LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
 
   if (!loadInputValue(QueuePtr, B, AMDGPUFunctionArgInfo::QUEUE_PTR))
     return Register();
 
   // Offset into amd_queue_t for group_segment_aperture_base_hi /
   // private_segment_aperture_base_hi.
   uint32_t StructOffset = (AS == AMDGPUAS::LOCAL_ADDRESS) ? 0x40 : 0x44;
 
-  // TODO: can we be smarter about machine pointer info?
-  MachinePointerInfo PtrInfo(AMDGPUAS::CONSTANT_ADDRESS);
   MachineMemOperand *MMO = MF.getMachineMemOperand(
       PtrInfo,
       MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
           MachineMemOperand::MOInvariant,
       LLT::scalar(32), commonAlignment(Align(64), StructOffset));
 
   Register LoadAddr;
 

     MachineInstr &MI, MachineRegisterInfo &MRI, MachineIRBuilder &B) const {
   B.buildInstr(AMDGPU::S_ENDPGM).addImm(0);
   MI.eraseFromParent();
   return true;
 }
 
 bool AMDGPULegalizerInfo::legalizeTrapHsaQueuePtr(
     MachineInstr &MI, MachineRegisterInfo &MRI, MachineIRBuilder &B) const {
+  MachineFunction &MF = B.getMF();
+  const GCNSubtarget &ST = MF.getSubtarget<GCNSubtarget>();
+  const LLT S64 = LLT::scalar(64);
+
+  Register SGPR01(AMDGPU::SGPR0_SGPR1);
+  // For code object version 5, queue_ptr is passed through implicit kernarg.
+  if (AMDGPU::getAmdhsaCodeObjectVersion() == 5) {
+    AMDGPUTargetLowering::ImplicitParameter Param =
+        AMDGPUTargetLowering::QUEUE_PTR;
+    uint64_t Offset =
+        ST.getTargetLowering()->getImplicitParameterOffset(B.getMF(), Param);
+
+    Register KernargPtrReg = MRI.createGenericVirtualRegister(
+        LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
+
+    if (!loadInputValue(KernargPtrReg, B,
+                        AMDGPUFunctionArgInfo::KERNARG_SEGMENT_PTR))
+      return false;
+
+    // TODO: can we be smarter about machine pointer info?
+    MachinePointerInfo PtrInfo(AMDGPUAS::CONSTANT_ADDRESS);
+    MachineMemOperand *MMO = MF.getMachineMemOperand(
+        PtrInfo,
+        MachineMemOperand::MOLoad | MachineMemOperand::MODereferenceable |
+            MachineMemOperand::MOInvariant,
+        LLT::scalar(64), commonAlignment(Align(64), Offset));
+
+    // Pointer address
+    Register LoadAddr;
+    B.materializePtrAdd(LoadAddr, KernargPtrReg, LLT::scalar(64), Offset);
+    // Load address
+    Register Temp = B.buildLoad(S64, LoadAddr, *MMO).getReg(0);
+    B.buildCopy(SGPR01, Temp);
+    B.buildInstr(AMDGPU::S_TRAP)
+        .addImm(static_cast<unsigned>(GCNSubtarget::TrapID::LLVMAMDHSATrap))
+        .addReg(SGPR01, RegState::Implicit);
+    MI.eraseFromParent();
+    return true;
+  }
+
   // Pass queue pointer to trap handler as input, and insert trap instruction
   // Reference: https://llvm.org/docs/AMDGPUUsage.html#trap-handler-abi
   Register LiveIn =
     MRI.createGenericVirtualRegister(LLT::pointer(AMDGPUAS::CONSTANT_ADDRESS, 64));
   if (!loadInputValue(LiveIn, B, AMDGPUFunctionArgInfo::QUEUE_PTR))
     return false;
 
-  Register SGPR01(AMDGPU::SGPR0_SGPR1);
   B.buildCopy(SGPR01, LiveIn);
   B.buildInstr(AMDGPU::S_TRAP)
       .addImm(static_cast<unsigned>(GCNSubtarget::TrapID::LLVMAMDHSATrap))
       .addReg(SGPR01, RegState::Implicit);
 
   MI.eraseFromParent();
   return true;
 }

llvm/lib/Target/AMDGPU/SIDefines.h

 enum OpSel : uint64_t {
   OP_SEL_HI_0 = UINT64_C(1) << 59,
   OP_SEL_HI_1 = UINT64_C(1) << 60,
   OP_SEL_HI_2 = UINT64_C(1) << 14,
 };
 
 } // namespace VOP3PEncoding
 
+// Implicit kernel argument offset for code object version 5.
+enum ImplicitKernargOffset : unsigned {
+  HOSTCALL_PTR_OFFSET = 80,
+  HEAP_PTR_OFFSET = 96,
+  PRIVATE_BASE_OFFSET = 192,
+  SHARED_BASE_OFFSET = 196,
+  QUEUE_PTR_OFFSET = 200,
+};
 } // namespace AMDGPU
 
 #define R_00B028_SPI_SHADER_PGM_RSRC1_PS                                0x00B028
 #define   S_00B028_VGPRS(x)                                           (((x) & 0x3F) << 0)
 #define   S_00B028_SGPRS(x)                                           (((x) & 0x0F) << 6)
 #define   S_00B028_MEM_ORDERED(x)                                     (((x) & 0x1) << 25)
 #define   G_00B028_MEM_ORDERED(x)                                     (((x) >> 25) & 0x1)
 #define   C_00B028_MEM_ORDERED                                        0xFDFFFFFF

llvm/lib/Target/AMDGPU/SIISelLowering.h

   SDValue lowerKernArgParameterPtr(SelectionDAG &DAG, const SDLoc &SL,
                                    SDValue Chain, uint64_t Offset) const;
   SDValue getImplicitArgPtr(SelectionDAG &DAG, const SDLoc &SL) const;
   SDValue lowerKernargMemParameter(SelectionDAG &DAG, EVT VT, EVT MemVT,
                                    const SDLoc &SL, SDValue Chain,
                                    uint64_t Offset, Align Alignment,
                                    bool Signed,
                                    const ISD::InputArg *Arg = nullptr) const;
+SDValue loadImplicitKernelArgument(SelectionDAG &DAG, MVT VT, const SDLoc &DL,
+                                     Align Alignment, ImplicitParameter Param) const;
 
   SDValue lowerStackParameter(SelectionDAG &DAG, CCValAssign &VA,
                               const SDLoc &SL, SDValue Chain,
                               const ISD::InputArg &Arg) const;
   SDValue getPreloadedValue(SelectionDAG &DAG,
                             const SIMachineFunctionInfo &MFI,
                             EVT VT,
                             AMDGPUFunctionArgInfo::PreloadedValue) const;

llvm/lib/Target/AMDGPU/SIISelLowering.cpp

 
 SDValue SITargetLowering::lowerTrapEndpgm(
     SDValue Op, SelectionDAG &DAG) const {
   SDLoc SL(Op);
   SDValue Chain = Op.getOperand(0);
   return DAG.getNode(AMDGPUISD::ENDPGM, SL, MVT::Other, Chain);
 }
 
+SDValue SITargetLowering::loadImplicitKernelArgument(SelectionDAG &DAG, MVT VT,
+    const SDLoc &DL, Align Alignment, ImplicitParameter Param) const {
+  MachineFunction &MF = DAG.getMachineFunction();
+  uint64_t Offset = getImplicitParameterOffset(MF, Param);
+  SDValue Ptr = lowerKernArgParameterPtr(DAG, DL, DAG.getEntryNode(), Offset);
+  MachinePointerInfo PtrInfo(AMDGPUAS::CONSTANT_ADDRESS);
+  return DAG.getLoad(
+      VT, DL, DAG.getEntryNode(), Ptr, PtrInfo, Alignment,
+      MachineMemOperand::MODereferenceable | MachineMemOperand::MOInvariant);
+}
+
 SDValue SITargetLowering::lowerTrapHsaQueuePtr(
     SDValue Op, SelectionDAG &DAG) const {
   SDLoc SL(Op);
   SDValue Chain = Op.getOperand(0);
 
-  MachineFunction &MF = DAG.getMachineFunction();
-  SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
-  Register UserSGPR = Info->getQueuePtrUserSGPR();
-
   SDValue QueuePtr;
-  if (UserSGPR == AMDGPU::NoRegister) {
-    // We probably are in a function incorrectly marked with
-    // amdgpu-no-queue-ptr. This is undefined. We don't want to delete the trap,
-    // so just use a null pointer.
-    QueuePtr = DAG.getConstant(0, SL, MVT::i64);
+  // For code object version 5, QueuePtr is passed through implicit kernarg.
+  if (AMDGPU::getAmdhsaCodeObjectVersion() == 5) {
+    QueuePtr = loadImplicitKernelArgument(DAG, MVT::i64, SL, Align(8),
+                                          QUEUE_PTR);
   } else {
-    QueuePtr = CreateLiveInRegister(
-      DAG, &AMDGPU::SReg_64RegClass, UserSGPR, MVT::i64);
+    MachineFunction &MF = DAG.getMachineFunction();
+    SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
+    Register UserSGPR = Info->getQueuePtrUserSGPR();
+
+    if (UserSGPR == AMDGPU::NoRegister) {
+      // We probably are in a function incorrectly marked with
+      // amdgpu-no-queue-ptr. This is undefined. We don't want to delete the trap,
+      // so just use a null pointer.
+      QueuePtr = DAG.getConstant(0, SL, MVT::i64);
+    } else {
+      QueuePtr = CreateLiveInRegister(DAG, &AMDGPU::SReg_64RegClass, UserSGPR,
+                                      MVT::i64);
+    }
   }
 
   SDValue SGPR01 = DAG.getRegister(AMDGPU::SGPR0_SGPR1, MVT::i64);
   SDValue ToReg = DAG.getCopyToReg(Chain, SL, SGPR01,
                                    QueuePtr, SDValue());
 
   uint64_t TrapID = static_cast<uint64_t>(GCNSubtarget::TrapID::LLVMAMDHSATrap);
   SDValue Ops[] = {
     ToReg,
     DAG.getTargetConstant(TrapID, SL, MVT::i16),
     SGPR01,
     ToReg.getValue(1)
   };

 
     SDValue EncodingImm = DAG.getTargetConstant(Encoding, DL, MVT::i16);
     SDValue ApertureReg = SDValue(
         DAG.getMachineNode(AMDGPU::S_GETREG_B32, DL, MVT::i32, EncodingImm), 0);
     SDValue ShiftAmount = DAG.getTargetConstant(WidthM1 + 1, DL, MVT::i32);
     return DAG.getNode(ISD::SHL, DL, MVT::i32, ApertureReg, ShiftAmount);
   }
 
+  // For code object version 5, private_base and shared_base are passed through
+  // implicit kernargs.
+  if (AMDGPU::getAmdhsaCodeObjectVersion() == 5) {
+    ImplicitParameter Param =
+        (AS == AMDGPUAS::LOCAL_ADDRESS) ? SHARED_BASE : PRIVATE_BASE;
+    return loadImplicitKernelArgument(DAG, MVT::i32, DL, Align(4), Param);
+  }
+
   MachineFunction &MF = DAG.getMachineFunction();
   SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
   Register UserSGPR = Info->getQueuePtrUserSGPR();
   if (UserSGPR == AMDGPU::NoRegister) {
     // We probably are in a function incorrectly marked with
     // amdgpu-no-queue-ptr. This is undefined.
     return DAG.getUNDEF(MVT::i32);
   }

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

 bool isHsaAbiVersion5(const MCSubtargetInfo *STI);
 /// \returns True if HSA OS ABI Version identification is 3 and above,
 /// false otherwise.
 bool isHsaAbiVersion3AndAbove(const MCSubtargetInfo *STI);
 
 /// \returns The offset of the hostcall pointer argument from implicitarg_ptr
 unsigned getHostcallImplicitArgPosition();
 
-/// \returns The offset of the heap ptr argument from implicitarg_ptr
-unsigned getHeapPtrImplicitArgPosition();
-
-/// \returns The offset of the queue ptr argument from implicitarg_ptr
-unsigned getQueuePtrImplicitArgPosition();
-
 /// \returns Code object version.
 unsigned getAmdhsaCodeObjectVersion();
 
 struct GcnBufferFormatInfo {
   unsigned Format;
   unsigned BitsPerComp;
   unsigned NumComponents;
   unsigned NumFormat;

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

 // central TD file.
 unsigned getHostcallImplicitArgPosition() {
   switch (AmdhsaCodeObjectVersion) {
   case 2:
   case 3:
   case 4:
     return 24;
   case 5:
-    return 80;
+    return AMDGPU::HOSTCALL_PTR_OFFSET;
   default:
     llvm_unreachable("Unexpected code object version");
     return 0;
   }
 }
 
-unsigned getHeapPtrImplicitArgPosition() {
-  if (AmdhsaCodeObjectVersion == 5)
-    return 96;
-  llvm_unreachable("hidden_heap is supported only by code object version 5");
-  return 0;
-}
-
-unsigned getQueuePtrImplicitArgPosition() {
-  if (AmdhsaCodeObjectVersion == 5)
-    return 200;
-  llvm_unreachable("queue_ptr is supported only by code object version 5");
-  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/test/CodeGen/AMDGPU/GlobalISel/implicit-kernarg-backend-usage-global-isel.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 --amdhsa-code-object-version=3 < %s | FileCheck --check-prefix=GFX8V3 %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 --amdhsa-code-object-version=4 < %s | FileCheck --check-prefix=GFX8V4 %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 --amdhsa-code-object-version=5 < %s | FileCheck --check-prefix=GFX8V5 %s
+
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdhsa-code-object-version=3 < %s | FileCheck --check-prefixes=GFX9V3 %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdhsa-code-object-version=4 < %s | FileCheck --check-prefixes=GFX9V4 %s
+; RUN: llc -global-isel -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdhsa-code-object-version=5 < %s | FileCheck --check-prefixes=GFX9V5 %s
+
+define amdgpu_kernel void @addrspacecast(i32 addrspace(5)* %ptr.private, i32 addrspace(3)* %ptr.local) {
+; GFX8V3-LABEL: addrspacecast:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V3-NEXT:    s_load_dword s3, s[4:5], 0x44
+; GFX8V3-NEXT:    s_load_dword s5, s[4:5], 0x40
+; GFX8V3-NEXT:    v_mov_b32_e32 v2, 1
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_mov_b32 s2, s0
+; GFX8V3-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX8V3-NEXT:    s_cselect_b64 s[2:3], s[2:3], 0
+; GFX8V3-NEXT:    s_mov_b32 s4, s1
+; GFX8V3-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s2
+; GFX8V3-NEXT:    s_cselect_b64 s[0:1], s[4:5], 0
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s3
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v2
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    v_mov_b32_e32 v2, 2
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v2
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: addrspacecast:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V4-NEXT:    s_load_dword s3, s[4:5], 0x44
+; GFX8V4-NEXT:    s_load_dword s5, s[4:5], 0x40
+; GFX8V4-NEXT:    v_mov_b32_e32 v2, 1
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_mov_b32 s2, s0
+; GFX8V4-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX8V4-NEXT:    s_cselect_b64 s[2:3], s[2:3], 0
+; GFX8V4-NEXT:    s_mov_b32 s4, s1
+; GFX8V4-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s2
+; GFX8V4-NEXT:    s_cselect_b64 s[0:1], s[4:5], 0
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s3
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v2
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    v_mov_b32_e32 v2, 2
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v2
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: addrspacecast:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX8V5-NEXT:    s_load_dword s3, s[4:5], 0xc8
+; GFX8V5-NEXT:    s_load_dword s5, s[4:5], 0xcc
+; GFX8V5-NEXT:    v_mov_b32_e32 v2, 1
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_mov_b32 s2, s0
+; GFX8V5-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX8V5-NEXT:    s_cselect_b64 s[2:3], s[2:3], 0
+; GFX8V5-NEXT:    s_mov_b32 s4, s1
+; GFX8V5-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s2
+; GFX8V5-NEXT:    s_cselect_b64 s[0:1], s[4:5], 0
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s3
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v2
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    v_mov_b32_e32 v2, 2
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v2
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: addrspacecast:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V3-NEXT:    s_getreg_b32 s2, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s3, s2, 16
+; GFX9V3-NEXT:    s_getreg_b32 s4, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V3-NEXT:    v_mov_b32_e32 v2, 1
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    s_mov_b32 s2, s0
+; GFX9V3-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX9V3-NEXT:    s_cselect_b64 s[2:3], s[2:3], 0
+; GFX9V3-NEXT:    s_lshl_b32 s5, s4, 16
+; GFX9V3-NEXT:    s_mov_b32 s4, s1
+; GFX9V3-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s2
+; GFX9V3-NEXT:    s_cselect_b64 s[0:1], s[4:5], 0
+; GFX9V3-NEXT:    v_mov_b32_e32 v1, s3
+; GFX9V3-NEXT:    flat_store_dword v[0:1], v2
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V3-NEXT:    v_mov_b32_e32 v2, 2
+; GFX9V3-NEXT:    v_mov_b32_e32 v1, s1
+; GFX9V3-NEXT:    flat_store_dword v[0:1], v2
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: addrspacecast:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V4-NEXT:    s_getreg_b32 s2, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s3, s2, 16
+; GFX9V4-NEXT:    s_getreg_b32 s4, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V4-NEXT:    v_mov_b32_e32 v2, 1
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    s_mov_b32 s2, s0
+; GFX9V4-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX9V4-NEXT:    s_cselect_b64 s[2:3], s[2:3], 0
+; GFX9V4-NEXT:    s_lshl_b32 s5, s4, 16
+; GFX9V4-NEXT:    s_mov_b32 s4, s1
+; GFX9V4-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s2
+; GFX9V4-NEXT:    s_cselect_b64 s[0:1], s[4:5], 0
+; GFX9V4-NEXT:    v_mov_b32_e32 v1, s3
+; GFX9V4-NEXT:    flat_store_dword v[0:1], v2
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V4-NEXT:    v_mov_b32_e32 v2, 2
+; GFX9V4-NEXT:    v_mov_b32_e32 v1, s1
+; GFX9V4-NEXT:    flat_store_dword v[0:1], v2
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: addrspacecast:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V5-NEXT:    s_getreg_b32 s2, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s3, s2, 16
+; GFX9V5-NEXT:    s_getreg_b32 s4, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V5-NEXT:    v_mov_b32_e32 v2, 1
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    s_mov_b32 s2, s0
+; GFX9V5-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX9V5-NEXT:    s_cselect_b64 s[2:3], s[2:3], 0
+; GFX9V5-NEXT:    s_lshl_b32 s5, s4, 16
+; GFX9V5-NEXT:    s_mov_b32 s4, s1
+; GFX9V5-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s2
+; GFX9V5-NEXT:    s_cselect_b64 s[0:1], s[4:5], 0
+; GFX9V5-NEXT:    v_mov_b32_e32 v1, s3
+; GFX9V5-NEXT:    flat_store_dword v[0:1], v2
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V5-NEXT:    v_mov_b32_e32 v2, 2
+; GFX9V5-NEXT:    v_mov_b32_e32 v1, s1
+; GFX9V5-NEXT:    flat_store_dword v[0:1], v2
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %flat.private = addrspacecast i32 addrspace(5)* %ptr.private to i32*
+  %flat.local = addrspacecast i32 addrspace(3)* %ptr.local to i32*
+  store volatile i32 1, i32* %flat.private
+  store volatile i32 2, i32* %flat.local
+  ret void
+}
+
+define amdgpu_kernel void @llvm_amdgcn_is_shared(i8* %ptr) {
+; GFX8V3-LABEL: llvm_amdgcn_is_shared:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_load_dword s0, s[4:5], 0x40
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V3-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: llvm_amdgcn_is_shared:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_load_dword s0, s[4:5], 0x40
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V4-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: llvm_amdgcn_is_shared:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_load_dword s0, s[4:5], 0xcc
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V5-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: llvm_amdgcn_is_shared:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V3-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX9V3-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V3-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: llvm_amdgcn_is_shared:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V4-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX9V4-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V4-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: llvm_amdgcn_is_shared:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V5-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX9V5-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V5-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %is.shared = call i1 @llvm.amdgcn.is.shared(i8* %ptr)
+  %zext = zext i1 %is.shared to i32
+  store volatile i32 %zext, i32 addrspace(1)* undef
+  ret void
+}
+
+define amdgpu_kernel void @llvm_amdgcn_is_private(i8* %ptr) {
+; GFX8V3-LABEL: llvm_amdgcn_is_private:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_load_dword s0, s[4:5], 0x44
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V3-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: llvm_amdgcn_is_private:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_load_dword s0, s[4:5], 0x44
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V4-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: llvm_amdgcn_is_private:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_load_dword s0, s[4:5], 0xc8
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V5-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: llvm_amdgcn_is_private:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V3-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX9V3-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V3-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: llvm_amdgcn_is_private:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V4-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX9V4-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V4-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: llvm_amdgcn_is_private:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V5-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX9V5-NEXT:    s_cselect_b32 s0, 1, 0
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V5-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %is.private = call i1 @llvm.amdgcn.is.private(i8* %ptr)
+  %zext = zext i1 %is.private to i32
+  store volatile i32 %zext, i32 addrspace(1)* undef
+  ret void
+}
+
+define amdgpu_kernel void @llvm_trap() {
+; GFX8V3-LABEL: llvm_trap:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_mov_b64 s[0:1], s[4:5]
+; GFX8V3-NEXT:    s_trap 2
+;
+; GFX8V4-LABEL: llvm_trap:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_mov_b64 s[0:1], s[4:5]
+; GFX8V4-NEXT:    s_trap 2
+;
+; GFX8V5-LABEL: llvm_trap:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0xc8
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_trap 2
+;
+; GFX9V3-LABEL: llvm_trap:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_mov_b64 s[0:1], s[4:5]
+; GFX9V3-NEXT:    s_trap 2
+;
+; GFX9V4-LABEL: llvm_trap:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_trap 2
+;
+; GFX9V5-LABEL: llvm_trap:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_trap 2
+  call void @llvm.trap()
+  unreachable
+}
+
+define amdgpu_kernel void @llvm_debugtrap() {
+; GFX8V3-LABEL: llvm_debugtrap:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_trap 3
+;
+; GFX8V4-LABEL: llvm_debugtrap:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_trap 3
+;
+; GFX8V5-LABEL: llvm_debugtrap:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_trap 3
+;
+; GFX9V3-LABEL: llvm_debugtrap:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_trap 3
+;
+; GFX9V4-LABEL: llvm_debugtrap:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_trap 3
+;
+; GFX9V5-LABEL: llvm_debugtrap:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_trap 3
+  call void @llvm.debugtrap()
+  unreachable
+}
+
+define amdgpu_kernel void @llvm_amdgcn_queue_ptr(i64 addrspace(1)* %ptr)  {
+; GFX8V3-LABEL: llvm_amdgcn_queue_ptr:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s6
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s7
+; GFX8V3-NEXT:    s_add_u32 s0, s8, 8
+; GFX8V3-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V3-NEXT:    s_addc_u32 s1, s9, 0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V3-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s4
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s5
+; GFX8V3-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V3-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s10
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s11
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v3, s1
+; GFX8V3-NEXT:    v_mov_b32_e32 v2, s0
+; GFX8V3-NEXT:    flat_store_dwordx2 v[2:3], v[0:1]
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: llvm_amdgcn_queue_ptr:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s6
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s7
+; GFX8V4-NEXT:    s_add_u32 s0, s8, 8
+; GFX8V4-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V4-NEXT:    s_addc_u32 s1, s9, 0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V4-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s4
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s5
+; GFX8V4-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V4-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s10
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s11
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v3, s1
+; GFX8V4-NEXT:    v_mov_b32_e32 v2, s0
+; GFX8V4-NEXT:    flat_store_dwordx2 v[2:3], v[0:1]
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: llvm_amdgcn_queue_ptr:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_add_u32 s0, s6, 8
+; GFX8V5-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V5-NEXT:    s_addc_u32 s1, s7, 0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V5-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s4
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s5
+; GFX8V5-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s8
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s9
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v3, s1
+; GFX8V5-NEXT:    v_mov_b32_e32 v2, s0
+; GFX8V5-NEXT:    flat_store_dwordx2 v[2:3], v[0:1]
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: llvm_amdgcn_queue_ptr:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    v_mov_b32_e32 v2, 0
+; GFX9V3-NEXT:    global_load_ubyte v0, v2, s[6:7] glc
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    global_load_ubyte v0, v2, s[8:9] offset:8 glc
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    global_load_ubyte v0, v2, s[4:5] glc
+; GFX9V3-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s10
+; GFX9V3-NEXT:    v_mov_b32_e32 v1, s11
+; GFX9V3-NEXT:    ; kill: killed $sgpr6_sgpr7
+; GFX9V3-NEXT:    ; kill: killed $sgpr4_sgpr5
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: llvm_amdgcn_queue_ptr:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    v_mov_b32_e32 v2, 0
+; GFX9V4-NEXT:    global_load_ubyte v0, v2, s[6:7] glc
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    global_load_ubyte v0, v2, s[8:9] offset:8 glc
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    global_load_ubyte v0, v2, s[4:5] glc
+; GFX9V4-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s10
+; GFX9V4-NEXT:    v_mov_b32_e32 v1, s11
+; GFX9V4-NEXT:    ; kill: killed $sgpr6_sgpr7
+; GFX9V4-NEXT:    ; kill: killed $sgpr4_sgpr5
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: llvm_amdgcn_queue_ptr:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    v_mov_b32_e32 v2, 0
+; GFX9V5-NEXT:    global_load_ubyte v0, v[0:1], off glc
+; GFX9V5-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    global_load_ubyte v0, v2, s[6:7] offset:8 glc
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    global_load_ubyte v0, v2, s[4:5] glc
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s8
+; GFX9V5-NEXT:    v_mov_b32_e32 v1, s9
+; GFX9V5-NEXT:    ; kill: killed $sgpr4_sgpr5
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %queue.ptr = call i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
+  %implicitarg.ptr = call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %dispatch.ptr = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+  %dispatch.id = call i64 @llvm.amdgcn.dispatch.id()
+  %queue.load = load volatile i8, i8 addrspace(4)* %queue.ptr
+  %implicitarg.load = load volatile i8, i8 addrspace(4)* %implicitarg.ptr
+  %dispatch.load = load volatile i8, i8 addrspace(4)* %dispatch.ptr
+  store volatile i64 %dispatch.id, i64 addrspace(1)* %ptr
+  ret void
+}
+
+declare noalias i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
+declare noalias i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+declare i64 @llvm.amdgcn.dispatch.id()
+declare noalias i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+declare i1 @llvm.amdgcn.is.shared(i8*)
+declare i1 @llvm.amdgcn.is.private(i8*)
+declare void @llvm.trap()
+declare void @llvm.debugtrap()

llvm/test/CodeGen/AMDGPU/implicit-kernarg-backend-usage.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 --amdhsa-code-object-version=3 < %s | FileCheck --check-prefix=GFX8V3 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 --amdhsa-code-object-version=4 < %s | FileCheck --check-prefix=GFX8V4 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx803 --amdhsa-code-object-version=5 < %s | FileCheck --check-prefix=GFX8V5 %s
+
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdhsa-code-object-version=3 < %s | FileCheck --check-prefixes=GFX9V3 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdhsa-code-object-version=4 < %s | FileCheck --check-prefixes=GFX9V4 %s
+; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx906 --amdhsa-code-object-version=5 < %s | FileCheck --check-prefixes=GFX9V5 %s
+
+define amdgpu_kernel void @addrspacecast(i32 addrspace(5)* %ptr.private, i32 addrspace(3)* %ptr.local) {
+; GFX8V3-LABEL: addrspacecast:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V3-NEXT:    s_load_dword s2, s[4:5], 0x44
+; GFX8V3-NEXT:    s_load_dword s3, s[4:5], 0x40
+; GFX8V3-NEXT:    v_mov_b32_e32 v4, 1
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s2
+; GFX8V3-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX8V3-NEXT:    v_cndmask_b32_e32 v1, 0, v0, vcc
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX8V3-NEXT:    v_cndmask_b32_e32 v0, 0, v0, vcc
+; GFX8V3-NEXT:    v_mov_b32_e32 v2, s3
+; GFX8V3-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX8V3-NEXT:    v_cndmask_b32_e32 v3, 0, v2, vcc
+; GFX8V3-NEXT:    v_mov_b32_e32 v2, s1
+; GFX8V3-NEXT:    v_cndmask_b32_e32 v2, 0, v2, vcc
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v4
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, 2
+; GFX8V3-NEXT:    flat_store_dword v[2:3], v0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: addrspacecast:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V4-NEXT:    s_load_dword s2, s[4:5], 0x44
+; GFX8V4-NEXT:    s_load_dword s3, s[4:5], 0x40
+; GFX8V4-NEXT:    v_mov_b32_e32 v4, 1
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s2
+; GFX8V4-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX8V4-NEXT:    v_cndmask_b32_e32 v1, 0, v0, vcc
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX8V4-NEXT:    v_cndmask_b32_e32 v0, 0, v0, vcc
+; GFX8V4-NEXT:    v_mov_b32_e32 v2, s3
+; GFX8V4-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX8V4-NEXT:    v_cndmask_b32_e32 v3, 0, v2, vcc
+; GFX8V4-NEXT:    v_mov_b32_e32 v2, s1
+; GFX8V4-NEXT:    v_cndmask_b32_e32 v2, 0, v2, vcc
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v4
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, 2
+; GFX8V4-NEXT:    flat_store_dword v[2:3], v0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: addrspacecast:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX8V5-NEXT:    s_load_dword s2, s[4:5], 0xc8
+; GFX8V5-NEXT:    s_load_dword s3, s[4:5], 0xcc
+; GFX8V5-NEXT:    v_mov_b32_e32 v4, 1
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s2
+; GFX8V5-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX8V5-NEXT:    v_cndmask_b32_e32 v1, 0, v0, vcc
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX8V5-NEXT:    v_cndmask_b32_e32 v0, 0, v0, vcc
+; GFX8V5-NEXT:    v_mov_b32_e32 v2, s3
+; GFX8V5-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX8V5-NEXT:    v_cndmask_b32_e32 v3, 0, v2, vcc
+; GFX8V5-NEXT:    v_mov_b32_e32 v2, s1
+; GFX8V5-NEXT:    v_cndmask_b32_e32 v2, 0, v2, vcc
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v4
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, 2
+; GFX8V5-NEXT:    flat_store_dword v[2:3], v0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: addrspacecast:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V3-NEXT:    s_getreg_b32 s2, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s2, s2, 16
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s2
+; GFX9V3-NEXT:    v_mov_b32_e32 v4, 1
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX9V3-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX9V3-NEXT:    v_cndmask_b32_e32 v1, 0, v0, vcc
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V3-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V3-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX9V3-NEXT:    v_cndmask_b32_e32 v0, 0, v0, vcc
+; GFX9V3-NEXT:    v_mov_b32_e32 v2, s0
+; GFX9V3-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX9V3-NEXT:    v_cndmask_b32_e32 v3, 0, v2, vcc
+; GFX9V3-NEXT:    v_mov_b32_e32 v2, s1
+; GFX9V3-NEXT:    v_cndmask_b32_e32 v2, 0, v2, vcc
+; GFX9V3-NEXT:    flat_store_dword v[0:1], v4
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, 2
+; GFX9V3-NEXT:    flat_store_dword v[2:3], v0
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: addrspacecast:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V4-NEXT:    s_getreg_b32 s2, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s2, s2, 16
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s2
+; GFX9V4-NEXT:    v_mov_b32_e32 v4, 1
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX9V4-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX9V4-NEXT:    v_cndmask_b32_e32 v1, 0, v0, vcc
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V4-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V4-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX9V4-NEXT:    v_cndmask_b32_e32 v0, 0, v0, vcc
+; GFX9V4-NEXT:    v_mov_b32_e32 v2, s0
+; GFX9V4-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX9V4-NEXT:    v_cndmask_b32_e32 v3, 0, v2, vcc
+; GFX9V4-NEXT:    v_mov_b32_e32 v2, s1
+; GFX9V4-NEXT:    v_cndmask_b32_e32 v2, 0, v2, vcc
+; GFX9V4-NEXT:    flat_store_dword v[0:1], v4
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, 2
+; GFX9V4-NEXT:    flat_store_dword v[2:3], v0
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: addrspacecast:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0x0
+; GFX9V5-NEXT:    s_getreg_b32 s2, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s2, s2, 16
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s2
+; GFX9V5-NEXT:    v_mov_b32_e32 v4, 1
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    s_cmp_lg_u32 s0, -1
+; GFX9V5-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX9V5-NEXT:    v_cndmask_b32_e32 v1, 0, v0, vcc
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX9V5-NEXT:    s_getreg_b32 s0, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s0, s0, 16
+; GFX9V5-NEXT:    s_cmp_lg_u32 s1, -1
+; GFX9V5-NEXT:    v_cndmask_b32_e32 v0, 0, v0, vcc
+; GFX9V5-NEXT:    v_mov_b32_e32 v2, s0
+; GFX9V5-NEXT:    s_cselect_b64 vcc, -1, 0
+; GFX9V5-NEXT:    v_cndmask_b32_e32 v3, 0, v2, vcc
+; GFX9V5-NEXT:    v_mov_b32_e32 v2, s1
+; GFX9V5-NEXT:    v_cndmask_b32_e32 v2, 0, v2, vcc
+; GFX9V5-NEXT:    flat_store_dword v[0:1], v4
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, 2
+; GFX9V5-NEXT:    flat_store_dword v[2:3], v0
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %flat.private = addrspacecast i32 addrspace(5)* %ptr.private to i32*
+  %flat.local = addrspacecast i32 addrspace(3)* %ptr.local to i32*
+  store volatile i32 1, i32* %flat.private
+  store volatile i32 2, i32* %flat.local
+  ret void
+}
+
+define amdgpu_kernel void @llvm_amdgcn_is_shared(i8* %ptr) {
+; GFX8V3-LABEL: llvm_amdgcn_is_shared:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_load_dword s0, s[4:5], 0x40
+; GFX8V3-NEXT:    s_load_dword s1, s[6:7], 0x4
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V3-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX8V3-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: llvm_amdgcn_is_shared:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_load_dword s0, s[4:5], 0x40
+; GFX8V4-NEXT:    s_load_dword s1, s[6:7], 0x4
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V4-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX8V4-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: llvm_amdgcn_is_shared:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dword s0, s[4:5], 0xcc
+; GFX8V5-NEXT:    s_load_dword s1, s[4:5], 0x4
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V5-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX8V5-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: llvm_amdgcn_is_shared:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_load_dword s0, s[4:5], 0x4
+; GFX9V3-NEXT:    s_getreg_b32 s1, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s1, s1, 16
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    s_cmp_eq_u32 s0, s1
+; GFX9V3-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX9V3-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX9V3-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: llvm_amdgcn_is_shared:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_load_dword s0, s[4:5], 0x4
+; GFX9V4-NEXT:    s_getreg_b32 s1, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s1, s1, 16
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    s_cmp_eq_u32 s0, s1
+; GFX9V4-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX9V4-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX9V4-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: llvm_amdgcn_is_shared:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_load_dword s0, s[4:5], 0x4
+; GFX9V5-NEXT:    s_getreg_b32 s1, hwreg(HW_REG_SH_MEM_BASES, 16, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s1, s1, 16
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    s_cmp_eq_u32 s0, s1
+; GFX9V5-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX9V5-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX9V5-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %is.shared = call i1 @llvm.amdgcn.is.shared(i8* %ptr)
+  %zext = zext i1 %is.shared to i32
+  store volatile i32 %zext, i32 addrspace(1)* undef
+  ret void
+}
+
+define amdgpu_kernel void @llvm_amdgcn_is_private(i8* %ptr) {
+; GFX8V3-LABEL: llvm_amdgcn_is_private:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_load_dword s0, s[4:5], 0x44
+; GFX8V3-NEXT:    s_load_dword s1, s[6:7], 0x4
+; GFX8V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V3-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V3-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX8V3-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX8V3-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: llvm_amdgcn_is_private:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_load_dword s0, s[4:5], 0x44
+; GFX8V4-NEXT:    s_load_dword s1, s[6:7], 0x4
+; GFX8V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V4-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V4-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX8V4-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX8V4-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: llvm_amdgcn_is_private:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dword s0, s[4:5], 0xc8
+; GFX8V5-NEXT:    s_load_dword s1, s[4:5], 0x4
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_cmp_eq_u32 s1, s0
+; GFX8V5-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX8V5-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX8V5-NEXT:    flat_store_dword v[0:1], v0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: llvm_amdgcn_is_private:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_load_dword s0, s[4:5], 0x4
+; GFX9V3-NEXT:    s_getreg_b32 s1, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V3-NEXT:    s_lshl_b32 s1, s1, 16
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    s_cmp_eq_u32 s0, s1
+; GFX9V3-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX9V3-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX9V3-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: llvm_amdgcn_is_private:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_load_dword s0, s[4:5], 0x4
+; GFX9V4-NEXT:    s_getreg_b32 s1, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V4-NEXT:    s_lshl_b32 s1, s1, 16
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    s_cmp_eq_u32 s0, s1
+; GFX9V4-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX9V4-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX9V4-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: llvm_amdgcn_is_private:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_load_dword s0, s[4:5], 0x4
+; GFX9V5-NEXT:    s_getreg_b32 s1, hwreg(HW_REG_SH_MEM_BASES, 0, 16)
+; GFX9V5-NEXT:    s_lshl_b32 s1, s1, 16
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    s_cmp_eq_u32 s0, s1
+; GFX9V5-NEXT:    s_cselect_b64 s[0:1], -1, 0
+; GFX9V5-NEXT:    v_cndmask_b32_e64 v0, 0, 1, s[0:1]
+; GFX9V5-NEXT:    global_store_dword v[0:1], v0, off
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %is.private = call i1 @llvm.amdgcn.is.private(i8* %ptr)
+  %zext = zext i1 %is.private to i32
+  store volatile i32 %zext, i32 addrspace(1)* undef
+  ret void
+}
+
+define amdgpu_kernel void @llvm_trap() {
+; GFX8V3-LABEL: llvm_trap:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_mov_b64 s[0:1], s[4:5]
+; GFX8V3-NEXT:    s_trap 2
+;
+; GFX8V4-LABEL: llvm_trap:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_mov_b64 s[0:1], s[4:5]
+; GFX8V4-NEXT:    s_trap 2
+;
+; GFX8V5-LABEL: llvm_trap:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[4:5], 0xc8
+; GFX8V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX8V5-NEXT:    s_trap 2
+;
+; GFX9V3-LABEL: llvm_trap:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_mov_b64 s[0:1], s[4:5]
+; GFX9V3-NEXT:    s_trap 2
+;
+; GFX9V4-LABEL: llvm_trap:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_trap 2
+;
+; GFX9V5-LABEL: llvm_trap:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_trap 2
+  call void @llvm.trap()
+  unreachable
+}
+
+define amdgpu_kernel void @llvm_debugtrap() {
+; GFX8V3-LABEL: llvm_debugtrap:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    s_trap 3
+;
+; GFX8V4-LABEL: llvm_debugtrap:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    s_trap 3
+;
+; GFX8V5-LABEL: llvm_debugtrap:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_trap 3
+;
+; GFX9V3-LABEL: llvm_debugtrap:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    s_trap 3
+;
+; GFX9V4-LABEL: llvm_debugtrap:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    s_trap 3
+;
+; GFX9V5-LABEL: llvm_debugtrap:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    s_trap 3
+  call void @llvm.debugtrap()
+  unreachable
+}
+
+define amdgpu_kernel void @llvm_amdgcn_queue_ptr(i64 addrspace(1)* %ptr)  {
+; GFX8V3-LABEL: llvm_amdgcn_queue_ptr:
+; GFX8V3:       ; %bb.0:
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s6
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s7
+; GFX8V3-NEXT:    s_add_u32 s0, s8, 8
+; GFX8V3-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V3-NEXT:    s_addc_u32 s1, s9, 0
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V3-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s4
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s5
+; GFX8V3-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V3-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX8V3-NEXT:    v_mov_b32_e32 v2, s10
+; GFX8V3-NEXT:    v_mov_b32_e32 v3, s11
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
+; GFX8V3-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V3-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V3-NEXT:    flat_store_dwordx2 v[0:1], v[2:3]
+; GFX8V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V3-NEXT:    s_endpgm
+;
+; GFX8V4-LABEL: llvm_amdgcn_queue_ptr:
+; GFX8V4:       ; %bb.0:
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s6
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s7
+; GFX8V4-NEXT:    s_add_u32 s0, s8, 8
+; GFX8V4-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V4-NEXT:    s_addc_u32 s1, s9, 0
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V4-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s4
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s5
+; GFX8V4-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V4-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX8V4-NEXT:    v_mov_b32_e32 v2, s10
+; GFX8V4-NEXT:    v_mov_b32_e32 v3, s11
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
+; GFX8V4-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V4-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V4-NEXT:    flat_store_dwordx2 v[0:1], v[2:3]
+; GFX8V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V4-NEXT:    s_endpgm
+;
+; GFX8V5-LABEL: llvm_amdgcn_queue_ptr:
+; GFX8V5:       ; %bb.0:
+; GFX8V5-NEXT:    s_add_u32 s0, s6, 8
+; GFX8V5-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V5-NEXT:    s_addc_u32 s1, s7, 0
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V5-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s4
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s5
+; GFX8V5-NEXT:    flat_load_ubyte v0, v[0:1] glc
+; GFX8V5-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX8V5-NEXT:    v_mov_b32_e32 v2, s8
+; GFX8V5-NEXT:    v_mov_b32_e32 v3, s9
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0) lgkmcnt(0)
+; GFX8V5-NEXT:    v_mov_b32_e32 v0, s0
+; GFX8V5-NEXT:    v_mov_b32_e32 v1, s1
+; GFX8V5-NEXT:    flat_store_dwordx2 v[0:1], v[2:3]
+; GFX8V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX8V5-NEXT:    s_endpgm
+;
+; GFX9V3-LABEL: llvm_amdgcn_queue_ptr:
+; GFX9V3:       ; %bb.0:
+; GFX9V3-NEXT:    v_mov_b32_e32 v2, 0
+; GFX9V3-NEXT:    global_load_ubyte v0, v2, s[6:7] glc
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    global_load_ubyte v0, v2, s[8:9] offset:8 glc
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    global_load_ubyte v0, v2, s[4:5] glc
+; GFX9V3-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    v_mov_b32_e32 v0, s10
+; GFX9V3-NEXT:    v_mov_b32_e32 v1, s11
+; GFX9V3-NEXT:    ; kill: killed $sgpr6_sgpr7
+; GFX9V3-NEXT:    ; kill: killed $sgpr4_sgpr5
+; GFX9V3-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V3-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX9V3-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V3-NEXT:    s_endpgm
+;
+; GFX9V4-LABEL: llvm_amdgcn_queue_ptr:
+; GFX9V4:       ; %bb.0:
+; GFX9V4-NEXT:    v_mov_b32_e32 v2, 0
+; GFX9V4-NEXT:    global_load_ubyte v0, v2, s[6:7] glc
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    global_load_ubyte v0, v2, s[8:9] offset:8 glc
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    global_load_ubyte v0, v2, s[4:5] glc
+; GFX9V4-NEXT:    s_load_dwordx2 s[0:1], s[8:9], 0x0
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    v_mov_b32_e32 v0, s10
+; GFX9V4-NEXT:    v_mov_b32_e32 v1, s11
+; GFX9V4-NEXT:    ; kill: killed $sgpr6_sgpr7
+; GFX9V4-NEXT:    ; kill: killed $sgpr4_sgpr5
+; GFX9V4-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V4-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX9V4-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V4-NEXT:    s_endpgm
+;
+; GFX9V5-LABEL: llvm_amdgcn_queue_ptr:
+; GFX9V5:       ; %bb.0:
+; GFX9V5-NEXT:    v_mov_b32_e32 v2, 0
+; GFX9V5-NEXT:    global_load_ubyte v0, v2, s[0:1] glc
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    global_load_ubyte v0, v2, s[6:7] offset:8 glc
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    global_load_ubyte v0, v2, s[4:5] glc
+; GFX9V5-NEXT:    ; kill: killed $sgpr0_sgpr1
+; GFX9V5-NEXT:    s_load_dwordx2 s[0:1], s[6:7], 0x0
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    v_mov_b32_e32 v0, s8
+; GFX9V5-NEXT:    v_mov_b32_e32 v1, s9
+; GFX9V5-NEXT:    ; kill: killed $sgpr4_sgpr5
+; GFX9V5-NEXT:    s_waitcnt lgkmcnt(0)
+; GFX9V5-NEXT:    global_store_dwordx2 v2, v[0:1], s[0:1]
+; GFX9V5-NEXT:    s_waitcnt vmcnt(0)
+; GFX9V5-NEXT:    s_endpgm
+  %queue.ptr = call i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
+  %implicitarg.ptr = call i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+  %dispatch.ptr = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+  %dispatch.id = call i64 @llvm.amdgcn.dispatch.id()
+  %queue.load = load volatile i8, i8 addrspace(4)* %queue.ptr
+  %implicitarg.load = load volatile i8, i8 addrspace(4)* %implicitarg.ptr
+  %dispatch.load = load volatile i8, i8 addrspace(4)* %dispatch.ptr
+  store volatile i64 %dispatch.id, i64 addrspace(1)* %ptr
+  ret void
+}
+
+declare noalias i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
+declare noalias i8 addrspace(4)* @llvm.amdgcn.implicitarg.ptr()
+declare i64 @llvm.amdgcn.dispatch.id()
+declare noalias i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+declare i1 @llvm.amdgcn.is.shared(i8*)
+declare i1 @llvm.amdgcn.is.private(i8*)
+declare void @llvm.trap()
+declare void @llvm.debugtrap()