[MLIR][SPIRV] Rename `spv.globalVariable` to `spv.GlobalVariable`.

[ergawy]

Summary

To unify the naming scheme across all ops in the SPIR-V dialect, we are
moving from spv.camelCase to spv.CamelCase everywhere.

mlir/docs/Dialects/SPIR-V.md

*   Various normal constant instructions are represented by the same
     transformations less tedious.
 *   Normal constants are not placed in `spv.module`'s region; they are localized
     into functions. This is to make functions in the SPIR-V dialect to be
     isolated and explicit capturing. Constants are cheap to duplicate given
     attributes are made unique in `MLIRContext`.
 
 #### Adopt symbol-based global variables and specialization constant
 
-*   Global variables are defined with the `spv.globalVariable` op. They do not
+*   Global variables are defined with the `spv.GlobalVariable` op. They do not
     generate SSA values. Instead they have symbols and should be referenced via
     symbols. To use global variables in a function block, `spv.mlir.addressof` is
     needed to turn the symbol into an SSA value.
 *   Specialization constants are defined with the `spv.specConstant` op. Similar
     to global variables, they do not generate SSA values and have symbols for
     reference, too. `spv.mlir.referenceof` is needed to turn the symbol into an SSA
     value for use in a function block.
 

 Similarly, a few transformations are performed during deserialization:
 
 *   Instructions for execution environment requirements (extensions,
     capabilities, extended instruction sets, etc.) will be placed as attributes
     on `spv.module`.
 *   `OpType*` instructions will be converted into proper `mlir::Type`s.
 *   `OpConstant*` instructions are materialized as `spv.constant` at each use
     site.
-*   `OpVariable` instructions will be converted to `spv.globalVariable` ops if
+*   `OpVariable` instructions will be converted to `spv.GlobalVariable` ops if
     in module-level; otherwise they will be converted into `spv.Variable` ops.
 *   Every use of a module-level `OpVariable` instruction will materialize a
     `spv.mlir.addressof` op to turn the symbol of the corresponding
-    `spv.globalVariable` into an SSA value.
+    `spv.GlobalVariable` into an SSA value.
 *   Every use of a `OpSpecConstant` instruction will materialize a
     `spv.mlir.referenceof` op to turn the symbol of the corresponding
     `spv.specConstant` into an SSA value.
 *   `OpPhi` instructions are converted to block arguments.
 *   Structured control flow are placed inside `spv.selection` and `spv.loop`.
 
 ## Conversions
 

 
 The method `mlir::spirv::setABIAttrs` allows setting the [shader interface
 attributes](#shader-interface-abi) for a function that is to be an entry
 point function within the `spv.module` on lowering. A later pass
 `mlir::spirv::LowerABIAttributesPass` uses this information to lower the entry
 point function and its ABI consistent with the Vulkan validation
 rules. Specifically,
 
-*   Creates `spv.globalVariable`s for the arguments, and replaces all uses of
+*   Creates `spv.GlobalVariable`s for the arguments, and replaces all uses of
     the argument with this variable. The SSA value used for replacement is
     obtained using the `spv.mlir.addressof` operation.
 *   Adds the `spv.EntryPoint` and `spv.ExecutionMode` operations into the
     `spv.module` for the entry function.
 
 #### Setting layout for shader interface variables
 
 SPIR-V validation rules for shaders require composite objects to be explicitly
-laid out. If a `spv.globalVariable` is not explicitly laid out, the utility
+laid out. If a `spv.GlobalVariable` is not explicitly laid out, the utility
 method `mlir::spirv::decorateType` implements a layout consistent with
 the [Vulkan shader requirements][VulkanShaderInterface].
 
 #### Creating builtin variables
 
-In SPIR-V dialect, builtins are represented using `spv.globalVariable`s, with
+In SPIR-V dialect, builtins are represented using `spv.GlobalVariable`s, with
 `spv.mlir.addressof` used to get a handle to the builtin as an SSA value.  The
-method `mlir::spirv::getBuiltinVariableValue` creates a `spv.globalVariable` for
+method `mlir::spirv::getBuiltinVariableValue` creates a `spv.GlobalVariable` for
 the builtin in the current `spv.module` if it does not exist already, and
 returns an SSA value generated from an `spv.mlir.addressof` operation.
 
 ### Current conversions to SPIR-V
 
 Using the above infrastructure, conversions are implemented from
 
 *   [Standard Dialect][MlirStandardDialect] : Only arithmetic and logical

mlir/docs/SPIRVToLLVMDialectConversion.md

*   **Nontemporal**: set `nontemporal` flag, for example: `mlir // %res =
     %ptr ["Nontemporal"] : f32`
 
 *   **Volatile**: mark the op as `volatile`, for example: `mlir // %res =
     llvm.load volatile %ptr : !llvm.ptr<f32> %res = spv.Load "Function" %ptr
     ["Volatile"] : f32` Otherwise the conversion fails as other cases
     (`MakePointerAvailable`, `MakePointerVisible`, `NonPrivatePointer`) are not
     supported yet.
 
-#### `spv.globalVariable` and `spv.mlir.addressof`
+#### `spv.GlobalVariable` and `spv.mlir.addressof`
 
-`spv.globalVariable` is modelled with `llvm.mlir.global` op. However, there
+`spv.GlobalVariable` is modelled with `llvm.mlir.global` op. However, there
 is a difference that has to be pointed out.
 
 In SPIR-V dialect, the global variable returns a pointer, whereas in LLVM
 dialect the global holds an actual value. This difference is handled by
 `spv.mlir.addressof` and `llvm.mlir.addressof` ops that both return a pointer and
 are used to reference the global.
 
 ```mlir
 // Original SPIR-V module
 spv.module Logical GLSL450 {
-  spv.globalVariable @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
+  spv.GlobalVariable @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
   spv.func @func() -> () "None" {
     %0 = spv.mlir.addressof @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
     spv.Return
   }
 }
 
 // Converted result
 module {

 > By default, functions and global variables are private to a module and cannot
 be accessed by other modules. However, a module may be written to export or
 import functions and global (module scope) variables.
 
 If the global variable's pointer has `Input` storage class, then a `constant`
 flag is added to LLVM op:
 
 ```mlir
-spv.globalVariable @var : !spv.ptr<f32, Input>    =>    llvm.mlir.global external constant @var() : f32
+spv.GlobalVariable @var : !spv.ptr<f32, Input>    =>    llvm.mlir.global external constant @var() : f32
 ```
 
 #### `spv.Variable`
 
 Per SPIR-V dialect spec, `spv.Variable` allocates an object in memory, resulting
 in a pointer to it, which can be used with `spv.Load` and `spv.Store`. It is
 also a function-level variable.
 
func @main() {
   }
 }
 ```
 
 Lowering `gpu` dialect to SPIR-V dialect results in
 
 ```mlir
 spv.module @__spv__foo /*VCE triple and other metadata here*/ {
-  spv.globalVariable @__spv__foo_arg bind(0,0) : ...
+  spv.GlobalVariable @__spv__foo_arg bind(0,0) : ...
   spv.func @bar() {
     // Kernel code.
   }
   spv.EntryPoint @bar, ...
 }
 
 func @main() {
   // Fill the buffer with some data.
   %buffer = alloc : memref<8xi32>
   %data = ...
   call fillBuffer(%buffer, %data)
 
   "gpu.launch_func"(/*grid dimensions*/, %buffer) {
     kernel = @foo::bar
   }
 }
 ```
 
 Then, the lowering from standard dialect to LLVM dialect is applied to the host
 code.
 
 ```mlir
 spv.module @__spv__foo /*VCE triple and other metadata here*/ {
-  spv.globalVariable @__spv__foo_arg bind(0,0) : ...
+  spv.GlobalVariable @__spv__foo_arg bind(0,0) : ...
   spv.func @bar() {
     // Kernel code.
   }
   spv.EntryPoint @bar, ...
 }
 
 // Kernel function declaration.
 llvm.func @__spv__foo_bar() : ...

mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td

let description = [{
     static call tree. See Execution Model.
 
     Entry Point must be the Result <id> of an OpFunction instruction.
 
     Name is a name string for the entry point. A module cannot have two
     OpEntryPoint instructions with the same Execution Model and the same
     Name string.
 
-    Interface is a list of symbol references to `spv.globalVariable`
+    Interface is a list of symbol references to `spv.GlobalVariable`
     operations. These declare the set of global variables from a
     module that form the interface of this entry point. The set of
     Interface symbols must be equal to or a superset of the
-    `spv.globalVariable`s referenced by the entry point’s static call
+    `spv.GlobalVariable`s referenced by the entry point’s static call
     tree, within the interface’s storage classes.  Before version 1.4,
     the interface’s storage classes are limited to the Input and
     Output storage classes. Starting with version 1.4, the interface’s
     storage classes are all storage classes used in declaring all
     global variables referenced by the entry point’s call tree.
 
     <!-- End of AutoGen section -->
 
private:
     /// type and the presence of the (potentially empty) function body.
     /// Ensures SPIR-V specific semantics.
     LogicalResult verifyBody();
   }];
 }
 
 // -----
 
-def SPV_GlobalVariableOp : SPV_Op<"globalVariable", [InModuleScope, Symbol]> {
+def SPV_GlobalVariableOp : SPV_Op<"GlobalVariable", [InModuleScope, Symbol]> {
   let summary = [{
     Allocate an object in memory at module scope. The object is
     referenced using a symbol name.
   }];
 
   let description = [{
     The variable type must be an OpTypePointer. Its type operand is the type of
     object in memory.
 
     Storage Class is the Storage Class of the memory holding the object. It
     cannot be Generic. It must be the same as the Storage Class operand of
     the variable types. Only those storage classes that are valid at module
     scope (like Input, Output, StorageBuffer, etc.) are valid.
 
     Initializer is optional.  If Initializer is present, it will be
     the initial value of the variable’s memory content. Initializer
     must be an symbol defined from a constant instruction or other
-    `spv.globalVariable` operation in module scope. Initializer must
+    `spv.GlobalVariable` operation in module scope. Initializer must
     have the same type as the type of the defined symbol.
 
     <!-- End of AutoGen section -->
 
     ```
-    variable-op ::= `spv.globalVariable` spirv-type symbol-ref-id
+    variable-op ::= `spv.GlobalVariable` spirv-type symbol-ref-id
                     (`initializer(` symbol-ref-id `)`)?
                     (`bind(` integer-literal, integer-literal `)`)?
                     (`built_in(` string-literal `)`)?
                     attribute-dict?
     ```
 
     where `initializer` specifies initializer and `bind` specifies the
     descriptor set and binding number. `built_in` specifies SPIR-V
     BuiltIn decoration associated with the op.
 
     #### Example:
 
     ```mlir
-    spv.globalVariable @var0 : !spv.ptr<f32, Input> @var0
-    spv.globalVariable @var1 initializer(@var0) : !spv.ptr<f32, Output>
-    spv.globalVariable @var2 bind(1, 2) : !spv.ptr<f32, Uniform>
-    spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+    spv.GlobalVariable @var0 : !spv.ptr<f32, Input> @var0
+    spv.GlobalVariable @var1 initializer(@var0) : !spv.ptr<f32, Output>
+    spv.GlobalVariable @var2 bind(1, 2) : !spv.ptr<f32, Uniform>
+    spv.GlobalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
     ```
   }];
 
   let arguments = (ins
     TypeAttr:$type,
     StrAttr:$sym_name,
     OptionalAttr<FlatSymbolRefAttr>:$initializer
   );

mlir/include/mlir/Dialect/SPIRV/Linking/ModuleCombiner.h

 ///   - If 2 spv.func's have the same symbol name, then rename one of the
 ///   functions.
 ///   - If an spv.func and another op have the same symbol name, then rename the
 ///   other symbol.
 ///   - If none of the 2 conflicting ops are spv.func, then rename either.
 ///
 /// For deduplication, the following 3 cases are taken into consideration:
 ///
-///   - If 2 spv.globalVariable's have either the same descriptor set + binding
+///   - If 2 spv.GlobalVariable's have either the same descriptor set + binding
 ///   or the same build_in attribute value, then replace one of them using the
 ///   other.
 ///   - If 2 spv.specConstant's have the same spec_id attribute value, then
 ///   replace one of them using the other.
 ///   - If 2 spv.func's are identical replace one of them using the other.
 ///
 /// In all cases, the references to the updated symbol (whether renamed or
 /// deduplicated) are also updated to reflect the change.

mlir/lib/Conversion/SPIRVToLLVM/ConvertLaunchFuncToLLVMCalls.cpp

static bool hasDescriptorSetAndBinding(spirv::GlobalVariableOp op) {
   IntegerAttr descriptorSet =
       op->getAttrOfType<IntegerAttr>(descriptorSetName());
   IntegerAttr binding = op->getAttrOfType<IntegerAttr>(bindingName());
   return descriptorSet && binding;
 }
 
 /// Fills `globalVariableMap` with SPIR-V global variables that represent kernel
 /// arguments from the given SPIR-V module. We assume that the module contains a
-/// single entry point function. Hence, all `spv.globalVariable`s with a bind
+/// single entry point function. Hence, all `spv.GlobalVariable`s with a bind
 /// attribute are kernel arguments.
 static LogicalResult getKernelGlobalVariables(
     spirv::ModuleOp module,
     DenseMap<uint32_t, spirv::GlobalVariableOp> &globalVariableMap) {
   auto entryPoints = module.getOps<spirv::EntryPointOp>();
   if (!llvm::hasSingleElement(entryPoints)) {
     return module.emitError(
         "The module must contain exactly one entry point function");

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp

for (unsigned i = 0, e = values.size(); i < e; ++i) {
                           rewriter.getIntegerAttr(rewriter.getI32Type(), i)}));
     }
     rewriter.create<LLVM::ReturnOp>(loc, ArrayRef<Value>({structValue}));
     rewriter.eraseOp(op);
     return success();
   }
 };
 
-/// Converts `spv.globalVariable` to `llvm.mlir.global`. Note that SPIR-V global
+/// Converts `spv.GlobalVariable` to `llvm.mlir.global`. Note that SPIR-V global
 /// returns a pointer, whereas in LLVM dialect the global holds an actual value.
 /// This difference is handled by `spv.mlir.addressof` and
 /// `llvm.mlir.addressof`ops that both return a pointer.
 class GlobalVariablePattern
     : public SPIRVToLLVMConversion<spirv::GlobalVariableOp> {
 public:
   using SPIRVToLLVMConversion<spirv::GlobalVariableOp>::SPIRVToLLVMConversion;
 

mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp

matchAndRewrite(AllocOp operation, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     MemRefType allocType = operation.getType();
     if (!isAllocationSupported(allocType))
       return operation.emitError("unhandled allocation type");
 
     // Get the SPIR-V type for the allocation.
     Type spirvType = getTypeConverter()->convertType(allocType);
 
-    // Insert spv.globalVariable for this allocation.
+    // Insert spv.GlobalVariable for this allocation.
     Operation *parent =
         SymbolTable::getNearestSymbolTable(operation->getParentOp());
     if (!parent)
       return failure();
     Location loc = operation.getLoc();
     spirv::GlobalVariableOp varOp;
     {
       OpBuilder::InsertionGuard guard(rewriter);

mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp

void spirv::AddressOfOp::build(OpBuilder &builder, OperationState &state,
   build(builder, state, var.type(), builder.getSymbolRefAttr(var));
 }
 
 static LogicalResult verify(spirv::AddressOfOp addressOfOp) {
   auto varOp = dyn_cast_or_null<spirv::GlobalVariableOp>(
       SymbolTable::lookupNearestSymbolFrom(addressOfOp->getParentOp(),
                                            addressOfOp.variable()));
   if (!varOp) {
-    return addressOfOp.emitOpError("expected spv.globalVariable symbol");
+    return addressOfOp.emitOpError("expected spv.GlobalVariable symbol");
   }
   if (addressOfOp.pointer().getType() != varOp.type()) {
     return addressOfOp.emitOpError(
         "result type mismatch with the referenced global variable's type");
   }
   return success();
 }
 
CallInterfaceCallable spirv::FunctionCallOp::getCallableForCallee() {
   return (*this)->getAttrOfType<SymbolRefAttr>(kCallee);
 }
 
 Operation::operand_range spirv::FunctionCallOp::getArgOperands() {
   return arguments();
 }
 
 //===----------------------------------------------------------------------===//
-// spv.globalVariable
+// spv.GlobalVariable
 //===----------------------------------------------------------------------===//
 
 void spirv::GlobalVariableOp::build(OpBuilder &builder, OperationState &state,
                                     Type type, StringRef name,
                                     unsigned descriptorSet, unsigned binding) {
   build(builder, state, TypeAttr::get(type), builder.getStringAttr(name),
         nullptr);
   state.addAttribute(
if (auto init =
     Operation *initOp = SymbolTable::lookupNearestSymbolFrom(
         varOp->getParentOp(), init.getValue());
     // TODO: Currently only variable initialization with specialization
     // constants and other variables is supported. They could be normal
     // constants in the module scope as well.
     if (!initOp ||
         !isa<spirv::GlobalVariableOp, spirv::SpecConstantOp>(initOp)) {
       return varOp.emitOpError("initializer must be result of a "
-                               "spv.specConstant or spv.globalVariable op");
+                               "spv.specConstant or spv.GlobalVariable op");
     }
   }
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.GroupBroadcast
if (auto entryPointOp = dyn_cast<spirv::EntryPointOp>(op)) {
             return entryPointOp.emitError(
                        "expected symbol reference for interface "
                        "specification instead of '")
                    << varRef;
           }
           auto variableOp =
               table.lookup<spirv::GlobalVariableOp>(varSymRef.getValue());
           if (!variableOp) {
-            return entryPointOp.emitError("expected spv.globalVariable "
+            return entryPointOp.emitError("expected spv.GlobalVariable "
                                           "symbol reference instead of'")
                    << varSymRef << "'";
           }
         }
       }
 
       auto key = std::pair<spirv::FuncOp, spirv::ExecutionModel>(
           funcOp, entryPointOp.execution_model());

 
 static LogicalResult verify(spirv::VariableOp varOp) {
   // SPIR-V spec: "Storage Class is the Storage Class of the memory holding the
   // object. It cannot be Generic. It must be the same as the Storage Class
   // operand of the Result Type."
   if (varOp.storage_class() != spirv::StorageClass::Function) {
     return varOp.emitOpError(
         "can only be used to model function-level variables. Use "
-        "spv.globalVariable for module-level variables.");
+        "spv.GlobalVariable for module-level variables.");
   }
 
   auto pointerType = varOp.pointer().getType().cast<spirv::PointerType>();
   if (varOp.storage_class() != pointerType.getStorageClass())
     return varOp.emitOpError(
         "storage class must match result pointer's storage class");
 
   if (varOp.getNumOperands() != 0) {
     // SPIR-V spec: "Initializer must be an <id> from a constant instruction or
     // a global (module scope) OpVariable instruction".
     auto *initOp = varOp.getOperand(0).getDefiningOp();
     if (!initOp || !isa<spirv::ConstantOp,    // for normal constant
                         spirv::ReferenceOfOp, // for spec constant
                         spirv::AddressOfOp>(initOp))
       return varOp.emitOpError("initializer must be the result of a "
-                               "constant or spv.globalVariable op");
+                               "constant or spv.GlobalVariable op");
   }
 
   // TODO: generate these strings using ODS.
   auto *op = varOp.getOperation();
   auto descriptorSetName = llvm::convertToSnakeFromCamelCase(
       stringifyDecoration(spirv::Decoration::DescriptorSet));
   auto bindingName = llvm::convertToSnakeFromCamelCase(
       stringifyDecoration(spirv::Decoration::Binding));
   auto builtInName = llvm::convertToSnakeFromCamelCase(
       stringifyDecoration(spirv::Decoration::BuiltIn));
 
   for (const auto &attr : {descriptorSetName, bindingName, builtInName}) {
     if (op->getAttr(attr))
       return varOp.emitOpError("cannot have '")
-             << attr << "' attribute (only allowed in spv.globalVariable)";
+             << attr << "' attribute (only allowed in spv.GlobalVariable)";
   }
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // spv.VectorShuffle
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/SPIRV/Transforms/SPIRVConversion.cpp

 
 //===----------------------------------------------------------------------===//
 // Builtin Variables
 //===----------------------------------------------------------------------===//
 
 static spirv::GlobalVariableOp getBuiltinVariable(Block &body,
                                                   spirv::BuiltIn builtin) {
   // Look through all global variables in the given `body` block and check if
-  // there is a spv.globalVariable that has the same `builtin` attribute.
+  // there is a spv.GlobalVariable that has the same `builtin` attribute.
   for (auto varOp : body.getOps<spirv::GlobalVariableOp>()) {
     if (auto builtinAttr = varOp->getAttrOfType<StringAttr>(
             spirv::SPIRVDialect::getAttributeName(
                 spirv::Decoration::BuiltIn))) {
       auto varBuiltIn = spirv::symbolizeBuiltIn(builtinAttr.getValue());
       if (varBuiltIn && varBuiltIn.getValue() == builtin) {
         return varOp;
       }

mlir/lib/Target/SPIRV/Deserialization/Deserializer.h

private:
   }
 
   /// Creates a spirv::SpecConstantOp.
   spirv::SpecConstantOp createSpecConstant(Location loc, uint32_t resultID,
                                            Attribute defaultValue);
 
   /// Processes the OpVariable instructions at current `offset` into `binary`.
   /// It is expected that this method is used for variables that are to be
-  /// defined at module scope and will be deserialized into a spv.globalVariable
+  /// defined at module scope and will be deserialized into a spv.GlobalVariable
   /// instruction.
   LogicalResult processGlobalVariable(ArrayRef<uint32_t> operands);
 
   /// Gets the global variable associated with a result <id> of OpVariable.
   spirv::GlobalVariableOp getGlobalVariable(uint32_t id) {
     return globalVariableMap.lookup(id);
   }
 

mlir/test/Conversion/GPUToSPIRV/builtins.mlir

 // RUN: mlir-opt -split-input-file -convert-gpu-to-spirv %s -o - | FileCheck %s
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_y
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_y() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}1 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "y"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_workgroup_id_z
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  // CHECK: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
   gpu.module @kernels {
     gpu.func @builtin_workgroup_id_z() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[WORKGROUPID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}2 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "z"} : () -> index
       gpu.return
module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_local_id_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
+  // CHECK: spv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
   gpu.module @kernels {
     gpu.func @builtin_local_id_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[LOCALINVOCATIONID]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.thread_id"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   func @builtin() {
     %c0 = constant 1 : index
     gpu.launch_func @kernels::@builtin_num_workgroups_x
         blocks in (%c0, %c0, %c0) threads in (%c0, %c0, %c0)
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
+  // CHECK: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_workgroups_x() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMWORKGROUPS]]
       // CHECK-NEXT: [[VEC:%.*]] = spv.Load "Input" [[ADDRESS]]
       // CHECK-NEXT: {{%.*}} = spv.CompositeExtract [[VEC]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.grid_dim"() {dimension = "x"} : () -> index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
+  // CHECK: spv.GlobalVariable [[SUBGROUPID:@.*]] built_in("SubgroupId")
   gpu.module @kernels {
     gpu.func @builtin_subgroup_id() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPID]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_id : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
+  // CHECK: spv.GlobalVariable [[NUMSUBGROUPS:@.*]] built_in("NumSubgroups")
   gpu.module @kernels {
     gpu.func @builtin_num_subgroups() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[NUMSUBGROUPS]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.num_subgroups : index
       gpu.return
     }
   }
 }
 
 // -----
 
 module attributes {gpu.container_module} {
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
-  // CHECK: spv.globalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
+  // CHECK: spv.GlobalVariable [[SUBGROUPSIZE:@.*]] built_in("SubgroupSize")
   gpu.module @kernels {
     gpu.func @builtin_subgroup_size() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
       // CHECK: [[ADDRESS:%.*]] = spv.mlir.addressof [[SUBGROUPSIZE]]
       // CHECK-NEXT: {{%.*}} = spv.Load "Input" [[ADDRESS]]
       %0 = gpu.subgroup_size : index
       gpu.return
     }
   }
 }

mlir/test/Conversion/GPUToSPIRV/load-store.mlir

gpu.launch_func @kernels::@load_store_kernel
         blocks in (%0, %c1_2, %c1_2) threads in (%1, %c1_2, %c1_2)
         args(%arg0 : memref<12x4xf32>, %arg1 : memref<12x4xf32>, %arg2 : memref<12x4xf32>, 
              %c0 : index, %c0_0 : index, %c1 : index, %c1_1 : index)
     return
   }
 
   // CHECK-LABEL: spv.module @{{.*}} Logical GLSL450
   gpu.module @kernels {
-    // CHECK-DAG: spv.globalVariable @[[NUMWORKGROUPSVAR:.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
-    // CHECK-DAG: spv.globalVariable @[[$LOCALINVOCATIONIDVAR:.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
-    // CHECK-DAG: spv.globalVariable @[[$WORKGROUPIDVAR:.*]] built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
+    // CHECK-DAG: spv.GlobalVariable @[[NUMWORKGROUPSVAR:.*]] built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
+    // CHECK-DAG: spv.GlobalVariable @[[$LOCALINVOCATIONIDVAR:.*]] built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+    // CHECK-DAG: spv.GlobalVariable @[[$WORKGROUPIDVAR:.*]] built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
     // CHECK-LABEL:    spv.func @load_store_kernel
     // CHECK-SAME: %[[ARG0:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 0)>}
     // CHECK-SAME: %[[ARG1:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}
     // CHECK-SAME: %[[ARG2:.*]]: !spv.ptr<!spv.struct<(!spv.array<48 x f32, stride=4> [0])>, StorageBuffer> {spv.interface_var_abi = #spv.interface_var_abi<(0, 2)>}
     // CHECK-SAME: %[[ARG3:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 3), StorageBuffer>}
     // CHECK-SAME: %[[ARG4:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 4), StorageBuffer>}
     // CHECK-SAME: %[[ARG5:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>}
     // CHECK-SAME: %[[ARG6:.*]]: i32 {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}

mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir

 // RUN: mlir-opt %s -convert-gpu-launch-to-vulkan-launch | FileCheck %s
 
 // CHECK: %[[resource:.*]] = alloc() : memref<12xf32>
 // CHECK: %[[index:.*]] = constant 1 : index
 // CHECK: call @vulkanLaunch(%[[index]], %[[index]], %[[index]], %[[resource]]) {spirv_blob = "{{.*}}", spirv_entry_point = "kernel"}
 
 module attributes {gpu.container_module} {
   spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]> {
-    spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+    spv.GlobalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
     spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
       %2 = spv.constant 0 : i32
       %3 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
       %4 = spv.AccessChain %0[%2, %2] : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
       %5 = spv.Load "StorageBuffer" %4 : f32
       spv.Return
     }

mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir

#single_workgroup_reduction_trait = {
   ]
 }
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.3, [Shader, GroupNonUniformArithmetic], []>, {}>
 } {
 
-// CHECK:      spv.globalVariable
+// CHECK:      spv.GlobalVariable
 // CHECK-SAME: built_in("LocalInvocationId")
 
 // CHECK:      @single_workgroup_reduction
 // CHECK-SAME: (%[[INPUT:.+]]: !spv.ptr{{.+}}, %[[OUTPUT:.+]]: !spv.ptr{{.+}})
 
 // CHECK:        %[[ZERO:.+]] = spv.constant 0 : i32
 // CHECK:        %[[ID:.+]] = spv.Load "Input" %{{.+}} : vector<3xi32>
 // CHECK:        %[[X:.+]] = spv.CompositeExtract %[[ID]][0 : i32]

mlir/test/Conversion/SPIRVToLLVM/lower-host-to-llvm-calls.mlir

 // RUN: mlir-opt --lower-host-to-llvm %s | FileCheck %s
   
 module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]>, {max_compute_workgroup_invocations = 128 : i32, max_compute_workgroup_size = dense<[128, 128, 64]> : vector<3xi32>}>} {
 
   //       CHECK: llvm.mlir.global linkonce @__spv__foo_bar_arg_0_descriptor_set0_binding0() : !llvm.struct<(array<6 x i32>)>
   //       CHECK: llvm.func @__spv__foo_bar()
 
   //       CHECK: spv.module @__spv__foo
-  //       CHECK:   spv.globalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
+  //       CHECK:   spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
   //       CHECK:   spv.func @__spv__foo_bar
   
   //       CHECK:   spv.EntryPoint "GLCompute" @__spv__foo_bar
   //       CHECK:   spv.ExecutionMode @__spv__foo_bar "LocalSize", 1, 1, 1
 
   // CHECK-LABEL: @main
   //       CHECK:   %[[SRC:.*]] = llvm.extractvalue %{{.*}}[0] : !llvm.struct<(ptr<i32>, ptr<i32>, i64, array<1 x i64>, array<1 x i64>)>
   //  CHECK-NEXT:   %[[DEST:.*]] = llvm.mlir.addressof @__spv__foo_bar_arg_0_descriptor_set0_binding0 : !llvm.ptr<struct<(array<6 x i32>)>>
   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr<struct<(array<6 x i32>)>>, !llvm.ptr<i32>, i64, i1) -> ()
   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
   //  CHECK-NEXT:   llvm.mlir.constant(false) : i1
   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr<i32>, !llvm.ptr<struct<(array<6 x i32>)>>, i64, i1) -> ()
 
   spv.module @__spv__foo Logical GLSL450 requires #spv.vce<v1.0, [Shader], [SPV_KHR_variable_pointers]> {
-    spv.globalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
+    spv.GlobalVariable @bar_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
     spv.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv.mlir.addressof @bar_arg_0 : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
       spv.Return
     }
     spv.EntryPoint "GLCompute" @bar
     spv.ExecutionMode @bar "LocalSize", 1, 1, 1
   }
 

mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir

spv.func @access_chain_array(%arg0 : i32) "None" {
   // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
   // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %{{.*}}] : (!llvm.ptr<array<4 x array<4 x f32>>>, i32, i32) -> !llvm.ptr<array<4 x f32>>
   %1 = spv.AccessChain %0[%arg0] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, i32
   %2 = spv.Load "Function" %1 ["Volatile"] : !spv.array<4xf32>
   spv.Return
 }
 
 //===----------------------------------------------------------------------===//
-// spv.globalVariable and spv.mlir.addressof
+// spv.GlobalVariable and spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
   // CHECK: llvm.mlir.global external constant @var() : f32
-  spv.globalVariable @var : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var : !spv.ptr<f32, Input>
 }
 
 spv.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global private @struct() : !llvm.struct<packed (f32, array<10 x f32>)>
   // CHECK-LABEL: @func
   //       CHECK:   llvm.mlir.addressof @struct : !llvm.ptr<struct<packed (f32, array<10 x f32>)>>
-  spv.globalVariable @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
+  spv.GlobalVariable @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
   spv.func @func() "None" {
     %0 = spv.mlir.addressof @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
     spv.Return
   }
 }
 
 spv.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global external @bar_descriptor_set0_binding0() : i32
   // CHECK-LABEL: @foo
   //       CHECK:   llvm.mlir.addressof @bar_descriptor_set0_binding0 : !llvm.ptr<i32>
-  spv.globalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
+  spv.GlobalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
 }
 
 spv.module @name Logical GLSL450 {
   //       CHECK: llvm.mlir.global external @name_bar_descriptor_set0_binding0() : i32
   // CHECK-LABEL: @foo
   //       CHECK:   llvm.mlir.addressof @name_bar_descriptor_set0_binding0 : !llvm.ptr<i32>
-  spv.globalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
+  spv.GlobalVariable @bar bind(0, 0) : !spv.ptr<i32, StorageBuffer>
   spv.func @foo() "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
 }
 
 //===----------------------------------------------------------------------===//
 // spv.Load

mlir/test/Conversion/StandardToSPIRV/alloc.mlir

module attributes {
   func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
     %0 = alloc() : memref<4x5xf32, 3>
     %1 = load %0[%arg0, %arg1] : memref<4x5xf32, 3>
     store %1, %0[%arg0, %arg1] : memref<4x5xf32, 3>
     dealloc %0 : memref<4x5xf32, 3>
     return
   }
 }
-//     CHECK: spv.globalVariable @[[VAR:.+]] : !spv.ptr<!spv.struct<(!spv.array<20 x f32, stride=4>)>, Workgroup>
+//     CHECK: spv.GlobalVariable @[[VAR:.+]] : !spv.ptr<!spv.struct<(!spv.array<20 x f32, stride=4>)>, Workgroup>
 //     CHECK: func @alloc_dealloc_workgroup_mem
 // CHECK-NOT:   alloc
 //     CHECK:   %[[PTR:.+]] = spv.mlir.addressof @[[VAR]]
 //     CHECK:   %[[LOADPTR:.+]] = spv.AccessChain %[[PTR]]
 //     CHECK:   %[[VAL:.+]] = spv.Load "Workgroup" %[[LOADPTR]] : f32
 //     CHECK:   %[[STOREPTR:.+]] = spv.AccessChain %[[PTR]]
 //     CHECK:   spv.Store "Workgroup" %[[STOREPTR]], %[[VAL]] : f32
 // CHECK-NOT:   dealloc
func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
     %0 = alloc() : memref<4x5xi16, 3>
     %1 = load %0[%arg0, %arg1] : memref<4x5xi16, 3>
     store %1, %0[%arg0, %arg1] : memref<4x5xi16, 3>
     dealloc %0 : memref<4x5xi16, 3>
     return
   }
 }
 
-//       CHECK: spv.globalVariable @__workgroup_mem__{{[0-9]+}}
+//       CHECK: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 //  CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<20 x i32, stride=4>)>, Workgroup>
 // CHECK_LABEL: spv.func @alloc_dealloc_workgroup_mem
 //       CHECK:   %[[VAR:.+]] = spv.mlir.addressof @__workgroup_mem__0
 //       CHECK:   %[[LOC:.+]] = spv.SDiv
 //       CHECK:   %[[PTR:.+]] = spv.AccessChain %[[VAR]][%{{.+}}, %[[LOC]]]
 //       CHECK:   %{{.+}} = spv.Load "Workgroup" %[[PTR]] : i32
 //       CHECK:   %[[LOC:.+]] = spv.SDiv
 //       CHECK:   %[[PTR:.+]] = spv.AccessChain %[[VAR]][%{{.+}}, %[[LOC]]]

 {
   func @two_allocs() {
     %0 = alloc() : memref<4x5xf32, 3>
     %1 = alloc() : memref<2x3xi32, 3>
     return
   }
 }
 
-//  CHECK-DAG: spv.globalVariable @__workgroup_mem__{{[0-9]+}}
+//  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4>)>, Workgroup>
-//  CHECK-DAG: spv.globalVariable @__workgroup_mem__{{[0-9]+}}
+//  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<20 x f32, stride=4>)>, Workgroup>
 //      CHECK: spv.func @two_allocs()
 //      CHECK: spv.Return
 
 // -----
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
   }
 {
   func @two_allocs_vector() {
     %0 = alloc() : memref<4xvector<4xf32>, 3>
     %1 = alloc() : memref<2xvector<2xi32>, 3>
     return
   }
 }
 
-//  CHECK-DAG: spv.globalVariable @__workgroup_mem__{{[0-9]+}}
+//  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<2 x vector<2xi32>, stride=8>)>, Workgroup>
-//  CHECK-DAG: spv.globalVariable @__workgroup_mem__{{[0-9]+}}
+//  CHECK-DAG: spv.GlobalVariable @__workgroup_mem__{{[0-9]+}}
 // CHECK-SAME:   !spv.ptr<!spv.struct<(!spv.array<4 x vector<4xf32>, stride=16>)>, Workgroup>
 //      CHECK: spv.func @two_allocs_vector()
 //      CHECK: spv.Return
 
 
 // -----
 
 module attributes {

mlir/test/Dialect/SPIRV/IR/memory-ops.mlir

func @aligned_load_incorrect_attributes() -> () {
   // expected-error @+1 {{expected ']'}}
   %1 = spv.Load "Function" %0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @var0 : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
   // CHECK_LABEL: @simple_load
   spv.func @simple_load() -> () "None" {
     // CHECK: spv.Load "Input" {{%.*}} : f32
     %0 = spv.mlir.addressof @var0 : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.Return
   }
 }
func @aligned_store_incorrect_attributes(%arg0 : f32) -> () {
   // expected-error @+1 {{expected ']'}}
   spv.Store  "Function" %0, %arg0 ["Aligned", 4, 23] : f32
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @var0 : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
   spv.func @simple_store(%arg0 : f32) -> () "None" {
     %0 = spv.mlir.addressof @var0 : !spv.ptr<f32, Input>
     // CHECK: spv.Store  "Input" {{%.*}}, {{%.*}} : f32
     spv.Store  "Input" %0, %arg0 : f32
     spv.Return
   }
 }
 
func @variable_init_normal_constant() -> () {
   // CHECK: spv.Variable init(%0) : !spv.ptr<f32, Function>
   %1 = spv.Variable init(%0) : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @global : !spv.ptr<f32, Workgroup>
+  spv.GlobalVariable @global : !spv.ptr<f32, Workgroup>
   spv.func @variable_init_global_variable() -> () "None" {
     %0 = spv.mlir.addressof @global : !spv.ptr<f32, Workgroup>
     // CHECK: spv.Variable init({{.*}}) : !spv.ptr<!spv.ptr<f32, Workgroup>, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<!spv.ptr<f32, Workgroup>, Function>
     spv.Return
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.specConstant @sc = 42 : i32
   // CHECK-LABEL: @variable_init_spec_constant
   spv.func @variable_init_spec_constant() -> () "None" {
     %0 = spv.mlir.referenceof @sc : i32
     // CHECK: spv.Variable init(%0) : !spv.ptr<i32, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<i32, Function>
     spv.Return
   }
 }
 
 // -----
 
 func @variable_bind() -> () {
-  // expected-error @+1 {{cannot have 'descriptor_set' attribute (only allowed in spv.globalVariable)}}
+  // expected-error @+1 {{cannot have 'descriptor_set' attribute (only allowed in spv.GlobalVariable)}}
   %0 = spv.Variable bind(1, 2) : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
 func @variable_init_bind() -> () {
   %0 = spv.constant 4.0 : f32
-  // expected-error @+1 {{cannot have 'binding' attribute (only allowed in spv.globalVariable)}}
+  // expected-error @+1 {{cannot have 'binding' attribute (only allowed in spv.GlobalVariable)}}
   %1 = spv.Variable init(%0) {binding = 5 : i32} : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
 func @variable_builtin() -> () {
-  // expected-error @+1 {{cannot have 'built_in' attribute (only allowed in spv.globalVariable)}}
+  // expected-error @+1 {{cannot have 'built_in' attribute (only allowed in spv.GlobalVariable)}}
   %1 = spv.Variable built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Function>
   return
 }
 
 // -----
 
 func @expect_ptr_result_type(%arg0: f32) -> () {
   // expected-error @+1 {{expected spv.ptr type}}
   %0 = spv.Variable : f32
   return
 }
 
 // -----
 
 func @variable_init(%arg0: f32) -> () {
-  // expected-error @+1 {{op initializer must be the result of a constant or spv.globalVariable op}}
+  // expected-error @+1 {{op initializer must be the result of a constant or spv.GlobalVariable op}}
   %0 = spv.Variable init(%arg0) : !spv.ptr<f32, Function>
   return
 }
 
 // -----
 
 func @cannot_be_generic_storage_class(%arg0: f32) -> () {
-  // expected-error @+1 {{op can only be used to model function-level variables. Use spv.globalVariable for module-level variables}}
+  // expected-error @+1 {{op can only be used to model function-level variables. Use spv.GlobalVariable for module-level variables}}
   %0 = spv.Variable : !spv.ptr<f32, Generic>
   return
 }
 
 // -----
 
 func @copy_memory_incompatible_ptrs() {
   %0 = spv.Variable : !spv.ptr<f32, Function>

mlir/test/Dialect/SPIRV/IR/structure-ops.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
 // spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @access_chain() -> () "None" {
     %0 = spv.constant 1: i32
     // CHECK: [[VAR1:%.*]] = spv.mlir.addressof @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
     // CHECK-NEXT: spv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
     %1 = spv.mlir.addressof @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
     %2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>, i32, i32
     spv.Return
   }
 }
 
 // -----
 
 // Allow taking address of global variables in other module-like ops
-spv.globalVariable @var : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
+spv.GlobalVariable @var : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
 func @addressof() -> () {
   // CHECK: spv.mlir.addressof @var
   %1 = spv.mlir.addressof @var : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @foo() -> () "None" {
-    // expected-error @+1 {{expected spv.globalVariable symbol}}
+    // expected-error @+1 {{expected spv.GlobalVariable symbol}}
     %0 = spv.mlir.addressof @var2 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced global variable's type}}
     %0 = spv.mlir.addressof @var1 : !spv.ptr<f32, Input>
   }
 }
 
 // -----
 
spv.module Logical GLSL450 {
    spv.func @do_nothing() -> () "None" {
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_nothing
    spv.EntryPoint "GLCompute" @do_nothing
 }
 
 spv.module Logical GLSL450 {
-   spv.globalVariable @var2 : !spv.ptr<f32, Input>
-   spv.globalVariable @var3 : !spv.ptr<f32, Output>
+   spv.GlobalVariable @var2 : !spv.ptr<f32, Input>
+   spv.GlobalVariable @var3 : !spv.ptr<f32, Output>
    spv.func @do_something(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
      %1 = spv.Load "Input" %arg0 : f32
      spv.Store "Output" %arg1, %1 : f32
      spv.Return
    }
    // CHECK: spv.EntryPoint "GLCompute" @do_something, @var2, @var3
    spv.EntryPoint "GLCompute" @do_something, @var2, @var3
 }
spv.func @outer_func() -> () "None" {
     }
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
-// spv.globalVariable
+// spv.GlobalVariable
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var0 : !spv.ptr<f32, Input>
-  spv.globalVariable @var0 : !spv.ptr<f32, Input>
+  // CHECK: spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
 }
 
 // TODO: Fix test case after initialization with normal constant is addressed
 // spv.module Logical GLSL450 {
 //   %0 = spv.constant 4.0 : f32
 //   // CHECK1: spv.Variable init(%0) : !spv.ptr<f32, Private>
-//   spv.globalVariable @var1 init(%0) : !spv.ptr<f32, Private>
+//   spv.GlobalVariable @var1 init(%0) : !spv.ptr<f32, Private>
 // }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.specConstant @sc = 4.0 : f32
-  // CHECK: spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
-  spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
+  // CHECK: spv.GlobalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
+  spv.GlobalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
 }
 
 // -----
 
 // Allow initializers coming from other module-like ops
 spv.specConstant @sc = 4.0 : f32
-// CHECK: spv.globalVariable @var initializer(@sc)
-spv.globalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
+// CHECK: spv.GlobalVariable @var initializer(@sc)
+spv.GlobalVariable @var initializer(@sc) : !spv.ptr<f32, Private>
 
 // -----
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
-  spv.globalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
+  // CHECK: spv.GlobalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
+  spv.GlobalVariable @var0 bind(1, 2) : !spv.ptr<f32, Uniform>
 }
 
 // TODO: Fix test case after initialization with constant is addressed
 // spv.module Logical GLSL450 {
 //   %0 = spv.constant 4.0 : f32
-//   // CHECK1: spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
-//   spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
+//   // CHECK1: spv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
+//   spv.GlobalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 // }
 
 // -----
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
-  spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
-  // CHECK: spv.globalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
-  spv.globalVariable @var2 {built_in = "GlobalInvocationID"} : !spv.ptr<vector<3xi32>, Input>
+  // CHECK: spv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
+  // CHECK: spv.GlobalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @var2 {built_in = "GlobalInvocationID"} : !spv.ptr<vector<3xi32>, Input>
 }
 
 // -----
 
 // Allow in other module-like ops
 module {
-  // CHECK: spv.globalVariable
-  spv.globalVariable @var0 : !spv.ptr<f32, Input>
+  // CHECK: spv.GlobalVariable
+  spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   // expected-error @+1 {{expected spv.ptr type}}
-  spv.globalVariable @var0 : f32
+  spv.GlobalVariable @var0 : f32
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  // expected-error @+1 {{op initializer must be result of a spv.specConstant or spv.globalVariable op}}
-  spv.globalVariable @var0 initializer(@var1) : !spv.ptr<f32, Private>
+  // expected-error @+1 {{op initializer must be result of a spv.specConstant or spv.GlobalVariable op}}
+  spv.GlobalVariable @var0 initializer(@var1) : !spv.ptr<f32, Private>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   // expected-error @+1 {{storage class cannot be 'Generic'}}
-  spv.globalVariable @var0 : !spv.ptr<f32, Generic>
+  spv.GlobalVariable @var0 : !spv.ptr<f32, Generic>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   // expected-error @+1 {{storage class cannot be 'Function'}}
-  spv.globalVariable @var0 : !spv.ptr<f32, Function>
+  spv.GlobalVariable @var0 : !spv.ptr<f32, Function>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.func @foo() "None" {
     // expected-error @+1 {{op must appear in a module-like op's block}}
-    spv.globalVariable @var0 : !spv.ptr<f32, Input>
+    spv.GlobalVariable @var0 : !spv.ptr<f32, Input>
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.module

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict-resolution.mlir

 // Resolve conflicting funcOp and globalVariableOp.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
-// CHECK-NEXT:     spv.globalVariable @foo_1
+// CHECK-NEXT:     spv.GlobalVariable @foo_1
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 }
 
 // -----
 
 // Resolve conflicting funcOp and globalVariableOp and update the global variable's
 // references.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
-// CHECK-NEXT:     spv.globalVariable @foo_1
+// CHECK-NEXT:     spv.GlobalVariable @foo_1
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.mlir.addressof @foo_1
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @bar() -> f32 "None" {
     %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and funcOp and update the global variable's
 // references.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo_1
+// CHECK-NEXT:     spv.GlobalVariable @foo_1
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.mlir.addressof @foo_1
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @bar() -> f32 "None" {
     %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 

 }
 
 // -----
 
 // Resolve conflicting globalVariableOps.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo_1 bind(1, 0)
+// CHECK-NEXT:     spv.GlobalVariable @foo_1 bind(1, 0)
 
-// CHECK-NEXT:     spv.globalVariable @foo bind(2, 0)
+// CHECK-NEXT:     spv.GlobalVariable @foo bind(2, 0)
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(2, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(2, 0) : !spv.ptr<f32, Input>
 }
 }
 
 // -----
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo_1 built_in("GlobalInvocationId")
+// CHECK-NEXT:     spv.GlobalVariable @foo_1 built_in("GlobalInvocationId")
 
-// CHECK-NEXT:     spv.globalVariable @foo built_in("LocalInvocationId")
+// CHECK-NEXT:     spv.GlobalVariable @foo built_in("LocalInvocationId")
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @foo built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @foo built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and specConstantOp.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo_1
+// CHECK-NEXT:     spv.GlobalVariable @foo_1
 
 // CHECK-NEXT:     spv.specConstant @foo
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 spv.module Logical GLSL450 {
   spv.specConstant @foo = -5 : i32
 }
 }
 
 // -----
 
 // Resolve conflicting specConstantOp and globalVariableOp.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.specConstant @foo_1
 
-// CHECK-NEXT:     spv.globalVariable @foo
+// CHECK-NEXT:     spv.GlobalVariable @foo
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
   spv.specConstant @foo = -5 : i32
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and specConstantCompositeOp.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo_1
+// CHECK-NEXT:     spv.GlobalVariable @foo_1
 
 // CHECK-NEXT:     spv.specConstant @bar
 // CHECK-NEXT:     spv.specConstantComposite @foo (@bar, @bar)
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 spv.module Logical GLSL450 {
   spv.specConstant @bar = -5 : i32
   spv.specConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 }
 }
 
 // -----
 
 // Resolve conflicting globalVariableOp and specConstantComposite.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.specConstant @bar
 // CHECK-NEXT:     spv.specConstantComposite @foo_1 (@bar, @bar)
 
-// CHECK-NEXT:     spv.globalVariable @foo
+// CHECK-NEXT:     spv.GlobalVariable @foo
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
   spv.specConstant @bar = -5 : i32
   spv.specConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 }
 }

mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/deduplication.mlir

 // RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // Deduplicate 2 global variables with the same descriptor set and binding.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo
+// CHECK-NEXT:     spv.GlobalVariable @foo
 
 // CHECK-NEXT:     spv.func @use_foo
 // CHECK-NEXT:       spv.mlir.addressof @foo
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @use_bar
 // CHECK-NEXT:       spv.mlir.addressof @foo
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.FAdd
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @use_foo() -> f32 "None" {
     %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @bar bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @bar bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @use_bar() -> f32 "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     %2 = spv.FAdd %1, %1 : f32
     spv.ReturnValue %2 : f32
   }
 }
 }
 
 // -----
 
 // Deduplicate 2 global variables with the same descriptor set and binding but different types.
 
 // CHECK:      module {
 // CHECK-NEXT: spv.module Logical GLSL450 {
-// CHECK-NEXT:   spv.globalVariable @foo bind(1, 0)
+// CHECK-NEXT:   spv.GlobalVariable @foo bind(1, 0)
 
-// CHECK-NEXT:   spv.globalVariable @bar bind(1, 0)
+// CHECK-NEXT:   spv.GlobalVariable @bar bind(1, 0)
 
 // CHECK-NEXT:   spv.func @use_bar
 // CHECK-NEXT:     spv.mlir.addressof @bar
 // CHECK-NEXT:     spv.Load
 // CHECK-NEXT:     spv.ReturnValue
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo bind(1, 0) : !spv.ptr<i32, Input>
+  spv.GlobalVariable @foo bind(1, 0) : !spv.ptr<i32, Input>
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @bar bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @bar bind(1, 0) : !spv.ptr<f32, Input>
 
   spv.func @use_bar() -> f32 "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 }
 
 // -----
 
 // Deduplicate 2 global variables with the same built-in attribute.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
-// CHECK-NEXT:     spv.globalVariable @foo built_in("GlobalInvocationId")
+// CHECK-NEXT:     spv.GlobalVariable @foo built_in("GlobalInvocationId")
 // CHECK-NEXT:     spv.func @use_bar
 // CHECK-NEXT:       spv.mlir.addressof @foo
 // CHECK-NEXT:       spv.Load
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
-  spv.globalVariable @foo built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @foo built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @bar built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @bar built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 
   spv.func @use_bar() -> vector<3xi32> "None" {
     %0 = spv.mlir.addressof @bar : !spv.ptr<vector<3xi32>, Input>
     %1 = spv.Load "Input" %0 : vector<3xi32>
     spv.ReturnValue %1 : vector<3xi32>
   }
 }
 }

mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
 } {
 
 // CHECK-LABEL: spv.module
 spv.module Logical GLSL450 {
-  // CHECK-DAG:    spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<(f32 [0])>, StorageBuffer>
-  // CHECK-DAG:    spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+  // CHECK-DAG:    spv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<(f32 [0])>, StorageBuffer>
+  // CHECK-DAG:    spv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
   // CHECK:    spv.func [[FN:@.*]]()
   spv.func @kernel(
     %arg0: f32
            {spv.interface_var_abi = #spv.interface_var_abi<(0, 0), StorageBuffer>},
     %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x f32>)>, StorageBuffer>
            {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>}) "None"
   attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]

mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir

 // RUN: mlir-opt -spirv-lower-abi-attrs -verify-diagnostics %s -o - | FileCheck %s
 
 module attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.0, [Shader], [SPV_KHR_storage_buffer_storage_class]>, {}>
 } {
 
 // CHECK-LABEL: spv.module
 spv.module Logical GLSL450 {
-  // CHECK-DAG: spv.globalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
-  spv.globalVariable @__builtin_var_WorkgroupSize__ built_in("WorkgroupSize") : !spv.ptr<vector<3xi32>, Input>
-  // CHECK-DAG: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
-  spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
-  // CHECK-DAG: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
-  spv.globalVariable @__builtin_var_LocalInvocationId__ built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
-  // CHECK-DAG: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
-  spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
-  // CHECK-DAG: spv.globalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR2:@.*]] bind(0, 2) : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR3:@.*]] bind(0, 3) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR4:@.*]] bind(0, 4) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR5:@.*]] bind(0, 5) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
-  // CHECK-DAG: spv.globalVariable [[VAR6:@.*]] bind(0, 6) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize")
+  spv.GlobalVariable @__builtin_var_WorkgroupSize__ built_in("WorkgroupSize") : !spv.ptr<vector<3xi32>, Input>
+  // CHECK-DAG: spv.GlobalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups")
+  spv.GlobalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
+  // CHECK-DAG: spv.GlobalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId")
+  spv.GlobalVariable @__builtin_var_LocalInvocationId__ built_in("LocalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  // CHECK-DAG: spv.GlobalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId")
+  spv.GlobalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
+  // CHECK-DAG: spv.GlobalVariable [[VAR0:@.*]] bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[VAR1:@.*]] bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[VAR2:@.*]] bind(0, 2) : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32, stride=4>, stride=16> [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[VAR3:@.*]] bind(0, 3) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[VAR4:@.*]] bind(0, 4) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[VAR5:@.*]] bind(0, 5) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+  // CHECK-DAG: spv.GlobalVariable [[VAR6:@.*]] bind(0, 6) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
   // CHECK: spv.func [[FN:@.*]]()
   spv.func @load_store_kernel(
     %arg0: !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32>>)>, StorageBuffer>
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 0)>},
     %arg1: !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32>>)>, StorageBuffer>
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 1)>},
     %arg2: !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32>>)>, StorageBuffer>
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 2)>},

mlir/test/Dialect/SPIRV/Transforms/inlining.mlir

spv.func @calling_single_block_retval_func() -> i32 "None" {
     // CHECK-NEXT: spv.ReturnValue %[[CST]]
     spv.ReturnValue %0 : i32
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
+  spv.GlobalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
   spv.func @callee() "None" {
     %0 = spv.mlir.addressof @data : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
     %1 = spv.constant 0: i32
     %2 = spv.AccessChain %0[%1, %1] : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>, i32, i32
     spv.Branch ^next
 
   ^next:
     %3 = spv.constant 42: i32
spv.func @calling_loop_no_ret_func() "None" {
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  spv.globalVariable @arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
-  spv.globalVariable @arg_1 bind(0, 1) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+  spv.GlobalVariable @arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+  spv.GlobalVariable @arg_1 bind(0, 1) : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
 
   // CHECK: @inline_into_selection_region
   spv.func @inline_into_selection_region() "None" {
     %1 = spv.constant 0 : i32
     // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv.mlir.addressof @arg_0
     // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv.mlir.addressof @arg_1
     // CHECK-DAG: [[LOADPTR:%.*]] = spv.AccessChain [[ADDRESS_ARG0]]
     // CHECK: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOADPTR]]

mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir

 // RUN: mlir-opt -decorate-spirv-composite-type-layout -split-input-file -verify-diagnostics %s -o - | FileCheck %s
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
-  spv.globalVariable @var0 bind(0,1) : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
+  // CHECK: spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
+  spv.GlobalVariable @var0 bind(0,1) : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
 
-  // CHECK: spv.globalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<(!spv.array<64 x i32, stride=4> [0], f32 [256])>, StorageBuffer>
-  spv.globalVariable @var1 bind(0,2) : !spv.ptr<!spv.struct<(!spv.array<64xi32>, f32)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<(!spv.array<64 x i32, stride=4> [0], f32 [256])>, StorageBuffer>
+  spv.GlobalVariable @var1 bind(0,2) : !spv.ptr<!spv.struct<(!spv.array<64xi32>, f32)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var2 bind(1, 0) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64 x i32, stride=4> [0], f32 [256])> [0], i32 [260])>, StorageBuffer>
-  spv.globalVariable @var2 bind(1,0) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64xi32>, f32)>, i32)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var2 bind(1, 0) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64 x i32, stride=4> [0], f32 [256])> [0], i32 [260])>, StorageBuffer>
+  spv.GlobalVariable @var2 bind(1,0) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64xi32>, f32)>, i32)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var3 : !spv.ptr<!spv.struct<(!spv.array<16 x !spv.struct<(f32 [0], f32 [4], !spv.array<16 x f32, stride=4> [8])>, stride=72> [0], f32 [1152])>, StorageBuffer>
-  spv.globalVariable @var3 : !spv.ptr<!spv.struct<(!spv.array<16x!spv.struct<(f32, f32, !spv.array<16xf32>)>>, f32)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var3 : !spv.ptr<!spv.struct<(!spv.array<16 x !spv.struct<(f32 [0], f32 [4], !spv.array<16 x f32, stride=4> [8])>, stride=72> [0], f32 [1152])>, StorageBuffer>
+  spv.GlobalVariable @var3 : !spv.ptr<!spv.struct<(!spv.array<16x!spv.struct<(f32, f32, !spv.array<16xf32>)>>, f32)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var4 bind(1, 2) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(i1 [0], i8 [1], i16 [2], i32 [4], i64 [8])> [0], f32 [16], i1 [20])> [0], i1 [24])>, StorageBuffer>
-  spv.globalVariable @var4 bind(1,2) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(i1, i8, i16, i32, i64)>, f32, i1)>, i1)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var4 bind(1, 2) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(i1 [0], i8 [1], i16 [2], i32 [4], i64 [8])> [0], f32 [16], i1 [20])> [0], i1 [24])>, StorageBuffer>
+  spv.GlobalVariable @var4 bind(1,2) : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(i1, i8, i16, i32, i64)>, f32, i1)>, i1)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var5 bind(1, 3) : !spv.ptr<!spv.struct<(!spv.array<256 x f32, stride=4> [0])>, StorageBuffer>
-  spv.globalVariable @var5 bind(1,3) : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var5 bind(1, 3) : !spv.ptr<!spv.struct<(!spv.array<256 x f32, stride=4> [0])>, StorageBuffer>
+  spv.GlobalVariable @var5 bind(1,3) : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, StorageBuffer>
 
   spv.func @kernel() -> () "None" {
     %c0 = spv.constant 0 : i32
     // CHECK: {{%.*}} = spv.mlir.addressof @var0 : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
     %0 = spv.mlir.addressof @var0 : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
     // CHECK:  {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
     %1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>, i32
     spv.Return
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(i1 [0], i1 [1], f64 [8])> [0], i1 [16])> [0], i1 [24])> [0], i1 [32])> [0], i1 [40])>, Uniform>
-  spv.globalVariable @var0 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(i1, i1, f64)>, i1)>, i1)>, i1)>, i1)>, Uniform>
+  // CHECK: spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(i1 [0], i1 [1], f64 [8])> [0], i1 [16])> [0], i1 [24])> [0], i1 [32])> [0], i1 [40])>, Uniform>
+  spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(!spv.struct<(i1, i1, f64)>, i1)>, i1)>, i1)>, i1)>, Uniform>
 
-  // CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<(!spv.struct<(i16 [0], !spv.struct<(i1 [0], f64 [8])> [8], f32 [24])> [0], f32 [32])>, Uniform>
-  spv.globalVariable @var1 : !spv.ptr<!spv.struct<(!spv.struct<(i16, !spv.struct<(i1, f64)>, f32)>, f32)>, Uniform>
+  // CHECK: spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(!spv.struct<(i16 [0], !spv.struct<(i1 [0], f64 [8])> [8], f32 [24])> [0], f32 [32])>, Uniform>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(!spv.struct<(i16, !spv.struct<(i1, f64)>, f32)>, f32)>, Uniform>
 
-  // CHECK: spv.globalVariable @var2 : !spv.ptr<!spv.struct<(!spv.struct<(i16 [0], !spv.struct<(i1 [0], !spv.array<16 x !spv.array<16 x i64, stride=8>, stride=128> [8])> [8], f32 [2064])> [0], f32 [2072])>, Uniform>
-  spv.globalVariable @var2 : !spv.ptr<!spv.struct<(!spv.struct<(i16, !spv.struct<(i1, !spv.array<16x!spv.array<16xi64>>)>, f32)>, f32)>, Uniform>
+  // CHECK: spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(!spv.struct<(i16 [0], !spv.struct<(i1 [0], !spv.array<16 x !spv.array<16 x i64, stride=8>, stride=128> [8])> [8], f32 [2064])> [0], f32 [2072])>, Uniform>
+  spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(!spv.struct<(i16, !spv.struct<(i1, !spv.array<16x!spv.array<16xi64>>)>, f32)>, f32)>, Uniform>
 
-  // CHECK: spv.globalVariable @var3 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64 x i64, stride=8> [0], i1 [512])> [0], i1 [520])>, Uniform>
-  spv.globalVariable @var3 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64xi64>, i1)>, i1)>, Uniform>
+  // CHECK: spv.GlobalVariable @var3 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64 x i64, stride=8> [0], i1 [512])> [0], i1 [520])>, Uniform>
+  spv.GlobalVariable @var3 : !spv.ptr<!spv.struct<(!spv.struct<(!spv.array<64xi64>, i1)>, i1)>, Uniform>
 
-  // CHECK: spv.globalVariable @var4 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(i64 [0], i1 [8], i1 [9], i1 [10], i1 [11])> [8], i1 [24])>, Uniform>
-  spv.globalVariable @var4 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i64, i1, i1, i1, i1)>, i1)>, Uniform>
+  // CHECK: spv.GlobalVariable @var4 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(i64 [0], i1 [8], i1 [9], i1 [10], i1 [11])> [8], i1 [24])>, Uniform>
+  spv.GlobalVariable @var4 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i64, i1, i1, i1, i1)>, i1)>, Uniform>
 
-  // CHECK: spv.globalVariable @var5 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(i1 [0], i1 [1], i1 [2], i1 [3], i64 [8])> [8], i1 [24])>, Uniform>
-  spv.globalVariable @var5 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i1, i1, i1, i1, i64)>, i1)>, Uniform>
+  // CHECK: spv.GlobalVariable @var5 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(i1 [0], i1 [1], i1 [2], i1 [3], i64 [8])> [8], i1 [24])>, Uniform>
+  spv.GlobalVariable @var5 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i1, i1, i1, i1, i64)>, i1)>, Uniform>
 
-  // CHECK: spv.globalVariable @var6 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(i64 [0], i32 [8], i16 [12], i8 [14], i1 [15])> [8], i1 [24])>, Uniform>
-  spv.globalVariable @var6 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i64, i32, i16, i8, i1)>, i1)>, Uniform>
+  // CHECK: spv.GlobalVariable @var6 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(i64 [0], i32 [8], i16 [12], i8 [14], i1 [15])> [8], i1 [24])>, Uniform>
+  spv.GlobalVariable @var6 : !spv.ptr<!spv.struct<(i1, !spv.struct<(i64, i32, i16, i8, i1)>, i1)>, Uniform>
 
-  // CHECK: spv.globalVariable @var7 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(!spv.struct<(i1 [0], i64 [8])> [0], i1 [16])> [8], i1 [32])>, Uniform>
-  spv.globalVariable @var7 : !spv.ptr<!spv.struct<(i1, !spv.struct<(!spv.struct<(i1, i64)>, i1)>, i1)>, Uniform>
+  // CHECK: spv.GlobalVariable @var7 : !spv.ptr<!spv.struct<(i1 [0], !spv.struct<(!spv.struct<(i1 [0], i64 [8])> [0], i1 [16])> [8], i1 [32])>, Uniform>
+  spv.GlobalVariable @var7 : !spv.ptr<!spv.struct<(i1, !spv.struct<(!spv.struct<(i1, i64)>, i1)>, i1)>, Uniform>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<(vector<2xi32> [0], f32 [8])>, StorageBuffer>
-  spv.globalVariable @var0 : !spv.ptr<!spv.struct<(vector<2xi32>, f32)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(vector<2xi32> [0], f32 [8])>, StorageBuffer>
+  spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(vector<2xi32>, f32)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<(vector<3xi32> [0], f32 [12])>, StorageBuffer>
-  spv.globalVariable @var1 : !spv.ptr<!spv.struct<(vector<3xi32>, f32)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(vector<3xi32> [0], f32 [12])>, StorageBuffer>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(vector<3xi32>, f32)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @var2 : !spv.ptr<!spv.struct<(vector<4xi32> [0], f32 [16])>, StorageBuffer>
-  spv.globalVariable @var2 : !spv.ptr<!spv.struct<(vector<4xi32>, f32)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(vector<4xi32> [0], f32 [16])>, StorageBuffer>
+  spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(vector<4xi32>, f32)>, StorageBuffer>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<(!spv.struct<()> [0])>, StorageBuffer>
-  spv.globalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<(!spv.struct<()>)>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<(!spv.struct<()> [0])>, StorageBuffer>
+  spv.GlobalVariable @emptyStructAsMember : !spv.ptr<!spv.struct<(!spv.struct<()>)>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @arrayType : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, StorageBuffer>
-  spv.globalVariable @arrayType : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, StorageBuffer>
+  // CHECK: spv.GlobalVariable @arrayType : !spv.ptr<!spv.array<4 x !spv.array<4 x f32>>, StorageBuffer>
+  spv.GlobalVariable @arrayType : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, StorageBuffer>
 
-  // CHECK: spv.globalVariable @InputStorage : !spv.ptr<!spv.struct<(!spv.array<256 x f32>)>, Input>
-  spv.globalVariable @InputStorage : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, Input>
+  // CHECK: spv.GlobalVariable @InputStorage : !spv.ptr<!spv.struct<(!spv.array<256 x f32>)>, Input>
+  spv.GlobalVariable @InputStorage : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, Input>
 
-  // CHECK: spv.globalVariable @customLayout : !spv.ptr<!spv.struct<(f32 [256], i32 [512])>, Uniform>
-  spv.globalVariable @customLayout : !spv.ptr<!spv.struct<(f32 [256], i32 [512])>, Uniform>
+  // CHECK: spv.GlobalVariable @customLayout : !spv.ptr<!spv.struct<(f32 [256], i32 [512])>, Uniform>
+  spv.GlobalVariable @customLayout : !spv.ptr<!spv.struct<(f32 [256], i32 [512])>, Uniform>
 
-  // CHECK:  spv.globalVariable @emptyStruct : !spv.ptr<!spv.struct<()>, Uniform>
-  spv.globalVariable @emptyStruct : !spv.ptr<!spv.struct<()>, Uniform>
+  // CHECK:  spv.GlobalVariable @emptyStruct : !spv.ptr<!spv.struct<()>, Uniform>
+  spv.GlobalVariable @emptyStruct : !spv.ptr<!spv.struct<()>, Uniform>
 }
 
 // -----
 
 spv.module Logical GLSL450 {
-  // CHECK: spv.globalVariable @var0 : !spv.ptr<!spv.struct<(i32 [0])>, PushConstant>
-  spv.globalVariable @var0 : !spv.ptr<!spv.struct<(i32)>, PushConstant>
-  // CHECK: spv.globalVariable @var1 : !spv.ptr<!spv.struct<(i32 [0])>, PhysicalStorageBuffer>
-  spv.globalVariable @var1 : !spv.ptr<!spv.struct<(i32)>, PhysicalStorageBuffer>
+  // CHECK: spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(i32 [0])>, PushConstant>
+  spv.GlobalVariable @var0 : !spv.ptr<!spv.struct<(i32)>, PushConstant>
+  // CHECK: spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(i32 [0])>, PhysicalStorageBuffer>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.struct<(i32)>, PhysicalStorageBuffer>
 }

mlir/test/Dialect/SPIRV/Transforms/vce-deduction.mlir

 // * Buffer requires ImageBuffer or SampledBuffer.
 // * Rg32f requires StorageImageExtendedFormats.
 // CHECK: requires #spv.vce<v1.0, [UniformAndStorageBuffer8BitAccess, StorageUniform16, Shader, ImageBuffer, StorageImageExtendedFormats], [SPV_KHR_8bit_storage, SPV_KHR_16bit_storage]>
 spv.module Logical GLSL450 attributes {
   spv.target_env = #spv.target_env<
     #spv.vce<v1.5, [Shader, UniformAndStorageBuffer8BitAccess, StorageBuffer16BitAccess, StorageUniform16, Int16, ImageBuffer, StorageImageExtendedFormats], []>,
     {}>
 } {
-  spv.globalVariable @data : !spv.ptr<!spv.struct<(i8 [0], f16 [2], i64 [4])>, Uniform>
-  spv.globalVariable @img  : !spv.ptr<!spv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Rg32f>, UniformConstant>
+  spv.GlobalVariable @data : !spv.ptr<!spv.struct<(i8 [0], f16 [2], i64 [4])>, Uniform>
+  spv.GlobalVariable @img  : !spv.ptr<!spv.image<f32, Buffer, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Rg32f>, UniformConstant>
 }

mlir/test/Target/SPIRV/array.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @array_stride(%arg0 : !spv.ptr<!spv.array<4x!spv.array<4xf32, stride=4>, stride=128>, StorageBuffer>, %arg1 : i32, %arg2 : i32) "None" {
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.array<4 x !spv.array<4 x f32, stride=4>, stride=128>, StorageBuffer>, i32, i32
     %2 = spv.AccessChain %arg0[%arg1, %arg2] : !spv.ptr<!spv.array<4x!spv.array<4xf32, stride=4>, stride=128>, StorageBuffer>, i32, i32
     spv.Return
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<f32, stride=4>, StorageBuffer>
-  spv.globalVariable @var0 : !spv.ptr<!spv.rtarray<f32, stride=4>, StorageBuffer>
-  // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
-  spv.globalVariable @var1 : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
+  // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.rtarray<f32, stride=4>, StorageBuffer>
+  spv.GlobalVariable @var0 : !spv.ptr<!spv.rtarray<f32, stride=4>, StorageBuffer>
+  // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.rtarray<vector<4xf16>>, Input>
 }

mlir/test/Target/SPIRV/debug.mlir

 // RUN: mlir-translate -test-spirv-roundtrip-debug -mlir-print-debuginfo -mlir-print-local-scope %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: loc({{".*debug.mlir"}}:5:3)
-  spv.globalVariable @var0 bind(0, 1) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<f32, Input>
   spv.func @arithmetic(%arg0 : vector<4xf32>, %arg1 : vector<4xf32>) "None" {
     // CHECK: loc({{".*debug.mlir"}}:8:10)
     %0 = spv.FAdd %arg0, %arg1 : vector<4xf32>
     // CHECK: loc({{".*debug.mlir"}}:10:10)
     %1 = spv.FNegate %arg0 : vector<4xf32>
     spv.Return
   }
 

mlir/test/Target/SPIRV/decorations.mlir

 // RUN: mlir-translate -split-input-file -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: location = 0 : i32
-  spv.globalVariable @var {location = 0 : i32} : !spv.ptr<vector<4xf32>, Input>
+  spv.GlobalVariable @var {location = 0 : i32} : !spv.ptr<vector<4xf32>, Input>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: no_perspective
-  spv.globalVariable @var {no_perspective} : !spv.ptr<vector<4xf32>, Input>
+  spv.GlobalVariable @var {no_perspective} : !spv.ptr<vector<4xf32>, Input>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: flat
-  spv.globalVariable @var {flat} : !spv.ptr<si32, Input>
+  spv.GlobalVariable @var {flat} : !spv.ptr<si32, Input>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: aliased
   // CHECK: aliased
-  spv.globalVariable @var1 bind(0, 0) {aliased} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
-  spv.globalVariable @var2 bind(0, 0) {aliased} : !spv.ptr<!spv.struct<(vector<4xf32>[0])>, StorageBuffer>
+  spv.GlobalVariable @var1 bind(0, 0) {aliased} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
+  spv.GlobalVariable @var2 bind(0, 0) {aliased} : !spv.ptr<!spv.struct<(vector<4xf32>[0])>, StorageBuffer>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: non_readable
-  spv.globalVariable @var bind(0, 0) {non_readable} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
+  spv.GlobalVariable @var bind(0, 0) {non_readable} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: non_writable
-  spv.globalVariable @var bind(0, 0) {non_writable} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
+  spv.GlobalVariable @var bind(0, 0) {non_writable} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: restrict
-  spv.globalVariable @var bind(0, 0) {restrict} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
+  spv.GlobalVariable @var bind(0, 0) {restrict} : !spv.ptr<!spv.struct<(!spv.array<4xf32, stride=4>[0])>, StorageBuffer>
 }
 

mlir/test/Target/SPIRV/entry-point.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @noop() -> () "None" {
     spv.Return
   }
   // CHECK:      spv.EntryPoint "GLCompute" @noop
   // CHECK-NEXT: spv.ExecutionMode @noop "ContractionOff"
   spv.EntryPoint "GLCompute" @noop
   spv.ExecutionMode @noop "ContractionOff"
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  // CHECK:       spv.globalVariable @var2 : !spv.ptr<f32, Input>
-  // CHECK-NEXT:  spv.globalVariable @var3 : !spv.ptr<f32, Output>
+  // CHECK:       spv.GlobalVariable @var2 : !spv.ptr<f32, Input>
+  // CHECK-NEXT:  spv.GlobalVariable @var3 : !spv.ptr<f32, Output>
   // CHECK-NEXT:  spv.func @noop({{%.*}}: !spv.ptr<f32, Input>, {{%.*}}: !spv.ptr<f32, Output>) "None"
   // CHECK:       spv.EntryPoint "GLCompute" @noop, @var2, @var3
-  spv.globalVariable @var2 : !spv.ptr<f32, Input>
-  spv.globalVariable @var3 : !spv.ptr<f32, Output>
+  spv.GlobalVariable @var2 : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var3 : !spv.ptr<f32, Output>
   spv.func @noop(%arg0 : !spv.ptr<f32, Input>, %arg1 : !spv.ptr<f32, Output>) -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @noop, @var2, @var3
   spv.ExecutionMode @noop "ContractionOff"
 }

mlir/test/Target/SPIRV/function-call.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  spv.globalVariable @var1 : !spv.ptr<!spv.array<4xf32>, Input>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.array<4xf32>, Input>
   spv.func @fmain() -> i32 "None" {
     %0 = spv.constant 16 : i32
     %1 = spv.mlir.addressof @var1 : !spv.ptr<!spv.array<4xf32>, Input>
     // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}) : (i32) -> i32
     %3 = spv.FunctionCall @f_0(%0) : (i32) -> i32
     // CHECK: spv.FunctionCall @f_1({{%.*}}, {{%.*}}) : (i32, !spv.ptr<!spv.array<4 x f32>, Input>) -> ()
     spv.FunctionCall @f_1(%3, %1) : (i32, !spv.ptr<!spv.array<4xf32>, Input>) ->  ()
     // CHECK: {{%.*}} =  spv.FunctionCall @f_2({{%.*}}) : (!spv.ptr<!spv.array<4 x f32>, Input>) -> !spv.ptr<!spv.array<4 x f32>, Input>

mlir/test/Target/SPIRV/global-variable.mlir

 // RUN: mlir-translate -test-spirv-roundtrip -split-input-file %s | FileCheck %s
 
-// CHECK:      spv.globalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
-// CHECK-NEXT: spv.globalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
-// CHECK-NEXT: spv.globalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
-// CHECK-NEXT: spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+// CHECK:      spv.GlobalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
+// CHECK-NEXT: spv.GlobalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
+// CHECK-NEXT: spv.GlobalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+// CHECK-NEXT: spv.GlobalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  spv.globalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
-  spv.globalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
-  spv.globalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr<vector<3xi32>, Input>
-  spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @var0 bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var1 bind(0, 1) : !spv.ptr<f32, Output>
+  spv.GlobalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  // CHECK:         spv.globalVariable @var1 : !spv.ptr<f32, Input>
-  // CHECK-NEXT:    spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
-  spv.globalVariable @var1 : !spv.ptr<f32, Input>
-  spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
+  // CHECK:         spv.GlobalVariable @var1 : !spv.ptr<f32, Input>
+  // CHECK-NEXT:    spv.GlobalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var1 : !spv.ptr<f32, Input>
+  spv.GlobalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr<f32, Input>
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr<vector<3xi32>, Input>
   spv.func @foo() "None" {
     // CHECK: %[[ADDR:.*]] = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %0 = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %1 = spv.constant 0: i32
     // CHECK: spv.AccessChain %[[ADDR]]
     %2 = spv.AccessChain %0[%1] : !spv.ptr<vector<3xi32>, Input>, i32
     spv.Return
   }
 }

mlir/test/Target/SPIRV/image.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>, UniformConstant>
-  spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>, UniformConstant>
+  spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, SamplerUnknown, Unknown>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.image<si32, Cube, IsDepth, NonArrayed, SingleSampled, NeedSampler, R8ui>, UniformConstant>
-  spv.globalVariable @var1 : !spv.ptr<!spv.image<si32, Cube, IsDepth, NonArrayed, SingleSampled, NeedSampler, R8ui>, UniformConstant>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.image<si32, Cube, IsDepth, NonArrayed, SingleSampled, NeedSampler, R8ui>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.image<i32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>, UniformConstant>
-  spv.globalVariable @var2 : !spv.ptr<!spv.image<i32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>, UniformConstant>
+  spv.GlobalVariable @var2 : !spv.ptr<!spv.image<i32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>, UniformConstant>
 }

mlir/test/Target/SPIRV/loop.mlir

spv.func @main() -> () "None" {
     spv.Return
   }
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  spv.globalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
-  spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
+  spv.GlobalVariable @GV1 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
+  spv.GlobalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
   spv.func @loop_kernel() "None" {
     %0 = spv.mlir.addressof @GV1 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
     %1 = spv.constant 0 : i32
     %2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
     %3 = spv.mlir.addressof @GV2 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
     %5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
     %6 = spv.constant 4 : i32
     %7 = spv.constant 42 : i32

mlir/test/Target/SPIRV/matrix.mlir

spv.func @matrix_times_matrix_2(%arg0: !spv.matrix<3 x vector<2xf32>>, %arg1: !spv.matrix<2 x vector<3xf32>>) -> !spv.matrix<2 x vector<2xf32>> "None"{
     %result = spv.MatrixTimesMatrix %arg0, %arg1 : !spv.matrix<3 x vector<2xf32>>, !spv.matrix<2 x vector<3xf32>> -> !spv.matrix<2 x vector<2xf32>>
     spv.ReturnValue %result : !spv.matrix<2 x vector<2xf32>>
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, StorageBuffer>
-  spv.globalVariable @var0 : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, StorageBuffer>
+  // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, StorageBuffer>
+  spv.GlobalVariable @var0 : !spv.ptr<!spv.matrix<3 x vector<3xf32>>, StorageBuffer>
 
-  // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.matrix<2 x vector<3xf32>>, StorageBuffer>
-  spv.globalVariable @var1 : !spv.ptr<!spv.matrix<2 x vector<3xf32>>, StorageBuffer>
+  // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.matrix<2 x vector<3xf32>>, StorageBuffer>
+  spv.GlobalVariable @var1 : !spv.ptr<!spv.matrix<2 x vector<3xf32>>, StorageBuffer>
 
-  // CHECK: spv.globalVariable {{@.*}} : !spv.ptr<!spv.matrix<4 x vector<4xf16>>, StorageBuffer>
-  spv.globalVariable @var2 : !spv.ptr<!spv.matrix<4 x vector<4xf16>>, StorageBuffer>
+  // CHECK: spv.GlobalVariable {{@.*}} : !spv.ptr<!spv.matrix<4 x vector<4xf16>>, StorageBuffer>
+  spv.GlobalVariable @var2 : !spv.ptr<!spv.matrix<4 x vector<4xf16>>, StorageBuffer>
 }

mlir/test/Target/SPIRV/phi.mlir

// CHECK-NEXT:   spv.mlir.merge
   spv.EntryPoint "GLCompute" @main
 }
 
 // -----
 
 // Test nested loops with block arguments
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
-  spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
-  spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
+  spv.GlobalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
   spv.func @fmul_kernel() "None" {
     %3 = spv.constant 12 : i32
     %4 = spv.constant 32 : i32
     %5 = spv.constant 4 : i32
     %6 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %7 = spv.Load "Input" %6 : vector<3xi32>
     %8 = spv.CompositeExtract %7[0 : i32] : vector<3xi32>
     %9 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>

mlir/test/Target/SPIRV/sampled-image.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.sampled_image<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, UniformConstant>
-  spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.sampled_image<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, UniformConstant>
+  spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.sampled_image<!spv.image<f32, Dim1D, NoDepth, NonArrayed, SingleSampled, NoSampler, Unknown>>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.sampled_image<!spv.image<si32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>>, UniformConstant>
-  spv.globalVariable @var1 bind(0, 0) : !spv.ptr<!spv.sampled_image<!spv.image<si32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>>, UniformConstant>
+  spv.GlobalVariable @var1 bind(0, 0) : !spv.ptr<!spv.sampled_image<!spv.image<si32, SubpassData, DepthUnknown, Arrayed, MultiSampled, NoSampler, Unknown>>, UniformConstant>
 
   // CHECK: !spv.ptr<!spv.sampled_image<!spv.image<i32, Rect, DepthUnknown, Arrayed, MultiSampled, NeedSampler, R8ui>>, UniformConstant>
-  spv.globalVariable @var2 bind(0, 0) : !spv.ptr<!spv.sampled_image<!spv.image<i32, Rect, DepthUnknown, Arrayed, MultiSampled, NeedSampler, R8ui>>, UniformConstant>
+  spv.GlobalVariable @var2 bind(0, 0) : !spv.ptr<!spv.sampled_image<!spv.image<i32, Rect, DepthUnknown, Arrayed, MultiSampled, NeedSampler, R8ui>>, UniformConstant>
 }

mlir/test/Target/SPIRV/struct.mlir

 // RUN: mlir-translate -test-spirv-roundtrip %s | FileCheck %s
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>
-  spv.globalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>
+  spv.GlobalVariable @var0 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [0], !spv.struct<(f32 [0], !spv.array<16 x f32, stride=4> [4])> [4])>, Input>
-  spv.globalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<(f32 [0], !spv.struct<(f32 [0], !spv.array<16 x f32, stride=4> [4])> [4])>, Input>
+  spv.GlobalVariable @var1 bind(0, 2) : !spv.ptr<!spv.struct<(f32 [0], !spv.struct<(f32 [0], !spv.array<16 x f32, stride=4> [4])> [4])>, Input>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38])>, StorageBuffer>
-  spv.globalVariable @var2 : !spv.ptr<!spv.struct<(f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38])>, StorageBuffer>
+  spv.GlobalVariable @var2 : !spv.ptr<!spv.struct<(f32 [0], i32 [4], f64 [8], i64 [16], f32 [24], i32 [30], f32 [34], i32 [38])>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<(!spv.array<128 x !spv.struct<(!spv.array<128 x f32, stride=4> [0])>, stride=512> [0])>, StorageBuffer>
-  spv.globalVariable @var3 : !spv.ptr<!spv.struct<(!spv.array<128 x !spv.struct<(!spv.array<128 x f32, stride=4> [0])>, stride=512> [0])>, StorageBuffer>
+  spv.GlobalVariable @var3 : !spv.ptr<!spv.struct<(!spv.array<128 x !spv.struct<(!spv.array<128 x f32, stride=4> [0])>, stride=512> [0])>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [0, NonWritable], i32 [4])>, StorageBuffer>
-  spv.globalVariable @var4 : !spv.ptr<!spv.struct<(f32 [0, NonWritable], i32 [4])>, StorageBuffer>
+  spv.GlobalVariable @var4 : !spv.ptr<!spv.struct<(f32 [0, NonWritable], i32 [4])>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [NonWritable], i32 [NonWritable, NonReadable])>, StorageBuffer>
-  spv.globalVariable @var5 : !spv.ptr<!spv.struct<(f32 [NonWritable], i32 [NonWritable, NonReadable])>, StorageBuffer>
+  spv.GlobalVariable @var5 : !spv.ptr<!spv.struct<(f32 [NonWritable], i32 [NonWritable, NonReadable])>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<(f32 [0, NonWritable], i32 [4, NonWritable, NonReadable])>, StorageBuffer>
-  spv.globalVariable @var6 : !spv.ptr<!spv.struct<(f32 [0, NonWritable], i32 [4, NonWritable, NonReadable])>, StorageBuffer>
+  spv.GlobalVariable @var6 : !spv.ptr<!spv.struct<(f32 [0, NonWritable], i32 [4, NonWritable, NonReadable])>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<(!spv.matrix<3 x vector<3xf32>> [0, ColMajor, MatrixStride=16])>, StorageBuffer>
-  spv.globalVariable @var7 : !spv.ptr<!spv.struct<(!spv.matrix<3 x vector<3xf32>> [0, ColMajor, MatrixStride=16])>, StorageBuffer>
+  spv.GlobalVariable @var7 : !spv.ptr<!spv.struct<(!spv.matrix<3 x vector<3xf32>> [0, ColMajor, MatrixStride=16])>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<()>, StorageBuffer>
-  spv.globalVariable @empty : !spv.ptr<!spv.struct<()>, StorageBuffer>
+  spv.GlobalVariable @empty : !spv.ptr<!spv.struct<()>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<empty_struct, ()>, StorageBuffer>
-  spv.globalVariable @id_empty : !spv.ptr<!spv.struct<empty_struct, ()>, StorageBuffer>
+  spv.GlobalVariable @id_empty : !spv.ptr<!spv.struct<empty_struct, ()>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<test_id, (!spv.array<128 x f32, stride=4> [0])>, Input>
-  spv.globalVariable @id_var0 : !spv.ptr<!spv.struct<test_id, (!spv.array<128 x f32, stride=4> [0])>, Input>
+  spv.GlobalVariable @id_var0 : !spv.ptr<!spv.struct<test_id, (!spv.array<128 x f32, stride=4> [0])>, Input>
 
 
   // CHECK: !spv.ptr<!spv.struct<rec, (!spv.ptr<!spv.struct<rec>, StorageBuffer>)>, StorageBuffer>
-  spv.globalVariable @recursive_simple : !spv.ptr<!spv.struct<rec, (!spv.ptr<!spv.struct<rec>, StorageBuffer>)>, StorageBuffer>
+  spv.GlobalVariable @recursive_simple : !spv.ptr<!spv.struct<rec, (!spv.ptr<!spv.struct<rec>, StorageBuffer>)>, StorageBuffer>
 
   // CHECK: !spv.ptr<!spv.struct<a, (!spv.ptr<!spv.struct<b, (!spv.ptr<!spv.struct<a>, Uniform>)>, Uniform>)>, Uniform>
-  spv.globalVariable @recursive_2 : !spv.ptr<!spv.struct<a, (!spv.ptr<!spv.struct<b, (!spv.ptr<!spv.struct<a>, Uniform>)>, Uniform>)>, Uniform>
+  spv.GlobalVariable @recursive_2 : !spv.ptr<!spv.struct<a, (!spv.ptr<!spv.struct<b, (!spv.ptr<!spv.struct<a>, Uniform>)>, Uniform>)>, Uniform>
 
   // CHECK: !spv.ptr<!spv.struct<axx, (!spv.ptr<!spv.struct<bxx, (!spv.ptr<!spv.struct<axx>, Uniform>, !spv.ptr<!spv.struct<bxx>, Uniform>)>, Uniform>)>, Uniform>
-  spv.globalVariable @recursive_3 : !spv.ptr<!spv.struct<axx, (!spv.ptr<!spv.struct<bxx, (!spv.ptr<!spv.struct<axx>, Uniform>, !spv.ptr<!spv.struct<bxx>, Uniform>)>, Uniform>)>, Uniform>
+  spv.GlobalVariable @recursive_3 : !spv.ptr<!spv.struct<axx, (!spv.ptr<!spv.struct<bxx, (!spv.ptr<!spv.struct<axx>, Uniform>, !spv.ptr<!spv.struct<bxx>, Uniform>)>, Uniform>)>, Uniform>
 
   // CHECK: !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>,
   // CHECK-SAME: !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Output>
   spv.func @kernel(%arg0: !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Input>, %arg1: !spv.ptr<!spv.struct<(!spv.array<128 x f32, stride=4> [0])>, Output>) -> () "None" {
     spv.Return
   }
 }