[MLIR][SPIRV] Rename `spv._reference_of` to `spv.mlir.referenceof`

[ergawy]

Summary

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

mlir/docs/Dialects/SPIR-V.md

 #### 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.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
+    reference, too. `spv.mlir.referenceof` 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.
 
 #### Disallow implicit capturing in functions
 
 *   In SPIR-V specification, functions support implicit capturing: they can

 *   `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.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.mlir.referenceof` op to turn the symbol of the corresponding
     `spv.specConstant` into an SSA value.
 *   `OpPhi` instructions are converted to block arguments.
 *   Structured control flow are placed inside `spv.selection` and `spv.loop`.
 
 ## Conversions
 
 One of the main features of MLIR is the ability to progressively lower from
 dialects that capture programmer abstraction into dialects that are closer to a

mlir/docs/SPIRVToLLVMDialectConversion.md

 *   spv.CompositeConstruct
 *   spv.ControlBarrier
 *   spv.CopyMemory
 *   spv.FMod
 *   spv.GLSL.SAbs
 *   spv.GLSL.SSign
 *   spv.GLSL.FSign
 *   spv.MemoryBarrier
-*   spv._reference_of
+*   spv.mlir.referenceof
 *   spv.SMod
 *   spv.specConstant
 *   spv.SubgroupBallotKHR
 *   spv.Unreachable
 
 ## Control flow conversion
 
 ### Branch ops

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

def SPV_ModuleEndOp : SPV_Op<"_module_end", [InModuleScope, Terminator]> {
 
   let hasOpcode = 0;
 
   let autogenSerialization = 0;
 }
 
 // -----
 
-def SPV_ReferenceOfOp : SPV_Op<"_reference_of", [NoSideEffect]> {
+def SPV_ReferenceOfOp : SPV_Op<"mlir.referenceof", [NoSideEffect]> {
   let summary = "Reference a specialization constant.";
 
   let description = [{
     Specialization constants 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. This op's return type is
     the same as the specialization constant.
 
     <!-- End of AutoGen section -->
 
     ```
-    spv-reference-of-op ::= ssa-id `=` `spv._reference_of` symbol-ref-id
+    spv-reference-of-op ::= ssa-id `=` `spv.mlir.referenceof` symbol-ref-id
                                        `:` spirv-scalar-type
     ```
 
     #### Example:
 
     ```mlir
-    %0 = spv._reference_of @spec_const : f32
+    %0 = spv.mlir.referenceof @spec_const : f32
     ```
 
     TODO Add support for composite specialization constants.
   }];
 
   let arguments = (ins
     FlatSymbolRefAttr:$spec_const
   );

mlir/lib/Dialect/SPIRV/SPIRVOps.cpp

if (auto entryPointOp = dyn_cast<spirv::EntryPointOp>(op)) {
         }
     }
   }
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
-// spv._reference_of
+// spv.mlir.referenceof
 //===----------------------------------------------------------------------===//
 
 static LogicalResult verify(spirv::ReferenceOfOp referenceOfOp) {
   auto *specConstSym = SymbolTable::lookupNearestSymbolFrom(
       referenceOfOp.getParentOp(), referenceOfOp.spec_const());
   Type constType;
 
   auto specConstOp = dyn_cast_or_null<spirv::SpecConstantOp>(specConstSym);

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.mlir.addressof` and `spv._reference_of`) to turn an symbol into a
+  /// (`spv.mlir.addressof` and `spv.mlir.referenceof`) 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.

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

 // RUN: mlir-opt -test-spirv-module-combiner -split-input-file -verify-diagnostics %s | FileCheck %s
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.specConstant @m1_sc
 // CHECK-NEXT:     spv.specConstant @m2_sc
 // CHECK-NEXT:     spv.func @variable_init_spec_constant
-// CHECK-NEXT:       spv._reference_of @m2_sc
+// CHECK-NEXT:       spv.mlir.referenceof @m2_sc
 // CHECK-NEXT:       spv.Variable init
 // CHECK-NEXT:       spv.Return
 // CHECK-NEXT:     }
 // CHECK-NEXT:   }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
   spv.specConstant @m1_sc = 42.42 : f32
 }
 
 spv.module Logical GLSL450 {
   spv.specConstant @m2_sc = 42 : i32
   spv.func @variable_init_spec_constant() -> () "None" {
-    %0 = spv._reference_of @m2_sc : i32
+    %0 = spv.mlir.referenceof @m2_sc : i32
     %1 = spv.Variable init(%0) : !spv.ptr<i32, Function>
     spv.Return
   }
 }
 }
 
 // -----
 

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.specConstant @foo_1
 // CHECK-NEXT:     spv.func @bar
-// CHECK-NEXT:       spv._reference_of @foo_1
+// CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // 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.specConstant @foo = -5 : i32
 
   spv.func @bar() -> i32 "None" {
-    %0 = spv._reference_of @foo : i32 
+    %0 = spv.mlir.referenceof @foo : i32 
     spv.ReturnValue %0 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting specConstantOp and funcOp and update the spec constant's
 // references.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.specConstant @foo_1
 // CHECK-NEXT:     spv.func @bar
-// CHECK-NEXT:       spv._reference_of @foo_1
+// CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // 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.specConstant @foo = -5 : i32
 
   spv.func @bar() -> i32 "None" {
-    %0 = spv._reference_of @foo : i32
+    %0 = spv.mlir.referenceof @foo : i32
     spv.ReturnValue %0 : i32
   }
 }
 
 spv.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }

 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.specConstant @bar
 // CHECK-NEXT:     spv.specConstantComposite @foo_1 (@bar, @bar)
 // CHECK-NEXT:     spv.func @baz
-// CHECK-NEXT:       spv._reference_of @foo_1
+// CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // CHECK-NEXT:       spv.CompositeExtract
 // 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.specConstant @bar = -5 : i32
   spv.specConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 
   spv.func @baz() -> i32 "None" {
-    %0 = spv._reference_of @foo : !spv.array<2 x i32>
+    %0 = spv.mlir.referenceof @foo : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     spv.ReturnValue %1 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting specConstantCompositeOp and funcOp and update the spec
 // constant's references.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.specConstant @bar
 // CHECK-NEXT:     spv.specConstantComposite @foo_1 (@bar, @bar)
 // CHECK-NEXT:     spv.func @baz
-// CHECK-NEXT:       spv._reference_of @foo_1
+// CHECK-NEXT:       spv.mlir.referenceof @foo_1
 // CHECK-NEXT:       spv.CompositeExtract
 // 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.specConstant @bar = -5 : i32
   spv.specConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 
   spv.func @baz() -> i32 "None" {
-    %0 = spv._reference_of @foo : !spv.array<2 x i32>
+    %0 = spv.mlir.referenceof @foo : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     spv.ReturnValue %1 : i32
   }
 }
 
 spv.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32
   }
 }
 }
 
 // -----
 
 // Resolve conflicting spec constants and funcOps and update the spec constant's
 // references.
 
 // CHECK:      module {
 // CHECK-NEXT:   spv.module Logical GLSL450 {
 // CHECK-NEXT:     spv.specConstant @bar_1
 // CHECK-NEXT:     spv.specConstantComposite @foo_2 (@bar_1, @bar_1)
 // CHECK-NEXT:     spv.func @baz
-// CHECK-NEXT:       spv._reference_of @foo_2
+// CHECK-NEXT:       spv.mlir.referenceof @foo_2
 // CHECK-NEXT:       spv.CompositeExtract
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @foo
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 
 // CHECK-NEXT:     spv.func @bar
 // CHECK-NEXT:       spv.ReturnValue
 // CHECK-NEXT:     }
 // CHECK-NEXT: }
 
 module {
 spv.module Logical GLSL450 {
   spv.specConstant @bar = -5 : i32
   spv.specConstantComposite @foo (@bar, @bar) : !spv.array<2 x i32>
 
   spv.func @baz() -> i32 "None" {
-    %0 = spv._reference_of @foo : !spv.array<2 x i32>
+    %0 = spv.mlir.referenceof @foo : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     spv.ReturnValue %1 : i32
   }
 }
 
 spv.module Logical GLSL450 {
   spv.func @foo(%arg0 : i32) -> i32 "None" {
     spv.ReturnValue %arg0 : i32

mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir

 
 // -----
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.specConstant @condition_scalar = true
   spv.func @select() -> () "None" {
     %0 = spv.constant 4.0 : f32
     %1 = spv.constant 5.0 : f32
-    %2 = spv._reference_of @condition_scalar : i1
+    %2 = spv.mlir.referenceof @condition_scalar : i1
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, f32
     %3 = spv.Select %2, %0, %1 : i1, f32
     %4 = spv.constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32>
     %5 = spv.constant dense<[6.0, 7.0, 8.0, 9.0]> : vector<4xf32>
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, vector<4xf32>
     %6 = spv.Select %2, %4, %5 : i1, vector<4xf32>
     %7 = spv.constant dense<[true, true, true, true]> : vector<4xi1>
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : vector<4xi1>, vector<4xf32>
     %8 = spv.Select %7, %4, %5 : vector<4xi1>, vector<4xf32>
     spv.Return
   }
 }

mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir

spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK: spv.specConstant @sc_float spec_id(5) = 1.000000e+00 : f32
   spv.specConstant @sc_float spec_id(5) = 1. : f32
 
   // CHECK: spv.specConstantComposite @scc (@sc_int, @sc_int) : !spv.array<2 x i32>
   spv.specConstantComposite @scc (@sc_int, @sc_int) : !spv.array<2 x i32>
 
   // CHECK-LABEL: @use
   spv.func @use() -> (i32) "None" {
-    // We materialize a `spv._reference_of` op at every use of a
+    // We materialize a `spv.mlir.referenceof` op at every use of a
     // specialization constant in the deserializer. So two ops here.
-    // CHECK: %[[USE1:.*]] = spv._reference_of @sc_int : i32
-    // CHECK: %[[USE2:.*]] = spv._reference_of @sc_int : i32
+    // CHECK: %[[USE1:.*]] = spv.mlir.referenceof @sc_int : i32
+    // CHECK: %[[USE2:.*]] = spv.mlir.referenceof @sc_int : i32
     // CHECK: spv.IAdd %[[USE1]], %[[USE2]]
 
-    %0 = spv._reference_of @sc_int : i32
+    %0 = spv.mlir.referenceof @sc_int : i32
     %1 = spv.IAdd %0, %0 : i32
     spv.ReturnValue %1 : i32
   }
 
   // CHECK-LABEL: @use
   spv.func @use_composite() -> (i32) "None" {
-    // We materialize a `spv._reference_of` op at every use of a
+    // We materialize a `spv.mlir.referenceof` op at every use of a
     // specialization constant in the deserializer. So two ops here.
-    // CHECK: %[[USE1:.*]] = spv._reference_of @scc : !spv.array<2 x i32>
+    // CHECK: %[[USE1:.*]] = spv.mlir.referenceof @scc : !spv.array<2 x i32>
     // CHECK: %[[ITM0:.*]] = spv.CompositeExtract %[[USE1]][0 : i32] : !spv.array<2 x i32>
-    // CHECK: %[[USE2:.*]] = spv._reference_of @scc : !spv.array<2 x i32>
+    // CHECK: %[[USE2:.*]] = spv.mlir.referenceof @scc : !spv.array<2 x i32>
     // CHECK: %[[ITM1:.*]] = spv.CompositeExtract %[[USE2]][1 : i32] : !spv.array<2 x i32>
     // CHECK: spv.IAdd %[[ITM0]], %[[ITM1]]
 
-    %0 = spv._reference_of @scc : !spv.array<2 x i32>
+    %0 = spv.mlir.referenceof @scc : !spv.array<2 x i32>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<2 x i32>
     %2 = spv.CompositeExtract %0[1 : i32] : !spv.array<2 x i32>
     %3 = spv.IAdd %1, %2 : i32
     spv.ReturnValue %3 : i32
   }
 }
 
 // -----

mlir/test/Dialect/SPIRV/ops.mlir

 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.specConstant @sc = 42 : i32
   // CHECK-LABEL: @variable_init_spec_constant
   spv.func @variable_init_spec_constant() -> () "None" {
-    %0 = spv._reference_of @sc : i32
+    %0 = spv.mlir.referenceof @sc : i32
     // CHECK: spv.Variable init(%0) : !spv.ptr<i32, Function>
     %1 = spv.Variable init(%0) : !spv.ptr<i32, Function>
     spv.Return
   }
 }
 
 // -----
 

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

 func @module_end_not_in_module() -> () {
   // expected-error @+1 {{op must appear in a module-like op's block}}
   spv._module_end
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
-// spv._reference_of
+// spv.mlir.referenceof
 //===----------------------------------------------------------------------===//
 
 spv.module Logical GLSL450 {
   spv.specConstant @sc1 = false
   spv.specConstant @sc2 = 42 : i64
   spv.specConstant @sc3 = 1.5 : f32
 
   spv.specConstantComposite @scc (@sc1, @sc2, @sc3) : !spv.struct<(i1, i64, f32)>
 
   // CHECK-LABEL: @reference
   spv.func @reference() -> i1 "None" {
-    // CHECK: spv._reference_of @sc1 : i1
-    %0 = spv._reference_of @sc1 : i1
+    // CHECK: spv.mlir.referenceof @sc1 : i1
+    %0 = spv.mlir.referenceof @sc1 : i1
     spv.ReturnValue %0 : i1
   }
 
   // CHECK-LABEL: @reference_composite
   spv.func @reference_composite() -> i1 "None" {
-    // CHECK: spv._reference_of @scc : !spv.struct<(i1, i64, f32)>
-    %0 = spv._reference_of @scc : !spv.struct<(i1, i64, f32)>
+    // CHECK: spv.mlir.referenceof @scc : !spv.struct<(i1, i64, f32)>
+    %0 = spv.mlir.referenceof @scc : !spv.struct<(i1, i64, f32)>
     %1 = spv.CompositeExtract %0[0 : i32] : !spv.struct<(i1, i64, f32)>
     spv.ReturnValue %1 : i1
   }
 
   // CHECK-LABEL: @initialize
   spv.func @initialize() -> i64 "None" {
-    // CHECK: spv._reference_of @sc2 : i64
-    %0 = spv._reference_of @sc2 : i64
+    // CHECK: spv.mlir.referenceof @sc2 : i64
+    %0 = spv.mlir.referenceof @sc2 : i64
     %1 = spv.Variable init(%0) : !spv.ptr<i64, Function>
     %2 = spv.Load "Function" %1 : i64
     spv.ReturnValue %2 : i64
   }
 
   // CHECK-LABEL: @compute
   spv.func @compute() -> f32 "None" {
-    // CHECK: spv._reference_of @sc3 : f32
-    %0 = spv._reference_of @sc3 : f32
+    // CHECK: spv.mlir.referenceof @sc3 : f32
+    %0 = spv.mlir.referenceof @sc3 : f32
     %1 = spv.constant 6.0 : f32
     %2 = spv.FAdd %0, %1 : f32
     spv.ReturnValue %2 : f32
   }
 }
 
 // -----
 
 // Allow taking reference of spec constant in other module-like ops
 spv.specConstant @sc = 5 : i32
 func @reference_of() {
-  // CHECK: spv._reference_of @sc
-  %0 = spv._reference_of @sc : i32
+  // CHECK: spv.mlir.referenceof @sc
+  %0 = spv.mlir.referenceof @sc : i32
   return
 }
 
 // -----
 
 spv.specConstant @sc = 5 : i32
 spv.specConstantComposite @scc (@sc) : !spv.array<1 x i32>
 
 func @reference_of_composite() {
-  // CHECK: spv._reference_of @scc : !spv.array<1 x i32>
-  %0 = spv._reference_of @scc : !spv.array<1 x i32>
+  // CHECK: spv.mlir.referenceof @scc : !spv.array<1 x i32>
+  %0 = spv.mlir.referenceof @scc : !spv.array<1 x i32>
   %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<1 x i32>
   return
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{expected spv.specConstant or spv.SpecConstantComposite symbol}}
-    %0 = spv._reference_of @sc : i32
+    %0 = spv.mlir.referenceof @sc : i32
     spv.Return
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.specConstant @sc = 42 : i32
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}}
-    %0 = spv._reference_of @sc : f32
+    %0 = spv.mlir.referenceof @sc : f32
     spv.Return
   }
 }
 
 // -----
 
 spv.module Logical GLSL450 {
   spv.specConstant @sc = 42 : i32
   spv.specConstantComposite @scc (@sc) : !spv.array<1 x i32>
   spv.func @foo() -> () "None" {
     // expected-error @+1 {{result type mismatch with the referenced specialization constant's type}}
-    %0 = spv._reference_of @scc : f32
+    %0 = spv.mlir.referenceof @scc : f32
     spv.Return
   }
 }
 
 // -----
 
 //===----------------------------------------------------------------------===//
 // spv.specConstant