diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md --- a/mlir/docs/Dialects/SPIR-V.md +++ b/mlir/docs/Dialects/SPIR-V.md @@ -182,7 +182,7 @@ 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 @@ `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 --- a/mlir/docs/SPIRVToLLVMDialectConversion.md +++ b/mlir/docs/SPIRVToLLVMDialectConversion.md @@ -618,7 +618,7 @@ * 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 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVStructureOps.td @@ -468,7 +468,7 @@ // ----- -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 @@ ``` - 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 --- a/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp +++ b/mlir/lib/Dialect/SPIRV/SPIRVOps.cpp @@ -2571,7 +2571,7 @@ } //===----------------------------------------------------------------------===// -// 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 --- a/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp +++ b/mlir/lib/Dialect/SPIRV/Serialization/Deserializer.cpp @@ -399,7 +399,7 @@ /// Get the Value associated with a result . /// /// 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 --- 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.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 spv.Return } diff --git a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir --- a/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir +++ b/mlir/test/Dialect/SPIRV/Linking/ModuleCombiner/conflict_resolution.mlir @@ -363,7 +363,7 @@ // 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.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 @@ // 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.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 @@ // 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.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 @@ // 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.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 @@ // 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.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 --- a/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/logical-ops.mlir @@ -87,7 +87,7 @@ 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 --- a/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/spec-constant.mlir @@ -17,28 +17,28 @@ // 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 --- a/mlir/test/Dialect/SPIRV/ops.mlir +++ b/mlir/test/Dialect/SPIRV/ops.mlir @@ -1237,7 +1237,7 @@ 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 %1 = spv.Variable init(%0) : !spv.ptr spv.Return diff --git a/mlir/test/Dialect/SPIRV/structure-ops.mlir b/mlir/test/Dialect/SPIRV/structure-ops.mlir --- a/mlir/test/Dialect/SPIRV/structure-ops.mlir +++ b/mlir/test/Dialect/SPIRV/structure-ops.mlir @@ -488,7 +488,7 @@ // ----- //===----------------------------------------------------------------------===// -// spv._reference_of +// spv.mlir.referenceof //===----------------------------------------------------------------------===// spv.module Logical GLSL450 { @@ -500,23 +500,23 @@ // 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 %2 = spv.Load "Function" %1 : i64 spv.ReturnValue %2 : i64 @@ -524,8 +524,8 @@ // 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 @@ // 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.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 @@ 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.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.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 } }