[SPIR-V] Rename SPIRVGlobalRegistry to SPIRVGlobalObjectRegistry

[mpaszkowski]

Summary

This patch renames SPIRVGlobalRegistry to SPIRVGlobalObjectRegistry -- in preparation for adding SPIRVGlobalInstrRegistry (will create a separate revision and link below). This patch also marks GlobalObjectRegistry find/add method arguments as const.

[iliya-diyachkov]

llvm/lib/Target/SPIRV/Registries/SPIRVGlobalObjectRegistry.cpp

1–2

The same suggestion as for SPIRVGlobalObjectRegistry.h.

8

The same issue.

16

The same issue.

llvm/lib/Target/SPIRV/Registries/SPIRVGlobalObjectRegistry.h

1–2

This line break doesn't look very nice. There are such examples in llvm, but much more often this comment is written in one line. We could use something like this

//==- SPIRVGlobalObjectRegistry.h - SPIR-V Global Object Registry -*- C++ -*-=//

On other hand, you are going to introduce SPIRVGlobalInstrRegistry files in the next step, reducing Instruction to Instr. So it may be a suitable solution (in the same way) to use Obj instead of Object in the file/class name. It's still meaningful, and the names won't be so long.

8

The preferred string length is 80 characters.

15

The same issue

llvm/lib/Target/SPIRV/CMakeLists.txt

 add_llvm_component_group(SPIRV)
 
 set(LLVM_TARGET_DEFINITIONS SPIRV.td)
 
 tablegen(LLVM SPIRVGenAsmWriter.inc -gen-asm-writer)
 tablegen(LLVM SPIRVGenGlobalISel.inc -gen-global-isel)
 tablegen(LLVM SPIRVGenInstrInfo.inc -gen-instr-info)
 tablegen(LLVM SPIRVGenMCCodeEmitter.inc -gen-emitter)
 tablegen(LLVM SPIRVGenRegisterBank.inc -gen-register-bank)
 tablegen(LLVM SPIRVGenRegisterInfo.inc -gen-register-info)
 tablegen(LLVM SPIRVGenSubtargetInfo.inc -gen-subtarget)
 tablegen(LLVM SPIRVGenTables.inc -gen-searchable-tables)
 
 add_public_tablegen_target(SPIRVCommonTableGen)
 
 add_llvm_target(SPIRVCodeGen
+  Registries/SPIRVGlobalObjectRegistry.cpp  
   SPIRVAsmPrinter.cpp
   SPIRVBuiltins.cpp
   SPIRVCallLowering.cpp
   SPIRVDuplicatesTracker.cpp
   SPIRVEmitIntrinsics.cpp
-  SPIRVGlobalRegistry.cpp
   SPIRVInstrInfo.cpp
   SPIRVInstructionSelector.cpp
   SPIRVISelLowering.cpp
   SPIRVLegalizerInfo.cpp
   SPIRVMCInstLower.cpp
   SPIRVModuleAnalysis.cpp
   SPIRVPreLegalizer.cpp
   SPIRVPrepareFunctions.cpp

llvm/lib/Target/SPIRV/Registries/SPIRVGlobalObjectRegistry.h

-//===-- SPIRVGlobalRegistry.h - SPIR-V Global Registry ----------*- C++ -*-===//
+//===-- SPIRVGlobalObjectRegistry.h - SPIR-V Global Object Registry -*- C++
+//-*-===//
 //
 // 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
 //
-//===----------------------------------------------------------------------===//
+//===--------------------------------------------------------------------------===//
 //
-// SPIRVGlobalRegistry is used to maintain rich type information required for
-// SPIR-V even after lowering from LLVM IR to GMIR. It can convert an llvm::Type
-// into an OpTypeXXX instruction, and map it to a virtual register. Also it
-// builds and supports consistency of constants and global variables.
+// SPIRVGlobalObjectRegistry is used to maintain rich type information required
+// for SPIR-V even after lowering from LLVM IR to GMIR. It can convert an
+// llvm::Type into an OpTypeXXX instruction, and map it to a virtual register.
+// Also it builds and supports consistency of constants and global variables.
 //
-//===----------------------------------------------------------------------===//
+//===--------------------------------------------------------------------------===//
 
-#ifndef LLVM_LIB_TARGET_SPIRV_SPIRVTYPEMANAGER_H
-#define LLVM_LIB_TARGET_SPIRV_SPIRVTYPEMANAGER_H
+#ifndef LLVM_LIB_TARGET_SPIRV_SPIRVGLOBALOBJECTREGISTRY_H
+#define LLVM_LIB_TARGET_SPIRV_SPIRVGLOBALOBJECTREGISTRY_H
 
 #include "MCTargetDesc/SPIRVBaseInfo.h"
 #include "SPIRVDuplicatesTracker.h"
 #include "SPIRVInstrInfo.h"
 #include "llvm/CodeGen/GlobalISel/MachineIRBuilder.h"
 
 namespace llvm {
 using SPIRVType = const MachineInstr;
 
-class SPIRVGlobalRegistry {
+class SPIRVGlobalObjectRegistry {
   // Registers holding values which have types associated with them.
   // Initialized upon VReg definition in IRTranslator.
   // Do not confuse this with DuplicatesTracker as DT maps Type* to <MF, Reg>
   // where Reg = OpType...
   // while VRegToTypeMap tracks SPIR-V type assigned to other regs (i.e. not
   // type-declaring ones).
   DenseMap<const MachineFunction *, DenseMap<Register, SPIRVType *>>
       VRegToTypeMap;
SPIRVType *findSPIRVType(const Type *Ty, MachineIRBuilder &MIRBuilder,
                                SPIRV::AccessQualifier::ReadWrite,
                            bool EmitIR = true);
   SPIRVType *
   restOfCreateSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder,
                         SPIRV::AccessQualifier::AccessQualifier AccessQual,
                         bool EmitIR);
 
 public:
-  SPIRVGlobalRegistry(unsigned PointerSize);
+  SPIRVGlobalObjectRegistry(unsigned PointerSize);
 
   MachineFunction *CurMF;
 
-  void add(const Constant *C, MachineFunction *MF, Register R) {
+  void add(const Constant *C, const MachineFunction *MF, const Register R) {
     DT.add(C, MF, R);
   }
 
-  void add(const GlobalVariable *GV, MachineFunction *MF, Register R) {
+  void add(const GlobalVariable *GV, const MachineFunction *MF,
+           const Register R) {
     DT.add(GV, MF, R);
   }
 
-  void add(const Function *F, MachineFunction *MF, Register R) {
+  void add(const Function *F, const MachineFunction *MF, const Register R) {
     DT.add(F, MF, R);
   }
 
-  void add(const Argument *Arg, MachineFunction *MF, Register R) {
+  void add(const Argument *Arg, const MachineFunction *MF, const Register R) {
     DT.add(Arg, MF, R);
   }
 
-  Register find(const Constant *C, MachineFunction *MF) {
+  Register find(const Constant *C, const MachineFunction *MF) {
     return DT.find(C, MF);
   }
 
-  Register find(const GlobalVariable *GV, MachineFunction *MF) {
+  Register find(const GlobalVariable *GV, const MachineFunction *MF) {
     return DT.find(GV, MF);
   }
 
-  Register find(const Function *F, MachineFunction *MF) {
+  Register find(const Function *F, const MachineFunction *MF) {
     return DT.find(F, MF);
   }
 
   void buildDepsGraph(std::vector<SPIRV::DTSortableEntry *> &Graph,
                       MachineModuleInfo *MMI = nullptr) {
     DT.buildDepsGraph(Graph, MMI);
   }
 
SPIRVType *getOrCreateOpTypeFunctionWithArgs(
       const Type *Ty, SPIRVType *RetType,
       const SmallVectorImpl<SPIRVType *> &ArgTypes,
       MachineIRBuilder &MIRBuilder);
   SPIRVType *getOrCreateOpTypeByOpcode(const Type *Ty,
                                        MachineIRBuilder &MIRBuilder,
                                        unsigned Opcode);
 };
 } // end namespace llvm
-#endif // LLLVM_LIB_TARGET_SPIRV_SPIRVTYPEMANAGER_H
+#endif // LLVM_LIB_TARGET_SPIRV_SPIRVGLOBALOBJECTREGISTRY_H

llvm/lib/Target/SPIRV/Registries/SPIRVGlobalObjectRegistry.cpp

-//===-- SPIRVGlobalRegistry.cpp - SPIR-V Global Registry --------*- C++ -*-===//
+//===-- SPIRVGlobalObjectRegistry.cpp - SPIR-V Global Object Registry --*- C++
+//-*-===//
 //
 // 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
 //
-//===----------------------------------------------------------------------===//
+//===-----------------------------------------------------------------------------===//
 //
-// This file contains the implementation of the SPIRVGlobalRegistry class,
+// This file contains the implementation of the SPIRVGlobalObjectRegistry class,
 // which is used to maintain rich type information required for SPIR-V even
-// after lowering from LLVM IR to GMIR. It can convert an llvm::Type into
-// an OpTypeXXX instruction, and map it to a virtual register. Also it builds
-// and supports consistency of constants and global variables.
+// after lowering from LLVM IR to GMIR. It can convert an llvm::Type into an
+// OpTypeXXX instruction, and map it to a virtual register. Also it builds and
+// supports consistency of constants and global variables.
 //
-//===----------------------------------------------------------------------===//
+//===-----------------------------------------------------------------------------===//
 
-#include "SPIRVGlobalRegistry.h"
+#include "SPIRVGlobalObjectRegistry.h"
 #include "SPIRV.h"
 #include "SPIRVBuiltins.h"
 #include "SPIRVSubtarget.h"
 #include "SPIRVTargetMachine.h"
 #include "SPIRVUtils.h"
 
 using namespace llvm;
-SPIRVGlobalRegistry::SPIRVGlobalRegistry(unsigned PointerSize)
+
+SPIRVGlobalObjectRegistry::SPIRVGlobalObjectRegistry(unsigned PointerSize)
     : PointerSize(PointerSize) {}
 
-SPIRVType *SPIRVGlobalRegistry::assignIntTypeToVReg(unsigned BitWidth,
-                                                    Register VReg,
-                                                    MachineInstr &I,
-                                                    const SPIRVInstrInfo &TII) {
+SPIRVType *
+SPIRVGlobalObjectRegistry::assignIntTypeToVReg(unsigned BitWidth, Register VReg,
   SPIRVType *SpirvType = getOrCreateSPIRVIntegerType(BitWidth, I, TII);
   assignSPIRVTypeToVReg(SpirvType, VReg, *CurMF);
   return SpirvType;
 }
 
-SPIRVType *SPIRVGlobalRegistry::assignVectTypeToVReg(
+SPIRVType *SPIRVGlobalObjectRegistry::assignVectTypeToVReg(
     SPIRVType *BaseType, unsigned NumElements, Register VReg, MachineInstr &I,
     const SPIRVInstrInfo &TII) {
   SPIRVType *SpirvType =
       getOrCreateSPIRVVectorType(BaseType, NumElements, I, TII);
   assignSPIRVTypeToVReg(SpirvType, VReg, *CurMF);
   return SpirvType;
 }
 
-SPIRVType *SPIRVGlobalRegistry::assignTypeToVReg(
+SPIRVType *SPIRVGlobalObjectRegistry::assignTypeToVReg(
     const Type *Type, Register VReg, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccessQual, bool EmitIR) {
 
   SPIRVType *SpirvType =
       getOrCreateSPIRVType(Type, MIRBuilder, AccessQual, EmitIR);
   assignSPIRVTypeToVReg(SpirvType, VReg, MIRBuilder.getMF());
   return SpirvType;
 }
 
-void SPIRVGlobalRegistry::assignSPIRVTypeToVReg(SPIRVType *SpirvType,
+void SPIRVGlobalObjectRegistry::assignSPIRVTypeToVReg(SPIRVType *SpirvType,
                                                       Register VReg,
                                                       MachineFunction &MF) {
   VRegToTypeMap[&MF][VReg] = SpirvType;
 }
 
 static Register createTypeVReg(MachineIRBuilder &MIRBuilder) {
   auto &MRI = MIRBuilder.getMF().getRegInfo();
   auto Res = MRI.createGenericVirtualRegister(LLT::scalar(32));
   MRI.setRegClass(Res, &SPIRV::TYPERegClass);
   return Res;
 }
 
 static Register createTypeVReg(MachineRegisterInfo &MRI) {
   auto Res = MRI.createGenericVirtualRegister(LLT::scalar(32));
   MRI.setRegClass(Res, &SPIRV::TYPERegClass);
   return Res;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeBool(MachineIRBuilder &MIRBuilder) {
+SPIRVType *
+SPIRVGlobalObjectRegistry::getOpTypeBool(MachineIRBuilder &MIRBuilder) {
   return MIRBuilder.buildInstr(SPIRV::OpTypeBool)
       .addDef(createTypeVReg(MIRBuilder));
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeInt(uint32_t Width,
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypeInt(uint32_t Width,
                                                    MachineIRBuilder &MIRBuilder,
                                                    bool IsSigned) {
   assert(Width <= 64 && "Unsupported integer width!");
   if (Width <= 8)
     Width = 8;
   else if (Width <= 16)
     Width = 16;
   else if (Width <= 32)
     Width = 32;
   else if (Width <= 64)
     Width = 64;
 
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeInt)
                  .addDef(createTypeVReg(MIRBuilder))
                  .addImm(Width)
                  .addImm(IsSigned ? 1 : 0);
   return MIB;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeFloat(uint32_t Width,
-                                               MachineIRBuilder &MIRBuilder) {
+SPIRVType *
+SPIRVGlobalObjectRegistry::getOpTypeFloat(uint32_t Width,
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeFloat)
                  .addDef(createTypeVReg(MIRBuilder))
                  .addImm(Width);
   return MIB;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeVoid(MachineIRBuilder &MIRBuilder) {
+SPIRVType *
+SPIRVGlobalObjectRegistry::getOpTypeVoid(MachineIRBuilder &MIRBuilder) {
   return MIRBuilder.buildInstr(SPIRV::OpTypeVoid)
       .addDef(createTypeVReg(MIRBuilder));
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeVector(uint32_t NumElems,
-                                                SPIRVType *ElemType,
-                                                MachineIRBuilder &MIRBuilder) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypeVector(
+    uint32_t NumElems, SPIRVType *ElemType, MachineIRBuilder &MIRBuilder) {
   auto EleOpc = ElemType->getOpcode();
   assert((EleOpc == SPIRV::OpTypeInt || EleOpc == SPIRV::OpTypeFloat ||
           EleOpc == SPIRV::OpTypeBool) &&
          "Invalid vector element type");
 
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeVector)
                  .addDef(createTypeVReg(MIRBuilder))
                  .addUse(getSPIRVTypeID(ElemType))
                  .addImm(NumElems);
   return MIB;
 }
 
 std::tuple<Register, ConstantInt *, bool>
-SPIRVGlobalRegistry::getOrCreateConstIntReg(uint64_t Val, SPIRVType *SpvType,
+SPIRVGlobalObjectRegistry::getOrCreateConstIntReg(uint64_t Val,
+                                                  SPIRVType *SpvType,
                                                   MachineIRBuilder *MIRBuilder,
                                                   MachineInstr *I,
                                                   const SPIRVInstrInfo *TII) {
   const IntegerType *LLVMIntTy;
   if (SpvType)
     LLVMIntTy = cast<IntegerType>(getTypeForSPIRVType(SpvType));
   else
     LLVMIntTy = IntegerType::getInt32Ty(CurMF->getFunction().getContext());
   bool NewInstr = false;
   // Find a constant in DT or build a new one.
   ConstantInt *CI = ConstantInt::get(const_cast<IntegerType *>(LLVMIntTy), Val);
   Register Res = DT.find(CI, CurMF);
   if (!Res.isValid()) {
     unsigned BitWidth = SpvType ? getScalarOrVectorBitWidth(SpvType) : 32;
     LLT LLTy = LLT::scalar(32);
     Res = CurMF->getRegInfo().createGenericVirtualRegister(LLTy);
     CurMF->getRegInfo().setRegClass(Res, &SPIRV::IDRegClass);
     if (MIRBuilder)
       assignTypeToVReg(LLVMIntTy, Res, *MIRBuilder);
     else
       assignIntTypeToVReg(BitWidth, Res, *I, *TII);
     DT.add(CI, CurMF, Res);
     NewInstr = true;
   }
   return std::make_tuple(Res, CI, NewInstr);
 }
 
-Register SPIRVGlobalRegistry::getOrCreateConstInt(uint64_t Val, MachineInstr &I,
-                                                  SPIRVType *SpvType,
-                                                  const SPIRVInstrInfo &TII) {
+Register
+SPIRVGlobalObjectRegistry::getOrCreateConstInt(uint64_t Val, MachineInstr &I,
   assert(SpvType);
   ConstantInt *CI;
   Register Res;
   bool New;
   std::tie(Res, CI, New) =
       getOrCreateConstIntReg(Val, SpvType, nullptr, &I, &TII);
   // If we have found Res register which is defined by the passed G_CONSTANT
   // machine instruction, a new constant instruction should be created.
MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
               .addUse(getSPIRVTypeID(SpvType));
   }
   const auto &ST = CurMF->getSubtarget();
   constrainSelectedInstRegOperands(*MIB, *ST.getInstrInfo(),
                                    *ST.getRegisterInfo(), *ST.getRegBankInfo());
   return Res;
 }
 
-Register SPIRVGlobalRegistry::buildConstantInt(uint64_t Val,
-                                               MachineIRBuilder &MIRBuilder,
-                                               SPIRVType *SpvType,
-                                               bool EmitIR) {
+Register
+SPIRVGlobalObjectRegistry::buildConstantInt(uint64_t Val,
+                                            SPIRVType *SpvType, bool EmitIR) {
   auto &MF = MIRBuilder.getMF();
   const IntegerType *LLVMIntTy;
   if (SpvType)
     LLVMIntTy = cast<IntegerType>(getTypeForSPIRVType(SpvType));
   else
     LLVMIntTy = IntegerType::getInt32Ty(MF.getFunction().getContext());
   // Find a constant in DT or build a new one.
   const auto ConstInt =
if (EmitIR) {
       constrainSelectedInstRegOperands(*MIB, *Subtarget.getInstrInfo(),
                                        *Subtarget.getRegisterInfo(),
                                        *Subtarget.getRegBankInfo());
     }
   }
   return Res;
 }
 
-Register SPIRVGlobalRegistry::buildConstantFP(APFloat Val,
-                                              MachineIRBuilder &MIRBuilder,
-                                              SPIRVType *SpvType) {
+Register SPIRVGlobalObjectRegistry::buildConstantFP(
+    APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType) {
   auto &MF = MIRBuilder.getMF();
   const Type *LLVMFPTy;
   if (SpvType) {
     LLVMFPTy = getTypeForSPIRVType(SpvType);
     assert(LLVMFPTy->isFloatingPointTy());
   } else {
     LLVMFPTy = IntegerType::getFloatTy(MF.getFunction().getContext());
   }
   // Find a constant in DT or build a new one.
   const auto ConstFP = ConstantFP::get(LLVMFPTy->getContext(), Val);
   Register Res = DT.find(ConstFP, &MF);
   if (!Res.isValid()) {
     unsigned BitWidth = SpvType ? getScalarOrVectorBitWidth(SpvType) : 32;
     Res = MF.getRegInfo().createGenericVirtualRegister(LLT::scalar(BitWidth));
     MF.getRegInfo().setRegClass(Res, &SPIRV::IDRegClass);
     assignTypeToVReg(LLVMFPTy, Res, MIRBuilder);
     DT.add(ConstFP, &MF, Res);
     MIRBuilder.buildFConstant(Res, *ConstFP);
   }
   return Res;
 }
 
-Register SPIRVGlobalRegistry::getOrCreateIntCompositeOrNull(
+Register SPIRVGlobalObjectRegistry::getOrCreateIntCompositeOrNull(
     uint64_t Val, MachineInstr &I, SPIRVType *SpvType,
     const SPIRVInstrInfo &TII, Constant *CA, unsigned BitWidth,
     unsigned ElemCnt) {
   // Find a constant vector in DT or build a new one.
   Register Res = DT.find(CA, CurMF);
   if (!Res.isValid()) {
     SPIRVType *SpvBaseType = getOrCreateSPIRVIntegerType(BitWidth, I, TII);
     // SpvScalConst should be created before SpvVecConst to avoid undefined ID
if (!Res.isValid()) {
     constrainSelectedInstRegOperands(*MIB, *Subtarget.getInstrInfo(),
                                      *Subtarget.getRegisterInfo(),
                                      *Subtarget.getRegBankInfo());
     return SpvVecConst;
   }
   return Res;
 }
 
-Register
-SPIRVGlobalRegistry::getOrCreateConsIntVector(uint64_t Val, MachineInstr &I,
-                                              SPIRVType *SpvType,
-                                              const SPIRVInstrInfo &TII) {
+Register SPIRVGlobalObjectRegistry::getOrCreateConsIntVector(
+    uint64_t Val, MachineInstr &I, SPIRVType *SpvType,
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);
   assert(LLVMTy->isVectorTy());
   const FixedVectorType *LLVMVecTy = cast<FixedVectorType>(LLVMTy);
   Type *LLVMBaseTy = LLVMVecTy->getElementType();
   const auto ConstInt = ConstantInt::get(LLVMBaseTy, Val);
   auto ConstVec =
       ConstantVector::getSplat(LLVMVecTy->getElementCount(), ConstInt);
   unsigned BW = getScalarOrVectorBitWidth(SpvType);
   return getOrCreateIntCompositeOrNull(Val, I, SpvType, TII, ConstVec, BW,
                                        SpvType->getOperand(2).getImm());
 }
 
-Register
-SPIRVGlobalRegistry::getOrCreateConsIntArray(uint64_t Val, MachineInstr &I,
-                                             SPIRVType *SpvType,
-                                             const SPIRVInstrInfo &TII) {
+Register SPIRVGlobalObjectRegistry::getOrCreateConsIntArray(
+    uint64_t Val, MachineInstr &I, SPIRVType *SpvType,
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);
   assert(LLVMTy->isArrayTy());
   const ArrayType *LLVMArrTy = cast<ArrayType>(LLVMTy);
   Type *LLVMBaseTy = LLVMArrTy->getElementType();
   const auto ConstInt = ConstantInt::get(LLVMBaseTy, Val);
   auto ConstArr =
       ConstantArray::get(const_cast<ArrayType *>(LLVMArrTy), {ConstInt});
   SPIRVType *SpvBaseTy = getSPIRVTypeForVReg(SpvType->getOperand(1).getReg());
   unsigned BW = getScalarOrVectorBitWidth(SpvBaseTy);
   return getOrCreateIntCompositeOrNull(Val, I, SpvType, TII, ConstArr, BW,
                                        LLVMArrTy->getNumElements());
 }
 
-Register SPIRVGlobalRegistry::getOrCreateIntCompositeOrNull(
+Register SPIRVGlobalObjectRegistry::getOrCreateIntCompositeOrNull(
     uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType, bool EmitIR,
     Constant *CA, unsigned BitWidth, unsigned ElemCnt) {
   Register Res = DT.find(CA, CurMF);
   if (!Res.isValid()) {
     Register SpvScalConst;
     if (Val || EmitIR) {
       SPIRVType *SpvBaseType =
           getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
if (EmitIR) {
             .addUse(getSPIRVTypeID(SpvType));
       }
     }
     return SpvVecConst;
   }
   return Res;
 }
 
-Register
-SPIRVGlobalRegistry::getOrCreateConsIntVector(uint64_t Val,
-                                              MachineIRBuilder &MIRBuilder,
-                                              SPIRVType *SpvType, bool EmitIR) {
+Register SPIRVGlobalObjectRegistry::getOrCreateConsIntVector(
+    uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType,
+    bool EmitIR) {
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);
   assert(LLVMTy->isVectorTy());
   const FixedVectorType *LLVMVecTy = cast<FixedVectorType>(LLVMTy);
   Type *LLVMBaseTy = LLVMVecTy->getElementType();
   const auto ConstInt = ConstantInt::get(LLVMBaseTy, Val);
   auto ConstVec =
       ConstantVector::getSplat(LLVMVecTy->getElementCount(), ConstInt);
   unsigned BW = getScalarOrVectorBitWidth(SpvType);
   return getOrCreateIntCompositeOrNull(Val, MIRBuilder, SpvType, EmitIR,
                                        ConstVec, BW,
                                        SpvType->getOperand(2).getImm());
 }
 
-Register
-SPIRVGlobalRegistry::getOrCreateConsIntArray(uint64_t Val,
-                                             MachineIRBuilder &MIRBuilder,
-                                             SPIRVType *SpvType, bool EmitIR) {
+Register SPIRVGlobalObjectRegistry::getOrCreateConsIntArray(
+    uint64_t Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType,
+    bool EmitIR) {
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);
   assert(LLVMTy->isArrayTy());
   const ArrayType *LLVMArrTy = cast<ArrayType>(LLVMTy);
   Type *LLVMBaseTy = LLVMArrTy->getElementType();
   const auto ConstInt = ConstantInt::get(LLVMBaseTy, Val);
   auto ConstArr =
       ConstantArray::get(const_cast<ArrayType *>(LLVMArrTy), {ConstInt});
   SPIRVType *SpvBaseTy = getSPIRVTypeForVReg(SpvType->getOperand(1).getReg());
   unsigned BW = getScalarOrVectorBitWidth(SpvBaseTy);
   return getOrCreateIntCompositeOrNull(Val, MIRBuilder, SpvType, EmitIR,
                                        ConstArr, BW,
                                        LLVMArrTy->getNumElements());
 }
 
 Register
-SPIRVGlobalRegistry::getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder,
+SPIRVGlobalObjectRegistry::getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder,
                                                    SPIRVType *SpvType) {
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);
   const PointerType *LLVMPtrTy = cast<PointerType>(LLVMTy);
   // Find a constant in DT or build a new one.
   Constant *CP = ConstantPointerNull::get(const_cast<PointerType *>(LLVMPtrTy));
   Register Res = DT.find(CP, CurMF);
   if (!Res.isValid()) {
     LLT LLTy = LLT::pointer(LLVMPtrTy->getAddressSpace(), PointerSize);
     Res = CurMF->getRegInfo().createGenericVirtualRegister(LLTy);
     CurMF->getRegInfo().setRegClass(Res, &SPIRV::IDRegClass);
     assignSPIRVTypeToVReg(SpvType, Res, *CurMF);
     MIRBuilder.buildInstr(SPIRV::OpConstantNull)
         .addDef(Res)
         .addUse(getSPIRVTypeID(SpvType));
     DT.add(CP, CurMF, Res);
   }
   return Res;
 }
 
-Register SPIRVGlobalRegistry::buildConstantSampler(
+Register SPIRVGlobalObjectRegistry::buildConstantSampler(
     Register ResReg, unsigned AddrMode, unsigned Param, unsigned FilerMode,
     MachineIRBuilder &MIRBuilder, SPIRVType *SpvType) {
   SPIRVType *SampTy;
   if (SpvType)
     SampTy = getOrCreateSPIRVType(getTypeForSPIRVType(SpvType), MIRBuilder);
   else
     SampTy = getOrCreateSPIRVTypeByName("opencl.sampler_t", MIRBuilder);
 
   auto Sampler =
       ResReg.isValid()
           ? ResReg
           : MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
   auto Res = MIRBuilder.buildInstr(SPIRV::OpConstantSampler)
                  .addDef(Sampler)
                  .addUse(getSPIRVTypeID(SampTy))
                  .addImm(AddrMode)
                  .addImm(Param)
                  .addImm(FilerMode);
   assert(Res->getOperand(0).isReg());
   return Res->getOperand(0).getReg();
 }
 
-Register SPIRVGlobalRegistry::buildGlobalVariable(
+Register SPIRVGlobalObjectRegistry::buildGlobalVariable(
     Register ResVReg, SPIRVType *BaseType, StringRef Name,
     const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage,
     const MachineInstr *Init, bool IsConst, bool HasLinkageTy,
     SPIRV::LinkageType::LinkageType LinkageType, MachineIRBuilder &MIRBuilder,
     bool IsInstSelector) {
   const GlobalVariable *GVar = nullptr;
   if (GV)
     GVar = cast<const GlobalVariable>(GV);
Register SPIRVGlobalObjectRegistry::buildGlobalVariable(
   SPIRV::BuiltIn::BuiltIn BuiltInId;
   if (getSpirvBuiltInIdByName(Name, BuiltInId))
     buildOpDecorate(Reg, MIRBuilder, SPIRV::Decoration::BuiltIn,
                     {static_cast<uint32_t>(BuiltInId)});
 
   return Reg;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeArray(uint32_t NumElems,
-                                               SPIRVType *ElemType,
-                                               MachineIRBuilder &MIRBuilder,
-                                               bool EmitIR) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypeArray(
+    uint32_t NumElems, SPIRVType *ElemType, MachineIRBuilder &MIRBuilder,
   assert((ElemType->getOpcode() != SPIRV::OpTypeVoid) &&
          "Invalid array element type");
   Register NumElementsVReg =
       buildConstantInt(NumElems, MIRBuilder, nullptr, EmitIR);
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeArray)
                  .addDef(createTypeVReg(MIRBuilder))
                  .addUse(getSPIRVTypeID(ElemType))
                  .addUse(NumElementsVReg);
   return MIB;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeOpaque(const StructType *Ty,
-                                                MachineIRBuilder &MIRBuilder) {
+SPIRVType *
+SPIRVGlobalObjectRegistry::getOpTypeOpaque(const StructType *Ty,
   assert(Ty->hasName());
   const StringRef Name = Ty->hasName() ? Ty->getName() : "";
   Register ResVReg = createTypeVReg(MIRBuilder);
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeOpaque).addDef(ResVReg);
   addStringImm(Name, MIB);
   buildOpName(ResVReg, Name, MIRBuilder);
   return MIB;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeStruct(const StructType *Ty,
-                                                MachineIRBuilder &MIRBuilder,
-                                                bool EmitIR) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypeStruct(
+    const StructType *Ty, MachineIRBuilder &MIRBuilder, bool EmitIR) {
   SmallVector<Register, 4> FieldTypes;
   for (const auto &Elem : Ty->elements()) {
     SPIRVType *ElemTy = findSPIRVType(Elem, MIRBuilder);
     assert(ElemTy && ElemTy->getOpcode() != SPIRV::OpTypeVoid &&
            "Invalid struct element type");
     FieldTypes.push_back(getSPIRVTypeID(ElemTy));
   }
   Register ResVReg = createTypeVReg(MIRBuilder);
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeStruct).addDef(ResVReg);
   for (const auto &Ty : FieldTypes)
     MIB.addUse(Ty);
   if (Ty->hasName())
     buildOpName(ResVReg, Ty->getName(), MIRBuilder);
   if (Ty->isPacked())
     buildOpDecorate(ResVReg, MIRBuilder, SPIRV::Decoration::CPacked, {});
   return MIB;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSpecialType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSpecialType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccQual) {
   // Some OpenCL and SPIRV builtins like image2d_t are passed in as
   // pointers, but should be treated as custom types like OpTypeImage.
   if (auto PType = dyn_cast<PointerType>(Ty)) {
     assert(!PType->isOpaque());
     Ty = PType->getNonOpaquePointerElementType();
   }
   assert(isSpecialOpaqueType(Ty) && "Not a special opaque builtin type");
   return SPIRV::lowerBuiltinType(Ty, AccQual, MIRBuilder, this);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypePointer(
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypePointer(
     SPIRV::StorageClass::StorageClass SC, SPIRVType *ElemType,
     MachineIRBuilder &MIRBuilder, Register Reg) {
   if (!Reg.isValid())
     Reg = createTypeVReg(MIRBuilder);
   return MIRBuilder.buildInstr(SPIRV::OpTypePointer)
       .addDef(Reg)
       .addImm(static_cast<uint32_t>(SC))
       .addUse(getSPIRVTypeID(ElemType));
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeForwardPointer(
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypeForwardPointer(
     SPIRV::StorageClass::StorageClass SC, MachineIRBuilder &MIRBuilder) {
   return MIRBuilder.buildInstr(SPIRV::OpTypeForwardPointer)
       .addUse(createTypeVReg(MIRBuilder))
       .addImm(static_cast<uint32_t>(SC));
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOpTypeFunction(
+SPIRVType *SPIRVGlobalObjectRegistry::getOpTypeFunction(
     SPIRVType *RetType, const SmallVectorImpl<SPIRVType *> &ArgTypes,
     MachineIRBuilder &MIRBuilder) {
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpTypeFunction)
                  .addDef(createTypeVReg(MIRBuilder))
                  .addUse(getSPIRVTypeID(RetType));
   for (const SPIRVType *ArgType : ArgTypes)
     MIB.addUse(getSPIRVTypeID(ArgType));
   return MIB;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeFunctionWithArgs(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypeFunctionWithArgs(
     const Type *Ty, SPIRVType *RetType,
     const SmallVectorImpl<SPIRVType *> &ArgTypes,
     MachineIRBuilder &MIRBuilder) {
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   SPIRVType *SpirvType = getOpTypeFunction(RetType, ArgTypes, MIRBuilder);
   return finishCreatingSPIRVType(Ty, SpirvType);
 }
 
-SPIRVType *SPIRVGlobalRegistry::findSPIRVType(
+SPIRVType *SPIRVGlobalObjectRegistry::findSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccQual, bool EmitIR) {
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   if (ForwardPointerTypes.find(Ty) != ForwardPointerTypes.end())
     return ForwardPointerTypes[Ty];
   return restOfCreateSPIRVType(Ty, MIRBuilder, AccQual, EmitIR);
 }
 
-Register SPIRVGlobalRegistry::getSPIRVTypeID(const SPIRVType *SpirvType) const {
+Register
+SPIRVGlobalObjectRegistry::getSPIRVTypeID(const SPIRVType *SpirvType) const {
   assert(SpirvType && "Attempting to get type id for nullptr type.");
   if (SpirvType->getOpcode() == SPIRV::OpTypeForwardPointer)
     return SpirvType->uses().begin()->getReg();
   return SpirvType->defs().begin()->getReg();
 }
 
-SPIRVType *SPIRVGlobalRegistry::createSPIRVType(
+SPIRVType *SPIRVGlobalObjectRegistry::createSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccQual, bool EmitIR) {
   if (isSpecialOpaqueType(Ty))
     return getOrCreateSpecialType(Ty, MIRBuilder, AccQual);
   auto &TypeToSPIRVTypeMap = DT.getTypes()->getAllUses();
   auto t = TypeToSPIRVTypeMap.find(Ty);
   if (t != TypeToSPIRVTypeMap.end()) {
     auto tt = t->second.find(&MIRBuilder.getMF());
if (auto PType = dyn_cast<PointerType>(Ty)) {
     if (ForwardPointerTypes.find(PType) != ForwardPointerTypes.end())
       Reg = getSPIRVTypeID(ForwardPointerTypes[PType]);
 
     return getOpTypePointer(SC, SpvElementType, MIRBuilder, Reg);
   }
   llvm_unreachable("Unable to convert LLVM type to SPIRVType");
 }
 
-SPIRVType *SPIRVGlobalRegistry::restOfCreateSPIRVType(
+SPIRVType *SPIRVGlobalObjectRegistry::restOfCreateSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccessQual, bool EmitIR) {
   if (TypesInProcessing.count(Ty) && !Ty->isPointerTy())
     return nullptr;
   TypesInProcessing.insert(Ty);
   SPIRVType *SpirvType = createSPIRVType(Ty, MIRBuilder, AccessQual, EmitIR);
   TypesInProcessing.erase(Ty);
   VRegToTypeMap[&MIRBuilder.getMF()][getSPIRVTypeID(SpirvType)] = SpirvType;
   SPIRVToLLVMType[SpirvType] = Ty;
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   // Do not add OpTypeForwardPointer to DT, a corresponding normal pointer type
   // will be added later. For special types it is already added to DT.
   if (SpirvType->getOpcode() != SPIRV::OpTypeForwardPointer && !Reg.isValid() &&
       !isSpecialOpaqueType(Ty))
     DT.add(Ty, &MIRBuilder.getMF(), getSPIRVTypeID(SpirvType));
 
   return SpirvType;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getSPIRVTypeForVReg(Register VReg) const {
+SPIRVType *SPIRVGlobalObjectRegistry::getSPIRVTypeForVReg(Register VReg) const {
   auto t = VRegToTypeMap.find(CurMF);
   if (t != VRegToTypeMap.end()) {
     auto tt = t->second.find(VReg);
     if (tt != t->second.end())
       return tt->second;
   }
   return nullptr;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVType(
     const Type *Ty, MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccessQual, bool EmitIR) {
   Register Reg = DT.find(Ty, &MIRBuilder.getMF());
   if (Reg.isValid() && !isSpecialOpaqueType(Ty))
     return getSPIRVTypeForVReg(Reg);
   TypesInProcessing.clear();
   SPIRVType *STy = restOfCreateSPIRVType(Ty, MIRBuilder, AccessQual, EmitIR);
   // Create normal pointer types for the corresponding OpTypeForwardPointers.
   for (auto &CU : ForwardPointerTypes) {
     const Type *Ty2 = CU.first;
     SPIRVType *STy2 = CU.second;
     if ((Reg = DT.find(Ty2, &MIRBuilder.getMF())).isValid())
       STy2 = getSPIRVTypeForVReg(Reg);
     else
       STy2 = restOfCreateSPIRVType(Ty2, MIRBuilder, AccessQual, EmitIR);
     if (Ty == Ty2)
       STy = STy2;
   }
   ForwardPointerTypes.clear();
   return STy;
 }
 
-bool SPIRVGlobalRegistry::isScalarOfType(Register VReg,
+bool SPIRVGlobalObjectRegistry::isScalarOfType(Register VReg,
                                                unsigned TypeOpcode) const {
   SPIRVType *Type = getSPIRVTypeForVReg(VReg);
   assert(Type && "isScalarOfType VReg has no type assigned");
   return Type->getOpcode() == TypeOpcode;
 }
 
-bool SPIRVGlobalRegistry::isScalarOrVectorOfType(Register VReg,
-                                                 unsigned TypeOpcode) const {
+bool SPIRVGlobalObjectRegistry::isScalarOrVectorOfType(
+    Register VReg, unsigned TypeOpcode) const {
   SPIRVType *Type = getSPIRVTypeForVReg(VReg);
   assert(Type && "isScalarOrVectorOfType VReg has no type assigned");
   if (Type->getOpcode() == TypeOpcode)
     return true;
   if (Type->getOpcode() == SPIRV::OpTypeVector) {
     Register ScalarTypeVReg = Type->getOperand(1).getReg();
     SPIRVType *ScalarType = getSPIRVTypeForVReg(ScalarTypeVReg);
     return ScalarType->getOpcode() == TypeOpcode;
   }
   return false;
 }
 
-unsigned
-SPIRVGlobalRegistry::getScalarOrVectorBitWidth(const SPIRVType *Type) const {
+unsigned SPIRVGlobalObjectRegistry::getScalarOrVectorBitWidth(
+    const SPIRVType *Type) const {
   assert(Type && "Invalid Type pointer");
   if (Type->getOpcode() == SPIRV::OpTypeVector) {
     auto EleTypeReg = Type->getOperand(1).getReg();
     Type = getSPIRVTypeForVReg(EleTypeReg);
   }
   if (Type->getOpcode() == SPIRV::OpTypeInt ||
       Type->getOpcode() == SPIRV::OpTypeFloat)
     return Type->getOperand(1).getImm();
   if (Type->getOpcode() == SPIRV::OpTypeBool)
     return 1;
   llvm_unreachable("Attempting to get bit width of non-integer/float type.");
 }
 
-bool SPIRVGlobalRegistry::isScalarOrVectorSigned(const SPIRVType *Type) const {
+bool SPIRVGlobalObjectRegistry::isScalarOrVectorSigned(
+    const SPIRVType *Type) const {
   assert(Type && "Invalid Type pointer");
   if (Type->getOpcode() == SPIRV::OpTypeVector) {
     auto EleTypeReg = Type->getOperand(1).getReg();
     Type = getSPIRVTypeForVReg(EleTypeReg);
   }
   if (Type->getOpcode() == SPIRV::OpTypeInt)
     return Type->getOperand(2).getImm() != 0;
   llvm_unreachable("Attempting to get sign of non-integer type.");
 }
 
 SPIRV::StorageClass::StorageClass
-SPIRVGlobalRegistry::getPointerStorageClass(Register VReg) const {
+SPIRVGlobalObjectRegistry::getPointerStorageClass(Register VReg) const {
   SPIRVType *Type = getSPIRVTypeForVReg(VReg);
   assert(Type && Type->getOpcode() == SPIRV::OpTypePointer &&
          Type->getOperand(1).isImm() && "Pointer type is expected");
   return static_cast<SPIRV::StorageClass::StorageClass>(
       Type->getOperand(1).getImm());
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeImage(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypeImage(
     MachineIRBuilder &MIRBuilder, SPIRVType *SampledType, SPIRV::Dim::Dim Dim,
     uint32_t Depth, uint32_t Arrayed, uint32_t Multisampled, uint32_t Sampled,
     SPIRV::ImageFormat::ImageFormat ImageFormat,
     SPIRV::AccessQualifier::AccessQualifier AccessQual) {
   SPIRV::ImageTypeDescriptor TD(SPIRVToLLVMType.lookup(SampledType), Dim, Depth,
                                 Arrayed, Multisampled, Sampled, ImageFormat,
                                 AccessQual);
   if (auto *Res = checkSpecialInstr(TD, MIRBuilder))
     return Res;
   Register ResVReg = createTypeVReg(MIRBuilder);
   DT.add(TD, &MIRBuilder.getMF(), ResVReg);
   return MIRBuilder.buildInstr(SPIRV::OpTypeImage)
       .addDef(ResVReg)
       .addUse(getSPIRVTypeID(SampledType))
       .addImm(Dim)
       .addImm(Depth)        // Depth (whether or not it is a Depth image).
       .addImm(Arrayed)      // Arrayed.
       .addImm(Multisampled) // Multisampled (0 = only single-sample).
       .addImm(Sampled)      // Sampled (0 = usage known at runtime).
       .addImm(ImageFormat)
       .addImm(AccessQual);
 }
 
-SPIRVType *
-SPIRVGlobalRegistry::getOrCreateOpTypeSampler(MachineIRBuilder &MIRBuilder) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypeSampler(
+    MachineIRBuilder &MIRBuilder) {
   SPIRV::SamplerTypeDescriptor TD;
   if (auto *Res = checkSpecialInstr(TD, MIRBuilder))
     return Res;
   Register ResVReg = createTypeVReg(MIRBuilder);
   DT.add(TD, &MIRBuilder.getMF(), ResVReg);
   return MIRBuilder.buildInstr(SPIRV::OpTypeSampler).addDef(ResVReg);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypePipe(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypePipe(
     MachineIRBuilder &MIRBuilder,
     SPIRV::AccessQualifier::AccessQualifier AccessQual) {
   SPIRV::PipeTypeDescriptor TD(AccessQual);
   if (auto *Res = checkSpecialInstr(TD, MIRBuilder))
     return Res;
   Register ResVReg = createTypeVReg(MIRBuilder);
   DT.add(TD, &MIRBuilder.getMF(), ResVReg);
   return MIRBuilder.buildInstr(SPIRV::OpTypePipe)
       .addDef(ResVReg)
       .addImm(AccessQual);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeDeviceEvent(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypeDeviceEvent(
     MachineIRBuilder &MIRBuilder) {
   SPIRV::DeviceEventTypeDescriptor TD;
   if (auto *Res = checkSpecialInstr(TD, MIRBuilder))
     return Res;
   Register ResVReg = createTypeVReg(MIRBuilder);
   DT.add(TD, &MIRBuilder.getMF(), ResVReg);
   return MIRBuilder.buildInstr(SPIRV::OpTypeDeviceEvent).addDef(ResVReg);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeSampledImage(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypeSampledImage(
     SPIRVType *ImageType, MachineIRBuilder &MIRBuilder) {
   SPIRV::SampledImageTypeDescriptor TD(
       SPIRVToLLVMType.lookup(MIRBuilder.getMF().getRegInfo().getVRegDef(
           ImageType->getOperand(1).getReg())),
       ImageType);
   if (auto *Res = checkSpecialInstr(TD, MIRBuilder))
     return Res;
   Register ResVReg = createTypeVReg(MIRBuilder);
   DT.add(TD, &MIRBuilder.getMF(), ResVReg);
   return MIRBuilder.buildInstr(SPIRV::OpTypeSampledImage)
       .addDef(ResVReg)
       .addUse(getSPIRVTypeID(ImageType));
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateOpTypeByOpcode(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateOpTypeByOpcode(
     const Type *Ty, MachineIRBuilder &MIRBuilder, unsigned Opcode) {
   Register ResVReg = DT.find(Ty, &MIRBuilder.getMF());
   if (ResVReg.isValid())
     return MIRBuilder.getMF().getRegInfo().getUniqueVRegDef(ResVReg);
   ResVReg = createTypeVReg(MIRBuilder);
   DT.add(Ty, &MIRBuilder.getMF(), ResVReg);
   return MIRBuilder.buildInstr(Opcode).addDef(ResVReg);
 }
 
-const MachineInstr *
-SPIRVGlobalRegistry::checkSpecialInstr(const SPIRV::SpecialTypeDescriptor &TD,
-                                       MachineIRBuilder &MIRBuilder) {
+const MachineInstr *SPIRVGlobalObjectRegistry::checkSpecialInstr(
+    const SPIRV::SpecialTypeDescriptor &TD, MachineIRBuilder &MIRBuilder) {
   Register Reg = DT.find(TD, &MIRBuilder.getMF());
   if (Reg.isValid())
     return MIRBuilder.getMF().getRegInfo().getUniqueVRegDef(Reg);
   return nullptr;
 }
 
 // TODO: maybe use tablegen to implement this.
-SPIRVType *
-SPIRVGlobalRegistry::getOrCreateSPIRVTypeByName(StringRef TypeStr,
-                                                MachineIRBuilder &MIRBuilder) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVTypeByName(
+    StringRef TypeStr, MachineIRBuilder &MIRBuilder) {
   unsigned VecElts = 0;
   auto &Ctx = MIRBuilder.getMF().getFunction().getContext();
 
   // Parse type name in either "typeN" or "type vector[N]" format, where
   // N is the number of elements of the vector.
   Type *Type;
   if (TypeStr.startswith("void")) {
     Type = Type::getVoidTy(Ctx);
SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVTypeByName(
   }
   TypeStr.getAsInteger(10, VecElts);
   auto SpirvTy = getOrCreateSPIRVType(Type, MIRBuilder);
   if (VecElts > 0)
     SpirvTy = getOrCreateSPIRVVectorType(SpirvTy, VecElts, MIRBuilder);
   return SpirvTy;
 }
 
-SPIRVType *
-SPIRVGlobalRegistry::getOrCreateSPIRVIntegerType(unsigned BitWidth,
-                                                 MachineIRBuilder &MIRBuilder) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVIntegerType(
+    unsigned BitWidth, MachineIRBuilder &MIRBuilder) {
   return getOrCreateSPIRVType(
       IntegerType::get(MIRBuilder.getMF().getFunction().getContext(), BitWidth),
       MIRBuilder);
 }
 
-SPIRVType *SPIRVGlobalRegistry::finishCreatingSPIRVType(const Type *LLVMTy,
-                                                        SPIRVType *SpirvType) {
+SPIRVType *
+SPIRVGlobalObjectRegistry::finishCreatingSPIRVType(const Type *LLVMTy,
   assert(CurMF == SpirvType->getMF());
   VRegToTypeMap[CurMF][getSPIRVTypeID(SpirvType)] = SpirvType;
   SPIRVToLLVMType[SpirvType] = LLVMTy;
   DT.add(LLVMTy, CurMF, getSPIRVTypeID(SpirvType));
   return SpirvType;
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVIntegerType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVIntegerType(
     unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII) {
   Type *LLVMTy = IntegerType::get(CurMF->getFunction().getContext(), BitWidth);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeInt))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addImm(BitWidth)
                  .addImm(0);
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
-SPIRVType *
-SPIRVGlobalRegistry::getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder) {
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVBoolType(
+    MachineIRBuilder &MIRBuilder) {
   return getOrCreateSPIRVType(
       IntegerType::get(MIRBuilder.getMF().getFunction().getContext(), 1),
       MIRBuilder);
 }
 
 SPIRVType *
-SPIRVGlobalRegistry::getOrCreateSPIRVBoolType(MachineInstr &I,
+SPIRVGlobalObjectRegistry::getOrCreateSPIRVBoolType(MachineInstr &I,
                                                     const SPIRVInstrInfo &TII) {
   Type *LLVMTy = IntegerType::get(CurMF->getFunction().getContext(), 1);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeBool))
                  .addDef(createTypeVReg(CurMF->getRegInfo()));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVVectorType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVVectorType(
     SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder) {
   return getOrCreateSPIRVType(
       FixedVectorType::get(const_cast<Type *>(getTypeForSPIRVType(BaseType)),
                            NumElements),
       MIRBuilder);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVVectorType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVVectorType(
     SPIRVType *BaseType, unsigned NumElements, MachineInstr &I,
     const SPIRVInstrInfo &TII) {
   Type *LLVMTy = FixedVectorType::get(
       const_cast<Type *>(getTypeForSPIRVType(BaseType)), NumElements);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeVector))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addUse(getSPIRVTypeID(BaseType))
                  .addImm(NumElements);
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVArrayType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVArrayType(
     SPIRVType *BaseType, unsigned NumElements, MachineInstr &I,
     const SPIRVInstrInfo &TII) {
   Type *LLVMTy = ArrayType::get(
       const_cast<Type *>(getTypeForSPIRVType(BaseType)), NumElements);
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   SPIRVType *SpirvType = getOrCreateSPIRVIntegerType(32, I, TII);
   Register Len = getOrCreateConstInt(NumElements, I, SpirvType, TII);
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypeArray))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addUse(getSPIRVTypeID(BaseType))
                  .addUse(Len);
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVPointerType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVPointerType(
     SPIRVType *BaseType, MachineIRBuilder &MIRBuilder,
     SPIRV::StorageClass::StorageClass SClass) {
   return getOrCreateSPIRVType(
       PointerType::get(const_cast<Type *>(getTypeForSPIRVType(BaseType)),
                        storageClassToAddressSpace(SClass)),
       MIRBuilder);
 }
 
-SPIRVType *SPIRVGlobalRegistry::getOrCreateSPIRVPointerType(
+SPIRVType *SPIRVGlobalObjectRegistry::getOrCreateSPIRVPointerType(
     SPIRVType *BaseType, MachineInstr &I, const SPIRVInstrInfo &TII,
     SPIRV::StorageClass::StorageClass SC) {
   Type *LLVMTy =
       PointerType::get(const_cast<Type *>(getTypeForSPIRVType(BaseType)),
                        storageClassToAddressSpace(SC));
   Register Reg = DT.find(LLVMTy, CurMF);
   if (Reg.isValid())
     return getSPIRVTypeForVReg(Reg);
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpTypePointer))
                  .addDef(createTypeVReg(CurMF->getRegInfo()))
                  .addImm(static_cast<uint32_t>(SC))
                  .addUse(getSPIRVTypeID(BaseType));
   return finishCreatingSPIRVType(LLVMTy, MIB);
 }
 
-Register SPIRVGlobalRegistry::getOrCreateUndef(MachineInstr &I,
-                                               SPIRVType *SpvType,
-                                               const SPIRVInstrInfo &TII) {
+Register
+SPIRVGlobalObjectRegistry::getOrCreateUndef(MachineInstr &I, SPIRVType *SpvType,
   assert(SpvType);
   const Type *LLVMTy = getTypeForSPIRVType(SpvType);
   assert(LLVMTy);
   // Find a constant in DT or build a new one.
   UndefValue *UV = UndefValue::get(const_cast<Type *>(LLVMTy));
   Register Res = DT.find(UV, CurMF);
   if (Res.isValid())
     return Res;

llvm/lib/Target/SPIRV/SPIRVBuiltins.h

 //===-- SPIRVBuiltins.h - SPIR-V Built-in Functions -------------*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Lowering builtin function calls and types using their demangled names.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_SPIRV_SPIRVBUILTINS_H
 #define LLVM_LIB_TARGET_SPIRV_SPIRVBUILTINS_H
 
-#include "SPIRVGlobalRegistry.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "llvm/CodeGen/GlobalISel/CallLowering.h"
 #include "llvm/CodeGen/GlobalISel/MachineIRBuilder.h"
 
 namespace llvm {
 namespace SPIRV {
 /// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
 /// and external instruction \p Set.
 ///
 /// \return the lowering success status if the called function is a recognized
 /// builtin, std::nullopt otherwise.
 ///
 /// \p DemangledCall is the skeleton of the lowered builtin function call.
 /// \p Set is the external instruction set containing the given builtin.
 /// \p OrigRet is the single original virtual return register if defined,
 /// Register(0) otherwise.
 /// \p OrigRetTy is the type of the \p OrigRet.
 /// \p Args are the arguments of the lowered builtin call.
 std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
                                  InstructionSet::InstructionSet Set,
                                  MachineIRBuilder &MIRBuilder,
                                  const Register OrigRet, const Type *OrigRetTy,
                                  const SmallVectorImpl<Register> &Args,
-                                 SPIRVGlobalRegistry *GR);
+                                 SPIRVGlobalObjectRegistry *GOR);
 /// Handles the translation of the provided special opaque/builtin type \p Type
 /// to SPIR-V type. Generates the corresponding machine instructions for the
 /// target type or gets the already existing OpType<...> register from the
-/// global registry \p GR.
+/// global registry \p GOR.
 ///
 /// \return A machine instruction representing the OpType<...> SPIR-V type.
 ///
 /// \p Type is the special opaque/builtin type to be lowered.
 SPIRVType *lowerBuiltinType(const Type *Type,
                             AccessQualifier::AccessQualifier AccessQual,
                             MachineIRBuilder &MIRBuilder,
-                            SPIRVGlobalRegistry *GR);
+                            SPIRVGlobalObjectRegistry *GOR);
 } // namespace SPIRV
 } // namespace llvm
 #endif // LLVM_LIB_TARGET_SPIRV_SPIRVBUILTINS_H

llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp

 //===----------------------------------------------------------------------===//
 
 /// Helper function building either a resulting scalar or vector bool register
 /// depending on the expected \p ResultType.
 ///
 /// \returns Tuple of the resulting register and its type.
 static std::tuple<Register, SPIRVType *>
 buildBoolRegister(MachineIRBuilder &MIRBuilder, const SPIRVType *ResultType,
-                  SPIRVGlobalRegistry *GR) {
+                  SPIRVGlobalObjectRegistry *GOR) {
   LLT Type;
-  SPIRVType *BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
+  SPIRVType *BoolType = GOR->getOrCreateSPIRVBoolType(MIRBuilder);
 
   if (ResultType->getOpcode() == SPIRV::OpTypeVector) {
     unsigned VectorElements = ResultType->getOperand(2).getImm();
     BoolType =
-        GR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);
+        GOR->getOrCreateSPIRVVectorType(BoolType, VectorElements, MIRBuilder);
     const FixedVectorType *LLVMVectorType =
-        cast<FixedVectorType>(GR->getTypeForSPIRVType(BoolType));
+        cast<FixedVectorType>(GOR->getTypeForSPIRVType(BoolType));
     Type = LLT::vector(LLVMVectorType->getElementCount(), 1);
   } else {
     Type = LLT::scalar(1);
   }
 
   Register ResultRegister =
       MIRBuilder.getMRI()->createGenericVirtualRegister(Type);
   MIRBuilder.getMRI()->setRegClass(ResultRegister, &SPIRV::IDRegClass);
-  GR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
+  GOR->assignSPIRVTypeToVReg(BoolType, ResultRegister, MIRBuilder.getMF());
   return std::make_tuple(ResultRegister, BoolType);
 }
 
 /// Helper function for building either a vector or scalar select instruction
 /// depending on the expected \p ResultType.
 static bool buildSelectInst(MachineIRBuilder &MIRBuilder,
                             Register ReturnRegister, Register SourceRegister,
                             const SPIRVType *ReturnType,
-                            SPIRVGlobalRegistry *GR) {
+                            SPIRVGlobalObjectRegistry *GOR) {
   Register TrueConst, FalseConst;
 
   if (ReturnType->getOpcode() == SPIRV::OpTypeVector) {
-    unsigned Bits = GR->getScalarOrVectorBitWidth(ReturnType);
+    unsigned Bits = GOR->getScalarOrVectorBitWidth(ReturnType);
     uint64_t AllOnes = APInt::getAllOnes(Bits).getZExtValue();
-    TrueConst = GR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);
-    FalseConst = GR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);
+    TrueConst = GOR->getOrCreateConsIntVector(AllOnes, MIRBuilder, ReturnType);
+    FalseConst = GOR->getOrCreateConsIntVector(0, MIRBuilder, ReturnType);
   } else {
-    TrueConst = GR->buildConstantInt(1, MIRBuilder, ReturnType);
-    FalseConst = GR->buildConstantInt(0, MIRBuilder, ReturnType);
+    TrueConst = GOR->buildConstantInt(1, MIRBuilder, ReturnType);
+    FalseConst = GOR->buildConstantInt(0, MIRBuilder, ReturnType);
   }
   return MIRBuilder.buildSelect(ReturnRegister, SourceRegister, TrueConst,
                                 FalseConst);
 }
 
 /// Helper function for building a load instruction loading into the
 /// \p DestinationReg.
 static Register buildLoadInst(SPIRVType *BaseType, Register PtrRegister,
                               MachineIRBuilder &MIRBuilder,
-                              SPIRVGlobalRegistry *GR, LLT LowLevelType,
+                              SPIRVGlobalObjectRegistry *GOR, LLT LowLevelType,
                               Register DestinationReg = Register(0)) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   if (!DestinationReg.isValid()) {
     DestinationReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
     MRI->setType(DestinationReg, LLT::scalar(32));
-    GR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF());
+    GOR->assignSPIRVTypeToVReg(BaseType, DestinationReg, MIRBuilder.getMF());
   }
   // TODO: consider using correct address space and alignment (p0 is canonical
   // type for selection though).
   MachinePointerInfo PtrInfo = MachinePointerInfo();
   MIRBuilder.buildLoad(DestinationReg, PtrRegister, PtrInfo, Align());
   return DestinationReg;
 }
 
 /// Helper function for building a load instruction for loading a builtin global
 /// variable of \p BuiltinValue value.
 static Register buildBuiltinVariableLoad(MachineIRBuilder &MIRBuilder,
                                          SPIRVType *VariableType,
-                                         SPIRVGlobalRegistry *GR,
+                                         SPIRVGlobalObjectRegistry *GOR,
                                          SPIRV::BuiltIn::BuiltIn BuiltinValue,
                                          LLT LLType,
                                          Register Reg = Register(0)) {
   Register NewRegister =
       MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
   MIRBuilder.getMRI()->setType(NewRegister,
-                               LLT::pointer(0, GR->getPointerSize()));
-  SPIRVType *PtrType = GR->getOrCreateSPIRVPointerType(
+                               LLT::pointer(0, GOR->getPointerSize()));
+  SPIRVType *PtrType = GOR->getOrCreateSPIRVPointerType(
       VariableType, MIRBuilder, SPIRV::StorageClass::Input);
-  GR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
+  GOR->assignSPIRVTypeToVReg(PtrType, NewRegister, MIRBuilder.getMF());
 
   // Set up the global OpVariable with the necessary builtin decorations.
-  Register Variable = GR->buildGlobalVariable(
+  Register Variable = GOR->buildGlobalVariable(
       NewRegister, PtrType, getLinkStringForBuiltIn(BuiltinValue), nullptr,
       SPIRV::StorageClass::Input, nullptr, true, true,
       SPIRV::LinkageType::Import, MIRBuilder, false);
 
   // Load the value from the global variable.
   Register LoadedRegister =
-      buildLoadInst(VariableType, Variable, MIRBuilder, GR, LLType, Reg);
+      buildLoadInst(VariableType, Variable, MIRBuilder, GOR, LLType, Reg);
   MIRBuilder.getMRI()->setType(LoadedRegister, LLType);
   return LoadedRegister;
 }
 
 /// Helper external function for inserting ASSIGN_TYPE instuction between \p Reg
 /// and its definition, set the new register as a destination of the definition,
 /// assign SPIRVType to both registers. If SpirvTy is provided, use it as
 /// SPIRVType in ASSIGN_TYPE, otherwise create it from \p Ty. Defined in
 /// SPIRVPreLegalizer.cpp.
 extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
-                                  SPIRVGlobalRegistry *GR,
+                                  SPIRVGlobalObjectRegistry *GOR,
                                   MachineIRBuilder &MIB,
                                   MachineRegisterInfo &MRI);
 
 // TODO: Move to TableGen.
 static SPIRV::MemorySemantics::MemorySemantics
 getSPIRVMemSemantics(std::memory_order MemOrder) {
   switch (MemOrder) {
   case std::memory_order::memory_order_relaxed:
case SPIRV::CLMemoryScope::memory_scope_all_svm_devices:
     return SPIRV::Scope::CrossDevice;
   case SPIRV::CLMemoryScope::memory_scope_sub_group:
     return SPIRV::Scope::Subgroup;
   }
   llvm_unreachable("Unknown CL memory scope");
 }
 
 static Register buildConstantIntReg(uint64_t Val, MachineIRBuilder &MIRBuilder,
-                                    SPIRVGlobalRegistry *GR,
+                                    SPIRVGlobalObjectRegistry *GOR,
                                     unsigned BitWidth = 32) {
-  SPIRVType *IntType = GR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
-  return GR->buildConstantInt(Val, MIRBuilder, IntType);
+  SPIRVType *IntType = GOR->getOrCreateSPIRVIntegerType(BitWidth, MIRBuilder);
+  return GOR->buildConstantInt(Val, MIRBuilder, IntType);
 }
 
 static Register buildScopeReg(Register CLScopeRegister,
                               SPIRV::Scope::Scope Scope,
                               MachineIRBuilder &MIRBuilder,
-                              SPIRVGlobalRegistry *GR,
+                              SPIRVGlobalObjectRegistry *GOR,
                               MachineRegisterInfo *MRI) {
   if (CLScopeRegister.isValid()) {
     auto CLScope =
         static_cast<SPIRV::CLMemoryScope>(getIConstVal(CLScopeRegister, MRI));
     Scope = getSPIRVScope(CLScope);
 
     if (CLScope == static_cast<unsigned>(Scope)) {
       MRI->setRegClass(CLScopeRegister, &SPIRV::IDRegClass);
       return CLScopeRegister;
     }
   }
-  return buildConstantIntReg(Scope, MIRBuilder, GR);
+  return buildConstantIntReg(Scope, MIRBuilder, GOR);
 }
 
 static Register buildMemSemanticsReg(Register SemanticsRegister,
                                      Register PtrRegister, unsigned &Semantics,
                                      MachineIRBuilder &MIRBuilder,
-                                     SPIRVGlobalRegistry *GR) {
+                                     SPIRVGlobalObjectRegistry *GOR) {
   if (SemanticsRegister.isValid()) {
     MachineRegisterInfo *MRI = MIRBuilder.getMRI();
     std::memory_order Order =
         static_cast<std::memory_order>(getIConstVal(SemanticsRegister, MRI));
-    Semantics =
-        getSPIRVMemSemantics(Order) |
-        getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
+    Semantics = getSPIRVMemSemantics(Order) |
+                getMemSemanticsForStorageClass(
+                    GOR->getPointerStorageClass(PtrRegister));
 
     if (Order == Semantics) {
       MRI->setRegClass(SemanticsRegister, &SPIRV::IDRegClass);
       return SemanticsRegister;
     }
   }
-  return buildConstantIntReg(Semantics, MIRBuilder, GR);
+  return buildConstantIntReg(Semantics, MIRBuilder, GOR);
 }
 
 /// Helper function for translating atomic init to OpStore.
 static bool buildAtomicInitInst(const SPIRV::IncomingCall *Call,
                                 MachineIRBuilder &MIRBuilder) {
   assert(Call->Arguments.size() == 2 &&
          "Need 2 arguments for atomic init translation");
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
   MIRBuilder.buildInstr(SPIRV::OpStore)
       .addUse(Call->Arguments[0])
       .addUse(Call->Arguments[1]);
   return true;
 }
 
 /// Helper function for building an atomic load instruction.
 static bool buildAtomicLoadInst(const SPIRV::IncomingCall *Call,
                                 MachineIRBuilder &MIRBuilder,
-                                SPIRVGlobalRegistry *GR) {
+                                SPIRVGlobalObjectRegistry *GOR) {
   Register PtrRegister = Call->Arguments[0];
   MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass);
   // TODO: if true insert call to __translate_ocl_memory_sccope before
   // OpAtomicLoad and the function implementation. We can use Translator's
   // output for transcoding/atomic_explicit_arguments.cl as an example.
   Register ScopeRegister;
   if (Call->Arguments.size() > 1) {
     ScopeRegister = Call->Arguments[1];
     MIRBuilder.getMRI()->setRegClass(ScopeRegister, &SPIRV::IDRegClass);
   } else
-    ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
+    ScopeRegister = buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GOR);
 
   Register MemSemanticsReg;
   if (Call->Arguments.size() > 2) {
     // TODO: Insert call to __translate_ocl_memory_order before OpAtomicLoad.
     MemSemanticsReg = Call->Arguments[2];
     MIRBuilder.getMRI()->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
   } else {
-    int Semantics =
-        SPIRV::MemorySemantics::SequentiallyConsistent |
-        getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
-    MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
+    int Semantics = SPIRV::MemorySemantics::SequentiallyConsistent |
+                    getMemSemanticsForStorageClass(
+                        GOR->getPointerStorageClass(PtrRegister));
+    MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GOR);
   }
 
   MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
       .addDef(Call->ReturnRegister)
-      .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+      .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
       .addUse(PtrRegister)
       .addUse(ScopeRegister)
       .addUse(MemSemanticsReg);
   return true;
 }
 
 /// Helper function for building an atomic store instruction.
 static bool buildAtomicStoreInst(const SPIRV::IncomingCall *Call,
                                  MachineIRBuilder &MIRBuilder,
-                                 SPIRVGlobalRegistry *GR) {
+                                 SPIRVGlobalObjectRegistry *GOR) {
   Register ScopeRegister =
-      buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GR);
+      buildConstantIntReg(SPIRV::Scope::Device, MIRBuilder, GOR);
   Register PtrRegister = Call->Arguments[0];
   MIRBuilder.getMRI()->setRegClass(PtrRegister, &SPIRV::IDRegClass);
   int Semantics =
       SPIRV::MemorySemantics::SequentiallyConsistent |
-      getMemSemanticsForStorageClass(GR->getPointerStorageClass(PtrRegister));
-  Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GR);
+      getMemSemanticsForStorageClass(GOR->getPointerStorageClass(PtrRegister));
+  Register MemSemanticsReg = buildConstantIntReg(Semantics, MIRBuilder, GOR);
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
   MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
       .addUse(PtrRegister)
       .addUse(ScopeRegister)
       .addUse(MemSemanticsReg)
       .addUse(Call->Arguments[1]);
   return true;
 }
 
 /// Helper function for building an atomic compare-exchange instruction.
 static bool buildAtomicCompareExchangeInst(const SPIRV::IncomingCall *Call,
                                            MachineIRBuilder &MIRBuilder,
-                                           SPIRVGlobalRegistry *GR) {
+                                           SPIRVGlobalObjectRegistry *GOR) {
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
   bool IsCmpxchg = Call->Builtin->Name.contains("cmpxchg");
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 
   Register ObjectPtr = Call->Arguments[0];   // Pointer (volatile A *object.)
   Register ExpectedArg = Call->Arguments[1]; // Comparator (C* expected).
   Register Desired = Call->Arguments[2];     // Value (C Desired).
   MRI->setRegClass(ObjectPtr, &SPIRV::IDRegClass);
   MRI->setRegClass(ExpectedArg, &SPIRV::IDRegClass);
   MRI->setRegClass(Desired, &SPIRV::IDRegClass);
-  SPIRVType *SpvDesiredTy = GR->getSPIRVTypeForVReg(Desired);
+  SPIRVType *SpvDesiredTy = GOR->getSPIRVTypeForVReg(Desired);
   LLT DesiredLLT = MRI->getType(Desired);
 
-  assert(GR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
+  assert(GOR->getSPIRVTypeForVReg(ObjectPtr)->getOpcode() ==
          SPIRV::OpTypePointer);
-  unsigned ExpectedType = GR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
+  unsigned ExpectedType = GOR->getSPIRVTypeForVReg(ExpectedArg)->getOpcode();
   assert(IsCmpxchg ? ExpectedType == SPIRV::OpTypeInt
                    : ExpectedType == SPIRV::OpTypePointer);
-  assert(GR->isScalarOfType(Desired, SPIRV::OpTypeInt));
+  assert(GOR->isScalarOfType(Desired, SPIRV::OpTypeInt));
 
-  SPIRVType *SpvObjectPtrTy = GR->getSPIRVTypeForVReg(ObjectPtr);
+  SPIRVType *SpvObjectPtrTy = GOR->getSPIRVTypeForVReg(ObjectPtr);
   assert(SpvObjectPtrTy->getOperand(2).isReg() && "SPIRV type is expected");
   auto StorageClass = static_cast<SPIRV::StorageClass::StorageClass>(
       SpvObjectPtrTy->getOperand(1).getImm());
   auto MemSemStorage = getMemSemanticsForStorageClass(StorageClass);
 
   Register MemSemEqualReg;
   Register MemSemUnequalReg;
   uint64_t MemSemEqual =
if (Call->Arguments.size() >= 4) {
     if (MemOrdEq == MemSemEqual)
       MemSemEqualReg = Call->Arguments[3];
     if (MemOrdNeq == MemSemEqual)
       MemSemUnequalReg = Call->Arguments[4];
     MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass);
     MRI->setRegClass(Call->Arguments[4], &SPIRV::IDRegClass);
   }
   if (!MemSemEqualReg.isValid())
-    MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GR);
+    MemSemEqualReg = buildConstantIntReg(MemSemEqual, MIRBuilder, GOR);
   if (!MemSemUnequalReg.isValid())
-    MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GR);
+    MemSemUnequalReg = buildConstantIntReg(MemSemUnequal, MIRBuilder, GOR);
 
   Register ScopeReg;
   auto Scope = IsCmpxchg ? SPIRV::Scope::Workgroup : SPIRV::Scope::Device;
   if (Call->Arguments.size() >= 6) {
     assert(Call->Arguments.size() == 6 &&
            "Extra args for explicit atomic cmpxchg");
     auto ClScope = static_cast<SPIRV::CLMemoryScope>(
         getIConstVal(Call->Arguments[5], MRI));
     Scope = getSPIRVScope(ClScope);
     if (ClScope == static_cast<unsigned>(Scope))
       ScopeReg = Call->Arguments[5];
     MRI->setRegClass(Call->Arguments[5], &SPIRV::IDRegClass);
   }
   if (!ScopeReg.isValid())
-    ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
+    ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GOR);
 
   Register Expected = IsCmpxchg
                           ? ExpectedArg
                           : buildLoadInst(SpvDesiredTy, ExpectedArg, MIRBuilder,
-                                          GR, LLT::scalar(32));
+                                          GOR, LLT::scalar(32));
   MRI->setType(Expected, DesiredLLT);
   Register Tmp = !IsCmpxchg ? MRI->createGenericVirtualRegister(DesiredLLT)
                             : Call->ReturnRegister;
   if (!MRI->getRegClassOrNull(Tmp))
     MRI->setRegClass(Tmp, &SPIRV::IDRegClass);
-  GR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
+  GOR->assignSPIRVTypeToVReg(SpvDesiredTy, Tmp, MIRBuilder.getMF());
 
-  SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
+  SPIRVType *IntTy = GOR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
   MIRBuilder.buildInstr(Opcode)
       .addDef(Tmp)
-      .addUse(GR->getSPIRVTypeID(IntTy))
+      .addUse(GOR->getSPIRVTypeID(IntTy))
       .addUse(ObjectPtr)
       .addUse(ScopeReg)
       .addUse(MemSemEqualReg)
       .addUse(MemSemUnequalReg)
       .addUse(Desired)
       .addUse(Expected);
   if (!IsCmpxchg) {
     MIRBuilder.buildInstr(SPIRV::OpStore).addUse(ExpectedArg).addUse(Tmp);
     MIRBuilder.buildICmp(CmpInst::ICMP_EQ, Call->ReturnRegister, Tmp, Expected);
   }
   return true;
 }
 
 /// Helper function for building an atomic load instruction.
 static bool buildAtomicRMWInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
                                MachineIRBuilder &MIRBuilder,
-                               SPIRVGlobalRegistry *GR) {
+                               SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   Register ScopeRegister =
       Call->Arguments.size() >= 4 ? Call->Arguments[3] : Register();
 
   assert(Call->Arguments.size() <= 4 &&
          "Too many args for explicit atomic RMW");
   ScopeRegister = buildScopeReg(ScopeRegister, SPIRV::Scope::Workgroup,
-                                MIRBuilder, GR, MRI);
+                                MIRBuilder, GOR, MRI);
 
   Register PtrRegister = Call->Arguments[0];
   unsigned Semantics = SPIRV::MemorySemantics::None;
   MRI->setRegClass(PtrRegister, &SPIRV::IDRegClass);
   Register MemSemanticsReg =
       Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
   MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
-                                         Semantics, MIRBuilder, GR);
+                                         Semantics, MIRBuilder, GOR);
   MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
   MIRBuilder.buildInstr(Opcode)
       .addDef(Call->ReturnRegister)
-      .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+      .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
       .addUse(PtrRegister)
       .addUse(ScopeRegister)
       .addUse(MemSemanticsReg)
       .addUse(Call->Arguments[1]);
   return true;
 }
 
 /// Helper function for building atomic flag instructions (e.g.
 /// OpAtomicFlagTestAndSet).
 static bool buildAtomicFlagInst(const SPIRV::IncomingCall *Call,
                                 unsigned Opcode, MachineIRBuilder &MIRBuilder,
-                                SPIRVGlobalRegistry *GR) {
+                                SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   Register PtrRegister = Call->Arguments[0];
   unsigned Semantics = SPIRV::MemorySemantics::SequentiallyConsistent;
   Register MemSemanticsReg =
       Call->Arguments.size() >= 2 ? Call->Arguments[1] : Register();
   MemSemanticsReg = buildMemSemanticsReg(MemSemanticsReg, PtrRegister,
-                                         Semantics, MIRBuilder, GR);
+                                         Semantics, MIRBuilder, GOR);
 
   assert((Opcode != SPIRV::OpAtomicFlagClear ||
           (Semantics != SPIRV::MemorySemantics::Acquire &&
            Semantics != SPIRV::MemorySemantics::AcquireRelease)) &&
          "Invalid memory order argument!");
 
   Register ScopeRegister =
       Call->Arguments.size() >= 3 ? Call->Arguments[2] : Register();
   ScopeRegister =
-      buildScopeReg(ScopeRegister, SPIRV::Scope::Device, MIRBuilder, GR, MRI);
+      buildScopeReg(ScopeRegister, SPIRV::Scope::Device, MIRBuilder, GOR, MRI);
 
   auto MIB = MIRBuilder.buildInstr(Opcode);
   if (Opcode == SPIRV::OpAtomicFlagTestAndSet)
     MIB.addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType));
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType));
 
   MIB.addUse(PtrRegister).addUse(ScopeRegister).addUse(MemSemanticsReg);
   return true;
 }
 
 /// Helper function for building barriers, i.e., memory/control ordering
 /// operations.
 static bool buildBarrierInst(const SPIRV::IncomingCall *Call, unsigned Opcode,
                              MachineIRBuilder &MIRBuilder,
-                             SPIRVGlobalRegistry *GR) {
+                             SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   unsigned MemFlags = getIConstVal(Call->Arguments[0], MRI);
   unsigned MemSemantics = SPIRV::MemorySemantics::None;
 
   if (MemFlags & SPIRV::CLK_LOCAL_MEM_FENCE)
     MemSemantics |= SPIRV::MemorySemantics::WorkgroupMemory;
 
   if (MemFlags & SPIRV::CLK_GLOBAL_MEM_FENCE)
if (Opcode == SPIRV::OpMemoryBarrier) {
     MemSemantics |= SPIRV::MemorySemantics::SequentiallyConsistent;
   }
 
   Register MemSemanticsReg;
   if (MemFlags == MemSemantics) {
     MemSemanticsReg = Call->Arguments[0];
     MRI->setRegClass(MemSemanticsReg, &SPIRV::IDRegClass);
   } else
-    MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GR);
+    MemSemanticsReg = buildConstantIntReg(MemSemantics, MIRBuilder, GOR);
 
   Register ScopeReg;
   SPIRV::Scope::Scope Scope = SPIRV::Scope::Workgroup;
   SPIRV::Scope::Scope MemScope = Scope;
   if (Call->Arguments.size() >= 2) {
     assert(
         ((Opcode != SPIRV::OpMemoryBarrier && Call->Arguments.size() == 2) ||
          (Opcode == SPIRV::OpMemoryBarrier && Call->Arguments.size() == 3)) &&
if (Call->Arguments.size() >= 2) {
 
     if (CLScope == static_cast<unsigned>(Scope)) {
       ScopeReg = Call->Arguments[1];
       MRI->setRegClass(ScopeReg, &SPIRV::IDRegClass);
     }
   }
 
   if (!ScopeReg.isValid())
-    ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GR);
+    ScopeReg = buildConstantIntReg(Scope, MIRBuilder, GOR);
 
   auto MIB = MIRBuilder.buildInstr(Opcode).addUse(ScopeReg);
   if (Opcode != SPIRV::OpMemoryBarrier)
-    MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GR));
+    MIB.addUse(buildConstantIntReg(MemScope, MIRBuilder, GOR));
   MIB.addUse(MemSemanticsReg);
   return true;
 }
 
 static unsigned getNumComponentsForDim(SPIRV::Dim::Dim dim) {
   switch (dim) {
   case SPIRV::Dim::DIM_1D:
   case SPIRV::Dim::DIM_Buffer:

 }
 
 //===----------------------------------------------------------------------===//
 // Implementation functions for each builtin group
 //===----------------------------------------------------------------------===//
 
 static bool generateExtInst(const SPIRV::IncomingCall *Call,
                             MachineIRBuilder &MIRBuilder,
-                            SPIRVGlobalRegistry *GR) {
+                            SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the extended instruction number in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   uint32_t Number =
       SPIRV::lookupExtendedBuiltin(Builtin->Name, Builtin->Set)->Number;
 
   // Build extended instruction.
   auto MIB =
       MIRBuilder.buildInstr(SPIRV::OpExtInst)
           .addDef(Call->ReturnRegister)
-          .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+          .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
           .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
           .addImm(Number);
 
   for (auto Argument : Call->Arguments)
     MIB.addUse(Argument);
   return true;
 }
 
 static bool generateRelationalInst(const SPIRV::IncomingCall *Call,
                                    MachineIRBuilder &MIRBuilder,
-                                   SPIRVGlobalRegistry *GR) {
+                                   SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 
   Register CompareRegister;
   SPIRVType *RelationType;
   std::tie(CompareRegister, RelationType) =
-      buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
+      buildBoolRegister(MIRBuilder, Call->ReturnType, GOR);
 
   // Build relational instruction.
   auto MIB = MIRBuilder.buildInstr(Opcode)
                  .addDef(CompareRegister)
-                 .addUse(GR->getSPIRVTypeID(RelationType));
+                 .addUse(GOR->getSPIRVTypeID(RelationType));
 
   for (auto Argument : Call->Arguments)
     MIB.addUse(Argument);
 
   // Build select instruction.
   return buildSelectInst(MIRBuilder, Call->ReturnRegister, CompareRegister,
-                         Call->ReturnType, GR);
+                         Call->ReturnType, GOR);
 }
 
 static bool generateGroupInst(const SPIRV::IncomingCall *Call,
                               MachineIRBuilder &MIRBuilder,
-                              SPIRVGlobalRegistry *GR) {
+                              SPIRVGlobalObjectRegistry *GOR) {
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   const SPIRV::GroupBuiltin *GroupBuiltin =
       SPIRV::lookupGroupBuiltin(Builtin->Name);
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   Register Arg0;
   if (GroupBuiltin->HasBoolArg) {
     Register ConstRegister = Call->Arguments[0];
     auto ArgInstruction = getDefInstrMaybeConstant(ConstRegister, MRI);
     // TODO: support non-constant bool values.
     assert(ArgInstruction->getOpcode() == TargetOpcode::G_CONSTANT &&
            "Only constant bool value args are supported");
-    if (GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
+    if (GOR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode() !=
         SPIRV::OpTypeBool)
-      Arg0 = GR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
-                                  GR->getOrCreateSPIRVBoolType(MIRBuilder));
+      Arg0 = GOR->buildConstantInt(getIConstVal(ConstRegister, MRI), MIRBuilder,
+                                   GOR->getOrCreateSPIRVBoolType(MIRBuilder));
   }
 
   Register GroupResultRegister = Call->ReturnRegister;
   SPIRVType *GroupResultType = Call->ReturnType;
 
   // TODO: maybe we need to check whether the result type is already boolean
   // and in this case do not insert select instruction.
   const bool HasBoolReturnTy =
       GroupBuiltin->IsElect || GroupBuiltin->IsAllOrAny ||
       GroupBuiltin->IsAllEqual || GroupBuiltin->IsLogical ||
       GroupBuiltin->IsInverseBallot || GroupBuiltin->IsBallotBitExtract;
 
   if (HasBoolReturnTy)
     std::tie(GroupResultRegister, GroupResultType) =
-        buildBoolRegister(MIRBuilder, Call->ReturnType, GR);
+        buildBoolRegister(MIRBuilder, Call->ReturnType, GOR);
 
   auto Scope = Builtin->Name.startswith("sub_group") ? SPIRV::Scope::Subgroup
                                                      : SPIRV::Scope::Workgroup;
-  Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GR);
+  Register ScopeRegister = buildConstantIntReg(Scope, MIRBuilder, GOR);
 
   // Build work/sub group instruction.
   auto MIB = MIRBuilder.buildInstr(GroupBuiltin->Opcode)
                  .addDef(GroupResultRegister)
-                 .addUse(GR->getSPIRVTypeID(GroupResultType))
+                 .addUse(GOR->getSPIRVTypeID(GroupResultType))
                  .addUse(ScopeRegister);
 
   if (!GroupBuiltin->NoGroupOperation)
     MIB.addImm(GroupBuiltin->GroupOperation);
   if (Call->Arguments.size() > 0) {
     MIB.addUse(Arg0.isValid() ? Arg0 : Call->Arguments[0]);
     MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
     for (unsigned i = 1; i < Call->Arguments.size(); i++) {
       MIB.addUse(Call->Arguments[i]);
       MRI->setRegClass(Call->Arguments[i], &SPIRV::IDRegClass);
     }
   }
 
   // Build select instruction.
   if (HasBoolReturnTy)
     buildSelectInst(MIRBuilder, Call->ReturnRegister, GroupResultRegister,
-                    Call->ReturnType, GR);
+                    Call->ReturnType, GOR);
   return true;
 }
 
 // These queries ask for a single size_t result for a given dimension index, e.g
 // size_t get_global_id(uint dimindex). In SPIR-V, the builtins corresonding to
 // these values are all vec3 types, so we need to extract the correct index or
 // return defaultVal (0 or 1 depending on the query). We also handle extending
 // or tuncating in case size_t does not match the expected result type's

 //    %tmp = OpVectorExtractDynamic %SizeT %loadedVec %idx
 //    %cmp = OpULessThan %bool %idx %const_3
 //    %res = OpSelect %SizeT %cmp %tmp %const_0
 //
 //  If the bitwidth of %res does not match the expected return type, we add an
 //  extend or truncate.
 static bool genWorkgroupQuery(const SPIRV::IncomingCall *Call,
                               MachineIRBuilder &MIRBuilder,
-                              SPIRVGlobalRegistry *GR,
+                              SPIRVGlobalObjectRegistry *GOR,
                               SPIRV::BuiltIn::BuiltIn BuiltinValue,
                               uint64_t DefaultValue) {
   Register IndexRegister = Call->Arguments[0];
   const unsigned ResultWidth = Call->ReturnType->getOperand(1).getImm();
-  const unsigned PointerSize = GR->getPointerSize();
+  const unsigned PointerSize = GOR->getPointerSize();
   const SPIRVType *PointerSizeType =
-      GR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
+      GOR->getOrCreateSPIRVIntegerType(PointerSize, MIRBuilder);
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   auto IndexInstruction = getDefInstrMaybeConstant(IndexRegister, MRI);
 
   // Set up the final register to do truncation or extension on at the end.
   Register ToTruncate = Call->ReturnRegister;
 
   // If the index is constant, we can statically determine if it is in range.
   bool IsConstantIndex =
       IndexInstruction->getOpcode() == TargetOpcode::G_CONSTANT;
 
   // If it's out of range (max dimension is 3), we can just return the constant
   // default value (0 or 1 depending on which query function).
   if (IsConstantIndex && getIConstVal(IndexRegister, MRI) >= 3) {
     Register DefaultReg = Call->ReturnRegister;
     if (PointerSize != ResultWidth) {
       DefaultReg = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
       MRI->setRegClass(DefaultReg, &SPIRV::IDRegClass);
-      GR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg,
+      GOR->assignSPIRVTypeToVReg(PointerSizeType, DefaultReg,
                                  MIRBuilder.getMF());
       ToTruncate = DefaultReg;
     }
     auto NewRegister =
-        GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
+        GOR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
     MIRBuilder.buildCopy(DefaultReg, NewRegister);
   } else { // If it could be in range, we need to load from the given builtin.
     auto Vec3Ty =
-        GR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
+        GOR->getOrCreateSPIRVVectorType(PointerSizeType, 3, MIRBuilder);
     Register LoadedVector =
-        buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GR, BuiltinValue,
+        buildBuiltinVariableLoad(MIRBuilder, Vec3Ty, GOR, BuiltinValue,
                                  LLT::fixed_vector(3, PointerSize));
     // Set up the vreg to extract the result to (possibly a new temporary one).
     Register Extracted = Call->ReturnRegister;
     if (!IsConstantIndex || PointerSize != ResultWidth) {
       Extracted = MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
       MRI->setRegClass(Extracted, &SPIRV::IDRegClass);
-      GR->assignSPIRVTypeToVReg(PointerSizeType, Extracted, MIRBuilder.getMF());
+      GOR->assignSPIRVTypeToVReg(PointerSizeType, Extracted,
+                                 MIRBuilder.getMF());
     }
     // Use Intrinsic::spv_extractelt so dynamic vs static extraction is
     // handled later: extr = spv_extractelt LoadedVector, IndexRegister.
     MachineInstrBuilder ExtractInst = MIRBuilder.buildIntrinsic(
         Intrinsic::spv_extractelt, ArrayRef<Register>{Extracted}, true);
     ExtractInst.addUse(LoadedVector).addUse(IndexRegister);
 
     // If the index is dynamic, need check if it's < 3, and then use a select.
     if (!IsConstantIndex) {
-      insertAssignInstr(Extracted, nullptr, PointerSizeType, GR, MIRBuilder,
+      insertAssignInstr(Extracted, nullptr, PointerSizeType, GOR, MIRBuilder,
                         *MRI);
 
-      auto IndexType = GR->getSPIRVTypeForVReg(IndexRegister);
-      auto BoolType = GR->getOrCreateSPIRVBoolType(MIRBuilder);
+      auto IndexType = GOR->getSPIRVTypeForVReg(IndexRegister);
+      auto BoolType = GOR->getOrCreateSPIRVBoolType(MIRBuilder);
 
       Register CompareRegister =
           MRI->createGenericVirtualRegister(LLT::scalar(1));
       MRI->setRegClass(CompareRegister, &SPIRV::IDRegClass);
-      GR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
+      GOR->assignSPIRVTypeToVReg(BoolType, CompareRegister, MIRBuilder.getMF());
 
       // Use G_ICMP to check if idxVReg < 3.
       MIRBuilder.buildICmp(CmpInst::ICMP_ULT, CompareRegister, IndexRegister,
-                           GR->buildConstantInt(3, MIRBuilder, IndexType));
+                           GOR->buildConstantInt(3, MIRBuilder, IndexType));
 
       // Get constant for the default value (0 or 1 depending on which
       // function).
       Register DefaultRegister =
-          GR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
+          GOR->buildConstantInt(DefaultValue, MIRBuilder, PointerSizeType);
 
       // Get a register for the selection result (possibly a new temporary one).
       Register SelectionResult = Call->ReturnRegister;
       if (PointerSize != ResultWidth) {
         SelectionResult =
             MRI->createGenericVirtualRegister(LLT::scalar(PointerSize));
         MRI->setRegClass(SelectionResult, &SPIRV::IDRegClass);
-        GR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
+        GOR->assignSPIRVTypeToVReg(PointerSizeType, SelectionResult,
                                    MIRBuilder.getMF());
       }
       // Create the final G_SELECT to return the extracted value or the default.
       MIRBuilder.buildSelect(SelectionResult, CompareRegister, Extracted,
                              DefaultRegister);
       ToTruncate = SelectionResult;
     } else {
       ToTruncate = Extracted;
     }
   }
   // Alter the result's bitwidth if it does not match the SizeT value extracted.
   if (PointerSize != ResultWidth)
     MIRBuilder.buildZExtOrTrunc(Call->ReturnRegister, ToTruncate);
   return true;
 }
 
 static bool generateBuiltinVar(const SPIRV::IncomingCall *Call,
                                MachineIRBuilder &MIRBuilder,
-                               SPIRVGlobalRegistry *GR) {
+                               SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the builtin variable record.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   SPIRV::BuiltIn::BuiltIn Value =
       SPIRV::lookupGetBuiltin(Builtin->Name, Builtin->Set)->Value;
 
   if (Value == SPIRV::BuiltIn::GlobalInvocationId)
-    return genWorkgroupQuery(Call, MIRBuilder, GR, Value, 0);
+    return genWorkgroupQuery(Call, MIRBuilder, GOR, Value, 0);
 
   // Build a load instruction for the builtin variable.
-  unsigned BitWidth = GR->getScalarOrVectorBitWidth(Call->ReturnType);
+  unsigned BitWidth = GOR->getScalarOrVectorBitWidth(Call->ReturnType);
   LLT LLType;
   if (Call->ReturnType->getOpcode() == SPIRV::OpTypeVector)
     LLType =
         LLT::fixed_vector(Call->ReturnType->getOperand(2).getImm(), BitWidth);
   else
     LLType = LLT::scalar(BitWidth);
 
-  return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GR, Value,
+  return buildBuiltinVariableLoad(MIRBuilder, Call->ReturnType, GOR, Value,
                                   LLType, Call->ReturnRegister);
 }
 
 static bool generateAtomicInst(const SPIRV::IncomingCall *Call,
                                MachineIRBuilder &MIRBuilder,
-                               SPIRVGlobalRegistry *GR) {
+                               SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 
   switch (Opcode) {
   case SPIRV::OpStore:
     return buildAtomicInitInst(Call, MIRBuilder);
   case SPIRV::OpAtomicLoad:
-    return buildAtomicLoadInst(Call, MIRBuilder, GR);
+    return buildAtomicLoadInst(Call, MIRBuilder, GOR);
   case SPIRV::OpAtomicStore:
-    return buildAtomicStoreInst(Call, MIRBuilder, GR);
+    return buildAtomicStoreInst(Call, MIRBuilder, GOR);
   case SPIRV::OpAtomicCompareExchange:
   case SPIRV::OpAtomicCompareExchangeWeak:
-    return buildAtomicCompareExchangeInst(Call, MIRBuilder, GR);
+    return buildAtomicCompareExchangeInst(Call, MIRBuilder, GOR);
   case SPIRV::OpAtomicIAdd:
   case SPIRV::OpAtomicISub:
   case SPIRV::OpAtomicOr:
   case SPIRV::OpAtomicXor:
   case SPIRV::OpAtomicAnd:
   case SPIRV::OpAtomicExchange:
-    return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GR);
+    return buildAtomicRMWInst(Call, Opcode, MIRBuilder, GOR);
   case SPIRV::OpMemoryBarrier:
-    return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GR);
+    return buildBarrierInst(Call, SPIRV::OpMemoryBarrier, MIRBuilder, GOR);
   case SPIRV::OpAtomicFlagTestAndSet:
   case SPIRV::OpAtomicFlagClear:
-    return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GR);
+    return buildAtomicFlagInst(Call, Opcode, MIRBuilder, GOR);
   default:
     return false;
   }
 }
 
 static bool generateBarrierInst(const SPIRV::IncomingCall *Call,
                                 MachineIRBuilder &MIRBuilder,
-                                SPIRVGlobalRegistry *GR) {
+                                SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 
-  return buildBarrierInst(Call, Opcode, MIRBuilder, GR);
+  return buildBarrierInst(Call, Opcode, MIRBuilder, GOR);
 }
 
 static bool generateDotOrFMulInst(const SPIRV::IncomingCall *Call,
                                   MachineIRBuilder &MIRBuilder,
-                                  SPIRVGlobalRegistry *GR) {
-  unsigned Opcode = GR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
+                                  SPIRVGlobalObjectRegistry *GOR) {
+  unsigned Opcode = GOR->getSPIRVTypeForVReg(Call->Arguments[0])->getOpcode();
   bool IsVec = Opcode == SPIRV::OpTypeVector;
   // Use OpDot only in case of vector args and OpFMul in case of scalar args.
   MIRBuilder.buildInstr(IsVec ? SPIRV::OpDot : SPIRV::OpFMulS)
       .addDef(Call->ReturnRegister)
-      .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+      .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
       .addUse(Call->Arguments[0])
       .addUse(Call->Arguments[1]);
   return true;
 }
 
 static bool generateGetQueryInst(const SPIRV::IncomingCall *Call,
                                  MachineIRBuilder &MIRBuilder,
-                                 SPIRVGlobalRegistry *GR) {
+                                 SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the builtin record.
   SPIRV::BuiltIn::BuiltIn Value =
       SPIRV::lookupGetBuiltin(Call->Builtin->Name, Call->Builtin->Set)->Value;
   uint64_t IsDefault = (Value == SPIRV::BuiltIn::GlobalSize ||
                         Value == SPIRV::BuiltIn::WorkgroupSize ||
                         Value == SPIRV::BuiltIn::EnqueuedWorkgroupSize);
-  return genWorkgroupQuery(Call, MIRBuilder, GR, Value, IsDefault ? 1 : 0);
+  return genWorkgroupQuery(Call, MIRBuilder, GOR, Value, IsDefault ? 1 : 0);
 }
 
 static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call,
                                        MachineIRBuilder &MIRBuilder,
-                                       SPIRVGlobalRegistry *GR) {
+                                       SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the image size query component number in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   uint32_t Component =
       SPIRV::lookupImageQueryBuiltin(Builtin->Name, Builtin->Set)->Component;
   // Query result may either be a vector or a scalar. If return type is not a
   // vector, expect only a single size component. Otherwise get the number of
   // expected components.
   SPIRVType *RetTy = Call->ReturnType;
   unsigned NumExpectedRetComponents = RetTy->getOpcode() == SPIRV::OpTypeVector
                                           ? RetTy->getOperand(2).getImm()
                                           : 1;
   // Get the actual number of query result/size components.
-  SPIRVType *ImgType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
+  SPIRVType *ImgType = GOR->getSPIRVTypeForVReg(Call->Arguments[0]);
   unsigned NumActualRetComponents = getNumSizeComponents(ImgType);
   Register QueryResult = Call->ReturnRegister;
   SPIRVType *QueryResultType = Call->ReturnType;
   if (NumExpectedRetComponents != NumActualRetComponents) {
     QueryResult = MIRBuilder.getMRI()->createGenericVirtualRegister(
         LLT::fixed_vector(NumActualRetComponents, 32));
     MIRBuilder.getMRI()->setRegClass(QueryResult, &SPIRV::IDRegClass);
-    SPIRVType *IntTy = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
-    QueryResultType = GR->getOrCreateSPIRVVectorType(
+    SPIRVType *IntTy = GOR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
+    QueryResultType = GOR->getOrCreateSPIRVVectorType(
         IntTy, NumActualRetComponents, MIRBuilder);
-    GR->assignSPIRVTypeToVReg(QueryResultType, QueryResult, MIRBuilder.getMF());
+    GOR->assignSPIRVTypeToVReg(QueryResultType, QueryResult,
+                               MIRBuilder.getMF());
   }
   bool IsDimBuf = ImgType->getOperand(2).getImm() == SPIRV::Dim::DIM_Buffer;
   unsigned Opcode =
       IsDimBuf ? SPIRV::OpImageQuerySize : SPIRV::OpImageQuerySizeLod;
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
   auto MIB = MIRBuilder.buildInstr(Opcode)
                  .addDef(QueryResult)
-                 .addUse(GR->getSPIRVTypeID(QueryResultType))
+                 .addUse(GOR->getSPIRVTypeID(QueryResultType))
                  .addUse(Call->Arguments[0]);
   if (!IsDimBuf)
-    MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Lod id.
+    MIB.addUse(buildConstantIntReg(0, MIRBuilder, GOR)); // Lod id.
   if (NumExpectedRetComponents == NumActualRetComponents)
     return true;
   if (NumExpectedRetComponents == 1) {
     // Only 1 component is expected, build OpCompositeExtract instruction.
     unsigned ExtractedComposite =
         Component == 3 ? NumActualRetComponents - 1 : Component;
     assert(ExtractedComposite < NumActualRetComponents &&
            "Invalid composite index!");
     MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
         .addUse(QueryResult)
         .addImm(ExtractedComposite);
   } else {
     // More than 1 component is expected, fill a new vector.
     auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle)
                    .addDef(Call->ReturnRegister)
-                   .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+                   .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
                    .addUse(QueryResult)
                    .addUse(QueryResult);
     for (unsigned i = 0; i < NumExpectedRetComponents; ++i)
       MIB.addImm(i < NumActualRetComponents ? i : 0xffffffff);
   }
   return true;
 }
 
 static bool generateImageMiscQueryInst(const SPIRV::IncomingCall *Call,
                                        MachineIRBuilder &MIRBuilder,
-                                       SPIRVGlobalRegistry *GR) {
+                                       SPIRVGlobalObjectRegistry *GOR) {
   assert(Call->ReturnType->getOpcode() == SPIRV::OpTypeInt &&
          "Image samples query result must be of int type!");
 
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 
   Register Image = Call->Arguments[0];
   MIRBuilder.getMRI()->setRegClass(Image, &SPIRV::IDRegClass);
   SPIRV::Dim::Dim ImageDimensionality = static_cast<SPIRV::Dim::Dim>(
-      GR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
+      GOR->getSPIRVTypeForVReg(Image)->getOperand(2).getImm());
 
   switch (Opcode) {
   case SPIRV::OpImageQuerySamples:
     assert(ImageDimensionality == SPIRV::Dim::DIM_2D &&
            "Image must be of 2D dimensionality");
     break;
   case SPIRV::OpImageQueryLevels:
     assert((ImageDimensionality == SPIRV::Dim::DIM_1D ||
             ImageDimensionality == SPIRV::Dim::DIM_2D ||
             ImageDimensionality == SPIRV::Dim::DIM_3D ||
             ImageDimensionality == SPIRV::Dim::DIM_Cube) &&
            "Image must be of 1D/2D/3D/Cube dimensionality");
     break;
   }
 
   MIRBuilder.buildInstr(Opcode)
       .addDef(Call->ReturnRegister)
-      .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+      .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
       .addUse(Image);
   return true;
 }
 
 // TODO: Move to TableGen.
 static SPIRV::SamplerAddressingMode::SamplerAddressingMode
 getSamplerAddressingModeFromBitmask(unsigned Bitmask) {
   switch (Bitmask & SPIRV::CLK_ADDRESS_MODE_MASK) {
getSamplerFilterModeFromBitmask(unsigned Bitmask) {
   if (Bitmask & SPIRV::CLK_FILTER_NEAREST)
     return SPIRV::SamplerFilterMode::Nearest;
   return SPIRV::SamplerFilterMode::Nearest;
 }
 
 static bool generateReadImageInst(const StringRef DemangledCall,
                                   const SPIRV::IncomingCall *Call,
                                   MachineIRBuilder &MIRBuilder,
-                                  SPIRVGlobalRegistry *GR) {
+                                  SPIRVGlobalObjectRegistry *GOR) {
   Register Image = Call->Arguments[0];
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   MRI->setRegClass(Image, &SPIRV::IDRegClass);
   MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
   bool HasOclSampler = DemangledCall.contains_insensitive("ocl_sampler");
   bool HasMsaa = DemangledCall.contains_insensitive("msaa");
   if (HasOclSampler || HasMsaa)
     MRI->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
   if (HasOclSampler) {
     Register Sampler = Call->Arguments[1];
 
-    if (!GR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
+    if (!GOR->isScalarOfType(Sampler, SPIRV::OpTypeSampler) &&
         getDefInstrMaybeConstant(Sampler, MRI)->getOperand(1).isCImm()) {
       uint64_t SamplerMask = getIConstVal(Sampler, MRI);
-      Sampler = GR->buildConstantSampler(
+      Sampler = GOR->buildConstantSampler(
           Register(), getSamplerAddressingModeFromBitmask(SamplerMask),
           getSamplerParamFromBitmask(SamplerMask),
           getSamplerFilterModeFromBitmask(SamplerMask), MIRBuilder,
-          GR->getSPIRVTypeForVReg(Sampler));
+          GOR->getSPIRVTypeForVReg(Sampler));
     }
-    SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
+    SPIRVType *ImageType = GOR->getSPIRVTypeForVReg(Image);
     SPIRVType *SampledImageType =
-        GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
+        GOR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
     Register SampledImage = MRI->createVirtualRegister(&SPIRV::IDRegClass);
 
     MIRBuilder.buildInstr(SPIRV::OpSampledImage)
         .addDef(SampledImage)
-        .addUse(GR->getSPIRVTypeID(SampledImageType))
+        .addUse(GOR->getSPIRVTypeID(SampledImageType))
         .addUse(Image)
         .addUse(Sampler);
 
-    Register Lod = GR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()),
+    Register Lod = GOR->buildConstantFP(APFloat::getZero(APFloat::IEEEsingle()),
                                         MIRBuilder);
     SPIRVType *TempType = Call->ReturnType;
     bool NeedsExtraction = false;
     if (TempType->getOpcode() != SPIRV::OpTypeVector) {
       TempType =
-          GR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
+          GOR->getOrCreateSPIRVVectorType(Call->ReturnType, 4, MIRBuilder);
       NeedsExtraction = true;
     }
-    LLT LLType = LLT::scalar(GR->getScalarOrVectorBitWidth(TempType));
+    LLT LLType = LLT::scalar(GOR->getScalarOrVectorBitWidth(TempType));
     Register TempRegister = MRI->createGenericVirtualRegister(LLType);
     MRI->setRegClass(TempRegister, &SPIRV::IDRegClass);
-    GR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
+    GOR->assignSPIRVTypeToVReg(TempType, TempRegister, MIRBuilder.getMF());
 
     MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
         .addDef(NeedsExtraction ? TempRegister : Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(TempType))
+        .addUse(GOR->getSPIRVTypeID(TempType))
         .addUse(SampledImage)
         .addUse(Call->Arguments[2]) // Coordinate.
         .addImm(SPIRV::ImageOperand::Lod)
         .addUse(Lod);
 
     if (NeedsExtraction)
       MIRBuilder.buildInstr(SPIRV::OpCompositeExtract)
           .addDef(Call->ReturnRegister)
-          .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+          .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
           .addUse(TempRegister)
           .addImm(0);
   } else if (HasMsaa) {
     MIRBuilder.buildInstr(SPIRV::OpImageRead)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
         .addUse(Image)
         .addUse(Call->Arguments[1]) // Coordinate.
         .addImm(SPIRV::ImageOperand::Sample)
         .addUse(Call->Arguments[2]);
   } else {
     MIRBuilder.buildInstr(SPIRV::OpImageRead)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
         .addUse(Image)
         .addUse(Call->Arguments[1]); // Coordinate.
   }
   return true;
 }
 
 static bool generateWriteImageInst(const SPIRV::IncomingCall *Call,
                                    MachineIRBuilder &MIRBuilder,
-                                   SPIRVGlobalRegistry *GR) {
+                                   SPIRVGlobalObjectRegistry *GOR) {
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
   MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
   MIRBuilder.buildInstr(SPIRV::OpImageWrite)
       .addUse(Call->Arguments[0])  // Image.
       .addUse(Call->Arguments[1])  // Coordinate.
       .addUse(Call->Arguments[2]); // Texel.
   return true;
 }
 
 static bool generateSampleImageInst(const StringRef DemangledCall,
                                     const SPIRV::IncomingCall *Call,
                                     MachineIRBuilder &MIRBuilder,
-                                    SPIRVGlobalRegistry *GR) {
+                                    SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   if (Call->Builtin->Name.contains_insensitive(
           "__translate_sampler_initializer")) {
     // Build sampler literal.
     uint64_t Bitmask = getIConstVal(Call->Arguments[0], MRI);
-    Register Sampler = GR->buildConstantSampler(
+    Register Sampler = GOR->buildConstantSampler(
         Call->ReturnRegister, getSamplerAddressingModeFromBitmask(Bitmask),
         getSamplerParamFromBitmask(Bitmask),
         getSamplerFilterModeFromBitmask(Bitmask), MIRBuilder, Call->ReturnType);
     return Sampler.isValid();
   } else if (Call->Builtin->Name.contains_insensitive("__spirv_SampledImage")) {
     // Create OpSampledImage.
     Register Image = Call->Arguments[0];
-    SPIRVType *ImageType = GR->getSPIRVTypeForVReg(Image);
+    SPIRVType *ImageType = GOR->getSPIRVTypeForVReg(Image);
     SPIRVType *SampledImageType =
-        GR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
+        GOR->getOrCreateOpTypeSampledImage(ImageType, MIRBuilder);
     Register SampledImage =
         Call->ReturnRegister.isValid()
             ? Call->ReturnRegister
             : MRI->createVirtualRegister(&SPIRV::IDRegClass);
     MIRBuilder.buildInstr(SPIRV::OpSampledImage)
         .addDef(SampledImage)
-        .addUse(GR->getSPIRVTypeID(SampledImageType))
+        .addUse(GOR->getSPIRVTypeID(SampledImageType))
         .addUse(Image)
         .addUse(Call->Arguments[1]); // Sampler.
     return true;
   } else if (Call->Builtin->Name.contains_insensitive(
                  "__spirv_ImageSampleExplicitLod")) {
     // Sample an image using an explicit level of detail.
     std::string ReturnType = DemangledCall.str();
     if (DemangledCall.contains("_R")) {
       ReturnType = ReturnType.substr(ReturnType.find("_R") + 2);
       ReturnType = ReturnType.substr(0, ReturnType.find('('));
     }
-    SPIRVType *Type = GR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);
+    SPIRVType *Type = GOR->getOrCreateSPIRVTypeByName(ReturnType, MIRBuilder);
     MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
     MRI->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
     MRI->setRegClass(Call->Arguments[3], &SPIRV::IDRegClass);
 
     MIRBuilder.buildInstr(SPIRV::OpImageSampleExplicitLod)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Type))
+        .addUse(GOR->getSPIRVTypeID(Type))
         .addUse(Call->Arguments[0]) // Image.
         .addUse(Call->Arguments[1]) // Coordinate.
         .addImm(SPIRV::ImageOperand::Lod)
         .addUse(Call->Arguments[3]);
     return true;
   }
   return false;
 }
 
 static bool generateSelectInst(const SPIRV::IncomingCall *Call,
                                MachineIRBuilder &MIRBuilder) {
   MIRBuilder.buildSelect(Call->ReturnRegister, Call->Arguments[0],
                          Call->Arguments[1], Call->Arguments[2]);
   return true;
 }
 
 static bool generateSpecConstantInst(const SPIRV::IncomingCall *Call,
                                      MachineIRBuilder &MIRBuilder,
-                                     SPIRVGlobalRegistry *GR) {
+                                     SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
   const MachineRegisterInfo *MRI = MIRBuilder.getMRI();
 
   switch (Opcode) {
   case SPIRV::OpSpecConstant: {
case SPIRV::OpSpecConstant: {
     if (Call->ReturnType->getOpcode() == SPIRV::OpTypeBool) {
       assert(ConstOperand.isCImm() && "Int constant operand is expected");
       Opcode = ConstOperand.getCImm()->getValue().getZExtValue()
                    ? SPIRV::OpSpecConstantTrue
                    : SPIRV::OpSpecConstantFalse;
     }
     auto MIB = MIRBuilder.buildInstr(Opcode)
                    .addDef(Call->ReturnRegister)
-                   .addUse(GR->getSPIRVTypeID(Call->ReturnType));
+                   .addUse(GOR->getSPIRVTypeID(Call->ReturnType));
 
     if (Call->ReturnType->getOpcode() != SPIRV::OpTypeBool) {
       if (Const->getOpcode() == TargetOpcode::G_CONSTANT)
         addNumImm(ConstOperand.getCImm()->getValue(), MIB);
       else
         addNumImm(ConstOperand.getFPImm()->getValueAPF().bitcastToAPInt(), MIB);
     }
     return true;
   }
   case SPIRV::OpSpecConstantComposite: {
     auto MIB = MIRBuilder.buildInstr(Opcode)
                    .addDef(Call->ReturnRegister)
-                   .addUse(GR->getSPIRVTypeID(Call->ReturnType));
+                   .addUse(GOR->getSPIRVTypeID(Call->ReturnType));
     for (unsigned i = 0; i < Call->Arguments.size(); i++)
       MIB.addUse(Call->Arguments[i]);
     return true;
   }
   default:
     return false;
   }
 }
 
 static bool buildNDRange(const SPIRV::IncomingCall *Call,
                          MachineIRBuilder &MIRBuilder,
-                         SPIRVGlobalRegistry *GR) {
+                         SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
-  SPIRVType *PtrType = GR->getSPIRVTypeForVReg(Call->Arguments[0]);
+  SPIRVType *PtrType = GOR->getSPIRVTypeForVReg(Call->Arguments[0]);
   assert(PtrType->getOpcode() == SPIRV::OpTypePointer &&
          PtrType->getOperand(2).isReg());
   Register TypeReg = PtrType->getOperand(2).getReg();
-  SPIRVType *StructType = GR->getSPIRVTypeForVReg(TypeReg);
+  SPIRVType *StructType = GOR->getSPIRVTypeForVReg(TypeReg);
   MachineFunction &MF = MIRBuilder.getMF();
   Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
-  GR->assignSPIRVTypeToVReg(StructType, TmpReg, MF);
+  GOR->assignSPIRVTypeToVReg(StructType, TmpReg, MF);
   // Skip the first arg, it's the destination pointer. OpBuildNDRange takes
   // three other arguments, so pass zero constant on absence.
   unsigned NumArgs = Call->Arguments.size();
   assert(NumArgs >= 2);
   Register GlobalWorkSize = Call->Arguments[NumArgs < 4 ? 1 : 2];
   MRI->setRegClass(GlobalWorkSize, &SPIRV::IDRegClass);
   Register LocalWorkSize =
       NumArgs == 2 ? Register(0) : Call->Arguments[NumArgs < 4 ? 2 : 3];
   if (LocalWorkSize.isValid())
     MRI->setRegClass(LocalWorkSize, &SPIRV::IDRegClass);
   Register GlobalWorkOffset = NumArgs <= 3 ? Register(0) : Call->Arguments[1];
   if (GlobalWorkOffset.isValid())
     MRI->setRegClass(GlobalWorkOffset, &SPIRV::IDRegClass);
   if (NumArgs < 4) {
     Register Const;
-    SPIRVType *SpvTy = GR->getSPIRVTypeForVReg(GlobalWorkSize);
+    SPIRVType *SpvTy = GOR->getSPIRVTypeForVReg(GlobalWorkSize);
     if (SpvTy->getOpcode() == SPIRV::OpTypePointer) {
       MachineInstr *DefInstr = MRI->getUniqueVRegDef(GlobalWorkSize);
       assert(DefInstr && isSpvIntrinsic(*DefInstr, Intrinsic::spv_gep) &&
              DefInstr->getOperand(3).isReg());
       Register GWSPtr = DefInstr->getOperand(3).getReg();
       if (!MRI->getRegClassOrNull(GWSPtr))
         MRI->setRegClass(GWSPtr, &SPIRV::IDRegClass);
       // TODO: Maybe simplify generation of the type of the fields.
       unsigned Size = Call->Builtin->Name.equals("ndrange_3D") ? 3 : 2;
-      unsigned BitWidth = GR->getPointerSize() == 64 ? 64 : 32;
+      unsigned BitWidth = GOR->getPointerSize() == 64 ? 64 : 32;
       Type *BaseTy = IntegerType::get(MF.getFunction().getContext(), BitWidth);
       Type *FieldTy = ArrayType::get(BaseTy, Size);
-      SPIRVType *SpvFieldTy = GR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
+      SPIRVType *SpvFieldTy = GOR->getOrCreateSPIRVType(FieldTy, MIRBuilder);
       GlobalWorkSize = MRI->createVirtualRegister(&SPIRV::IDRegClass);
-      GR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF);
+      GOR->assignSPIRVTypeToVReg(SpvFieldTy, GlobalWorkSize, MF);
       MIRBuilder.buildInstr(SPIRV::OpLoad)
           .addDef(GlobalWorkSize)
-          .addUse(GR->getSPIRVTypeID(SpvFieldTy))
+          .addUse(GOR->getSPIRVTypeID(SpvFieldTy))
           .addUse(GWSPtr);
-      Const = GR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
+      Const = GOR->getOrCreateConsIntArray(0, MIRBuilder, SpvFieldTy);
     } else {
-      Const = GR->buildConstantInt(0, MIRBuilder, SpvTy);
+      Const = GOR->buildConstantInt(0, MIRBuilder, SpvTy);
     }
     if (!LocalWorkSize.isValid())
       LocalWorkSize = Const;
     if (!GlobalWorkOffset.isValid())
       GlobalWorkOffset = Const;
   }
   assert(LocalWorkSize.isValid() && GlobalWorkOffset.isValid());
   MIRBuilder.buildInstr(SPIRV::OpBuildNDRange)
static const Type *getBlockStructType(Register ParamReg,
   assert(isSpvIntrinsic(*MI, Intrinsic::spv_alloca) &&
          "Blocks in OpenCL C must be traceable to allocation site");
   return getMachineInstrType(MI);
 }
 
 // TODO: maybe move to the global register.
 static SPIRVType *
 getOrCreateSPIRVDeviceEventPointer(MachineIRBuilder &MIRBuilder,
-                                   SPIRVGlobalRegistry *GR) {
+                                   SPIRVGlobalObjectRegistry *GOR) {
   LLVMContext &Context = MIRBuilder.getMF().getFunction().getContext();
   Type *OpaqueType = StructType::getTypeByName(Context, "spirv.DeviceEvent");
   if (!OpaqueType)
     OpaqueType = StructType::getTypeByName(Context, "opencl.clk_event_t");
   if (!OpaqueType)
     OpaqueType = StructType::create(Context, "spirv.DeviceEvent");
   unsigned SC0 = storageClassToAddressSpace(SPIRV::StorageClass::Function);
   unsigned SC1 = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
   Type *PtrType = PointerType::get(PointerType::get(OpaqueType, SC0), SC1);
-  return GR->getOrCreateSPIRVType(PtrType, MIRBuilder);
+  return GOR->getOrCreateSPIRVType(PtrType, MIRBuilder);
 }
 
 static bool buildEnqueueKernel(const SPIRV::IncomingCall *Call,
                                MachineIRBuilder &MIRBuilder,
-                               SPIRVGlobalRegistry *GR) {
+                               SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   const DataLayout &DL = MIRBuilder.getDataLayout();
   bool HasEvents = Call->Builtin->Name.find("events") != StringRef::npos;
-  const SPIRVType *Int32Ty = GR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
+  const SPIRVType *Int32Ty = GOR->getOrCreateSPIRVIntegerType(32, MIRBuilder);
 
   // Make vararg instructions before OpEnqueueKernel.
   // Local sizes arguments: Sizes of block invoke arguments. Clang generates
   // local size operands as an array, so we need to unpack them.
   SmallVector<Register, 16> LocalSizes;
   if (Call->Builtin->Name.find("_varargs") != StringRef::npos) {
     const unsigned LocalSizeArrayIdx = HasEvents ? 9 : 6;
     Register GepReg = Call->Arguments[LocalSizeArrayIdx];
     MachineInstr *GepMI = MRI->getUniqueVRegDef(GepReg);
     assert(isSpvIntrinsic(*GepMI, Intrinsic::spv_gep) &&
            GepMI->getOperand(3).isReg());
     Register ArrayReg = GepMI->getOperand(3).getReg();
     MachineInstr *ArrayMI = MRI->getUniqueVRegDef(ArrayReg);
     const Type *LocalSizeTy = getMachineInstrType(ArrayMI);
     assert(LocalSizeTy && "Local size type is expected");
     const uint64_t LocalSizeNum =
         cast<ArrayType>(LocalSizeTy)->getNumElements();
     unsigned SC = storageClassToAddressSpace(SPIRV::StorageClass::Generic);
-    const LLT LLType = LLT::pointer(SC, GR->getPointerSize());
-    const SPIRVType *PointerSizeTy = GR->getOrCreateSPIRVPointerType(
+    const LLT LLType = LLT::pointer(SC, GOR->getPointerSize());
+    const SPIRVType *PointerSizeTy = GOR->getOrCreateSPIRVPointerType(
         Int32Ty, MIRBuilder, SPIRV::StorageClass::Function);
     for (unsigned I = 0; I < LocalSizeNum; ++I) {
       Register Reg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
       MRI->setType(Reg, LLType);
-      GR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
+      GOR->assignSPIRVTypeToVReg(PointerSizeTy, Reg, MIRBuilder.getMF());
       auto GEPInst = MIRBuilder.buildIntrinsic(Intrinsic::spv_gep,
                                                ArrayRef<Register>{Reg}, true);
       GEPInst
-          .addImm(GepMI->getOperand(2).getImm())          // In bound.
-          .addUse(ArrayMI->getOperand(0).getReg())        // Alloca.
-          .addUse(buildConstantIntReg(0, MIRBuilder, GR)) // Indices.
-          .addUse(buildConstantIntReg(I, MIRBuilder, GR));
+          .addImm(GepMI->getOperand(2).getImm())           // In bound.
+          .addUse(ArrayMI->getOperand(0).getReg())         // Alloca.
+          .addUse(buildConstantIntReg(0, MIRBuilder, GOR)) // Indices.
+          .addUse(buildConstantIntReg(I, MIRBuilder, GOR));
       LocalSizes.push_back(Reg);
     }
   }
 
   // SPIRV OpEnqueueKernel instruction has 10+ arguments.
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpEnqueueKernel)
                  .addDef(Call->ReturnRegister)
-                 .addUse(GR->getSPIRVTypeID(Int32Ty));
+                 .addUse(GOR->getSPIRVTypeID(Int32Ty));
 
   // Copy all arguments before block invoke function pointer.
   const unsigned BlockFIdx = HasEvents ? 6 : 3;
   for (unsigned i = 0; i < BlockFIdx; i++)
     MIB.addUse(Call->Arguments[i]);
 
   // If there are no event arguments in the original call, add dummy ones.
   if (!HasEvents) {
-    MIB.addUse(buildConstantIntReg(0, MIRBuilder, GR)); // Dummy num events.
-    Register NullPtr = GR->getOrCreateConstNullPtr(
-        MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GR));
+    MIB.addUse(buildConstantIntReg(0, MIRBuilder, GOR)); // Dummy num events.
+    Register NullPtr = GOR->getOrCreateConstNullPtr(
+        MIRBuilder, getOrCreateSPIRVDeviceEventPointer(MIRBuilder, GOR));
     MIB.addUse(NullPtr); // Dummy wait events.
     MIB.addUse(NullPtr); // Dummy ret event.
   }
 
   MachineInstr *BlockMI = getBlockStructInstr(Call->Arguments[BlockFIdx], MRI);
   assert(BlockMI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
   // Invoke: Pointer to invoke function.
   MIB.addGlobalAddress(BlockMI->getOperand(1).getGlobal());
 
   Register BlockLiteralReg = Call->Arguments[BlockFIdx + 1];
   // Param: Pointer to block literal.
   MIB.addUse(BlockLiteralReg);
 
   Type *PType = const_cast<Type *>(getBlockStructType(BlockLiteralReg, MRI));
   // TODO: these numbers should be obtained from block literal structure.
   // Param Size: Size of block literal structure.
-  MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GR));
+  MIB.addUse(buildConstantIntReg(DL.getTypeStoreSize(PType), MIRBuilder, GOR));
   // Param Aligment: Aligment of block literal structure.
   MIB.addUse(
-      buildConstantIntReg(DL.getPrefTypeAlign(PType).value(), MIRBuilder, GR));
+      buildConstantIntReg(DL.getPrefTypeAlign(PType).value(), MIRBuilder, GOR));
 
   for (unsigned i = 0; i < LocalSizes.size(); i++)
     MIB.addUse(LocalSizes[i]);
   return true;
 }
 
 static bool generateEnqueueInst(const SPIRV::IncomingCall *Call,
                                 MachineIRBuilder &MIRBuilder,
-                                SPIRVGlobalRegistry *GR) {
+                                SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
 
   switch (Opcode) {
   case SPIRV::OpRetainEvent:
   case SPIRV::OpReleaseEvent:
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
     return MIRBuilder.buildInstr(Opcode).addUse(Call->Arguments[0]);
   case SPIRV::OpCreateUserEvent:
   case SPIRV::OpGetDefaultQueue:
     return MIRBuilder.buildInstr(Opcode)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType));
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType));
   case SPIRV::OpIsValidEvent:
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
     return MIRBuilder.buildInstr(Opcode)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
         .addUse(Call->Arguments[0]);
   case SPIRV::OpSetUserEventStatus:
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
     return MIRBuilder.buildInstr(Opcode)
         .addUse(Call->Arguments[0])
         .addUse(Call->Arguments[1]);
   case SPIRV::OpCaptureEventProfilingInfo:
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[1], &SPIRV::IDRegClass);
     MIRBuilder.getMRI()->setRegClass(Call->Arguments[2], &SPIRV::IDRegClass);
     return MIRBuilder.buildInstr(Opcode)
         .addUse(Call->Arguments[0])
         .addUse(Call->Arguments[1])
         .addUse(Call->Arguments[2]);
   case SPIRV::OpBuildNDRange:
-    return buildNDRange(Call, MIRBuilder, GR);
+    return buildNDRange(Call, MIRBuilder, GOR);
   case SPIRV::OpEnqueueKernel:
-    return buildEnqueueKernel(Call, MIRBuilder, GR);
+    return buildEnqueueKernel(Call, MIRBuilder, GOR);
   default:
     return false;
   }
 }
 
 static bool generateAsyncCopy(const SPIRV::IncomingCall *Call,
                               MachineIRBuilder &MIRBuilder,
-                              SPIRVGlobalRegistry *GR) {
+                              SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
-  auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR);
+  auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GOR);
 
   switch (Opcode) {
   case SPIRV::OpGroupAsyncCopy:
     return MIRBuilder.buildInstr(Opcode)
         .addDef(Call->ReturnRegister)
-        .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+        .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
         .addUse(Scope)
         .addUse(Call->Arguments[0])
         .addUse(Call->Arguments[1])
         .addUse(Call->Arguments[2])
-        .addUse(buildConstantIntReg(1, MIRBuilder, GR))
+        .addUse(buildConstantIntReg(1, MIRBuilder, GOR))
         .addUse(Call->Arguments[3]);
   case SPIRV::OpGroupWaitEvents:
     return MIRBuilder.buildInstr(Opcode)
         .addUse(Scope)
         .addUse(Call->Arguments[0])
         .addUse(Call->Arguments[1]);
   default:
     return false;
   }
 }
 
 static bool generateConvertInst(const StringRef DemangledCall,
                                 const SPIRV::IncomingCall *Call,
                                 MachineIRBuilder &MIRBuilder,
-                                SPIRVGlobalRegistry *GR) {
+                                SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the conversion builtin in the TableGen records.
   const SPIRV::ConvertBuiltin *Builtin =
       SPIRV::lookupConvertBuiltin(Call->Builtin->Name, Call->Builtin->Set);
 
   if (Builtin->IsSaturated)
     buildOpDecorate(Call->ReturnRegister, MIRBuilder,
                     SPIRV::Decoration::SaturatedConversion, {});
   if (Builtin->IsRounded)
     buildOpDecorate(Call->ReturnRegister, MIRBuilder,
                     SPIRV::Decoration::FPRoundingMode,
                     {(unsigned)Builtin->RoundingMode});
 
   unsigned Opcode = SPIRV::OpNop;
-  if (GR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
+  if (GOR->isScalarOrVectorOfType(Call->Arguments[0], SPIRV::OpTypeInt)) {
     // Int -> ...
-    if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
+    if (GOR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt)) {
       // Int -> Int
       if (Builtin->IsSaturated)
         Opcode = Builtin->IsDestinationSigned ? SPIRV::OpSatConvertUToS
                                               : SPIRV::OpSatConvertSToU;
       else
         Opcode = Builtin->IsDestinationSigned ? SPIRV::OpUConvert
                                               : SPIRV::OpSConvert;
-    } else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
+    } else if (GOR->isScalarOrVectorOfType(Call->ReturnRegister,
                                            SPIRV::OpTypeFloat)) {
       // Int -> Float
       bool IsSourceSigned =
           DemangledCall[DemangledCall.find_first_of('(') + 1] != 'u';
       Opcode = IsSourceSigned ? SPIRV::OpConvertSToF : SPIRV::OpConvertUToF;
     }
-  } else if (GR->isScalarOrVectorOfType(Call->Arguments[0],
+  } else if (GOR->isScalarOrVectorOfType(Call->Arguments[0],
                                          SPIRV::OpTypeFloat)) {
     // Float -> ...
-    if (GR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt))
+    if (GOR->isScalarOrVectorOfType(Call->ReturnRegister, SPIRV::OpTypeInt))
       // Float -> Int
       Opcode = Builtin->IsDestinationSigned ? SPIRV::OpConvertFToS
                                             : SPIRV::OpConvertFToU;
-    else if (GR->isScalarOrVectorOfType(Call->ReturnRegister,
+    else if (GOR->isScalarOrVectorOfType(Call->ReturnRegister,
                                          SPIRV::OpTypeFloat))
       // Float -> Float
       Opcode = SPIRV::OpFConvert;
   }
 
   assert(Opcode != SPIRV::OpNop &&
          "Conversion between the types not implemented!");
 
   MIRBuilder.buildInstr(Opcode)
       .addDef(Call->ReturnRegister)
-      .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+      .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
       .addUse(Call->Arguments[0]);
   return true;
 }
 
 static bool generateVectorLoadStoreInst(const SPIRV::IncomingCall *Call,
                                         MachineIRBuilder &MIRBuilder,
-                                        SPIRVGlobalRegistry *GR) {
+                                        SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the vector load/store builtin in the TableGen records.
   const SPIRV::VectorLoadStoreBuiltin *Builtin =
       SPIRV::lookupVectorLoadStoreBuiltin(Call->Builtin->Name,
                                           Call->Builtin->Set);
   // Build extended instruction.
   auto MIB =
       MIRBuilder.buildInstr(SPIRV::OpExtInst)
           .addDef(Call->ReturnRegister)
-          .addUse(GR->getSPIRVTypeID(Call->ReturnType))
+          .addUse(GOR->getSPIRVTypeID(Call->ReturnType))
           .addImm(static_cast<uint32_t>(SPIRV::InstructionSet::OpenCL_std))
           .addImm(Builtin->Number);
   for (auto Argument : Call->Arguments)
     MIB.addUse(Argument);
 
   // Rounding mode should be passed as a last argument in the MI for builtins
   // like "vstorea_halfn_r".
   if (Builtin->IsRounded)
     MIB.addImm(static_cast<uint32_t>(Builtin->RoundingMode));
   return true;
 }
 
 static bool generateLoadStoreInst(const SPIRV::IncomingCall *Call,
                                   MachineIRBuilder &MIRBuilder,
-                                  SPIRVGlobalRegistry *GR) {
+                                  SPIRVGlobalObjectRegistry *GOR) {
   // Lookup the instruction opcode in the TableGen records.
   const SPIRV::DemangledBuiltin *Builtin = Call->Builtin;
   unsigned Opcode =
       SPIRV::lookupNativeBuiltin(Builtin->Name, Builtin->Set)->Opcode;
   bool IsLoad = Opcode == SPIRV::OpLoad;
   // Build the instruction.
   auto MIB = MIRBuilder.buildInstr(Opcode);
   if (IsLoad) {
     MIB.addDef(Call->ReturnRegister);
-    MIB.addUse(GR->getSPIRVTypeID(Call->ReturnType));
+    MIB.addUse(GOR->getSPIRVTypeID(Call->ReturnType));
   }
   // Add a pointer to the value to load/store.
   MIB.addUse(Call->Arguments[0]);
   MachineRegisterInfo *MRI = MIRBuilder.getMRI();
   MRI->setRegClass(Call->Arguments[0], &SPIRV::IDRegClass);
   // Add a value to store.
   if (!IsLoad) {
     MIB.addUse(Call->Arguments[1]);

 /// Lowers a builtin funtion call using the provided \p DemangledCall skeleton
 /// and external instruction \p Set.
 namespace SPIRV {
 std::optional<bool> lowerBuiltin(const StringRef DemangledCall,
                                  SPIRV::InstructionSet::InstructionSet Set,
                                  MachineIRBuilder &MIRBuilder,
                                  const Register OrigRet, const Type *OrigRetTy,
                                  const SmallVectorImpl<Register> &Args,
-                                 SPIRVGlobalRegistry *GR) {
+                                 SPIRVGlobalObjectRegistry *GOR) {
   LLVM_DEBUG(dbgs() << "Lowering builtin call: " << DemangledCall << "\n");
 
   // SPIR-V type and return register.
   Register ReturnRegister = OrigRet;
   SPIRVType *ReturnType = nullptr;
   if (OrigRetTy && !OrigRetTy->isVoidTy()) {
-    ReturnType = GR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
+    ReturnType = GOR->assignTypeToVReg(OrigRetTy, OrigRet, MIRBuilder);
     if (!MIRBuilder.getMRI()->getRegClassOrNull(ReturnRegister))
       MIRBuilder.getMRI()->setRegClass(ReturnRegister, &SPIRV::IDRegClass);
   } else if (OrigRetTy && OrigRetTy->isVoidTy()) {
     ReturnRegister = MIRBuilder.getMRI()->createVirtualRegister(&IDRegClass);
     MIRBuilder.getMRI()->setType(ReturnRegister, LLT::scalar(32));
-    ReturnType = GR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
+    ReturnType = GOR->assignTypeToVReg(OrigRetTy, ReturnRegister, MIRBuilder);
   }
 
   // Lookup the builtin in the TableGen records.
   std::unique_ptr<const IncomingCall> Call =
       lookupBuiltin(DemangledCall, Set, ReturnRegister, ReturnType, Args);
 
   if (!Call) {
     LLVM_DEBUG(dbgs() << "Builtin record was not found!\n");
     return std::nullopt;
   }
 
   // TODO: check if the provided args meet the builtin requirments.
   assert(Args.size() >= Call->Builtin->MinNumArgs &&
          "Too few arguments to generate the builtin");
   if (Call->Builtin->MaxNumArgs && Args.size() > Call->Builtin->MaxNumArgs)
     LLVM_DEBUG(dbgs() << "More arguments provided than required!\n");
 
   // Match the builtin with implementation based on the grouping.
   switch (Call->Builtin->Group) {
   case SPIRV::Extended:
-    return generateExtInst(Call.get(), MIRBuilder, GR);
+    return generateExtInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::Relational:
-    return generateRelationalInst(Call.get(), MIRBuilder, GR);
+    return generateRelationalInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::Group:
-    return generateGroupInst(Call.get(), MIRBuilder, GR);
+    return generateGroupInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::Variable:
-    return generateBuiltinVar(Call.get(), MIRBuilder, GR);
+    return generateBuiltinVar(Call.get(), MIRBuilder, GOR);
   case SPIRV::Atomic:
-    return generateAtomicInst(Call.get(), MIRBuilder, GR);
+    return generateAtomicInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::Barrier:
-    return generateBarrierInst(Call.get(), MIRBuilder, GR);
+    return generateBarrierInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::Dot:
-    return generateDotOrFMulInst(Call.get(), MIRBuilder, GR);
+    return generateDotOrFMulInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::GetQuery:
-    return generateGetQueryInst(Call.get(), MIRBuilder, GR);
+    return generateGetQueryInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::ImageSizeQuery:
-    return generateImageSizeQueryInst(Call.get(), MIRBuilder, GR);
+    return generateImageSizeQueryInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::ImageMiscQuery:
-    return generateImageMiscQueryInst(Call.get(), MIRBuilder, GR);
+    return generateImageMiscQueryInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::ReadImage:
-    return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
+    return generateReadImageInst(DemangledCall, Call.get(), MIRBuilder, GOR);
   case SPIRV::WriteImage:
-    return generateWriteImageInst(Call.get(), MIRBuilder, GR);
+    return generateWriteImageInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::SampleImage:
-    return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GR);
+    return generateSampleImageInst(DemangledCall, Call.get(), MIRBuilder, GOR);
   case SPIRV::Select:
     return generateSelectInst(Call.get(), MIRBuilder);
   case SPIRV::SpecConstant:
-    return generateSpecConstantInst(Call.get(), MIRBuilder, GR);
+    return generateSpecConstantInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::Enqueue:
-    return generateEnqueueInst(Call.get(), MIRBuilder, GR);
+    return generateEnqueueInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::AsyncCopy:
-    return generateAsyncCopy(Call.get(), MIRBuilder, GR);
+    return generateAsyncCopy(Call.get(), MIRBuilder, GOR);
   case SPIRV::Convert:
-    return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GR);
+    return generateConvertInst(DemangledCall, Call.get(), MIRBuilder, GOR);
   case SPIRV::VectorLoadStore:
-    return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GR);
+    return generateVectorLoadStoreInst(Call.get(), MIRBuilder, GOR);
   case SPIRV::LoadStore:
-    return generateLoadStoreInst(Call.get(), MIRBuilder, GR);
+    return generateLoadStoreInst(Call.get(), MIRBuilder, GOR);
   }
   return false;
 }
 
 struct BuiltinType {
   StringRef Name;
   uint32_t Opcode;
 };

 
 //===----------------------------------------------------------------------===//
 // Implementation functions for builtin types.
 //===----------------------------------------------------------------------===//
 
 static SPIRVType *getNonParameterizedType(const TargetExtType *ExtensionType,
                                           const SPIRV::BuiltinType *TypeRecord,
                                           MachineIRBuilder &MIRBuilder,
-                                          SPIRVGlobalRegistry *GR) {
+                                          SPIRVGlobalObjectRegistry *GOR) {
   unsigned Opcode = TypeRecord->Opcode;
   // Create or get an existing type from GlobalRegistry.
-  return GR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode);
+  return GOR->getOrCreateOpTypeByOpcode(ExtensionType, MIRBuilder, Opcode);
 }
 
 static SPIRVType *getSamplerType(MachineIRBuilder &MIRBuilder,
-                                 SPIRVGlobalRegistry *GR) {
+                                 SPIRVGlobalObjectRegistry *GOR) {
   // Create or get an existing type from GlobalRegistry.
-  return GR->getOrCreateOpTypeSampler(MIRBuilder);
+  return GOR->getOrCreateOpTypeSampler(MIRBuilder);
 }
 
 static SPIRVType *getPipeType(const TargetExtType *ExtensionType,
                               MachineIRBuilder &MIRBuilder,
-                              SPIRVGlobalRegistry *GR) {
+                              SPIRVGlobalObjectRegistry *GOR) {
   assert(ExtensionType->getNumIntParameters() == 1 &&
          "Invalid number of parameters for SPIR-V pipe builtin!");
   // Create or get an existing type from GlobalRegistry.
-  return GR->getOrCreateOpTypePipe(MIRBuilder,
+  return GOR->getOrCreateOpTypePipe(MIRBuilder,
                                     SPIRV::AccessQualifier::AccessQualifier(
                                         ExtensionType->getIntParameter(0)));
 }
 
 static SPIRVType *
 getImageType(const TargetExtType *ExtensionType,
              const SPIRV::AccessQualifier::AccessQualifier Qualifier,
-             MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry *GR) {
+             MachineIRBuilder &MIRBuilder, SPIRVGlobalObjectRegistry *GOR) {
   assert(ExtensionType->getNumTypeParameters() == 1 &&
          "SPIR-V image builtin type must have sampled type parameter!");
   const SPIRVType *SampledType =
-      GR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder);
+      GOR->getOrCreateSPIRVType(ExtensionType->getTypeParameter(0), MIRBuilder);
   assert(ExtensionType->getNumIntParameters() == 7 &&
          "Invalid number of parameters for SPIR-V image builtin!");
   // Create or get an existing type from GlobalRegistry.
-  return GR->getOrCreateOpTypeImage(
+  return GOR->getOrCreateOpTypeImage(
       MIRBuilder, SampledType,
       SPIRV::Dim::Dim(ExtensionType->getIntParameter(0)),
       ExtensionType->getIntParameter(1), ExtensionType->getIntParameter(2),
       ExtensionType->getIntParameter(3), ExtensionType->getIntParameter(4),
       SPIRV::ImageFormat::ImageFormat(ExtensionType->getIntParameter(5)),
       Qualifier == SPIRV::AccessQualifier::WriteOnly
           ? SPIRV::AccessQualifier::WriteOnly
           : SPIRV::AccessQualifier::AccessQualifier(
                 ExtensionType->getIntParameter(6)));
 }
 
 static SPIRVType *getSampledImageType(const TargetExtType *OpaqueType,
                                       MachineIRBuilder &MIRBuilder,
-                                      SPIRVGlobalRegistry *GR) {
+                                      SPIRVGlobalObjectRegistry *GOR) {
   SPIRVType *OpaqueImageType = getImageType(
-      OpaqueType, SPIRV::AccessQualifier::ReadOnly, MIRBuilder, GR);
+      OpaqueType, SPIRV::AccessQualifier::ReadOnly, MIRBuilder, GOR);
   // Create or get an existing type from GlobalRegistry.
-  return GR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder);
+  return GOR->getOrCreateOpTypeSampledImage(OpaqueImageType, MIRBuilder);
 }
 
 namespace SPIRV {
 SPIRVType *lowerBuiltinType(const Type *OpaqueType,
                             SPIRV::AccessQualifier::AccessQualifier AccessQual,
                             MachineIRBuilder &MIRBuilder,
-                            SPIRVGlobalRegistry *GR) {
+                            SPIRVGlobalObjectRegistry *GOR) {
   // In LLVM IR, SPIR-V and OpenCL builtin types are represented as either
   // target(...) target extension types or pointers-to-opaque-structs. The
   // approach relying on structs is deprecated and works only in the non-opaque
   // pointer mode (-opaque-pointers=0).
   // In order to maintain compatibility with LLVM IR generated by older versions
   // of Clang and LLVM/SPIR-V Translator, the pointers-to-opaque-structs are
   // "translated" to target extension types. This translation is temporary and
   // will be removed in the future release of LLVM.
SPIRVType *lowerBuiltinType(const Type *OpaqueType,
 
   // "Lower" the BuiltinType into TargetType. The following get<...>Type methods
   // use the implementation details from TableGen records or TargetExtType
   // parameters to either create a new OpType<...> machine instruction or get an
   // existing equivalent SPIRVType from GlobalRegistry.
   SPIRVType *TargetType;
   switch (TypeRecord->Opcode) {
   case SPIRV::OpTypeImage:
-    TargetType = getImageType(BuiltinType, AccessQual, MIRBuilder, GR);
+    TargetType = getImageType(BuiltinType, AccessQual, MIRBuilder, GOR);
     break;
   case SPIRV::OpTypePipe:
-    TargetType = getPipeType(BuiltinType, MIRBuilder, GR);
+    TargetType = getPipeType(BuiltinType, MIRBuilder, GOR);
     break;
   case SPIRV::OpTypeDeviceEvent:
-    TargetType = GR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
+    TargetType = GOR->getOrCreateOpTypeDeviceEvent(MIRBuilder);
     break;
   case SPIRV::OpTypeSampler:
-    TargetType = getSamplerType(MIRBuilder, GR);
+    TargetType = getSamplerType(MIRBuilder, GOR);
     break;
   case SPIRV::OpTypeSampledImage:
-    TargetType = getSampledImageType(BuiltinType, MIRBuilder, GR);
+    TargetType = getSampledImageType(BuiltinType, MIRBuilder, GOR);
     break;
   default:
     TargetType =
-        getNonParameterizedType(BuiltinType, TypeRecord, MIRBuilder, GR);
+        getNonParameterizedType(BuiltinType, TypeRecord, MIRBuilder, GOR);
     break;
   }
 
   // Emit OpName instruction if a new OpType<...> instruction was added
   // (equivalent type was not found in GlobalRegistry).
   if (NumStartingVRegs < MIRBuilder.getMRI()->getNumVirtRegs())
-    buildOpName(GR->getSPIRVTypeID(TargetType), Name, MIRBuilder);
+    buildOpName(GOR->getSPIRVTypeID(TargetType), Name, MIRBuilder);
 
   return TargetType;
 }
 } // namespace SPIRV
 } // namespace llvm

llvm/lib/Target/SPIRV/SPIRVCallLowering.h

 //===--- SPIRVCallLowering.h - Call lowering --------------------*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file describes how to lower LLVM calls to machine code calls.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_SPIRV_SPIRVCALLLOWERING_H
 #define LLVM_LIB_TARGET_SPIRV_SPIRVCALLLOWERING_H
 
-#include "SPIRVGlobalRegistry.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "llvm/CodeGen/GlobalISel/CallLowering.h"
 
 namespace llvm {
 
-class SPIRVGlobalRegistry;
+class SPIRVGlobalObjectRegistry;
 class SPIRVTargetLowering;
 
 class SPIRVCallLowering : public CallLowering {
 private:
   // Used to create and assign function, argument, and return type information.
-  SPIRVGlobalRegistry *GR;
+  SPIRVGlobalObjectRegistry *GOR;
 
 public:
-  SPIRVCallLowering(const SPIRVTargetLowering &TLI, SPIRVGlobalRegistry *GR);
+  SPIRVCallLowering(const SPIRVTargetLowering &TLI, SPIRVGlobalObjectRegistry *GOR);
 
   // Built OpReturn or OpReturnValue.
   bool lowerReturn(MachineIRBuilder &MIRBuiler, const Value *Val,
                    ArrayRef<Register> VRegs, FunctionLoweringInfo &FLI,
                    Register SwiftErrorVReg) const override;
 
   // Build OpFunction, OpFunctionParameter, and any EntryPoint or Linkage data.
   bool lowerFormalArguments(MachineIRBuilder &MIRBuilder, const Function &F,

llvm/lib/Target/SPIRV/SPIRVCallLowering.cpp

 //===--- SPIRVCallLowering.cpp - Call lowering ------------------*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the lowering of LLVM calls to machine code calls for
 // GlobalISel.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIRVCallLowering.h"
 #include "MCTargetDesc/SPIRVBaseInfo.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "SPIRV.h"
 #include "SPIRVBuiltins.h"
-#include "SPIRVGlobalRegistry.h"
 #include "SPIRVISelLowering.h"
 #include "SPIRVRegisterInfo.h"
 #include "SPIRVSubtarget.h"
 #include "SPIRVUtils.h"
 #include "llvm/CodeGen/FunctionLoweringInfo.h"
 #include "llvm/Support/ModRef.h"
 
 using namespace llvm;
 
 SPIRVCallLowering::SPIRVCallLowering(const SPIRVTargetLowering &TLI,
-                                     SPIRVGlobalRegistry *GR)
-    : CallLowering(&TLI), GR(GR) {}
+                                     SPIRVGlobalObjectRegistry *GOR)
+    : CallLowering(&TLI), GOR(GOR) {}
 
 bool SPIRVCallLowering::lowerReturn(MachineIRBuilder &MIRBuilder,
                                     const Value *Val, ArrayRef<Register> VRegs,
                                     FunctionLoweringInfo &FLI,
                                     Register SwiftErrorVReg) const {
   // Currently all return types should use a single register.
   // TODO: handle the case of multiple registers.
   if (VRegs.size() > 1)
return ExistingOpaqueType
              ? ExistingOpaqueType
              : StructType::create(F.getContext(), KernelArgTypeStr);
 }
 
 bool SPIRVCallLowering::lowerFormalArguments(MachineIRBuilder &MIRBuilder,
                                              const Function &F,
                                              ArrayRef<ArrayRef<Register>> VRegs,
                                              FunctionLoweringInfo &FLI) const {
-  assert(GR && "Must initialize the SPIRV type registry before lowering args.");
-  GR->setCurrentFunc(MIRBuilder.getMF());
+  assert(GOR &&
+         "Must initialize the SPIRV type registry before lowering args.");
+  GOR->setCurrentFunc(MIRBuilder.getMF());
 
   // Assign types and names to all args, and store their types for later.
   FunctionType *FTy = getOriginalFunctionType(F);
   SmallVector<SPIRVType *, 4> ArgTypeVRegs;
   if (VRegs.size() > 0) {
     unsigned i = 0;
     for (const auto &Arg : F.args()) {
       // Currently formal args should use single registers.
       // TODO: handle the case of multiple registers.
       if (VRegs[i].size() > 1)
         return false;
       SPIRV::AccessQualifier::AccessQualifier ArgAccessQual =
           getArgAccessQual(F, i);
-      auto *SpirvTy = GR->assignTypeToVReg(getArgType(F, i), VRegs[i][0],
+      auto *SpirvTy = GOR->assignTypeToVReg(getArgType(F, i), VRegs[i][0],
                                             MIRBuilder, ArgAccessQual);
       ArgTypeVRegs.push_back(SpirvTy);
 
       if (Arg.hasName())
         buildOpName(VRegs[i][0], Arg.getName(), MIRBuilder);
       if (Arg.getType()->isPointerTy()) {
         auto DerefBytes = static_cast<unsigned>(Arg.getDereferenceableBytes());
         if (DerefBytes != 0)
           buildOpDecorate(VRegs[i][0], MIRBuilder,
if (VRegs.size() > 0) {
     }
   }
 
   // Generate a SPIR-V type for the function.
   auto MRI = MIRBuilder.getMRI();
   Register FuncVReg = MRI->createGenericVirtualRegister(LLT::scalar(32));
   MRI->setRegClass(FuncVReg, &SPIRV::IDRegClass);
   if (F.isDeclaration())
-    GR->add(&F, &MIRBuilder.getMF(), FuncVReg);
-  SPIRVType *RetTy = GR->getOrCreateSPIRVType(FTy->getReturnType(), MIRBuilder);
-  SPIRVType *FuncTy = GR->getOrCreateOpTypeFunctionWithArgs(
+    GOR->add(&F, &MIRBuilder.getMF(), FuncVReg);
+  SPIRVType *RetTy =
+      GOR->getOrCreateSPIRVType(FTy->getReturnType(), MIRBuilder);
+  SPIRVType *FuncTy = GOR->getOrCreateOpTypeFunctionWithArgs(
       FTy, RetTy, ArgTypeVRegs, MIRBuilder);
 
   // Build the OpTypeFunction declaring it.
   uint32_t FuncControl = getFunctionControl(F);
 
   MIRBuilder.buildInstr(SPIRV::OpFunction)
       .addDef(FuncVReg)
-      .addUse(GR->getSPIRVTypeID(RetTy))
+      .addUse(GOR->getSPIRVTypeID(RetTy))
       .addImm(FuncControl)
-      .addUse(GR->getSPIRVTypeID(FuncTy));
+      .addUse(GOR->getSPIRVTypeID(FuncTy));
 
   // Add OpFunctionParameters.
   int i = 0;
   for (const auto &Arg : F.args()) {
     assert(VRegs[i].size() == 1 && "Formal arg has multiple vregs");
     MRI->setRegClass(VRegs[i][0], &SPIRV::IDRegClass);
     MIRBuilder.buildInstr(SPIRV::OpFunctionParameter)
         .addDef(VRegs[i][0])
-        .addUse(GR->getSPIRVTypeID(ArgTypeVRegs[i]));
+        .addUse(GOR->getSPIRVTypeID(ArgTypeVRegs[i]));
     if (F.isDeclaration())
-      GR->add(&Arg, &MIRBuilder.getMF(), VRegs[i][0]);
+      GOR->add(&Arg, &MIRBuilder.getMF(), VRegs[i][0]);
     i++;
   }
   // Name the function.
   if (F.hasName())
     buildOpName(FuncVReg, F.getName(), MIRBuilder);
 
   // Handle entry points and function linkage.
   if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {

 
 bool SPIRVCallLowering::lowerCall(MachineIRBuilder &MIRBuilder,
                                   CallLoweringInfo &Info) const {
   // Currently call returns should have single vregs.
   // TODO: handle the case of multiple registers.
   if (Info.OrigRet.Regs.size() > 1)
     return false;
   MachineFunction &MF = MIRBuilder.getMF();
-  GR->setCurrentFunc(MF);
+  GOR->setCurrentFunc(MF);
   FunctionType *FTy = nullptr;
   const Function *CF = nullptr;
 
   // Emit a regular OpFunctionCall. If it's an externally declared function,
   // be sure to emit its type and function declaration here. It will be hoisted
   // globally later.
   if (Info.Callee.isGlobal()) {
     CF = dyn_cast_or_null<const Function>(Info.Callee.getGlobal());
if (!DemangledName.empty() && CF && CF->isDeclaration() &&
       ST->canUseExtInstSet(SPIRV::InstructionSet::OpenCL_std)) {
     const Type *OrigRetTy = Info.OrigRet.Ty;
     if (FTy)
       OrigRetTy = FTy->getReturnType();
     SmallVector<Register, 8> ArgVRegs;
     for (auto Arg : Info.OrigArgs) {
       assert(Arg.Regs.size() == 1 && "Call arg has multiple VRegs");
       ArgVRegs.push_back(Arg.Regs[0]);
-      SPIRVType *SPIRVTy = GR->getOrCreateSPIRVType(Arg.Ty, MIRBuilder);
-      GR->assignSPIRVTypeToVReg(SPIRVTy, Arg.Regs[0], MIRBuilder.getMF());
+      SPIRVType *SPIRVTy = GOR->getOrCreateSPIRVType(Arg.Ty, MIRBuilder);
+      GOR->assignSPIRVTypeToVReg(SPIRVTy, Arg.Regs[0], MIRBuilder.getMF());
     }
     if (auto Res = SPIRV::lowerBuiltin(
             DemangledName, SPIRV::InstructionSet::OpenCL_std, MIRBuilder,
-            ResVReg, OrigRetTy, ArgVRegs, GR))
+            ResVReg, OrigRetTy, ArgVRegs, GOR))
       return *Res;
   }
   if (CF && CF->isDeclaration() &&
-      !GR->find(CF, &MIRBuilder.getMF()).isValid()) {
+      !GOR->find(CF, &MIRBuilder.getMF()).isValid()) {
     // Emit the type info and forward function declaration to the first MBB
     // to ensure VReg definition dependencies are valid across all MBBs.
     MachineIRBuilder FirstBlockBuilder;
     FirstBlockBuilder.setMF(MF);
     FirstBlockBuilder.setMBB(*MF.getBlockNumbered(0));
 
     SmallVector<ArrayRef<Register>, 8> VRegArgs;
     SmallVector<SmallVector<Register, 1>, 8> ToInsert;
if (CF && CF->isDeclaration() &&
     FunctionLoweringInfo FuncInfo;
     lowerFormalArguments(FirstBlockBuilder, *CF, VRegArgs, FuncInfo);
   }
 
   // Make sure there's a valid return reg, even for functions returning void.
   if (!ResVReg.isValid())
     ResVReg = MIRBuilder.getMRI()->createVirtualRegister(&SPIRV::IDRegClass);
   SPIRVType *RetType =
-      GR->assignTypeToVReg(FTy->getReturnType(), ResVReg, MIRBuilder);
+      GOR->assignTypeToVReg(FTy->getReturnType(), ResVReg, MIRBuilder);
 
   // Emit the OpFunctionCall and its args.
   auto MIB = MIRBuilder.buildInstr(SPIRV::OpFunctionCall)
                  .addDef(ResVReg)
-                 .addUse(GR->getSPIRVTypeID(RetType))
+                 .addUse(GOR->getSPIRVTypeID(RetType))
                  .add(Info.Callee);
 
   for (const auto &Arg : Info.OrigArgs) {
     // Currently call args should have single vregs.
     if (Arg.Regs.size() > 1)
       return false;
     MIB.addUse(Arg.Regs[0]);
   }
   return MIB.constrainAllUses(MIRBuilder.getTII(), *ST->getRegisterInfo(),
                               *ST->getRegBankInfo());
 }

llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp

 //===- SPIRVInstructionSelector.cpp ------------------------------*- C++ -*-==//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the targeting of the InstructionSelector class for
 // SPIRV.
 // TODO: This should be generated by TableGen.
 //
 //===----------------------------------------------------------------------===//
 
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "SPIRV.h"
-#include "SPIRVGlobalRegistry.h"
 #include "SPIRVInstrInfo.h"
 #include "SPIRVRegisterBankInfo.h"
 #include "SPIRVRegisterInfo.h"
 #include "SPIRVTargetMachine.h"
 #include "SPIRVUtils.h"
 #include "llvm/ADT/APFloat.h"
 #include "llvm/CodeGen/GlobalISel/InstructionSelector.h"
 #include "llvm/CodeGen/GlobalISel/InstructionSelectorImpl.h"

 #include "SPIRVGenGlobalISel.inc"
 #undef GET_GLOBALISEL_PREDICATE_BITSET
 
 class SPIRVInstructionSelector : public InstructionSelector {
   const SPIRVSubtarget &STI;
   const SPIRVInstrInfo &TII;
   const SPIRVRegisterInfo &TRI;
   const RegisterBankInfo &RBI;
-  SPIRVGlobalRegistry &GR;
+  SPIRVGlobalObjectRegistry &GOR;
   MachineRegisterInfo *MRI;
 
 public:
   SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                            const SPIRVSubtarget &ST,
                            const RegisterBankInfo &RBI);
   void setupMF(MachineFunction &MF, GISelKnownBits *KB,
                CodeGenCoverage &CoverageInfo, ProfileSummaryInfo *PSI,

 #define GET_GLOBALISEL_IMPL
 #include "SPIRVGenGlobalISel.inc"
 #undef GET_GLOBALISEL_IMPL
 
 SPIRVInstructionSelector::SPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                                                    const SPIRVSubtarget &ST,
                                                    const RegisterBankInfo &RBI)
     : InstructionSelector(), STI(ST), TII(*ST.getInstrInfo()),
-      TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
+      TRI(*ST.getRegisterInfo()), RBI(RBI),
+      GOR(*ST.getSPIRVGlobalObjectRegistry()),
 #define GET_GLOBALISEL_PREDICATES_INIT
 #include "SPIRVGenGlobalISel.inc"
 #undef GET_GLOBALISEL_PREDICATES_INIT
 #define GET_GLOBALISEL_TEMPORARIES_INIT
 #include "SPIRVGenGlobalISel.inc"
 #undef GET_GLOBALISEL_TEMPORARIES_INIT
 {
 }
 
 void SPIRVInstructionSelector::setupMF(MachineFunction &MF, GISelKnownBits *KB,
                                        CodeGenCoverage &CoverageInfo,
                                        ProfileSummaryInfo *PSI,
                                        BlockFrequencyInfo *BFI) {
   MRI = &MF.getRegInfo();
-  GR.setCurrentFunc(MF);
+  GOR.setCurrentFunc(MF);
   InstructionSelector::setupMF(MF, KB, CoverageInfo, PSI, BFI);
 }
 
 static bool isImm(const MachineOperand &MO, MachineRegisterInfo *MRI);
 
 // Defined in SPIRVLegalizerInfo.cpp.
 extern bool isTypeFoldingSupported(unsigned Opcode);
 
if (I.getNumOperands() != I.getNumExplicitOperands()) {
     LLVM_DEBUG(errs() << "Generic instr has unexpected implicit operands\n");
     return false;
   }
 
   // Common code for getting return reg+type, and removing selected instr
   // from parent occurs here. Instr-specific selection happens in spvSelect().
   bool HasDefs = I.getNumDefs() > 0;
   Register ResVReg = HasDefs ? I.getOperand(0).getReg() : Register(0);
-  SPIRVType *ResType = HasDefs ? GR.getSPIRVTypeForVReg(ResVReg) : nullptr;
+  SPIRVType *ResType = HasDefs ? GOR.getSPIRVTypeForVReg(ResVReg) : nullptr;
   assert(!HasDefs || ResType || I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE);
   if (spvSelect(ResVReg, ResType, I)) {
     if (HasDefs) // Make all vregs 32 bits (for SPIR-V IDs).
       MRI->setType(ResVReg, LLT::scalar(32));
     I.removeFromParent();
     return true;
   }
   return false;
bool SPIRVInstructionSelector::spvSelect(Register ResVReg,
 
   case TargetOpcode::G_BUILD_VECTOR:
     return selectConstVector(ResVReg, ResType, I);
 
   case TargetOpcode::G_SHUFFLE_VECTOR: {
     MachineBasicBlock &BB = *I.getParent();
     auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorShuffle))
                    .addDef(ResVReg)
-                   .addUse(GR.getSPIRVTypeID(ResType))
+                   .addUse(GOR.getSPIRVTypeID(ResType))
                    .addUse(I.getOperand(1).getReg())
                    .addUse(I.getOperand(2).getReg());
     for (auto V : I.getOperand(3).getShuffleMask())
       MIB.addImm(V);
     return MIB.constrainAllUses(TII, TRI, RBI);
   }
   case TargetOpcode::G_MEMMOVE:
   case TargetOpcode::G_MEMCPY:
case TargetOpcode::G_PTR_ADD: {
     // TODO: extend the handler once we have other cases.
     assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
     Register GV = I.getOperand(1).getReg();
     MachineRegisterInfo::def_instr_iterator II = MRI->def_instr_begin(GV);
     assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
             (*II).getOpcode() == TargetOpcode::COPY ||
             (*II).getOpcode() == SPIRV::OpVariable) &&
            isImm(I.getOperand(2), MRI));
-    Register Idx = buildZerosVal(GR.getOrCreateSPIRVIntegerType(32, I, TII), I);
+    Register Idx =
+        buildZerosVal(GOR.getOrCreateSPIRVIntegerType(32, I, TII), I);
     MachineBasicBlock &BB = *I.getParent();
     auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
                    .addDef(ResVReg)
-                   .addUse(GR.getSPIRVTypeID(ResType))
+                   .addUse(GOR.getSPIRVTypeID(ResType))
                    .addImm(static_cast<uint32_t>(
                        SPIRV::Opcode::InBoundsPtrAccessChain))
                    .addUse(GV)
                    .addUse(Idx)
                    .addUse(I.getOperand(2).getReg());
     return MIB.constrainAllUses(TII, TRI, RBI);
   }
 
bool SPIRVInstructionSelector::selectExtInst(Register ResVReg,
 
   for (const auto &Ex : Insts) {
     SPIRV::InstructionSet::InstructionSet Set = Ex.first;
     uint32_t Opcode = Ex.second;
     if (STI.canUseExtInstSet(Set)) {
       MachineBasicBlock &BB = *I.getParent();
       auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpExtInst))
                      .addDef(ResVReg)
-                     .addUse(GR.getSPIRVTypeID(ResType))
+                     .addUse(GOR.getSPIRVTypeID(ResType))
                      .addImm(static_cast<uint32_t>(Set))
                      .addImm(Opcode);
       const unsigned NumOps = I.getNumOperands();
       for (unsigned i = 1; i < NumOps; ++i)
         MIB.add(I.getOperand(i));
       return MIB.constrainAllUses(TII, TRI, RBI);
     }
   }
   return false;
 }
 
 bool SPIRVInstructionSelector::selectUnOpWithSrc(Register ResVReg,
                                                  const SPIRVType *ResType,
                                                  MachineInstr &I,
                                                  Register SrcReg,
                                                  unsigned Opcode) const {
   return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(SrcReg)
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectUnOp(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I,
                                           unsigned Opcode) const {

 bool SPIRVInstructionSelector::selectLoad(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I) const {
   unsigned OpOffset =
       I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ? 1 : 0;
   Register Ptr = I.getOperand(1 + OpOffset).getReg();
   auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpLoad))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType))
+                 .addUse(GOR.getSPIRVTypeID(ResType))
                  .addUse(Ptr);
   if (!I.getNumMemOperands()) {
     assert(I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS);
     addMemoryOperands(I.getOperand(2 + OpOffset).getImm(), MIB);
   } else {
     addMemoryOperands(*I.memoperands_begin(), MIB);
   }
   return MIB.constrainAllUses(TII, TRI, RBI);

 bool SPIRVInstructionSelector::selectMemOperation(Register ResVReg,
                                                   MachineInstr &I) const {
   MachineBasicBlock &BB = *I.getParent();
   Register SrcReg = I.getOperand(1).getReg();
   if (I.getOpcode() == TargetOpcode::G_MEMSET) {
     assert(I.getOperand(1).isReg() && I.getOperand(2).isReg());
     unsigned Val = getIConstVal(I.getOperand(1).getReg(), MRI);
     unsigned Num = getIConstVal(I.getOperand(2).getReg(), MRI);
-    SPIRVType *ValTy = GR.getOrCreateSPIRVIntegerType(8, I, TII);
-    SPIRVType *ArrTy = GR.getOrCreateSPIRVArrayType(ValTy, Num, I, TII);
-    Register Const = GR.getOrCreateConsIntArray(Val, I, ArrTy, TII);
-    SPIRVType *VarTy = GR.getOrCreateSPIRVPointerType(
+    SPIRVType *ValTy = GOR.getOrCreateSPIRVIntegerType(8, I, TII);
+    SPIRVType *ArrTy = GOR.getOrCreateSPIRVArrayType(ValTy, Num, I, TII);
+    Register Const = GOR.getOrCreateConsIntArray(Val, I, ArrTy, TII);
+    SPIRVType *VarTy = GOR.getOrCreateSPIRVPointerType(
         ArrTy, I, TII, SPIRV::StorageClass::UniformConstant);
     // TODO: check if we have such GV, add init, use buildGlobalVariable.
     Type *LLVMArrTy = ArrayType::get(
-        IntegerType::get(GR.CurMF->getFunction().getContext(), 8), Num);
+        IntegerType::get(GOR.CurMF->getFunction().getContext(), 8), Num);
     GlobalVariable *GV =
         new GlobalVariable(LLVMArrTy, true, GlobalValue::InternalLinkage);
     Register VarReg = MRI->createGenericVirtualRegister(LLT::scalar(32));
-    GR.add(GV, GR.CurMF, VarReg);
+    GOR.add(GV, GOR.CurMF, VarReg);
 
     buildOpDecorate(VarReg, I, TII, SPIRV::Decoration::Constant, {});
     BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
         .addDef(VarReg)
-        .addUse(GR.getSPIRVTypeID(VarTy))
+        .addUse(GOR.getSPIRVTypeID(VarTy))
         .addImm(SPIRV::StorageClass::UniformConstant)
         .addUse(Const)
         .constrainAllUses(TII, TRI, RBI);
-    SPIRVType *SourceTy = GR.getOrCreateSPIRVPointerType(
+    SPIRVType *SourceTy = GOR.getOrCreateSPIRVPointerType(
         ValTy, I, TII, SPIRV::StorageClass::UniformConstant);
     SrcReg = MRI->createGenericVirtualRegister(LLT::scalar(32));
     selectUnOpWithSrc(SrcReg, SourceTy, I, VarReg, SPIRV::OpBitcast);
   }
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCopyMemorySized))
                  .addUse(I.getOperand(0).getReg())
                  .addUse(SrcReg)
                  .addUse(I.getOperand(2).getReg());
bool SPIRVInstructionSelector::selectAtomicRMW(Register ResVReg,
   assert(I.hasOneMemOperand());
   const MachineMemOperand *MemOp = *I.memoperands_begin();
   uint32_t Scope = static_cast<uint32_t>(getScope(MemOp->getSyncScopeID()));
   Register ScopeReg = buildI32Constant(Scope, I);
 
   Register Ptr = I.getOperand(1).getReg();
   // TODO: Changed as it's implemented in the translator. See test/atomicrmw.ll
   // auto ScSem =
-  // getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr));
+  // getMemSemanticsForStorageClass(GOR.getPointerStorageClass(Ptr));
   AtomicOrdering AO = MemOp->getSuccessOrdering();
   uint32_t MemSem = static_cast<uint32_t>(getMemSemantics(AO));
   Register MemSemReg = buildI32Constant(MemSem /*| ScSem*/, I);
 
   return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(NewOpcode))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(Ptr)
       .addUse(ScopeReg)
       .addUse(MemSemReg)
       .addUse(I.getOperand(2).getReg())
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectFence(MachineInstr &I) const {
bool SPIRVInstructionSelector::selectAtomicCmpXchg(Register ResVReg,
   Register Ptr = I.getOperand(2).getReg();
   if (I.getOpcode() != TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS) {
     assert(I.hasOneMemOperand());
     const MachineMemOperand *MemOp = *I.memoperands_begin();
     unsigned Scope = static_cast<uint32_t>(getScope(MemOp->getSyncScopeID()));
     ScopeReg = buildI32Constant(Scope, I);
 
     unsigned ScSem = static_cast<uint32_t>(
-        getMemSemanticsForStorageClass(GR.getPointerStorageClass(Ptr)));
+        getMemSemanticsForStorageClass(GOR.getPointerStorageClass(Ptr)));
     AtomicOrdering AO = MemOp->getSuccessOrdering();
     unsigned MemSemEq = static_cast<uint32_t>(getMemSemantics(AO)) | ScSem;
     MemSemEqReg = buildI32Constant(MemSemEq, I);
     AtomicOrdering FO = MemOp->getFailureOrdering();
     unsigned MemSemNeq = static_cast<uint32_t>(getMemSemantics(FO)) | ScSem;
     MemSemNeqReg =
         MemSemEq == MemSemNeq ? MemSemEqReg : buildI32Constant(MemSemNeq, I);
   } else {
     ScopeReg = I.getOperand(5).getReg();
     MemSemEqReg = I.getOperand(6).getReg();
     MemSemNeqReg = I.getOperand(7).getReg();
   }
 
   Register Cmp = I.getOperand(3).getReg();
   Register Val = I.getOperand(4).getReg();
-  SPIRVType *SpvValTy = GR.getSPIRVTypeForVReg(Val);
+  SPIRVType *SpvValTy = GOR.getSPIRVTypeForVReg(Val);
   Register ACmpRes = MRI->createVirtualRegister(&SPIRV::IDRegClass);
   const DebugLoc &DL = I.getDebugLoc();
   bool Result =
       BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpAtomicCompareExchange))
           .addDef(ACmpRes)
-          .addUse(GR.getSPIRVTypeID(SpvValTy))
+          .addUse(GOR.getSPIRVTypeID(SpvValTy))
           .addUse(Ptr)
           .addUse(ScopeReg)
           .addUse(MemSemEqReg)
           .addUse(MemSemNeqReg)
           .addUse(Val)
           .addUse(Cmp)
           .constrainAllUses(TII, TRI, RBI);
   Register CmpSuccReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
-  SPIRVType *BoolTy = GR.getOrCreateSPIRVBoolType(I, TII);
+  SPIRVType *BoolTy = GOR.getOrCreateSPIRVBoolType(I, TII);
   Result |= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpIEqual))
                 .addDef(CmpSuccReg)
-                .addUse(GR.getSPIRVTypeID(BoolTy))
+                .addUse(GOR.getSPIRVTypeID(BoolTy))
                 .addUse(ACmpRes)
                 .addUse(Cmp)
                 .constrainAllUses(TII, TRI, RBI);
   Register TmpReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
   Result |= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
                 .addDef(TmpReg)
-                .addUse(GR.getSPIRVTypeID(ResType))
+                .addUse(GOR.getSPIRVTypeID(ResType))
                 .addUse(ACmpRes)
-                .addUse(GR.getOrCreateUndef(I, ResType, TII))
+                .addUse(GOR.getOrCreateUndef(I, ResType, TII))
                 .addImm(0)
                 .constrainAllUses(TII, TRI, RBI);
   Result |= BuildMI(*I.getParent(), I, DL, TII.get(SPIRV::OpCompositeInsert))
                 .addDef(ResVReg)
-                .addUse(GR.getSPIRVTypeID(ResType))
+                .addUse(GOR.getSPIRVTypeID(ResType))
                 .addUse(CmpSuccReg)
                 .addUse(TmpReg)
                 .addImm(1)
                 .constrainAllUses(TII, TRI, RBI);
   return Result;
 }
 
 static bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC) {
bool SPIRVInstructionSelector::selectAddrSpaceCast(Register ResVReg,
   // OpVariable, we should select OpSpecConstantOp.
   auto UIs = MRI->use_instructions(ResVReg);
   if (!UIs.empty() && ++UIs.begin() == UIs.end() &&
       (UIs.begin()->getOpcode() == SPIRV::OpConstantComposite ||
        UIs.begin()->getOpcode() == SPIRV::OpVariable ||
        isSpvIntrinsic(*UIs.begin(), Intrinsic::spv_init_global))) {
     Register NewReg = I.getOperand(1).getReg();
     MachineBasicBlock &BB = *I.getParent();
-    SPIRVType *SpvBaseTy = GR.getOrCreateSPIRVIntegerType(8, I, TII);
-    ResType = GR.getOrCreateSPIRVPointerType(SpvBaseTy, I, TII,
+    SPIRVType *SpvBaseTy = GOR.getOrCreateSPIRVIntegerType(8, I, TII);
+    ResType = GOR.getOrCreateSPIRVPointerType(SpvBaseTy, I, TII,
                                               SPIRV::StorageClass::Generic);
     bool Result =
         BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpSpecConstantOp))
             .addDef(ResVReg)
-            .addUse(GR.getSPIRVTypeID(ResType))
+            .addUse(GOR.getSPIRVTypeID(ResType))
             .addImm(static_cast<uint32_t>(SPIRV::Opcode::PtrCastToGeneric))
             .addUse(NewReg)
             .constrainAllUses(TII, TRI, RBI);
     return Result;
   }
   Register SrcPtr = I.getOperand(1).getReg();
-  SPIRVType *SrcPtrTy = GR.getSPIRVTypeForVReg(SrcPtr);
-  SPIRV::StorageClass::StorageClass SrcSC = GR.getPointerStorageClass(SrcPtr);
-  SPIRV::StorageClass::StorageClass DstSC = GR.getPointerStorageClass(ResVReg);
+  SPIRVType *SrcPtrTy = GOR.getSPIRVTypeForVReg(SrcPtr);
+  SPIRV::StorageClass::StorageClass SrcSC = GOR.getPointerStorageClass(SrcPtr);
+  SPIRV::StorageClass::StorageClass DstSC = GOR.getPointerStorageClass(ResVReg);
 
   // Casting from an eligable pointer to Generic.
   if (DstSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(SrcSC))
     return selectUnOp(ResVReg, ResType, I, SPIRV::OpPtrCastToGeneric);
   // Casting from Generic to an eligable pointer.
   if (SrcSC == SPIRV::StorageClass::Generic && isGenericCastablePtr(DstSC))
     return selectUnOp(ResVReg, ResType, I, SPIRV::OpGenericCastToPtr);
   // Casting between 2 eligable pointers using Generic as an intermediary.
   if (isGenericCastablePtr(SrcSC) && isGenericCastablePtr(DstSC)) {
     Register Tmp = MRI->createVirtualRegister(&SPIRV::IDRegClass);
-    SPIRVType *GenericPtrTy = GR.getOrCreateSPIRVPointerType(
+    SPIRVType *GenericPtrTy = GOR.getOrCreateSPIRVPointerType(
         SrcPtrTy, I, TII, SPIRV::StorageClass::Generic);
     MachineBasicBlock &BB = *I.getParent();
     const DebugLoc &DL = I.getDebugLoc();
     bool Success = BuildMI(BB, I, DL, TII.get(SPIRV::OpPtrCastToGeneric))
                        .addDef(Tmp)
-                       .addUse(GR.getSPIRVTypeID(GenericPtrTy))
+                       .addUse(GOR.getSPIRVTypeID(GenericPtrTy))
                        .addUse(SrcPtr)
                        .constrainAllUses(TII, TRI, RBI);
     return Success && BuildMI(BB, I, DL, TII.get(SPIRV::OpGenericCastToPtr))
                           .addDef(ResVReg)
-                          .addUse(GR.getSPIRVTypeID(ResType))
+                          .addUse(GOR.getSPIRVTypeID(ResType))
                           .addUse(Tmp)
                           .constrainAllUses(TII, TRI, RBI);
   }
   // TODO Should this case just be disallowed completely?
   // We're casting 2 other arbitrary address spaces, so have to bitcast.
   return selectUnOp(ResVReg, ResType, I, SPIRV::OpBitcast);
 }
 

 }
 
 bool SPIRVInstructionSelector::selectBitreverse(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
   MachineBasicBlock &BB = *I.getParent();
   return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpBitReverse))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(I.getOperand(1).getReg())
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectConstVector(Register ResVReg,
                                                  const SPIRVType *ResType,
                                                  MachineInstr &I) const {
   // TODO: only const case is supported for now.
assert(std::all_of(
         assert(Const);
         return (Const->getOpcode() == TargetOpcode::G_CONSTANT ||
                 Const->getOpcode() == TargetOpcode::G_FCONSTANT);
       }));
 
   auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(),
                      TII.get(SPIRV::OpConstantComposite))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType));
+                 .addUse(GOR.getSPIRVTypeID(ResType));
   for (unsigned i = I.getNumExplicitDefs(); i < I.getNumExplicitOperands(); ++i)
     MIB.addUse(I.getOperand(i).getReg());
   return MIB.constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectCmp(Register ResVReg,
                                          const SPIRVType *ResType,
                                          unsigned CmpOpc,
                                          MachineInstr &I) const {
   Register Cmp0 = I.getOperand(2).getReg();
   Register Cmp1 = I.getOperand(3).getReg();
-  assert(GR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==
-             GR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&
+  assert(GOR.getSPIRVTypeForVReg(Cmp0)->getOpcode() ==
+             GOR.getSPIRVTypeForVReg(Cmp1)->getOpcode() &&
          "CMP operands should have the same type");
   return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(CmpOpc))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(Cmp0)
       .addUse(Cmp1)
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectICmp(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I) const {
   auto Pred = I.getOperand(1).getPredicate();
   unsigned CmpOpc;
 
   Register CmpOperand = I.getOperand(2).getReg();
-  if (GR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))
+  if (GOR.isScalarOfType(CmpOperand, SPIRV::OpTypePointer))
     CmpOpc = getPtrCmpOpcode(Pred);
-  else if (GR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))
+  else if (GOR.isScalarOrVectorOfType(CmpOperand, SPIRV::OpTypeBool))
     CmpOpc = getBoolCmpOpcode(Pred);
   else
     CmpOpc = getICmpOpcode(Pred);
   return selectCmp(ResVReg, ResType, CmpOpc, I);
 }
 
 void SPIRVInstructionSelector::renderFImm32(MachineInstrBuilder &MIB,
                                             const MachineInstr &I,
void SPIRVInstructionSelector::renderImm32(MachineInstrBuilder &MIB,
   assert(I.getOpcode() == TargetOpcode::G_CONSTANT && OpIdx == -1 &&
          "Expected G_CONSTANT");
   addNumImm(I.getOperand(1).getCImm()->getValue(), MIB);
 }
 
 Register
 SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &I,
                                            const SPIRVType *ResType) const {
-  Type *LLVMTy = IntegerType::get(GR.CurMF->getFunction().getContext(), 32);
+  Type *LLVMTy = IntegerType::get(GOR.CurMF->getFunction().getContext(), 32);
   const SPIRVType *SpvI32Ty =
-      ResType ? ResType : GR.getOrCreateSPIRVIntegerType(32, I, TII);
+      ResType ? ResType : GOR.getOrCreateSPIRVIntegerType(32, I, TII);
   // Find a constant in DT or build a new one.
   auto ConstInt = ConstantInt::get(LLVMTy, Val);
-  Register NewReg = GR.find(ConstInt, GR.CurMF);
+  Register NewReg = GOR.find(ConstInt, GOR.CurMF);
   if (!NewReg.isValid()) {
     NewReg = MRI->createGenericVirtualRegister(LLT::scalar(32));
-    GR.add(ConstInt, GR.CurMF, NewReg);
+    GOR.add(ConstInt, GOR.CurMF, NewReg);
     MachineInstr *MI;
     MachineBasicBlock &BB = *I.getParent();
     if (Val == 0) {
       MI = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
                .addDef(NewReg)
-               .addUse(GR.getSPIRVTypeID(SpvI32Ty));
+               .addUse(GOR.getSPIRVTypeID(SpvI32Ty));
     } else {
       MI = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))
                .addDef(NewReg)
-               .addUse(GR.getSPIRVTypeID(SpvI32Ty))
+               .addUse(GOR.getSPIRVTypeID(SpvI32Ty))
                .addImm(APInt(32, Val).getZExtValue());
     }
     constrainSelectedInstRegOperands(*MI, TII, TRI, RBI);
   }
   return NewReg;
 }
 
 bool SPIRVInstructionSelector::selectFCmp(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I) const {
   unsigned CmpOp = getFCmpOpcode(I.getOperand(1).getPredicate());
   return selectCmp(ResVReg, ResType, CmpOp, I);
 }
 
 Register SPIRVInstructionSelector::buildZerosVal(const SPIRVType *ResType,
                                                  MachineInstr &I) const {
   if (ResType->getOpcode() == SPIRV::OpTypeVector)
-    return GR.getOrCreateConsIntVector(0, I, ResType, TII);
-  return GR.getOrCreateConstInt(0, I, ResType, TII);
+    return GOR.getOrCreateConsIntVector(0, I, ResType, TII);
+  return GOR.getOrCreateConstInt(0, I, ResType, TII);
 }
 
 Register SPIRVInstructionSelector::buildOnesVal(bool AllOnes,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
-  unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
+  unsigned BitWidth = GOR.getScalarOrVectorBitWidth(ResType);
   APInt One =
       AllOnes ? APInt::getAllOnes(BitWidth) : APInt::getOneBitSet(BitWidth, 0);
   if (ResType->getOpcode() == SPIRV::OpTypeVector)
-    return GR.getOrCreateConsIntVector(One.getZExtValue(), I, ResType, TII);
-  return GR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);
+    return GOR.getOrCreateConsIntVector(One.getZExtValue(), I, ResType, TII);
+  return GOR.getOrCreateConstInt(One.getZExtValue(), I, ResType, TII);
 }
 
 bool SPIRVInstructionSelector::selectSelect(Register ResVReg,
                                             const SPIRVType *ResType,
                                             MachineInstr &I,
                                             bool IsSigned) const {
   // To extend a bool, we need to use OpSelect between constants.
   Register ZeroReg = buildZerosVal(ResType, I);
   Register OneReg = buildOnesVal(IsSigned, ResType, I);
   bool IsScalarBool =
-      GR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
+      GOR.isScalarOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool);
   unsigned Opcode =
       IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectSIVCond;
   return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(I.getOperand(1).getReg())
       .addUse(OneReg)
       .addUse(ZeroReg)
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectIToF(Register ResVReg,
                                           const SPIRVType *ResType,
                                           MachineInstr &I, bool IsSigned,
                                           unsigned Opcode) const {
   Register SrcReg = I.getOperand(1).getReg();
   // We can convert bool value directly to float type without OpConvert*ToF,
   // however the translator generates OpSelect+OpConvert*ToF, so we do the same.
-  if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
-    unsigned BitWidth = GR.getScalarOrVectorBitWidth(ResType);
-    SPIRVType *TmpType = GR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
+  if (GOR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool)) {
+    unsigned BitWidth = GOR.getScalarOrVectorBitWidth(ResType);
+    SPIRVType *TmpType = GOR.getOrCreateSPIRVIntegerType(BitWidth, I, TII);
     if (ResType->getOpcode() == SPIRV::OpTypeVector) {
       const unsigned NumElts = ResType->getOperand(2).getImm();
-      TmpType = GR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);
+      TmpType = GOR.getOrCreateSPIRVVectorType(TmpType, NumElts, I, TII);
     }
     SrcReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
     selectSelect(SrcReg, TmpType, I, false);
   }
   return selectUnOpWithSrc(ResVReg, ResType, I, SrcReg, Opcode);
 }
 
 bool SPIRVInstructionSelector::selectExt(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I, bool IsSigned) const {
-  if (GR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool))
+  if (GOR.isScalarOrVectorOfType(I.getOperand(1).getReg(), SPIRV::OpTypeBool))
     return selectSelect(ResVReg, ResType, I, IsSigned);
   unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
   return selectUnOp(ResVReg, ResType, I, Opcode);
 }
 
 bool SPIRVInstructionSelector::selectIntToBool(Register IntReg,
                                                Register ResVReg,
                                                MachineInstr &I,
                                                const SPIRVType *IntTy,
                                                const SPIRVType *BoolTy) const {
   // To truncate to a bool, we use OpBitwiseAnd 1 and OpINotEqual to zero.
   Register BitIntReg = MRI->createVirtualRegister(&SPIRV::IDRegClass);
   bool IsVectorTy = IntTy->getOpcode() == SPIRV::OpTypeVector;
   unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
   Register Zero = buildZerosVal(IntTy, I);
   Register One = buildOnesVal(false, IntTy, I);
   MachineBasicBlock &BB = *I.getParent();
   BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
       .addDef(BitIntReg)
-      .addUse(GR.getSPIRVTypeID(IntTy))
+      .addUse(GOR.getSPIRVTypeID(IntTy))
       .addUse(IntReg)
       .addUse(One)
       .constrainAllUses(TII, TRI, RBI);
   return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpINotEqual))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(BoolTy))
+      .addUse(GOR.getSPIRVTypeID(BoolTy))
       .addUse(BitIntReg)
       .addUse(Zero)
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectTrunc(Register ResVReg,
                                            const SPIRVType *ResType,
                                            MachineInstr &I) const {
-  if (GR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool)) {
+  if (GOR.isScalarOrVectorOfType(ResVReg, SPIRV::OpTypeBool)) {
     Register IntReg = I.getOperand(1).getReg();
-    const SPIRVType *ArgType = GR.getSPIRVTypeForVReg(IntReg);
+    const SPIRVType *ArgType = GOR.getSPIRVTypeForVReg(IntReg);
     return selectIntToBool(IntReg, ResVReg, I, ArgType, ResType);
   }
-  bool IsSigned = GR.isScalarOrVectorSigned(ResType);
+  bool IsSigned = GOR.isScalarOrVectorSigned(ResType);
   unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
   return selectUnOp(ResVReg, ResType, I, Opcode);
 }
 
 bool SPIRVInstructionSelector::selectConst(Register ResVReg,
                                            const SPIRVType *ResType,
                                            const APInt &Imm,
                                            MachineInstr &I) const {
   unsigned TyOpcode = ResType->getOpcode();
   assert(TyOpcode != SPIRV::OpTypePointer || Imm.isZero());
   MachineBasicBlock &BB = *I.getParent();
   if ((TyOpcode == SPIRV::OpTypePointer || TyOpcode == SPIRV::OpTypeEvent) &&
       Imm.isZero())
     return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
         .addDef(ResVReg)
-        .addUse(GR.getSPIRVTypeID(ResType))
+        .addUse(GOR.getSPIRVTypeID(ResType))
         .constrainAllUses(TII, TRI, RBI);
   if (TyOpcode == SPIRV::OpTypeInt) {
     assert(Imm.getBitWidth() <= 64 && "Unsupported integer width!");
-    Register Reg = GR.getOrCreateConstInt(Imm.getZExtValue(), I, ResType, TII);
+    Register Reg = GOR.getOrCreateConstInt(Imm.getZExtValue(), I, ResType, TII);
     if (Reg == ResVReg)
       return true;
     return BuildMI(BB, I, I.getDebugLoc(), TII.get(TargetOpcode::COPY))
         .addDef(ResVReg)
         .addUse(Reg)
         .constrainAllUses(TII, TRI, RBI);
   }
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantI))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType));
+                 .addUse(GOR.getSPIRVTypeID(ResType));
   // <=32-bit integers should be caught by the sdag pattern.
   assert(Imm.getBitWidth() > 32);
   addNumImm(Imm, MIB);
   return MIB.constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectOpUndef(Register ResVReg,
                                              const SPIRVType *ResType,
                                              MachineInstr &I) const {
   return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .constrainAllUses(TII, TRI, RBI);
 }
 
 static bool isImm(const MachineOperand &MO, MachineRegisterInfo *MRI) {
   assert(MO.isReg());
   const SPIRVType *TypeInst = MRI->getVRegDef(MO.getReg());
   if (TypeInst->getOpcode() != SPIRV::ASSIGN_TYPE)
     return false;

 }
 
 bool SPIRVInstructionSelector::selectInsertVal(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeInsert))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType))
+                 .addUse(GOR.getSPIRVTypeID(ResType))
                  // object to insert
                  .addUse(I.getOperand(3).getReg())
                  // composite to insert into
                  .addUse(I.getOperand(2).getReg());
   for (unsigned i = 4; i < I.getNumOperands(); i++)
     MIB.addImm(foldImm(I.getOperand(i), MRI));
   return MIB.constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectExtractVal(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
   MachineBasicBlock &BB = *I.getParent();
   auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpCompositeExtract))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType))
+                 .addUse(GOR.getSPIRVTypeID(ResType))
                  .addUse(I.getOperand(2).getReg());
   for (unsigned i = 3; i < I.getNumOperands(); i++)
     MIB.addImm(foldImm(I.getOperand(i), MRI));
   return MIB.constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectInsertElt(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
   if (isImm(I.getOperand(4), MRI))
     return selectInsertVal(ResVReg, ResType, I);
   MachineBasicBlock &BB = *I.getParent();
   return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorInsertDynamic))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(I.getOperand(2).getReg())
       .addUse(I.getOperand(3).getReg())
       .addUse(I.getOperand(4).getReg())
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectExtractElt(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
   if (isImm(I.getOperand(3), MRI))
     return selectExtractVal(ResVReg, ResType, I);
   MachineBasicBlock &BB = *I.getParent();
   return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpVectorExtractDynamic))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addUse(I.getOperand(2).getReg())
       .addUse(I.getOperand(3).getReg())
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectGEP(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I) const {
   // In general we should also support OpAccessChain instrs here (i.e. not
   // PtrAccessChain) but SPIRV-LLVM Translator doesn't emit them at all and so
   // do we to stay compliant with its test and more importantly consumers.
   unsigned Opcode = I.getOperand(2).getImm() ? SPIRV::OpInBoundsPtrAccessChain
                                              : SPIRV::OpPtrAccessChain;
   auto Res = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(Opcode))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType))
+                 .addUse(GOR.getSPIRVTypeID(ResType))
                  // Object to get a pointer to.
                  .addUse(I.getOperand(3).getReg());
   // Adding indices.
   for (unsigned i = 4; i < I.getNumExplicitOperands(); ++i)
     Res.addUse(I.getOperand(i).getReg());
   return Res.constrainAllUses(TII, TRI, RBI);
 }
 
MachineInstr *Init = I.getNumExplicitOperands() > 2
                              ? MRI->getVRegDef(I.getOperand(2).getReg())
                              : nullptr;
     assert(MI);
     return selectGlobalValue(MI->getOperand(0).getReg(), *MI, Init);
   }
   case Intrinsic::spv_undef: {
     auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpUndef))
                    .addDef(ResVReg)
-                   .addUse(GR.getSPIRVTypeID(ResType));
+                   .addUse(GOR.getSPIRVTypeID(ResType));
     return MIB.constrainAllUses(TII, TRI, RBI);
   }
   case Intrinsic::spv_const_composite: {
     // If no values are attached, the composite is null constant.
     bool IsNull = I.getNumExplicitDefs() + 1 == I.getNumExplicitOperands();
     unsigned Opcode =
         IsNull ? SPIRV::OpConstantNull : SPIRV::OpConstantComposite;
     auto MIB = BuildMI(BB, I, I.getDebugLoc(), TII.get(Opcode))
                    .addDef(ResVReg)
-                   .addUse(GR.getSPIRVTypeID(ResType));
+                   .addUse(GOR.getSPIRVTypeID(ResType));
     // skip type MD node we already used when generated assign.type for this
     if (!IsNull) {
       for (unsigned i = I.getNumExplicitDefs() + 1;
            i < I.getNumExplicitOperands(); ++i) {
         MIB.addUse(I.getOperand(i).getReg());
       }
     }
     return MIB.constrainAllUses(TII, TRI, RBI);
bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
   return true;
 }
 
 bool SPIRVInstructionSelector::selectFrameIndex(Register ResVReg,
                                                 const SPIRVType *ResType,
                                                 MachineInstr &I) const {
   return BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpVariable))
       .addDef(ResVReg)
-      .addUse(GR.getSPIRVTypeID(ResType))
+      .addUse(GOR.getSPIRVTypeID(ResType))
       .addImm(static_cast<uint32_t>(SPIRV::StorageClass::Function))
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectBranch(MachineInstr &I) const {
   // InstructionSelector walks backwards through the instructions. We can use
   // both a G_BR and a G_BRCOND to create an OpBranchConditional. We hit G_BR
   // first, so can generate an OpBranchConditional here. If there is no
return BuildMI(MBB, I, I.getDebugLoc(), TII.get(SPIRV::OpBranchConditional))
       .constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectPhi(Register ResVReg,
                                          const SPIRVType *ResType,
                                          MachineInstr &I) const {
   auto MIB = BuildMI(*I.getParent(), I, I.getDebugLoc(), TII.get(SPIRV::OpPhi))
                  .addDef(ResVReg)
-                 .addUse(GR.getSPIRVTypeID(ResType));
+                 .addUse(GOR.getSPIRVTypeID(ResType));
   const unsigned NumOps = I.getNumOperands();
   for (unsigned i = 1; i < NumOps; i += 2) {
     MIB.addUse(I.getOperand(i + 0).getReg());
     MIB.addMBB(I.getOperand(i + 1).getMBB());
   }
   return MIB.constrainAllUses(TII, TRI, RBI);
 }
 
 bool SPIRVInstructionSelector::selectGlobalValue(
     Register ResVReg, MachineInstr &I, const MachineInstr *Init) const {
   // FIXME: don't use MachineIRBuilder here, replace it with BuildMI.
   MachineIRBuilder MIRBuilder(I);
   const GlobalValue *GV = I.getOperand(1).getGlobal();
-  SPIRVType *ResType = GR.getOrCreateSPIRVType(
+  SPIRVType *ResType = GOR.getOrCreateSPIRVType(
       GV->getType(), MIRBuilder, SPIRV::AccessQualifier::ReadWrite, false);
 
   std::string GlobalIdent = GV->getGlobalIdentifier();
   // We have functions as operands in tests with blocks of instruction e.g. in
   // transcoding/global_block.ll. These operands are not used and should be
   // substituted by zero constants. Their type is expected to be always
   // OpTypePointer Function %uchar.
   if (isa<Function>(GV)) {
     const Constant *ConstVal = GV;
     MachineBasicBlock &BB = *I.getParent();
-    Register NewReg = GR.find(ConstVal, GR.CurMF);
+    Register NewReg = GOR.find(ConstVal, GOR.CurMF);
     if (!NewReg.isValid()) {
-      SPIRVType *SpvBaseTy = GR.getOrCreateSPIRVIntegerType(8, I, TII);
-      ResType = GR.getOrCreateSPIRVPointerType(SpvBaseTy, I, TII);
+      SPIRVType *SpvBaseTy = GOR.getOrCreateSPIRVIntegerType(8, I, TII);
+      ResType = GOR.getOrCreateSPIRVPointerType(SpvBaseTy, I, TII);
       Register NewReg = ResVReg;
-      GR.add(ConstVal, GR.CurMF, NewReg);
+      GOR.add(ConstVal, GOR.CurMF, NewReg);
       return BuildMI(BB, I, I.getDebugLoc(), TII.get(SPIRV::OpConstantNull))
           .addDef(NewReg)
-          .addUse(GR.getSPIRVTypeID(ResType))
+          .addUse(GOR.getSPIRVTypeID(ResType))
           .constrainAllUses(TII, TRI, RBI);
     }
     assert(NewReg != ResVReg);
     return BuildMI(BB, I, I.getDebugLoc(), TII.get(TargetOpcode::COPY))
         .addDef(ResVReg)
         .addUse(NewReg)
         .constrainAllUses(TII, TRI, RBI);
   }
SPIRV::StorageClass::StorageClass Storage =
       addressSpaceToStorageClass(AddrSpace);
   bool HasLnkTy = GV->getLinkage() != GlobalValue::InternalLinkage &&
                   Storage != SPIRV::StorageClass::Function;
   SPIRV::LinkageType::LinkageType LnkType =
       (GV->isDeclaration() || GV->hasAvailableExternallyLinkage())
           ? SPIRV::LinkageType::Import
           : SPIRV::LinkageType::Export;
 
-  Register Reg = GR.buildGlobalVariable(ResVReg, ResType, GlobalIdent, GV,
+  Register Reg = GOR.buildGlobalVariable(ResVReg, ResType, GlobalIdent, GV,
                                          Storage, Init, GlobalVar->isConstant(),
                                          HasLnkTy, LnkType, MIRBuilder, true);
   return Reg.isValid();
 }
 
 namespace llvm {
 InstructionSelector *
 createSPIRVInstructionSelector(const SPIRVTargetMachine &TM,
                                const SPIRVSubtarget &Subtarget,
                                const RegisterBankInfo &RBI) {
   return new SPIRVInstructionSelector(TM, Subtarget, RBI);
 }
 } // namespace llvm

llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.h

 //===- SPIRVLegalizerInfo.h --- SPIR-V Legalization Rules --------*- C++ -*-==//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file declares the targeting of the MachineLegalizer class for SPIR-V.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_SPIRV_SPIRVMACHINELEGALIZER_H
 #define LLVM_LIB_TARGET_SPIRV_SPIRVMACHINELEGALIZER_H
 
-#include "SPIRVGlobalRegistry.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "llvm/CodeGen/GlobalISel/LegalizerInfo.h"
 
 bool isTypeFoldingSupported(unsigned Opcode);
 
 namespace llvm {
 
 class LLVMContext;
 class SPIRVSubtarget;
 
 // This class provides the information for legalizing SPIR-V instructions.
 class SPIRVLegalizerInfo : public LegalizerInfo {
   const SPIRVSubtarget *ST;
-  SPIRVGlobalRegistry *GR;
+  SPIRVGlobalObjectRegistry *GOR;
 
 public:
   bool legalizeCustom(LegalizerHelper &Helper, MachineInstr &MI) const override;
   SPIRVLegalizerInfo(const SPIRVSubtarget &ST);
 };
 } // namespace llvm
 #endif // LLVM_LIB_TARGET_SPIRV_SPIRVMACHINELEGALIZER_H

llvm/lib/Target/SPIRV/SPIRVLegalizerInfo.cpp

 //===- SPIRVLegalizerInfo.cpp --- SPIR-V Legalization Rules ------*- C++ -*-==//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the targeting of the Machinelegalizer class for SPIR-V.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIRVLegalizerInfo.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "SPIRV.h"
-#include "SPIRVGlobalRegistry.h"
 #include "SPIRVSubtarget.h"
 #include "llvm/CodeGen/GlobalISel/LegalizerHelper.h"
 #include "llvm/CodeGen/GlobalISel/MachineIRBuilder.h"
 #include "llvm/CodeGen/MachineInstr.h"
 #include "llvm/CodeGen/MachineRegisterInfo.h"
 #include "llvm/CodeGen/TargetOpcodes.h"
 
 using namespace llvm;

 bool isTypeFoldingSupported(unsigned Opcode) {
   return TypeFoldingSupportingOpcs.count(Opcode) > 0;
 }
 
 SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
   using namespace TargetOpcode;
 
   this->ST = &ST;
-  GR = ST.getSPIRVGlobalRegistry();
+  GOR = ST.getSPIRVGlobalObjectRegistry();
 
   const LLT s1 = LLT::scalar(1);
   const LLT s8 = LLT::scalar(8);
   const LLT s16 = LLT::scalar(16);
   const LLT s32 = LLT::scalar(32);
   const LLT s64 = LLT::scalar(64);
 
   const LLT v16s64 = LLT::fixed_vector(16, 64);
SPIRVLegalizerInfo::SPIRVLegalizerInfo(const SPIRVSubtarget &ST) {
 
   getLegacyLegalizerInfo().computeTables();
   verify(*ST.getInstrInfo());
 }
 
 static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVType *SpirvType,
                                 LegalizerHelper &Helper,
                                 MachineRegisterInfo &MRI,
-                                SPIRVGlobalRegistry *GR) {
+                                SPIRVGlobalObjectRegistry *GOR) {
   Register ConvReg = MRI.createGenericVirtualRegister(ConvTy);
-  GR->assignSPIRVTypeToVReg(SpirvType, ConvReg, Helper.MIRBuilder.getMF());
+  GOR->assignSPIRVTypeToVReg(SpirvType, ConvReg, Helper.MIRBuilder.getMF());
   Helper.MIRBuilder.buildInstr(TargetOpcode::G_PTRTOINT)
       .addDef(ConvReg)
       .addUse(Reg);
   return ConvReg;
 }
 
 bool SPIRVLegalizerInfo::legalizeCustom(LegalizerHelper &Helper,
                                         MachineInstr &MI) const {
   auto Opc = MI.getOpcode();
   MachineRegisterInfo &MRI = MI.getMF()->getRegInfo();
   if (!isTypeFoldingSupported(Opc)) {
     assert(Opc == TargetOpcode::G_ICMP);
-    assert(GR->getSPIRVTypeForVReg(MI.getOperand(0).getReg()));
+    assert(GOR->getSPIRVTypeForVReg(MI.getOperand(0).getReg()));
     auto &Op0 = MI.getOperand(2);
     auto &Op1 = MI.getOperand(3);
     Register Reg0 = Op0.getReg();
     Register Reg1 = Op1.getReg();
     CmpInst::Predicate Cond =
         static_cast<CmpInst::Predicate>(MI.getOperand(1).getPredicate());
     if ((!ST->canDirectlyComparePointers() ||
          (Cond != CmpInst::ICMP_EQ && Cond != CmpInst::ICMP_NE)) &&
         MRI.getType(Reg0).isPointer() && MRI.getType(Reg1).isPointer()) {
       LLT ConvT = LLT::scalar(ST->getPointerSize());
       Type *LLVMTy = IntegerType::get(MI.getMF()->getFunction().getContext(),
                                       ST->getPointerSize());
-      SPIRVType *SpirvTy = GR->getOrCreateSPIRVType(LLVMTy, Helper.MIRBuilder);
-      Op0.setReg(convertPtrToInt(Reg0, ConvT, SpirvTy, Helper, MRI, GR));
-      Op1.setReg(convertPtrToInt(Reg1, ConvT, SpirvTy, Helper, MRI, GR));
+      SPIRVType *SpirvTy = GOR->getOrCreateSPIRVType(LLVMTy, Helper.MIRBuilder);
+      Op0.setReg(convertPtrToInt(Reg0, ConvT, SpirvTy, Helper, MRI, GOR));
+      Op1.setReg(convertPtrToInt(Reg1, ConvT, SpirvTy, Helper, MRI, GOR));
     }
     return true;
   }
   // TODO: implement legalization for other opcodes.
   return true;
 }

llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h

 // and performs the global register numbering.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef LLVM_LIB_TARGET_SPIRV_SPIRVMODULEANALYSIS_H
 #define LLVM_LIB_TARGET_SPIRV_SPIRVMODULEANALYSIS_H
 
 #include "MCTargetDesc/SPIRVBaseInfo.h"
-#include "SPIRVGlobalRegistry.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "SPIRVUtils.h"
 #include "llvm/ADT/DenseMap.h"
 #include "llvm/ADT/SmallSet.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringMap.h"
 
 namespace llvm {
 class SPIRVSubtarget;
void collectGlobalEntities(
       std::function<bool(const SPIRV::DTSortableEntry *)> Pred,
       bool UsePreOrder);
   void processDefInstrs(const Module &M);
   void collectFuncNames(MachineInstr &MI, const Function *F);
   void processOtherInstrs(const Module &M);
   void numberRegistersGlobally(const Module &M);
 
   const SPIRVSubtarget *ST;
-  SPIRVGlobalRegistry *GR;
+  SPIRVGlobalObjectRegistry *GOR;
   const SPIRVInstrInfo *TII;
   MachineModuleInfo *MMI;
 };
 } // namespace llvm
 #endif // LLVM_LIB_TARGET_SPIRV_SPIRVMODULEANALYSIS_H

llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp

 
 // The function initializes global register alias table for types, consts,
 // global vars and func decls and collects these instruction for output
 // at module level. Also it collects explicit OpExtension/OpCapability
 // instructions.
 void SPIRVModuleAnalysis::processDefInstrs(const Module &M) {
   std::vector<SPIRV::DTSortableEntry *> DepsGraph;
 
-  GR->buildDepsGraph(DepsGraph, SPVDumpDeps ? MMI : nullptr);
+  GOR->buildDepsGraph(DepsGraph, SPVDumpDeps ? MMI : nullptr);
 
   collectGlobalEntities(
       DepsGraph, SPIRV::MB_TypeConstVars,
       [](const SPIRV::DTSortableEntry *E) { return !E->getIsFunc(); });
 
   for (auto F = M.begin(), E = M.end(); F != E; ++F) {
     MachineFunction *MF = MMI->getMachineFunction(*F);
     if (!MF)
void SPIRVModuleAnalysis::getAnalysisUsage(AnalysisUsage &AU) const {
   AU.addRequired<TargetPassConfig>();
   AU.addRequired<MachineModuleInfoWrapperPass>();
 }
 
 bool SPIRVModuleAnalysis::runOnModule(Module &M) {
   SPIRVTargetMachine &TM =
       getAnalysis<TargetPassConfig>().getTM<SPIRVTargetMachine>();
   ST = TM.getSubtargetImpl();
-  GR = ST->getSPIRVGlobalRegistry();
+  GOR = ST->getSPIRVGlobalObjectRegistry();
   TII = ST->getInstrInfo();
 
   MMI = &getAnalysis<MachineModuleInfoWrapperPass>().getMMI();
 
   setBaseInfo(M);
 
   addDecorations(M, *TII, MMI, *ST, MAI);
 

llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp

 //===-- SPIRVPreLegalizer.cpp - prepare IR for legalization -----*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // The pass prepares IR for legalization: it assigns SPIR-V types to registers
 // and removes intrinsics which holded these types during IR translation.
-// Also it processes constants and registers them in GR to avoid duplication.
+// Also it processes constants and registers them in GOR to avoid duplication.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIRV.h"
 #include "SPIRVSubtarget.h"
 #include "SPIRVUtils.h"
 #include "llvm/ADT/PostOrderIterator.h"
 #include "llvm/Analysis/OptimizationRemarkEmitter.h"
public:
   static char ID;
   SPIRVPreLegalizer() : MachineFunctionPass(ID) {
     initializeSPIRVPreLegalizerPass(*PassRegistry::getPassRegistry());
   }
   bool runOnMachineFunction(MachineFunction &MF) override;
 };
 } // namespace
 
-static void addConstantsToTrack(MachineFunction &MF, SPIRVGlobalRegistry *GR) {
+static void addConstantsToTrack(MachineFunction &MF,
+                                SPIRVGlobalObjectRegistry *GOR) {
   MachineRegisterInfo &MRI = MF.getRegInfo();
   DenseMap<MachineInstr *, Register> RegsAlreadyAddedToDT;
   SmallVector<MachineInstr *, 10> ToErase, ToEraseComposites;
   for (MachineBasicBlock &MBB : MF) {
     for (MachineInstr &MI : MBB) {
       if (!isSpvIntrinsic(MI, Intrinsic::spv_track_constant))
         continue;
       ToErase.push_back(&MI);
       auto *Const =
           cast<Constant>(cast<ConstantAsMetadata>(
                              MI.getOperand(3).getMetadata()->getOperand(0))
                              ->getValue());
       if (auto *GV = dyn_cast<GlobalValue>(Const)) {
-        Register Reg = GR->find(GV, &MF);
+        Register Reg = GOR->find(GV, &MF);
         if (!Reg.isValid())
-          GR->add(GV, &MF, MI.getOperand(2).getReg());
+          GOR->add(GV, &MF, MI.getOperand(2).getReg());
         else
           RegsAlreadyAddedToDT[&MI] = Reg;
       } else {
-        Register Reg = GR->find(Const, &MF);
+        Register Reg = GOR->find(Const, &MF);
         if (!Reg.isValid()) {
           if (auto *ConstVec = dyn_cast<ConstantDataVector>(Const)) {
             auto *BuildVec = MRI.getVRegDef(MI.getOperand(2).getReg());
             assert(BuildVec &&
                    BuildVec->getOpcode() == TargetOpcode::G_BUILD_VECTOR);
             for (unsigned i = 0; i < ConstVec->getNumElements(); ++i)
-              GR->add(ConstVec->getElementAsConstant(i), &MF,
+              GOR->add(ConstVec->getElementAsConstant(i), &MF,
                        BuildVec->getOperand(1 + i).getReg());
           }
-          GR->add(Const, &MF, MI.getOperand(2).getReg());
+          GOR->add(Const, &MF, MI.getOperand(2).getReg());
         } else {
           RegsAlreadyAddedToDT[&MI] = Reg;
           // This MI is unused and will be removed. If the MI uses
           // const_composite, it will be unused and should be removed too.
           assert(MI.getOperand(2).isReg() && "Reg operand is expected");
           MachineInstr *SrcMI = MRI.getVRegDef(MI.getOperand(2).getReg());
           if (SrcMI && isSpvIntrinsic(*SrcMI, Intrinsic::spv_const_composite))
             ToEraseComposites.push_back(SrcMI);
for (MachineInstr &MI : MBB) {
           ToErase.push_back(ConstMI);
       }
     }
   }
   for (MachineInstr *MI : ToErase)
     MI->eraseFromParent();
 }
 
-static void insertBitcasts(MachineFunction &MF, SPIRVGlobalRegistry *GR,
+static void insertBitcasts(MachineFunction &MF, SPIRVGlobalObjectRegistry *GOR,
                            MachineIRBuilder MIB) {
   SmallVector<MachineInstr *, 10> ToErase;
   for (MachineBasicBlock &MBB : MF) {
     for (MachineInstr &MI : MBB) {
       if (!isSpvIntrinsic(MI, Intrinsic::spv_bitcast))
         continue;
       assert(MI.getOperand(2).isReg());
       MIB.setInsertPt(*MI.getParent(), MI);
       MIB.buildBitcast(MI.getOperand(0).getReg(), MI.getOperand(2).getReg());
       ToErase.push_back(&MI);
     }
   }
   for (MachineInstr *MI : ToErase)
     MI->eraseFromParent();
 }
 
 // Translating GV, IRTranslator sometimes generates following IR:
 //   %1 = G_GLOBAL_VALUE
 //   %2 = COPY %1
 //   %3 = G_ADDRSPACE_CAST %2
 // New registers have no SPIRVType and no register class info.
 //
 // Set SPIRVType for GV, propagate it from GV to other instructions,
 // also set register classes.
-static SPIRVType *propagateSPIRVType(MachineInstr *MI, SPIRVGlobalRegistry *GR,
+static SPIRVType *propagateSPIRVType(MachineInstr *MI,
+                                     SPIRVGlobalObjectRegistry *GOR,
                                      MachineRegisterInfo &MRI,
                                      MachineIRBuilder &MIB) {
   SPIRVType *SpirvTy = nullptr;
   assert(MI && "Machine instr is expected");
   if (MI->getOperand(0).isReg()) {
     Register Reg = MI->getOperand(0).getReg();
-    SpirvTy = GR->getSPIRVTypeForVReg(Reg);
+    SpirvTy = GOR->getSPIRVTypeForVReg(Reg);
     if (!SpirvTy) {
       switch (MI->getOpcode()) {
       case TargetOpcode::G_CONSTANT: {
         MIB.setInsertPt(*MI->getParent(), MI);
         Type *Ty = MI->getOperand(1).getCImm()->getType();
-        SpirvTy = GR->getOrCreateSPIRVType(Ty, MIB);
+        SpirvTy = GOR->getOrCreateSPIRVType(Ty, MIB);
         break;
       }
       case TargetOpcode::G_GLOBAL_VALUE: {
         MIB.setInsertPt(*MI->getParent(), MI);
         Type *Ty = MI->getOperand(1).getGlobal()->getType();
-        SpirvTy = GR->getOrCreateSPIRVType(Ty, MIB);
+        SpirvTy = GOR->getOrCreateSPIRVType(Ty, MIB);
         break;
       }
       case TargetOpcode::G_TRUNC:
       case TargetOpcode::G_ADDRSPACE_CAST:
       case TargetOpcode::G_PTR_ADD:
       case TargetOpcode::COPY: {
         MachineOperand &Op = MI->getOperand(1);
         MachineInstr *Def = Op.isReg() ? MRI.getVRegDef(Op.getReg()) : nullptr;
         if (Def)
-          SpirvTy = propagateSPIRVType(Def, GR, MRI, MIB);
+          SpirvTy = propagateSPIRVType(Def, GOR, MRI, MIB);
         break;
       }
       default:
         break;
       }
       if (SpirvTy)
-        GR->assignSPIRVTypeToVReg(SpirvTy, Reg, MIB.getMF());
+        GOR->assignSPIRVTypeToVReg(SpirvTy, Reg, MIB.getMF());
       if (!MRI.getRegClassOrNull(Reg))
         MRI.setRegClass(Reg, &SPIRV::IDRegClass);
     }
   }
   return SpirvTy;
 }
 
 // Insert ASSIGN_TYPE instuction between Reg and its definition, set NewReg as
 // a dst of the definition, assign SPIRVType to both registers. If SpirvTy is
 // provided, use it as SPIRVType in ASSIGN_TYPE, otherwise create it from Ty.
 // It's used also in SPIRVBuiltins.cpp.
 // TODO: maybe move to SPIRVUtils.
 namespace llvm {
 Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy,
-                           SPIRVGlobalRegistry *GR, MachineIRBuilder &MIB,
-                           MachineRegisterInfo &MRI) {
+                           SPIRVGlobalObjectRegistry *GOR,
+                           MachineIRBuilder &MIB, MachineRegisterInfo &MRI) {
   MachineInstr *Def = MRI.getVRegDef(Reg);
   assert((Ty || SpirvTy) && "Either LLVM or SPIRV type is expected.");
   MIB.setInsertPt(*Def->getParent(),
                   (Def->getNextNode() ? Def->getNextNode()->getIterator()
                                       : Def->getParent()->end()));
   Register NewReg = MRI.createGenericVirtualRegister(MRI.getType(Reg));
   if (auto *RC = MRI.getRegClassOrNull(Reg)) {
     MRI.setRegClass(NewReg, RC);
   } else {
     MRI.setRegClass(NewReg, &SPIRV::IDRegClass);
     MRI.setRegClass(Reg, &SPIRV::IDRegClass);
   }
-  SpirvTy = SpirvTy ? SpirvTy : GR->getOrCreateSPIRVType(Ty, MIB);
-  GR->assignSPIRVTypeToVReg(SpirvTy, Reg, MIB.getMF());
+  SpirvTy = SpirvTy ? SpirvTy : GOR->getOrCreateSPIRVType(Ty, MIB);
+  GOR->assignSPIRVTypeToVReg(SpirvTy, Reg, MIB.getMF());
   // This is to make it convenient for Legalizer to get the SPIRVType
   // when processing the actual MI (i.e. not pseudo one).
-  GR->assignSPIRVTypeToVReg(SpirvTy, NewReg, MIB.getMF());
+  GOR->assignSPIRVTypeToVReg(SpirvTy, NewReg, MIB.getMF());
   // Copy MIFlags from Def to ASSIGN_TYPE instruction. It's required to keep
   // the flags after instruction selection.
   const uint16_t Flags = Def->getFlags();
   MIB.buildInstr(SPIRV::ASSIGN_TYPE)
       .addDef(Reg)
       .addUse(NewReg)
-      .addUse(GR->getSPIRVTypeID(SpirvTy))
+      .addUse(GOR->getSPIRVTypeID(SpirvTy))
       .setMIFlags(Flags);
   Def->getOperand(0).setReg(NewReg);
   return NewReg;
 }
 } // namespace llvm
 
-static void generateAssignInstrs(MachineFunction &MF, SPIRVGlobalRegistry *GR,
+static void generateAssignInstrs(MachineFunction &MF,
+                                 SPIRVGlobalObjectRegistry *GOR,
                                  MachineIRBuilder MIB) {
   MachineRegisterInfo &MRI = MF.getRegInfo();
   SmallVector<MachineInstr *, 10> ToErase;
 
   for (MachineBasicBlock *MBB : post_order(&MF)) {
     if (MBB->empty())
       continue;
 
     bool ReachedBegin = false;
     for (auto MII = std::prev(MBB->end()), Begin = MBB->begin();
          !ReachedBegin;) {
       MachineInstr &MI = *MII;
 
       if (isSpvIntrinsic(MI, Intrinsic::spv_assign_type)) {
         Register Reg = MI.getOperand(1).getReg();
         Type *Ty = getMDOperandAsType(MI.getOperand(2).getMetadata(), 0);
         MachineInstr *Def = MRI.getVRegDef(Reg);
         assert(Def && "Expecting an instruction that defines the register");
         // G_GLOBAL_VALUE already has type info.
         if (Def->getOpcode() != TargetOpcode::G_GLOBAL_VALUE)
-          insertAssignInstr(Reg, Ty, nullptr, GR, MIB, MF.getRegInfo());
+          insertAssignInstr(Reg, Ty, nullptr, GOR, MIB, MF.getRegInfo());
         ToErase.push_back(&MI);
       } else if (MI.getOpcode() == TargetOpcode::G_CONSTANT ||
                  MI.getOpcode() == TargetOpcode::G_FCONSTANT ||
                  MI.getOpcode() == TargetOpcode::G_BUILD_VECTOR) {
         // %rc = G_CONSTANT ty Val
         // ===>
         // %cty = OpType* ty
         // %rctmp = G_CONSTANT ty Val
for (auto MII = std::prev(MBB->end()), Begin = MBB->begin();
           else if (ElemMI->getOpcode() == TargetOpcode::G_FCONSTANT)
             ElemTy = ElemMI->getOperand(1).getFPImm()->getType();
           else
             llvm_unreachable("Unexpected opcode");
           unsigned NumElts =
               MI.getNumExplicitOperands() - MI.getNumExplicitDefs();
           Ty = VectorType::get(ElemTy, NumElts, false);
         }
-        insertAssignInstr(Reg, Ty, nullptr, GR, MIB, MRI);
+        insertAssignInstr(Reg, Ty, nullptr, GOR, MIB, MRI);
       } else if (MI.getOpcode() == TargetOpcode::G_TRUNC ||
                  MI.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
                  MI.getOpcode() == TargetOpcode::COPY ||
                  MI.getOpcode() == TargetOpcode::G_ADDRSPACE_CAST) {
-        propagateSPIRVType(&MI, GR, MRI, MIB);
+        propagateSPIRVType(&MI, GOR, MRI, MIB);
       }
 
       if (MII == Begin)
         ReachedBegin = true;
       else
         --MII;
     }
   }
   for (MachineInstr *MI : ToErase)
     MI->eraseFromParent();
 }
 
 static std::pair<Register, unsigned>
 createNewIdReg(Register ValReg, unsigned Opcode, MachineRegisterInfo &MRI,
-               const SPIRVGlobalRegistry &GR) {
+               const SPIRVGlobalObjectRegistry &GOR) {
   LLT NewT = LLT::scalar(32);
-  SPIRVType *SpvType = GR.getSPIRVTypeForVReg(ValReg);
+  SPIRVType *SpvType = GOR.getSPIRVTypeForVReg(ValReg);
   assert(SpvType && "VReg is expected to have SPIRV type");
   bool IsFloat = SpvType->getOpcode() == SPIRV::OpTypeFloat;
   bool IsVectorFloat =
       SpvType->getOpcode() == SPIRV::OpTypeVector &&
-      GR.getSPIRVTypeForVReg(SpvType->getOperand(1).getReg())->getOpcode() ==
+      GOR.getSPIRVTypeForVReg(SpvType->getOperand(1).getReg())->getOpcode() ==
           SPIRV::OpTypeFloat;
   IsFloat |= IsVectorFloat;
   auto GetIdOp = IsFloat ? SPIRV::GET_fID : SPIRV::GET_ID;
   auto DstClass = IsFloat ? &SPIRV::fIDRegClass : &SPIRV::IDRegClass;
   if (MRI.getType(ValReg).isPointer()) {
     NewT = LLT::pointer(0, 32);
     GetIdOp = SPIRV::GET_pID;
     DstClass = &SPIRV::pIDRegClass;
   } else if (MRI.getType(ValReg).isVector()) {
     NewT = LLT::fixed_vector(2, NewT);
     GetIdOp = IsFloat ? SPIRV::GET_vfID : SPIRV::GET_vID;
     DstClass = IsFloat ? &SPIRV::vfIDRegClass : &SPIRV::vIDRegClass;
   }
   Register IdReg = MRI.createGenericVirtualRegister(NewT);
   MRI.setRegClass(IdReg, DstClass);
   return {IdReg, GetIdOp};
 }
 
 static void processInstr(MachineInstr &MI, MachineIRBuilder &MIB,
-                         MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR) {
+                         MachineRegisterInfo &MRI,
+                         SPIRVGlobalObjectRegistry *GOR) {
   unsigned Opc = MI.getOpcode();
   assert(MI.getNumDefs() > 0 && MRI.hasOneUse(MI.getOperand(0).getReg()));
   MachineInstr &AssignTypeInst =
       *(MRI.use_instr_begin(MI.getOperand(0).getReg()));
-  auto NewReg = createNewIdReg(MI.getOperand(0).getReg(), Opc, MRI, *GR).first;
+  auto NewReg = createNewIdReg(MI.getOperand(0).getReg(), Opc, MRI, *GOR).first;
   AssignTypeInst.getOperand(1).setReg(NewReg);
   MI.getOperand(0).setReg(NewReg);
   MIB.setInsertPt(*MI.getParent(),
                   (MI.getNextNode() ? MI.getNextNode()->getIterator()
                                     : MI.getParent()->end()));
   for (auto &Op : MI.operands()) {
     if (!Op.isReg() || Op.isDef())
       continue;
-    auto IdOpInfo = createNewIdReg(Op.getReg(), Opc, MRI, *GR);
+    auto IdOpInfo = createNewIdReg(Op.getReg(), Opc, MRI, *GOR);
     MIB.buildInstr(IdOpInfo.second).addDef(IdOpInfo.first).addUse(Op.getReg());
     Op.setReg(IdOpInfo.first);
   }
 }
 
 // Defined in SPIRVLegalizerInfo.cpp.
 extern bool isTypeFoldingSupported(unsigned Opcode);
 
 static void processInstrsWithTypeFolding(MachineFunction &MF,
-                                         SPIRVGlobalRegistry *GR,
+                                         SPIRVGlobalObjectRegistry *GOR,
                                          MachineIRBuilder MIB) {
   MachineRegisterInfo &MRI = MF.getRegInfo();
   for (MachineBasicBlock &MBB : MF) {
     for (MachineInstr &MI : MBB) {
       if (isTypeFoldingSupported(MI.getOpcode()))
-        processInstr(MI, MIB, MRI, GR);
+        processInstr(MI, MIB, MRI, GOR);
     }
   }
   for (MachineBasicBlock &MBB : MF) {
     for (MachineInstr &MI : MBB) {
       // We need to rewrite dst types for ASSIGN_TYPE instrs to be able
       // to perform tblgen'erated selection and we can't do that on Legalizer
       // as it operates on gMIR only.
       if (MI.getOpcode() != SPIRV::ASSIGN_TYPE)
for (MachineInstr &MI : MBB) {
         if (UseMI.getOpcode() == TargetOpcode::G_ADDRSPACE_CAST)
           continue;
       }
       MRI.setType(DstReg, LLT::scalar(32));
     }
   }
 }
 
-static void processSwitches(MachineFunction &MF, SPIRVGlobalRegistry *GR,
+static void processSwitches(MachineFunction &MF, SPIRVGlobalObjectRegistry *GOR,
                             MachineIRBuilder MIB) {
   // Before IRTranslator pass, calls to spv_switch intrinsic are inserted before
   // each switch instruction. IRTranslator lowers switches to G_ICMP + G_BRCOND
   // + G_BR triples. A switch with two cases may be transformed to this MIR
   // sequence:
   //
   //   intrinsic(@llvm.spv.switch), %CmpReg, %Const0, %Const1
   //   %Dst0 = G_ICMP intpred(eq), %CmpReg, %Const0
if (ParentMBB->empty()) {
       ParentMBB->eraseFromParent();
     }
   }
 }
 
 bool SPIRVPreLegalizer::runOnMachineFunction(MachineFunction &MF) {
   // Initialize the type registry.
   const SPIRVSubtarget &ST = MF.getSubtarget<SPIRVSubtarget>();
-  SPIRVGlobalRegistry *GR = ST.getSPIRVGlobalRegistry();
-  GR->setCurrentFunc(MF);
+  SPIRVGlobalObjectRegistry *GOR = ST.getSPIRVGlobalObjectRegistry();
+  GOR->setCurrentFunc(MF);
   MachineIRBuilder MIB(MF);
-  addConstantsToTrack(MF, GR);
+  addConstantsToTrack(MF, GOR);
   foldConstantsIntoIntrinsics(MF);
-  insertBitcasts(MF, GR, MIB);
-  generateAssignInstrs(MF, GR, MIB);
-  processSwitches(MF, GR, MIB);
-  processInstrsWithTypeFolding(MF, GR, MIB);
+  insertBitcasts(MF, GOR, MIB);
+  generateAssignInstrs(MF, GOR, MIB);
+  processSwitches(MF, GOR, MIB);
+  processInstrsWithTypeFolding(MF, GOR, MIB);
 
   return true;
 }
 
 INITIALIZE_PASS(SPIRVPreLegalizer, DEBUG_TYPE, "SPIRV pre legalizer", false,
                 false)
 
 char SPIRVPreLegalizer::ID = 0;
 
 FunctionPass *llvm::createSPIRVPreLegalizerPass() {
   return new SPIRVPreLegalizer();
 }

llvm/lib/Target/SPIRV/SPIRVSubtarget.h

 class SPIRVSubtarget : public SPIRVGenSubtargetInfo {
 private:
   const unsigned PointerSize;
   uint32_t SPIRVVersion;
   uint32_t OpenCLVersion;
 
   SmallSet<SPIRV::Extension::Extension, 4> AvailableExtensions;
   SmallSet<SPIRV::InstructionSet::InstructionSet, 4> AvailableExtInstSets;
-  std::unique_ptr<SPIRVGlobalRegistry> GR;
+  std::unique_ptr<SPIRVGlobalObjectRegistry> GOR;
 
   SPIRVInstrInfo InstrInfo;
   SPIRVFrameLowering FrameLowering;
   SPIRVTargetLowering TLInfo;
 
   // GlobalISel related APIs.
   std::unique_ptr<CallLowering> CallLoweringInfo;
   std::unique_ptr<RegisterBankInfo> RegBankInfo;
public:
   bool isAtLeastSPIRVVer(uint32_t VerToCompareTo) const;
   bool isAtLeastOpenCLVer(uint32_t VerToCompareTo) const;
   // TODO: implement command line args or other ways to determine this.
   bool hasOpenCLFullProfile() const { return true; }
   bool hasOpenCLImageSupport() const { return true; }
   bool canUseExtension(SPIRV::Extension::Extension E) const;
   bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const;
 
-  SPIRVGlobalRegistry *getSPIRVGlobalRegistry() const { return GR.get(); }
+  SPIRVGlobalObjectRegistry *getSPIRVGlobalObjectRegistry() const {
+    return GOR.get();
+  }
 
   const CallLowering *getCallLowering() const override {
     return CallLoweringInfo.get();
   }
   const RegisterBankInfo *getRegBankInfo() const override {
     return RegBankInfo.get();
   }
   const LegalizerInfo *getLegalizerInfo() const override {

llvm/lib/Target/SPIRV/SPIRVSubtarget.cpp

 //===-- SPIRVSubtarget.cpp - SPIR-V Subtarget Information ------*- C++ -*--===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the SPIR-V specific subclass of TargetSubtargetInfo.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIRVSubtarget.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "SPIRV.h"
-#include "SPIRVGlobalRegistry.h"
 #include "SPIRVLegalizerInfo.h"
 #include "SPIRVRegisterBankInfo.h"
 #include "SPIRVTargetMachine.h"
 #include "llvm/MC/TargetRegistry.h"
 #include "llvm/TargetParser/Host.h"
 
 using namespace llvm;
 
SPIRVSubtarget::SPIRVSubtarget(const Triple &TT, const std::string &CPU,
     : SPIRVGenSubtargetInfo(TT, CPU, /*TuneCPU=*/CPU, FS),
       PointerSize(computePointerSize(TT)), SPIRVVersion(0), OpenCLVersion(0),
       InstrInfo(), FrameLowering(initSubtargetDependencies(CPU, FS)),
       TLInfo(TM, *this) {
   // The order of initialization is important.
   initAvailableExtensions();
   initAvailableExtInstSets();
 
-  GR = std::make_unique<SPIRVGlobalRegistry>(PointerSize);
-  CallLoweringInfo = std::make_unique<SPIRVCallLowering>(TLInfo, GR.get());
+  GOR = std::make_unique<SPIRVGlobalObjectRegistry>(PointerSize);
+  CallLoweringInfo = std::make_unique<SPIRVCallLowering>(TLInfo, GOR.get());
   Legalizer = std::make_unique<SPIRVLegalizerInfo>(*this);
   RegBankInfo = std::make_unique<SPIRVRegisterBankInfo>();
   InstSelector.reset(
       createSPIRVInstructionSelector(TM, *this, *RegBankInfo.get()));
 }
 
 SPIRVSubtarget &SPIRVSubtarget::initSubtargetDependencies(StringRef CPU,
                                                           StringRef FS) {

llvm/lib/Target/SPIRV/SPIRVTargetMachine.cpp

 //===- SPIRVTargetMachine.cpp - Define TargetMachine for SPIR-V -*- C++ -*-===//
 //
 // 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
 //
 //===----------------------------------------------------------------------===//
 //
 // Implements the info about SPIR-V target spec.
 //
 //===----------------------------------------------------------------------===//
 
 #include "SPIRVTargetMachine.h"
 #include "SPIRV.h"
 #include "SPIRVCallLowering.h"
-#include "SPIRVGlobalRegistry.h"
+#include "Registries/SPIRVGlobalObjectRegistry.h"
 #include "SPIRVLegalizerInfo.h"
 #include "SPIRVTargetObjectFile.h"
 #include "SPIRVTargetTransformInfo.h"
 #include "TargetInfo/SPIRVTargetInfo.h"
 #include "llvm/CodeGen/GlobalISel/IRTranslator.h"
 #include "llvm/CodeGen/GlobalISel/InstructionSelect.h"
 #include "llvm/CodeGen/GlobalISel/Legalizer.h"
 #include "llvm/CodeGen/GlobalISel/RegBankSelect.h"