[Mlir-commits] [mlir] adf9f64 - [MLIR][SPIRV] Rename `spv._reference_of` to `spv.mlir.referenceof`
Lei Zhang
llvmlistbot at llvm.org
Wed Nov 18 10:29:10 PST 2020
Author: ergawy
Date: 2020-11-18T13:27:29-05:00
New Revision: adf9f64a02838a16d0c6cadfbd089254b126bd22
URL: https://github.com/llvm/llvm-project/commit/adf9f64a02838a16d0c6cadfbd089254b126bd22
DIFF: https://github.com/llvm/llvm-project/commit/adf9f64a02838a16d0c6cadfbd089254b126bd22.diff
LOG: [MLIR][SPIRV] Rename `spv._reference_of` to `spv.mlir.referenceof`
This commit does the renaming mentioned in the title in order to bring
'spv' dialect closer to the MLIR naming conventions.
Reviewed By: antiagainst
Differential Revision: https://reviews.llvm.org/D91715
Added:
Modified:
mlir/docs/Dialects/SPIR-V.md
mlir/docs/SPIRVToLLVMDialectConversion.md
mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
mlir/lib/Dialect/SPIRV/SPIRVOps.cpp
mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp
mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir
mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir
mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir
mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir
mlir/test/Dialect/SPIRV/ops.mlir
mlir/test/Dialect/SPIRV/structure-ops.mlir
Removed:
################################################################################
diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md
index 65eac4164962..1627beded6d0 100644
--- a/mlir/docs/Dialects/SPIR-V.md
+++ b/mlir/docs/Dialects/SPIR-V.md
@@ -182,7 +182,7 @@ instructions are represented in the SPIR-V dialect:
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
@@ -971,7 +971,7 @@ Similarly, a few transformations are performed during deserialization:
`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`.
diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md
index 127663601501..903f53f8fd7b 100644
--- a/mlir/docs/SPIRVToLLVMDialectConversion.md
+++ b/mlir/docs/SPIRVToLLVMDialectConversion.md
@@ -618,7 +618,7 @@ As well as:
* spv.GLSL.SSign
* spv.GLSL.FSign
* spv.MemoryBarrier
-* spv._reference_of
+* spv.mlir.referenceof
* spv.SMod
* spv.specConstant
* spv.SubgroupBallotKHR
diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
index 68c32936e511..9d768f684279 100644
--- a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
+++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td
@@ -468,7 +468,7 @@ def SPV_ModuleEndOp : SPV_Op<"_module_end", [InModuleScope, Terminator]> {
// -----
-def SPV_ReferenceOfOp : SPV_Op<"_reference_of", [NoSideEffect]> {
+def SPV_ReferenceOfOp : SPV_Op<"mlir.referenceof", [NoSideEffect]> {
let summary = "Reference a specialization constant.";
let description = [{
@@ -482,14 +482,14 @@ def SPV_ReferenceOfOp : SPV_Op<"_reference_of", [NoSideEffect]> {
<!-- 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.
diff --git a/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp b/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp
index ad57e9f3fadb..887d662d0040 100644
--- a/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp
+++ b/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp
@@ -2571,7 +2571,7 @@ static LogicalResult verify(spirv::ModuleOp moduleOp) {
}
//===----------------------------------------------------------------------===//
-// spv._reference_of
+// spv.mlir.referenceof
//===----------------------------------------------------------------------===//
static LogicalResult verify(spirv::ReferenceOfOp referenceOfOp) {
diff --git a/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp b/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp
index fbdf667cdf91..ea67db378798 100644
--- a/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp
+++ b/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp
@@ -399,7 +399,7 @@ class Deserializer {
/// 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);
diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir
index 07fd41e4fe86..32b2fcd375d7 100644
--- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir
+++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/basic.mlir
@@ -5,7 +5,7 @@
// 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: }
@@ -20,7 +20,7 @@ spv.module Logical GLSL450 {
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
}
diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir
index 2a3fde717285..4d9727cc11b3 100644
--- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir
+++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir
@@ -363,7 +363,7 @@ 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: }
@@ -379,7 +379,7 @@ 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
}
}
@@ -394,7 +394,7 @@ spv.module Logical GLSL450 {
// 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: }
@@ -408,7 +408,7 @@ 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
}
}
@@ -461,7 +461,7 @@ 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: }
@@ -479,7 +479,7 @@ spv.module Logical GLSL450 {
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
}
@@ -496,7 +496,7 @@ 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: }
@@ -512,7 +512,7 @@ spv.module Logical GLSL450 {
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
}
@@ -535,7 +535,7 @@ 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: }
@@ -555,7 +555,7 @@ spv.module Logical GLSL450 {
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
}
diff --git a/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir
index 77251e358741..5ef6227ca8be 100644
--- a/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir
@@ -87,7 +87,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
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>
diff --git a/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir b/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir
index 9dd9bb158b78..54b6e2a2eb12 100644
--- a/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir
+++ b/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir
@@ -17,28 +17,28 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// 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
diff --git a/mlir/test/Dialect/SPIRV/ops.mlir b/mlir/test/Dialect/SPIRV/ops.mlir
index f459eb0279f2..92cca0122cf7 100644
--- a/mlir/test/Dialect/SPIRV/ops.mlir
+++ b/mlir/test/Dialect/SPIRV/ops.mlir
@@ -1237,7 +1237,7 @@ 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
diff --git a/mlir/test/Dialect/SPIRV/structure-ops.mlir b/mlir/test/Dialect/SPIRV/structure-ops.mlir
index 0bfa2210fe4d..af1f5484d23a 100644
--- a/mlir/test/Dialect/SPIRV/structure-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/structure-ops.mlir
@@ -488,7 +488,7 @@ func @module_end_not_in_module() -> () {
// -----
//===----------------------------------------------------------------------===//
-// spv._reference_of
+// spv.mlir.referenceof
//===----------------------------------------------------------------------===//
spv.module Logical GLSL450 {
@@ -500,23 +500,23 @@ spv.module Logical GLSL450 {
// 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
@@ -524,8 +524,8 @@ spv.module Logical GLSL450 {
// 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
@@ -537,8 +537,8 @@ spv.module Logical GLSL450 {
// 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
}
@@ -548,8 +548,8 @@ 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
}
@@ -559,7 +559,7 @@ func @reference_of_composite() {
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
}
}
@@ -570,7 +570,7 @@ 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
}
}
@@ -582,7 +582,7 @@ spv.module Logical GLSL450 {
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
}
}
More information about the Mlir-commits
mailing list