[MLIR][SPIRV] Towards a consistent assembly scheme: GlobalVariable.

[ergawy]

Summary

This commit works towards the effort of defining a consistent assembly
scheme for the spv dialect ops. Currently there are some
inconsistencies in things like: using : and -> to express the return
type of an op.

Following is an initial list of rules proposed by Lei, we will refine
this list as we discuss during reviews:

• Operands:
  • Use , to separate SSA values in a list
  • Tightly associate type with :directly after the SSA value. : and the type can be omitted if it's clear (i.e., easy to deduce from other types)
  • Use keyword and = as prefix to make the meaning of certain SSA value clear
  • Use [] to mark the range for indexing SSA values. With [], the type for the SSA value that are inserted into can be placed after [].
• Result: use -> to associate the result type
• Attributes
  • Attributes defined in the grammar should be placed in [] , preferably directly after the op.
  • Prefer to use unquoted keyword instead of quoted string when possible (For IntEnumAttr it should be possible; for BitEnum it might not be possible to drop the quote given we can have | in the middle.)
• Use keyword like from/ into/etc. to make the assembly read nicer when possible
• Exceptions can exist

[ergawy]

@antiagainst, 2 points:

• I adopted the set of rules you suggested, we can refine them as we discuss on the review.
• I tried to define GlobalVariable using the declarative format, but it seems that is not possible due to the optional directives: bind and built_in. Let me know if I misunderstood something.

[mehdi_amini]

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

188

Using -> for separating the operands instead of : for the type seems the opposite of everywhere else in MLIR, isn't it?

[antiagainst]

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

188

Oh we still want to use : to associate types with operands. So not trying to replace : with -> everywhere. -> is for specifying result types. (Like how we represent result types in functions, etc.).

One thing that I find misleading w.r.t. types in custom assembly format is that it can be hard to figure out what a type is associated with, if they are all put at the very end, with a single : separating them with operands, especially if we have those types both for operands and results and wanting to omit types for some of them. Like

%0 : = dialect.op %a, %b, %c : typeA, typeB

It's hard to figure out what typeA and typeB are for (%a, %b, %c, or the result)?

For such cases, I'd think using : to immediately attach the type to the operand and use -> to mark result types would be nicer.

WDYT?

[mehdi_amini]

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

188

I'd just like to keep consistency with MLIR, what you're mentioning here could be a reasonable convention but I don't think it has a precedent so far.
What about starting a thread on Discourse about the "good practices and general convention for IR textual syntax"?

[antiagainst]

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

188

Sure will do that.

[ergawy]

We will halt working on this until a discussion on a good assembly scheme is discussed with a wider audience.

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

let description = [{
     must be an symbol defined from a constant instruction or other
     `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
-                    (`initializer(` symbol-ref-id `)`)?
+                    (`init=` symbol-ref-id)?
                     (`bind(` integer-literal, integer-literal `)`)?
                     (`built_in(` string-literal `)`)?
                     attribute-dict?
     ```
 
-    where `initializer` specifies initializer and `bind` specifies the
+    where `init` 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 init=@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
+    OptionalAttr<FlatSymbolRefAttr>:$init
   );
 
   let results = (outs);
 
   let builders = [
     OpBuilder<(ins "TypeAttr":$type, "ArrayRef<NamedAttribute>":$namedAttrs),
     [{
       $_state.addAttribute("type", type);

mlir/lib/Conversion/SPIRVToLLVM/SPIRVToLLVM.cpp

 public:
   using SPIRVToLLVMConversion<spirv::GlobalVariableOp>::SPIRVToLLVMConversion;
 
   LogicalResult
   matchAndRewrite(spirv::GlobalVariableOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {
     // Currently, there is no support of initialization with a constant value in
     // SPIR-V dialect. Specialization constants are not considered as well.
-    if (op.initializer())
+    if (op.init())
       return failure();
 
     auto srcType = op.type().cast<spirv::PointerType>();
     auto dstType = typeConverter.convertType(srcType.getPointeeType());
     if (!dstType)
       return failure();
 
     // Limit conversion to the current invocation only or `StorageBuffer`

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

 static constexpr const char kClusterSize[] = "cluster_size";
 static constexpr const char kControl[] = "control";
 static constexpr const char kDefaultValueAttrName[] = "default_value";
 static constexpr const char kExecutionScopeAttrName[] = "execution_scope";
 static constexpr const char kEqualSemanticsAttrName[] = "equal_semantics";
 static constexpr const char kFnNameAttrName[] = "fn";
 static constexpr const char kGroupOperationAttrName[] = "group_operation";
 static constexpr const char kIndicesAttrName[] = "indices";
-static constexpr const char kInitializerAttrName[] = "initializer";
+static constexpr const char kInitializerAttrName[] = "init";
 static constexpr const char kInterfaceAttrName[] = "interface";
 static constexpr const char kMemoryScopeAttrName[] = "memory_scope";
 static constexpr const char kSemanticsAttrName[] = "semantics";
 static constexpr const char kSpecIdAttrName[] = "spec_id";
 static constexpr const char kTypeAttrName[] = "type";
 static constexpr const char kUnequalSemanticsAttrName[] = "unequal_semantics";
 static constexpr const char kValueAttrName[] = "value";
 static constexpr const char kValuesAttrName[] = "values";
static ParseResult parseGlobalVariableOp(OpAsmParser &parser,
   if (parser.parseSymbolName(nameAttr, SymbolTable::getSymbolAttrName(),
                              state.attributes)) {
     return failure();
   }
 
   // Parse optional initializer
   if (succeeded(parser.parseOptionalKeyword(kInitializerAttrName))) {
     FlatSymbolRefAttr initSymbol;
-    if (parser.parseLParen() ||
+    if (parser.parseEqual() ||
         parser.parseAttribute(initSymbol, Type(), kInitializerAttrName,
-                              state.attributes) ||
-        parser.parseRParen())
+                              state.attributes))
       return failure();
   }
 
   if (parseVariableDecorations(parser, state)) {
     return failure();
   }
 
   Type type;
   auto loc = parser.getCurrentLocation();
-  if (parser.parseColonType(type)) {
+  if (parser.parseArrow() || parser.parseType(type)) {
     return failure();
   }
   if (!type.isa<spirv::PointerType>()) {
     return parser.emitError(loc, "expected spv.ptr type");
   }
   state.addAttribute(kTypeAttrName, TypeAttr::get(type));
 
   return success();
 }
 
 static void print(spirv::GlobalVariableOp varOp, OpAsmPrinter &printer) {
   auto *op = varOp.getOperation();
   SmallVector<StringRef, 4> elidedAttrs{
       spirv::attributeName<spirv::StorageClass>()};
   printer << spirv::GlobalVariableOp::getOperationName();
 
   // Print variable name.
   printer << ' ';
   printer.printSymbolName(varOp.sym_name());
   elidedAttrs.push_back(SymbolTable::getSymbolAttrName());
 
   // Print optional initializer
-  if (auto initializer = varOp.initializer()) {
-    printer << " " << kInitializerAttrName << '(';
+  if (auto initializer = varOp.init()) {
+    printer << " " << kInitializerAttrName << '=';
     printer.printSymbolName(initializer.getValue());
-    printer << ')';
     elidedAttrs.push_back(kInitializerAttrName);
   }
 
   elidedAttrs.push_back(kTypeAttrName);
   printVariableDecorations(op, printer, elidedAttrs);
-  printer << " : " << varOp.type();
+  printer << " -> " << varOp.type();
 }
 
 static LogicalResult verify(spirv::GlobalVariableOp 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."
   // Also, Function storage class is reserved by spv.Variable.
   auto storageClass = varOp.storageClass();

mlir/lib/Target/SPIRV/Serialization/SerializeOps.cpp

Serializer::processGlobalVariableOp(spirv::GlobalVariableOp varOp) {
   }
   globalVarIDMap[varName] = resultID;
   operands.push_back(resultID);
 
   // Encode StorageClass.
   operands.push_back(static_cast<uint32_t>(varOp.storageClass()));
 
   // Encode initialization.
-  if (auto initializer = varOp.initializer()) {
+  if (auto initializer = varOp.init()) {
     auto initializerID = getVariableID(initializer.getValue());
     if (!initializerID) {
       return emitError(varOp.getLoc(),
                        "invalid usage of undefined variable as initializer");
     }
     operands.push_back(initializerID);
-    elidedAttrs.push_back("initializer");
+    elidedAttrs.push_back("init");
   }
 
   (void)emitDebugLine(typesGlobalValues, varOp.getLoc());
   if (failed(encodeInstructionInto(typesGlobalValues, spirv::Opcode::OpVariable,
                                    operands))) {
-    elidedAttrs.push_back("initializer");
     return failure();
   }
 
   // Encode decorations.
   for (auto attr : varOp->getAttrs()) {
     if (llvm::any_of(elidedAttrs,
                      [&](StringRef elided) { return attr.first == elided; })) {
       continue;

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/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.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

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
   }
 }
 

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}}
     %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.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.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 init=@sc -> !spv.ptr<f32, Private>
+  spv.GlobalVariable @var init=@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 init=@sc
+spv.GlobalVariable @var init=@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>
 // }
 
 // -----
 
 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>
+  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>
+  spv.GlobalVariable @var0 init=@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

 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.
 

 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
   }
 }
 }

 
 // 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
   }
 }
 

 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.GlobalVariable @foo_1 bind(1, 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 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.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: }
 
 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.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>
 }
 }
 

 
 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

 // 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
   }
 }

 // 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
   }
 }
 }

 // 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>
+  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>
+  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>
+  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>
+  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 init=@var1 bind(1, 0) -> !spv.ptr<f32, Input>
+  spv.GlobalVariable @var1 -> !spv.ptr<f32, Input>
+  spv.GlobalVariable @var2 init=@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
   }
 }