[MLIR][SPIRV] Rename `spv._address_of` to `spv.mlir.addressof`

[ergawy]

Summary

This commit does the renaming mentioned in the title in order to bring
spv dialect closer to the MLIR naming conventions.

[ergawy]

I think this should be actually renamed to spv.mlir.addressof to be truly consistent the style used in the llvm dialect. Are there any objections to this? If not, I will do as such.

[ergawy]

Rename to addressof instead of address_of.

[ergawy]

Fix mistake due to batch editing using sed.

[ergawy]

Fix mistake due to batch editing using sed.

[antiagainst]

Awesome, thanks for addressing this!!

mlir/docs/Dialects/SPIR-V.md

*   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
     generate SSA values. Instead they have symbols and should be referenced via
-    symbols. To use global variables in a function block, `spv._address_of` is
+    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._reference_of` is needed to turn the symbol into an SSA
     value for use in a function block.
 
 The above choices enables functions in the SPIR-V dialect to be isolated and
 explicit capturing.
*   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
     in module-level; otherwise they will be converted into `spv.Variable` ops.
 *   Every use of a module-level `OpVariable` instruction will materialize a
-    `spv._address_of` op to turn the symbol of the corresponding
+    `spv.mlir.addressof` op to turn the symbol of the corresponding
     `spv.globalVariable` into an SSA value.
 *   Every use of a `OpSpecConstant` instruction will materialize a
     `spv._reference_of` 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

 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
     the argument with this variable. The SSA value used for replacement is
-    obtained using the `spv._address_of` operation.
+    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
 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
-`spv._address_of` used to get a handle to the builtin as an SSA value.  The
+`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
 the builtin in the current `spv.module` if it does not exist already, and
-returns an SSA value generated from an `spv._address_of` operation.
+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
     operations conversions are implemented.
 *   [GPU Dialect][MlirGpuDialect] : A gpu.module is converted to a `spv.module`.

mlir/docs/SPIRVToLLVMDialectConversion.md

 * **Volatile**: mark the op as `volatile`, for example:
   ```mlir
   // %res = llvm.load volatile %ptr : !llvm.ptr<float>
   %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._address_of`
+#### `spv.globalVariable` and `spv.mlir.addressof`
 
 `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._address_of` and `llvm.mlir.addressof` ops that both return a pointer and
+`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.func @func() -> () "None" {
-    %0 = spv._address_of @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
+    %0 = spv.mlir.addressof @struct : !spv.ptr<!spv.struct<f32, !spv.array<10xf32>>, Private>
     spv.Return
   }
 }
 
 // Converted result
 module {
   llvm.mlir.global private @struct() : !llvm.struct<packed (float, [10 x float])>
   llvm.func @func() {

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

 
 include "mlir/Dialect/SPIRV/SPIRVBase.td"
 include "mlir/IR/SymbolInterfaces.td"
 include "mlir/Interfaces/CallInterfaces.td"
 include "mlir/Interfaces/SideEffectInterfaces.td"
 
 // -----
 
-def SPV_AddressOfOp : SPV_Op<"_address_of", [InFunctionScope, NoSideEffect]> {
+def SPV_AddressOfOp : SPV_Op<"mlir.addressof", [InFunctionScope, NoSideEffect]> {
   let summary = "Get the address of a global variable.";
 
   let description = [{
     Variables in module scope are defined using symbol names. This op generates
     an SSA value that can be used to refer to the symbol within function scope
     for use in ops that expect an SSA value. This operation has no corresponding
     SPIR-V instruction; it's merely used for modelling purpose in the SPIR-V
     dialect. Since variables in module scope in SPIR-V dialect are of pointer
     type, this op returns a pointer type as well, and the type is the same as
     the variable referenced.
 
     <!-- End of AutoGen section -->
 
     ```
-    spv-address-of-op ::= ssa-id `=` `spv._address_of` symbol-ref-id
+    spv-address-of-op ::= ssa-id `=` `spv.mlir.addressof` symbol-ref-id
                                      `:` spirv-pointer-type
     ```
 
     #### Example:
 
     ```mlir
-    %0 = spv._address_of @global_var : !spv.ptr<f32, Input>
+    %0 = spv.mlir.addressof @global_var : !spv.ptr<f32, Input>
     ```
   }];
 
   let arguments = (ins
     FlatSymbolRefAttr:$variable
   );
 
   let results = (outs

mlir/lib/Conversion/SPIRVToLLVM/ConvertSPIRVToLLVM.cpp

matchAndRewrite(spirv::ExecutionModeOp op, ArrayRef<Value> operands,
     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
 /// returns a pointer, whereas in LLVM dialect the global holds an actual value.
-/// This difference is handled by `spv._address_of` and `llvm.mlir.addressof`ops
-/// that both return a pointer.
+/// 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;
 
   LogicalResult
   matchAndRewrite(spirv::GlobalVariableOp op, ArrayRef<Value> operands,
                   ConversionPatternRewriter &rewriter) const override {

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

if (resultType != providedResultType) {
     return accessChainOp.emitOpError("invalid result type: expected ")
            << resultType << ", but provided " << providedResultType;
   }
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
-// spv._address_of
+// spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 void spirv::AddressOfOp::build(OpBuilder &builder, OperationState &state,
                                spirv::GlobalVariableOp var) {
   build(builder, state, var.type(), builder.getSymbolRefAttr(var));
 }
 
 static LogicalResult verify(spirv::AddressOfOp addressOfOp) {

mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp

private:
 
   //===--------------------------------------------------------------------===//
   // Instruction
   //===--------------------------------------------------------------------===//
 
   /// Get the Value associated with a result <id>.
   ///
   /// This method materializes normal constants and inserts "casting" ops
-  /// (`spv._address_of` and `spv._reference_of`) to turn an symbol into a SSA
-  /// value for handling uses of module scope constants/variables in functions.
+  /// (`spv.mlir.addressof` and `spv._reference_of`) to turn an symbol into a
+  /// SSA value for handling uses of module scope constants/variables in
+  /// functions.
   Value getValue(uint32_t id);
 
   /// Slices the first instruction out of `binary` and returns its opcode and
   /// operands via `opcode` and `operands` respectively. Returns failure if
   /// there is no more remaining instructions (`expectedOpcode` will be used to
   /// compose the error message) or the next instruction is malformed.
   LogicalResult
   sliceInstruction(spirv::Opcode &opcode, ArrayRef<uint32_t> &operands,

mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp-n

private:
 
   //===--------------------------------------------------------------------===//
   // Instruction
   //===--------------------------------------------------------------------===//
 
   /// Get the Value associated with a result <id>.
   ///
   /// This method materializes normal constants and inserts "casting" ops
-  /// (`spv._address_of` and `spv._reference_of`) to turn an symbol into a SSA
-  /// value for handling uses of module scope constants/variables in functions.
+  /// (`spv.mlir.addressof` and `spv._reference_of`) to turn an symbol into a
+  /// SSA value for handling uses of module scope constants/variables in
+  /// functions.
   Value getValue(uint32_t id);
 
   /// Slices the first instruction out of `binary` and returns its opcode and
   /// operands via `opcode` and `operands` respectively. Returns failure if
   /// there is no more remaining instructions (`expectedOpcode` will be used to
   /// compose the error message) or the next instruction is malformed.
   LogicalResult
   sliceInstruction(spirv::Opcode &opcode, ArrayRef<uint32_t> &operands,

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")
   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._address_of [[WORKGROUPID]]
+      // 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")
   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._address_of [[WORKGROUPID]]
+      // 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")
   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._address_of [[WORKGROUPID]]
+      // 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
     }
   }
 }
 
func @builtin() {
     return
   }
 
   // CHECK-LABEL:  spv.module @{{.*}} Logical GLSL450
   // 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._address_of [[LOCALINVOCATIONID]]
+      // 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")
   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._address_of [[NUMWORKGROUPS]]
+      // 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")
   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._address_of [[SUBGROUPID]]
+      // 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")
   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._address_of [[NUMSUBGROUPS]]
+      // 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")
   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._address_of [[SUBGROUPSIZE]]
+      // 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.module @kernels {
     // 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>}
     gpu.func @load_store_kernel(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref<12x4xf32>, %arg3: index, %arg4: index, %arg5: index, %arg6: index) kernel
       attributes {spv.entry_point_abi = {local_size = dense<[16, 1, 1]>: vector<3xi32>}} {
-      // CHECK: %[[ADDRESSWORKGROUPID:.*]] = spv._address_of @[[$WORKGROUPIDVAR]]
+      // CHECK: %[[ADDRESSWORKGROUPID:.*]] = spv.mlir.addressof @[[$WORKGROUPIDVAR]]
       // CHECK: %[[WORKGROUPID:.*]] = spv.Load "Input" %[[ADDRESSWORKGROUPID]]
       // CHECK: %[[WORKGROUPIDX:.*]] = spv.CompositeExtract %[[WORKGROUPID]]{{\[}}0 : i32{{\]}}
-      // CHECK: %[[ADDRESSLOCALINVOCATIONID:.*]] = spv._address_of @[[$LOCALINVOCATIONIDVAR]]
+      // CHECK: %[[ADDRESSLOCALINVOCATIONID:.*]] = spv.mlir.addressof @[[$LOCALINVOCATIONIDVAR]]
       // CHECK: %[[LOCALINVOCATIONID:.*]] = spv.Load "Input" %[[ADDRESSLOCALINVOCATIONID]]
       // CHECK: %[[LOCALINVOCATIONIDX:.*]] = spv.CompositeExtract %[[LOCALINVOCATIONID]]{{\[}}0 : i32{{\]}}
       %0 = "gpu.block_id"() {dimension = "x"} : () -> index
       %1 = "gpu.block_id"() {dimension = "y"} : () -> index
       %2 = "gpu.block_id"() {dimension = "z"} : () -> index
       %3 = "gpu.thread_id"() {dimension = "x"} : () -> index
       %4 = "gpu.thread_id"() {dimension = "y"} : () -> index
       %5 = "gpu.thread_id"() {dimension = "z"} : () -> index

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.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
-      %0 = spv._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+      %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._address_of @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
+      %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
     }
     spv.EntryPoint "GLCompute" @kernel
     spv.ExecutionMode @kernel "LocalSize", 1, 1, 1
   }
   gpu.module @kernels {

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

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-NEXT:   "llvm.intr.memcpy"(%[[DEST]], %[[SRC]], %[[SIZE:.*]], %{{.*}}) : (!llvm.ptr<struct<(array<6 x i32>)>>, !llvm.ptr<i32>, !llvm.i64, !llvm.i1) -> ()
   //  CHECK-NEXT:   llvm.call @__spv__foo_bar() : () -> ()
   //  CHECK-NEXT:   llvm.mlir.constant(false) : !llvm.i1
   //  CHECK-NEXT:   "llvm.intr.memcpy"(%[[SRC]], %[[DEST]], %[[SIZE]], %{{.*}}) : (!llvm.ptr<i32>, !llvm.ptr<struct<(array<6 x i32>)>>, !llvm.i64, !llvm.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.func @bar() "None" attributes {workgroup_attributions = 0 : i64} {
-      %0 = spv._address_of @bar_arg_0 : !spv.ptr<!spv.struct<(!spv.array<6 x i32, stride=4> [0])>, StorageBuffer>
+      %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
   }
 
   gpu.module @foo {
     gpu.func @bar(%arg0: memref<6xi32>) kernel attributes {spv.entry_point_abi = {local_size = dense<1> : vector<3xi32>}} {

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

spv.func @access_chain_array(%arg0 : i32) "None" {
   // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : !llvm.i32
   // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %{{.*}}] : (!llvm.ptr<array<4 x array<4 x float>>>, !llvm.i32, !llvm.i32) -> !llvm.ptr<array<4 x float>>
   %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._address_of
+// spv.globalVariable and spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
   // CHECK: llvm.mlir.global external constant @var() : !llvm.float
   spv.globalVariable @var : !spv.ptr<f32, Input>
 }
 
 spv.module Logical GLSL450 {
   //       CHECK: llvm.mlir.global private @struct() : !llvm.struct<packed (float, array<10 x float>)>
   // CHECK-LABEL: @func
   //       CHECK:   llvm.mlir.addressof @struct : !llvm.ptr<struct<packed (float, array<10 x float>)>>
   spv.globalVariable @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
   spv.func @func() "None" {
-    %0 = spv._address_of @struct : !spv.ptr<!spv.struct<(f32, !spv.array<10xf32>)>, Private>
+    %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() : !llvm.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.func @foo() "None" {
-    %0 = spv._address_of @bar : !spv.ptr<i32, StorageBuffer>
+    %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() : !llvm.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.func @foo() "None" {
-    %0 = spv._address_of @bar : !spv.ptr<i32, StorageBuffer>
+    %0 = spv.mlir.addressof @bar : !spv.ptr<i32, StorageBuffer>
     spv.Return
   }
 }
 
 //===----------------------------------------------------------------------===//
 // spv.Load
 //===----------------------------------------------------------------------===//
 

mlir/test/Conversion/StandardToSPIRV/alloc.mlir

func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
     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: func @alloc_dealloc_workgroup_mem
 // CHECK-NOT:   alloc
-//     CHECK:   %[[PTR:.+]] = spv._address_of @[[VAR]]
+//     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
 //     CHECK:   spv.Return
 
 // -----
func @alloc_dealloc_workgroup_mem(%arg0 : index, %arg1 : index) {
     dealloc %0 : memref<4x5xi16, 3>
     return
   }
 }
 
 //       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._address_of @__workgroup_mem__0
+//       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]]]
 //       CHECK:   %{{.+}} = spv.AtomicAnd "Workgroup" "AcquireRelease" %[[PTR]], %{{.+}} : !spv.ptr<i32, Workgroup>
 //       CHECK:   %{{.+}} = spv.AtomicOr "Workgroup" "AcquireRelease" %[[PTR]], %{{.+}} : !spv.ptr<i32, Workgroup>
 

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

 // 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.func @bar
-// CHECK-NEXT:       spv._address_of @foo_1
+// 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.func @bar() -> f32 "None" {
-    %0 = spv._address_of @foo : !spv.ptr<f32, Input>
+    %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.func @bar
-// CHECK-NEXT:       spv._address_of @foo_1
+// 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.func @bar() -> f32 "None" {
-    %0 = spv._address_of @foo : !spv.ptr<f32, Input>
+    %0 = spv.mlir.addressof @foo : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.ReturnValue %1 : f32
   }
 }
 
 spv.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32

mlir/test/Dialect/SPIRV/Serialization/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.func @fmain() -> i32 "None" {
     %0 = spv.constant 16 : i32
-    %1 = spv._address_of @var1 : !spv.ptr<!spv.array<4xf32>, Input>
+    %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>
     %4 = spv.FunctionCall @f_2(%1) : (!spv.ptr<!spv.array<4xf32>, Input>) -> !spv.ptr<!spv.array<4xf32>, Input>
     spv.ReturnValue %3 : i32
   }

mlir/test/Dialect/SPIRV/Serialization/global-variable.mlir

spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   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.func @foo() "None" {
-    // CHECK: %[[ADDR:.*]] = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
-    %0 = spv._address_of @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
+    // 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/Dialect/SPIRV/Serialization/loop.mlir

 }
 
 // -----
 
 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.func @loop_kernel() "None" {
-    %0 = spv._address_of @GV1 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
+    %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._address_of @GV2 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
+    %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
     %8 = spv.constant 2 : i32
 // CHECK:        spv.Branch ^bb1(%{{.*}} : i32)
 // CHECK-NEXT: ^bb1(%[[OUTARG:.*]]: i32):
 // CHECK-NEXT:   spv.loop {
     spv.loop {

mlir/test/Dialect/SPIRV/Serialization/phi.mlir

 
 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.func @fmul_kernel() "None" {
     %3 = spv.constant 12 : i32
     %4 = spv.constant 32 : i32
     %5 = spv.constant 4 : i32
-    %6 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
+    %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._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
+    %9 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %10 = spv.Load "Input" %9 : vector<3xi32>
     %11 = spv.CompositeExtract %10[1 : i32] : vector<3xi32>
-    %18 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
+    %18 = spv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
     %19 = spv.Load "Input" %18 : vector<3xi32>
     %20 = spv.CompositeExtract %19[0 : i32] : vector<3xi32>
-    %21 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
+    %21 = spv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
     %22 = spv.Load "Input" %21 : vector<3xi32>
     %23 = spv.CompositeExtract %22[1 : i32] : vector<3xi32>
     %30 = spv.IMul %11, %4 : i32
     %31 = spv.IMul %23, %4 : i32
 
 // CHECK:   spv.Branch ^[[FN_BB:.*]](%{{.*}} : i32)
 // CHECK: ^[[FN_BB]](%[[FN_BB_ARG:.*]]: i32):
 // CHECK:   spv.loop {

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

spv.module Logical GLSL450 {
   // 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._address_of [[VAR1]]
-    // CHECK: [[ADDRESSARG0:%.*]] = spv._address_of [[VAR0]]
+    // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]
+    // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]]
     // CHECK: [[CONST0:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG0PTR:%.*]] = spv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]]
     // CHECK: [[ARG0:%.*]] = spv.Load "StorageBuffer" [[ARG0PTR]]
     // CHECK: spv.Return
     spv.Return
   }
   // CHECK: spv.EntryPoint "GLCompute" [[FN]]
   // CHECK: spv.ExecutionMode [[FN]] "LocalSize", 32, 1, 1
 } // end spv.module
 
 } // end module

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

spv.func @load_store_kernel(
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 3), StorageBuffer>},
     %arg4: i32
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 4), StorageBuffer>},
     %arg5: i32
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 5), StorageBuffer>},
     %arg6: i32
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}) "None"
   attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
-    // CHECK: [[ADDRESSARG6:%.*]] = spv._address_of [[VAR6]]
+    // CHECK: [[ADDRESSARG6:%.*]] = spv.mlir.addressof [[VAR6]]
     // CHECK: [[CONST6:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]]
-    // CHECK: [[ADDRESSARG5:%.*]] = spv._address_of [[VAR5]]
+    // CHECK: [[ADDRESSARG5:%.*]] = spv.mlir.addressof [[VAR5]]
     // CHECK: [[CONST5:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG5PTR]]
-    // CHECK: [[ADDRESSARG4:%.*]] = spv._address_of [[VAR4]]
+    // CHECK: [[ADDRESSARG4:%.*]] = spv.mlir.addressof [[VAR4]]
     // CHECK: [[CONST4:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG4PTR:%.*]] = spv.AccessChain [[ADDRESSARG4]]{{\[}}[[CONST4]]
     // CHECK: [[ARG4:%.*]] = spv.Load "StorageBuffer" [[ARG4PTR]]
-    // CHECK: [[ADDRESSARG3:%.*]] = spv._address_of [[VAR3]]
+    // CHECK: [[ADDRESSARG3:%.*]] = spv.mlir.addressof [[VAR3]]
     // CHECK: [[CONST3:%.*]] = spv.constant 0 : i32
     // CHECK: [[ARG3PTR:%.*]] = spv.AccessChain [[ADDRESSARG3]]{{\[}}[[CONST3]]
     // CHECK: [[ARG3:%.*]] = spv.Load "StorageBuffer" [[ARG3PTR]]
-    // CHECK: [[ADDRESSARG2:%.*]] = spv._address_of [[VAR2]]
+    // CHECK: [[ADDRESSARG2:%.*]] = spv.mlir.addressof [[VAR2]]
     // CHECK: [[ARG2:%.*]] = spv.Bitcast [[ADDRESSARG2]]
-    // CHECK: [[ADDRESSARG1:%.*]] = spv._address_of [[VAR1]]
+    // CHECK: [[ADDRESSARG1:%.*]] = spv.mlir.addressof [[VAR1]]
     // CHECK: [[ARG1:%.*]] = spv.Bitcast [[ADDRESSARG1]]
-    // CHECK: [[ADDRESSARG0:%.*]] = spv._address_of [[VAR0]]
+    // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]]
     // CHECK: [[ARG0:%.*]] = spv.Bitcast [[ADDRESSARG0]]
-    %0 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
+    %0 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %1 = spv.Load "Input" %0 : vector<3xi32>
     %2 = spv.CompositeExtract %1[0 : i32] : vector<3xi32>
-    %3 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
+    %3 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %4 = spv.Load "Input" %3 : vector<3xi32>
     %5 = spv.CompositeExtract %4[1 : i32] : vector<3xi32>
-    %6 = spv._address_of @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
+    %6 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %7 = spv.Load "Input" %6 : vector<3xi32>
     %8 = spv.CompositeExtract %7[2 : i32] : vector<3xi32>
-    %9 = spv._address_of @__builtin_var_LocalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
+    %9 = spv.mlir.addressof @__builtin_var_LocalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
     %10 = spv.Load "Input" %9 : vector<3xi32>
     %11 = spv.CompositeExtract %10[0 : i32] : vector<3xi32>
-    %12 = spv._address_of @__builtin_var_LocalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
+    %12 = spv.mlir.addressof @__builtin_var_LocalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
     %13 = spv.Load "Input" %12 : vector<3xi32>
     %14 = spv.CompositeExtract %13[1 : i32] : vector<3xi32>
-    %15 = spv._address_of @__builtin_var_LocalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
+    %15 = spv.mlir.addressof @__builtin_var_LocalInvocationId__ : !spv.ptr<vector<3xi32>, Input>
     %16 = spv.Load "Input" %15 : vector<3xi32>
     %17 = spv.CompositeExtract %16[2 : i32] : vector<3xi32>
-    %18 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
+    %18 = spv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
     %19 = spv.Load "Input" %18 : vector<3xi32>
     %20 = spv.CompositeExtract %19[0 : i32] : vector<3xi32>
-    %21 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
+    %21 = spv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
     %22 = spv.Load "Input" %21 : vector<3xi32>
     %23 = spv.CompositeExtract %22[1 : i32] : vector<3xi32>
-    %24 = spv._address_of @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
+    %24 = spv.mlir.addressof @__builtin_var_NumWorkgroups__ : !spv.ptr<vector<3xi32>, Input>
     %25 = spv.Load "Input" %24 : vector<3xi32>
     %26 = spv.CompositeExtract %25[2 : i32] : vector<3xi32>
-    %27 = spv._address_of @__builtin_var_WorkgroupSize__ : !spv.ptr<vector<3xi32>, Input>
+    %27 = spv.mlir.addressof @__builtin_var_WorkgroupSize__ : !spv.ptr<vector<3xi32>, Input>
     %28 = spv.Load "Input" %27 : vector<3xi32>
     %29 = spv.CompositeExtract %28[0 : i32] : vector<3xi32>
-    %30 = spv._address_of @__builtin_var_WorkgroupSize__ : !spv.ptr<vector<3xi32>, Input>
+    %30 = spv.mlir.addressof @__builtin_var_WorkgroupSize__ : !spv.ptr<vector<3xi32>, Input>
     %31 = spv.Load "Input" %30 : vector<3xi32>
     %32 = spv.CompositeExtract %31[1 : i32] : vector<3xi32>
-    %33 = spv._address_of @__builtin_var_WorkgroupSize__ : !spv.ptr<vector<3xi32>, Input>
+    %33 = spv.mlir.addressof @__builtin_var_WorkgroupSize__ : !spv.ptr<vector<3xi32>, Input>
     %34 = spv.Load "Input" %33 : vector<3xi32>
     %35 = spv.CompositeExtract %34[2 : i32] : vector<3xi32>
     // CHECK: spv.IAdd [[ARG3]]
     %36 = spv.IAdd %arg3, %2 : i32
     // CHECK: spv.IAdd [[ARG4]]
     %37 = spv.IAdd %arg4, %11 : i32
     // CHECK: spv.AccessChain [[ARG0]]
     %c0 = spv.constant 0 : i32

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

spv.module Logical GLSL450 {
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.globalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
   spv.func @callee() "None" {
-    %0 = spv._address_of @data : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
+    %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.Store "StorageBuffer" %2, %3 : i32
     spv.Return
   }
 
   // CHECK-LABEL: @calling_multi_block_ret_func
   spv.func @calling_multi_block_ret_func() "None" {
-    // CHECK-NEXT:   spv._address_of
+    // CHECK-NEXT:   spv.mlir.addressof
     // CHECK-NEXT:   spv.constant 0
     // CHECK-NEXT:   spv.AccessChain
     // CHECK-NEXT:   spv.Branch ^bb1
     // CHECK-NEXT: ^bb1:
     // CHECK-NEXT:   spv.constant
     // CHECK-NEXT:   spv.Store
     // CHECK-NEXT:   spv.Branch ^bb2
     spv.FunctionCall @callee() : () -> ()

 
 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>
 
   // CHECK: @inline_into_selection_region
   spv.func @inline_into_selection_region() "None" {
     %1 = spv.constant 0 : i32
-    // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv._address_of @arg_0
-    // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv._address_of @arg_1
+    // 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]]
-    %2 = spv._address_of @arg_0 : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
-    %3 = spv._address_of @arg_1 : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+    %2 = spv.mlir.addressof @arg_0 : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
+    %3 = spv.mlir.addressof @arg_1 : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>
     %4 = spv.AccessChain %2[%1] : !spv.ptr<!spv.struct<(i32 [0])>, StorageBuffer>, i32
     %5 = spv.Load "StorageBuffer" %4 : i32
     %6 = spv.SGreaterThan %5, %1 : i32
     // CHECK: spv.selection
     spv.selection {
       spv.BranchConditional %6, ^bb1, ^bb2
     ^bb1: // pred: ^bb0
       // CHECK: [[STOREPTR:%.*]] = spv.AccessChain [[ADDRESS_ARG1]]

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

spv.module Logical GLSL450 {
   // 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>
 
   spv.func @kernel() -> () "None" {
     %c0 = spv.constant 0 : i32
-    // CHECK: {{%.*}} = spv._address_of @var0 : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
-    %0 = spv._address_of @var0 : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
+    // 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
   }
 }
 
 // -----
 

mlir/test/Dialect/SPIRV/ops.mlir

 
 // -----
 
 spv.module Logical GLSL450 {
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
   // CHECK_LABEL: @simple_load
   spv.func @simple_load() -> () "None" {
     // CHECK: spv.Load "Input" {{%.*}} : f32
-    %0 = spv._address_of @var0 : !spv.ptr<f32, Input>
+    %0 = spv.mlir.addressof @var0 : !spv.ptr<f32, Input>
     %1 = spv.Load "Input" %0 : f32
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
func @aligned_store_incorrect_attributes(%arg0 : f32) -> () {
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.globalVariable @var0 : !spv.ptr<f32, Input>
   spv.func @simple_store(%arg0 : f32) -> () "None" {
-    %0 = spv._address_of @var0 : !spv.ptr<f32, Input>
+    %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() -> () {
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.globalVariable @global : !spv.ptr<f32, Workgroup>
   spv.func @variable_init_global_variable() -> () "None" {
-    %0 = spv._address_of @global : !spv.ptr<f32, Workgroup>
+    %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/structure-ops.mlir

 // RUN: mlir-opt -allow-unregistered-dialect -split-input-file -verify-diagnostics %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
-// spv._address_of
+// spv.mlir.addressof
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
   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._address_of @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
+    // 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._address_of @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, 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>
-func @address_of() -> () {
-  // CHECK: spv._address_of @var
-  %1 = spv._address_of @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.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.globalVariable symbol}}
-    %0 = spv._address_of @var2 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
+    %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.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced global variable's type}}
-    %0 = spv._address_of @var1 : !spv.ptr<f32, Input>
+    %0 = spv.mlir.addressof @var1 : !spv.ptr<f32, Input>
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.constant
 //===----------------------------------------------------------------------===//