[Mlir-commits] [mlir] 5abdca4 - [MLIR][SPIRV] Rename `spv.constant` to `spv.Constant`.
Lei Zhang
llvmlistbot at llvm.org
Thu Mar 4 13:16:12 PST 2021
Author: KareemErgawy-TomTom
Date: 2021-03-04T16:15:56-05:00
New Revision: 5abdca47b3aa5518091c823db2be7b5a334d83fd
URL: https://github.com/llvm/llvm-project/commit/5abdca47b3aa5518091c823db2be7b5a334d83fd
DIFF: https://github.com/llvm/llvm-project/commit/5abdca47b3aa5518091c823db2be7b5a334d83fd.diff
LOG: [MLIR][SPIRV] Rename `spv.constant` to `spv.Constant`.
To unify the naming scheme across all ops in the SPIR-V dialect, we are
moving from `spv.camelCase` to `spv.CamelCase` everywhere.
Reviewed By: antiagainst
Differential Revision: https://reviews.llvm.org/D97917
Added:
Modified:
mlir/docs/Dialects/SPIR-V.md
mlir/docs/SPIRVToLLVMDialectConversion.md
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td
mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp
mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp
mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp
mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
mlir/test/Conversion/SCFToSPIRV/for.mlir
mlir/test/Conversion/SCFToSPIRV/if.mlir
mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
mlir/test/Dialect/SPIRV/IR/composite-ops.mlir
mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir
mlir/test/Dialect/SPIRV/IR/logical-ops.mlir
mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
mlir/test/Dialect/SPIRV/IR/structure-ops.mlir
mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir
mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
mlir/test/Target/SPIRV/constant.mlir
mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir
mlir/test/Target/SPIRV/debug.mlir
mlir/test/Target/SPIRV/function-call.mlir
mlir/test/Target/SPIRV/global-variable.mlir
mlir/test/Target/SPIRV/logical-ops.mlir
mlir/test/Target/SPIRV/loop.mlir
mlir/test/Target/SPIRV/memory-ops.mlir
mlir/test/Target/SPIRV/non-uniform-ops.mlir
mlir/test/Target/SPIRV/phi.mlir
mlir/test/Target/SPIRV/selection.mlir
mlir/test/Target/SPIRV/spec-constant.mlir
mlir/test/Target/SPIRV/undef.mlir
Removed:
################################################################################
diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md
index 295dad8102cf..2d9e6687b12a 100644
--- a/mlir/docs/Dialects/SPIR-V.md
+++ b/mlir/docs/Dialects/SPIR-V.md
@@ -88,7 +88,7 @@ The SPIR-V dialect adopts the following conventions for IR:
* Ops with `snake_case` names are those that have
diff erent representation
from corresponding instructions (or concepts) in the specification. These
ops are mostly for defining the SPIR-V structure. For example, `spv.module`
- and `spv.constant`. They may correspond to one or more instructions during
+ and `spv.Constant`. They may correspond to one or more instructions during
(de)serialization.
* Ops with `mlir.snake_case` names are those that have no corresponding
instructions (or concepts) in the binary format. They are introduced to
@@ -166,7 +166,7 @@ instructions are represented in the SPIR-V dialect:
#### Unify and localize constants
* Various normal constant instructions are represented by the same
- `spv.constant` op. Those instructions are just for constants of
diff erent
+ `spv.Constant` op. Those instructions are just for constants of
diff erent
types; using one op to represent them reduces IR verbosity and makes
transformations less tedious.
* Normal constants are not placed in `spv.module`'s region; they are localized
@@ -475,7 +475,7 @@ For example,
can be represented in the dialect as
```mlir
-%0 = "spv.constant"() { value = 42 : i32 } : () -> i32
+%0 = "spv.Constant"() { value = 42 : i32 } : () -> i32
%1 = "spv.Variable"(%0) { storage_class = "Function" } : (i32) -> !spv.ptr<i32, Function>
%2 = "spv.IAdd"(%0, %0) : (i32, i32) -> i32
```
@@ -581,9 +581,9 @@ It will be represented as
```mlir
func @selection(%cond: i1) -> () {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
- %two = spv.constant 2: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
+ %two = spv.Constant 2: i32
%x = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.selection {
@@ -669,8 +669,8 @@ It will be represented as
```mlir
func @loop(%count : i32) -> () {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.loop {
@@ -732,15 +732,15 @@ func @foo() -> () {
%var = spv.Variable : !spv.ptr<i32, Function>
spv.selection {
- %true = spv.constant true
+ %true = spv.Constant true
spv.BranchConditional %true, ^true, ^false
^true:
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
spv.Branch ^phi(%zero: i32)
^false:
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.Branch ^phi(%one: i32)
^phi(%arg: i32):
@@ -964,7 +964,7 @@ the representational
diff erences between SPIR-V dialect and binary format:
instructions.
* Types are serialized into `OpType*` instructions in the SPIR-V binary module
section for types, constants, and global variables.
-* `spv.constant`s are unified and placed in the SPIR-V binary module section
+* `spv.Constant`s are unified and placed in the SPIR-V binary module section
for types, constants, and global variables.
* Attributes on ops, if not part of the op's binary encoding, are emitted as
`OpDecorate*` instructions in the SPIR-V binary module section for
@@ -980,7 +980,7 @@ Similarly, a few transformations are performed during deserialization:
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
+* `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.
diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md
index f16bd46011a6..4623be9267d1 100644
--- a/mlir/docs/SPIRVToLLVMDialectConversion.md
+++ b/mlir/docs/SPIRVToLLVMDialectConversion.md
@@ -434,7 +434,7 @@ order to go through the pointer.
```mlir
// Access the 1st element of the array
-%i = spv.constant 1: i32
+%i = spv.Constant 1: i32
%var = spv.Variable : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>
%el = spv.AccessChain %var[%i, %i] : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>, i32, i32
@@ -529,7 +529,7 @@ also a function-level variable.
`spv.Variable` is modelled as `llvm.alloca` op. If initialized, an additional
store instruction is used. Note that there is no initialization for arrays and
structs since constants of these types are not supported in LLVM dialect (TODO).
-Also, at the moment initialization is only possible via `spv.constant`.
+Also, at the moment initialization is only possible via `spv.Constant`.
```mlir
// Conversion of VariableOp without initialization
@@ -538,7 +538,7 @@ Also, at the moment initialization is only possible via `spv.constant`.
// Conversion of VariableOp with initialization
%c = llvm.mlir.constant(0 : i64) : i64
-%c = spv.constant 0 : i64 %size = llvm.mlir.constant(1 : i32) : i32
+%c = spv.Constant 0 : i64 %size = llvm.mlir.constant(1 : i32) : i32
%res = spv.Variable init(%c) : !spv.ptr<i64, Function> => %res = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr<i64>
llvm.store %c, %res : !llvm.ptr<i64>
```
@@ -584,14 +584,14 @@ bitwidth. This leads to the following conversions:
%res1 = spv.ShiftRightArithmetic %0, %1 : i32, i16 => %res1 = llvm.ashr %0, %ext: i32
```
-### `spv.constant`
+### `spv.Constant`
-At the moment `spv.constant` conversion supports scalar and vector constants
+At the moment `spv.Constant` conversion supports scalar and vector constants
**only**.
#### Mapping
-`spv.constant` is mapped to `llvm.mlir.constant`. This is a straightforward
+`spv.Constant` is mapped to `llvm.mlir.constant`. This is a straightforward
conversion pattern with a special case when the argument is signed or unsigned.
#### Special case
@@ -608,10 +608,10 @@ cover all possible corner cases.
```mlir
// %0 = llvm.mlir.constant(0 : i8) : i8
-%0 = spv.constant 0 : i8
+%0 = spv.Constant 0 : i8
// %1 = llvm.mlir.constant(dense<[2, 3, 4]> : vector<3xi32>) : vector<3xi32>
-%1 = spv.constant dense<[2, 3, 4]> : vector<3xui32>
+%1 = spv.Constant dense<[2, 3, 4]> : vector<3xui32>
```
### Not implemented ops
@@ -672,7 +672,7 @@ blocks being reachable. Moreover, selection and loop control attributes (such as
```mlir
// Conversion of selection
-%cond = spv.constant true %cond = llvm.mlir.constant(true) : i1
+%cond = spv.Constant true %cond = llvm.mlir.constant(true) : i1
spv.selection {
spv.BranchConditional %cond, ^true, ^false llvm.cond_br %cond, ^true, ^false
@@ -693,7 +693,7 @@ spv.selection {
```mlir
// Conversion of loop
-%cond = spv.constant true %cond = llvm.mlir.constant(true) : i1
+%cond = spv.Constant true %cond = llvm.mlir.constant(true) : i1
spv.loop {
spv.Branch ^header llvm.br ^header
diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td
index 9f514a587e6e..0cffdb5a6d06 100644
--- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td
+++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td
@@ -58,7 +58,7 @@ def SPV_AccessChainOp : SPV_Op<"AccessChain", [NoSideEffect]> {
#### Example:
```mlir
- %0 = "spv.constant"() { value = 1: i32} : () -> i32
+ %0 = "spv.Constant"() { value = 1: i32} : () -> i32
%1 = spv.Variable : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>
%2 = spv.AccessChain %1[%0] : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>
%3 = spv.Load "Function" %2 ["Volatile"] : !spv.array<4xf32>
@@ -276,7 +276,7 @@ def SPV_VariableOp : SPV_Op<"Variable", []> {
#### Example:
```mlir
- %0 = spv.constant ...
+ %0 = spv.Constant ...
%1 = spv.Variable : !spv.ptr<f32, Function>
%2 = spv.Variable init(%0): !spv.ptr<f32, Function>
diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
index e7a9f44bdaeb..050c49294bd4 100644
--- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
+++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
@@ -254,7 +254,7 @@ def SPV_GroupNonUniformFAddOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFAdd "Workgroup" "Reduce" %scalar : f32
@@ -314,7 +314,7 @@ def SPV_GroupNonUniformFMaxOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFMax "Workgroup" "Reduce" %scalar : f32
@@ -374,7 +374,7 @@ def SPV_GroupNonUniformFMinOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFMin "Workgroup" "Reduce" %scalar : f32
@@ -431,7 +431,7 @@ def SPV_GroupNonUniformFMulOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFMul "Workgroup" "Reduce" %scalar : f32
@@ -486,7 +486,7 @@ def SPV_GroupNonUniformIAddOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformIAdd "Workgroup" "Reduce" %scalar : i32
@@ -541,7 +541,7 @@ def SPV_GroupNonUniformIMulOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformIMul "Workgroup" "Reduce" %scalar : i32
@@ -598,7 +598,7 @@ def SPV_GroupNonUniformSMaxOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformSMax "Workgroup" "Reduce" %scalar : i32
@@ -655,7 +655,7 @@ def SPV_GroupNonUniformSMinOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformSMin "Workgroup" "Reduce" %scalar : i32
@@ -713,7 +713,7 @@ def SPV_GroupNonUniformUMaxOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformUMax "Workgroup" "Reduce" %scalar : i32
@@ -771,7 +771,7 @@ def SPV_GroupNonUniformUMinOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformUMin "Workgroup" "Reduce" %scalar : i32
diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td
index bcc4acd83581..188b4d7ddbd8 100644
--- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td
+++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td
@@ -67,7 +67,7 @@ def SPV_AddressOfOp : SPV_Op<"mlir.addressof", [InFunctionScope, NoSideEffect]>
// -----
-def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> {
+def SPV_ConstantOp : SPV_Op<"Constant", [ConstantLike, NoSideEffect]> {
let summary = "The op that declares a SPIR-V normal constant";
let description = [{
@@ -81,7 +81,7 @@ def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> {
* ...
Having such a plethora of constant instructions renders IR transformations
- more tedious. Therefore, we use a single `spv.constant` op to represent
+ more tedious. Therefore, we use a single `spv.Constant` op to represent
them all. Note that conversion between those SPIR-V constant instructions
and this op is purely mechanical; so it can be scoped to the binary
(de)serialization process.
@@ -89,16 +89,16 @@ def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> {
<!-- End of AutoGen section -->
```
- spv-constant-op ::= ssa-id `=` `spv.constant` attribute-value
+ spv.Constant-op ::= ssa-id `=` `spv.Constant` attribute-value
(`:` spirv-type)?
```
#### Example:
```mlir
- %0 = spv.constant true
- %1 = spv.constant dense<[2, 3]> : vector<2xf32>
- %2 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
+ %0 = spv.Constant true
+ %1 = spv.Constant dense<[2, 3]> : vector<2xf32>
+ %2 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
```
TODO: support constant structs
@@ -572,7 +572,7 @@ def SPV_SpecConstantOp : SPV_Op<"SpecConstant", [InModuleScope, Symbol]> {
* `OpSpecConstantTrue` and `OpSpecConstantFalse` for boolean constants
* `OpSpecConstant` for scalar constants
- Similar as `spv.constant`, this op represents all of the above cases.
+ Similar as `spv.Constant`, this op represents all of the above cases.
`OpSpecConstantComposite` and `OpSpecConstantOp` are modelled with
separate ops.
@@ -731,8 +731,8 @@ def SPV_SpecConstantOperationOp : SPV_Op<"SpecConstantOperation", [
#### Example:
```mlir
- %0 = spv.constant 1: i32
- %1 = spv.constant 1: i32
+ %0 = spv.Constant 1: i32
+ %1 = spv.Constant 1: i32
%2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32
```
diff --git a/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp b/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp
index e07934b16fe3..6b115e69d4d5 100644
--- a/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp
+++ b/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp
@@ -356,7 +356,7 @@ class BitwiseOpPattern final : public OpConversionPattern<StdOp> {
}
};
-/// Converts composite std.constant operation to spv.constant.
+/// Converts composite std.constant operation to spv.Constant.
class ConstantCompositeOpPattern final
: public OpConversionPattern<ConstantOp> {
public:
@@ -367,7 +367,7 @@ class ConstantCompositeOpPattern final
ConversionPatternRewriter &rewriter) const override;
};
-/// Converts scalar std.constant operation to spv.constant.
+/// Converts scalar std.constant operation to spv.Constant.
class ConstantScalarOpPattern final : public OpConversionPattern<ConstantOp> {
public:
using OpConversionPattern<ConstantOp>::OpConversionPattern;
diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp
index 7f268ca92964..8c47a89d15a1 100644
--- a/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp
+++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp
@@ -138,11 +138,11 @@ OpFoldResult spirv::CompositeExtractOp::fold(ArrayRef<Attribute> operands) {
}
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
OpFoldResult spirv::ConstantOp::fold(ArrayRef<Attribute> operands) {
- assert(operands.empty() && "spv.constant has no operands");
+ assert(operands.empty() && "spv.Constant has no operands");
return value();
}
diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp
index cc5e2041f9b5..07fa245db91e 100644
--- a/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp
+++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp
@@ -941,7 +941,7 @@ static Type getElementPtrType(Type type, ValueRange indices, Location baseLoc) {
Operation *op = indexSSA.getDefiningOp();
if (!op) {
emitError(baseLoc, "'spv.AccessChain' op index must be an "
- "integer spv.constant to access "
+ "integer spv.Constant to access "
"element of spv.struct");
return nullptr;
}
@@ -950,7 +950,7 @@ static Type getElementPtrType(Type type, ValueRange indices, Location baseLoc) {
// integer literals of other bitwidths.
if (failed(extractValueFromConstOp(op, index))) {
emitError(baseLoc,
- "'spv.AccessChain' index must be an integer spv.constant to "
+ "'spv.AccessChain' index must be an integer spv.Constant to "
"access element of spv.struct, but provided ")
<< op->getName();
return nullptr;
@@ -1483,7 +1483,7 @@ static void print(spirv::CompositeInsertOp compositeInsertOp,
}
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
static ParseResult parseConstantOp(OpAsmParser &parser, OperationState &state) {
diff --git a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
index 7424f083d5b0..06e7f8136e96 100644
--- a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
+++ b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
@@ -38,7 +38,7 @@ static inline spirv::Opcode extractOpcode(uint32_t word) {
Value spirv::Deserializer::getValue(uint32_t id) {
if (auto constInfo = getConstant(id)) {
- // Materialize a `spv.constant` op at every use site.
+ // Materialize a `spv.Constant` op at every use site.
return opBuilder.create<spirv::ConstantOp>(unknownLoc, constInfo->second,
constInfo->first);
}
diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
index 34fc4af5e121..05adf87873ea 100644
--- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir
@@ -88,7 +88,7 @@ module attributes {gpu.container_module} {
// Note that this ignores the workgroup size specification in gpu.launch.
// We may want to define gpu.workgroup_size and convert it to the entry
// point ABI we want here.
- // CHECK: spv.constant 32 : i32
+ // CHECK: spv.Constant 32 : i32
%0 = "gpu.block_dim"() {dimension = "x"} : () -> index
gpu.return
}
@@ -110,7 +110,7 @@ module attributes {gpu.container_module} {
gpu.func @builtin_workgroup_size_y() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
// The constant value is obtained from the spv.entry_point_abi.
- // CHECK: spv.constant 4 : i32
+ // CHECK: spv.Constant 4 : i32
%0 = "gpu.block_dim"() {dimension = "y"} : () -> index
gpu.return
}
@@ -132,7 +132,7 @@ module attributes {gpu.container_module} {
gpu.func @builtin_workgroup_size_z() kernel
attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
// The constant value is obtained from the spv.entry_point_abi.
- // CHECK: spv.constant 1 : i32
+ // CHECK: spv.Constant 1 : i32
%0 = "gpu.block_dim"() {dimension = "z"} : () -> index
gpu.return
}
diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
index 40ce2d48c58f..de4af83b93ab 100644
--- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
+++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir
@@ -59,12 +59,12 @@ module attributes {
%12 = addi %arg3, %0 : index
// CHECK: %[[INDEX2:.*]] = spv.IAdd %[[ARG4]], %[[LOCALINVOCATIONIDX]]
%13 = addi %arg4, %3 : index
- // CHECK: %[[ZERO:.*]] = spv.constant 0 : i32
- // CHECK: %[[OFFSET1_0:.*]] = spv.constant 0 : i32
- // CHECK: %[[STRIDE1_1:.*]] = spv.constant 4 : i32
+ // CHECK: %[[ZERO:.*]] = spv.Constant 0 : i32
+ // CHECK: %[[OFFSET1_0:.*]] = spv.Constant 0 : i32
+ // CHECK: %[[STRIDE1_1:.*]] = spv.Constant 4 : i32
// CHECK: %[[UPDATE1_1:.*]] = spv.IMul %[[STRIDE1_1]], %[[INDEX1]] : i32
// CHECK: %[[OFFSET1_1:.*]] = spv.IAdd %[[OFFSET1_0]], %[[UPDATE1_1]] : i32
- // CHECK: %[[STRIDE1_2:.*]] = spv.constant 1 : i32
+ // CHECK: %[[STRIDE1_2:.*]] = spv.Constant 1 : i32
// CHECK: %[[UPDATE1_2:.*]] = spv.IMul %[[STRIDE1_2]], %[[INDEX2]] : i32
// CHECK: %[[OFFSET1_2:.*]] = spv.IAdd %[[OFFSET1_1]], %[[UPDATE1_2]] : i32
// CHECK: %[[PTR1:.*]] = spv.AccessChain %[[ARG0]]{{\[}}%[[ZERO]], %[[OFFSET1_2]]{{\]}}
diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
index d99d28e0a864..9b64d4937102 100644
--- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
+++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
@@ -9,7 +9,7 @@ module attributes {gpu.container_module} {
spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
%0 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
- %2 = spv.constant 0 : i32
+ %2 = spv.Constant 0 : i32
%3 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
%4 = spv.AccessChain %0[%2, %2] : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
%5 = spv.Load "StorageBuffer" %4 : f32
diff --git a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
index e81c4cd103a5..7a23e5389295 100644
--- a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
+++ b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
@@ -23,7 +23,7 @@ module attributes {
// CHECK: @single_workgroup_reduction
// CHECK-SAME: (%[[INPUT:.+]]: !spv.ptr{{.+}}, %[[OUTPUT:.+]]: !spv.ptr{{.+}})
-// CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
+// CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
// CHECK: %[[ID:.+]] = spv.Load "Input" %{{.+}} : vector<3xi32>
// CHECK: %[[X:.+]] = spv.CompositeExtract %[[ID]][0 : i32]
diff --git a/mlir/test/Conversion/SCFToSPIRV/for.mlir b/mlir/test/Conversion/SCFToSPIRV/for.mlir
index 3e4545b1b1b5..c0aa5ef6997e 100644
--- a/mlir/test/Conversion/SCFToSPIRV/for.mlir
+++ b/mlir/test/Conversion/SCFToSPIRV/for.mlir
@@ -6,11 +6,11 @@ module attributes {
} {
func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
- // CHECK: %[[LB:.*]] = spv.constant 4 : i32
+ // CHECK: %[[LB:.*]] = spv.Constant 4 : i32
%lb = constant 4 : index
- // CHECK: %[[UB:.*]] = spv.constant 42 : i32
+ // CHECK: %[[UB:.*]] = spv.Constant 42 : i32
%ub = constant 42 : index
- // CHECK: %[[STEP:.*]] = spv.constant 2 : i32
+ // CHECK: %[[STEP:.*]] = spv.Constant 2 : i32
%step = constant 2 : index
// CHECK: spv.loop {
// CHECK-NEXT: spv.Branch ^[[HEADER:.*]](%[[LB]] : i32)
@@ -18,15 +18,15 @@ func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
// CHECK: %[[CMP:.*]] = spv.SLessThan %[[INDVAR]], %[[UB]] : i32
// CHECK: spv.BranchConditional %[[CMP]], ^[[BODY:.*]], ^[[MERGE:.*]]
// CHECK: ^[[BODY]]:
- // CHECK: %[[ZERO1:.*]] = spv.constant 0 : i32
- // CHECK: %[[OFFSET1:.*]] = spv.constant 0 : i32
- // CHECK: %[[STRIDE1:.*]] = spv.constant 1 : i32
+ // CHECK: %[[ZERO1:.*]] = spv.Constant 0 : i32
+ // CHECK: %[[OFFSET1:.*]] = spv.Constant 0 : i32
+ // CHECK: %[[STRIDE1:.*]] = spv.Constant 1 : i32
// CHECK: %[[UPDATE1:.*]] = spv.IMul %[[STRIDE1]], %[[INDVAR]] : i32
// CHECK: %[[INDEX1:.*]] = spv.IAdd %[[OFFSET1]], %[[UPDATE1]] : i32
// CHECK: spv.AccessChain {{%.*}}{{\[}}%[[ZERO1]], %[[INDEX1]]{{\]}}
- // CHECK: %[[ZERO2:.*]] = spv.constant 0 : i32
- // CHECK: %[[OFFSET2:.*]] = spv.constant 0 : i32
- // CHECK: %[[STRIDE2:.*]] = spv.constant 1 : i32
+ // CHECK: %[[ZERO2:.*]] = spv.Constant 0 : i32
+ // CHECK: %[[OFFSET2:.*]] = spv.Constant 0 : i32
+ // CHECK: %[[STRIDE2:.*]] = spv.Constant 1 : i32
// CHECK: %[[UPDATE2:.*]] = spv.IMul %[[STRIDE2]], %[[INDVAR]] : i32
// CHECK: %[[INDEX2:.*]] = spv.IAdd %[[OFFSET2]], %[[UPDATE2]] : i32
// CHECK: spv.AccessChain {{%.*}}[%[[ZERO2]], %[[INDEX2]]]
@@ -44,15 +44,15 @@ func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
// CHECK-LABEL: @loop_yield
func @loop_yield(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
- // CHECK: %[[LB:.*]] = spv.constant 4 : i32
+ // CHECK: %[[LB:.*]] = spv.Constant 4 : i32
%lb = constant 4 : index
- // CHECK: %[[UB:.*]] = spv.constant 42 : i32
+ // CHECK: %[[UB:.*]] = spv.Constant 42 : i32
%ub = constant 42 : index
- // CHECK: %[[STEP:.*]] = spv.constant 2 : i32
+ // CHECK: %[[STEP:.*]] = spv.Constant 2 : i32
%step = constant 2 : index
- // CHECK: %[[INITVAR1:.*]] = spv.constant 0.000000e+00 : f32
+ // CHECK: %[[INITVAR1:.*]] = spv.Constant 0.000000e+00 : f32
%s0 = constant 0.0 : f32
- // CHECK: %[[INITVAR2:.*]] = spv.constant 1.000000e+00 : f32
+ // CHECK: %[[INITVAR2:.*]] = spv.Constant 1.000000e+00 : f32
%s1 = constant 1.0 : f32
// CHECK: %[[VAR1:.*]] = spv.Variable : !spv.ptr<f32, Function>
// CHECK: %[[VAR2:.*]] = spv.Variable : !spv.ptr<f32, Function>
diff --git a/mlir/test/Conversion/SCFToSPIRV/if.mlir b/mlir/test/Conversion/SCFToSPIRV/if.mlir
index d7c048f517ab..0db2dfcbfbd5 100644
--- a/mlir/test/Conversion/SCFToSPIRV/if.mlir
+++ b/mlir/test/Conversion/SCFToSPIRV/if.mlir
@@ -86,14 +86,14 @@ func @simple_if_yield(%arg2 : memref<10xf32>, %arg3 : i1) {
// CHECK: spv.selection {
// CHECK-NEXT: spv.BranchConditional {{%.*}}, [[TRUE:\^.*]], [[FALSE:\^.*]]
// CHECK-NEXT: [[TRUE]]:
- // CHECK: %[[RET1TRUE:.*]] = spv.constant 0.000000e+00 : f32
- // CHECK: %[[RET2TRUE:.*]] = spv.constant 1.000000e+00 : f32
+ // CHECK: %[[RET1TRUE:.*]] = spv.Constant 0.000000e+00 : f32
+ // CHECK: %[[RET2TRUE:.*]] = spv.Constant 1.000000e+00 : f32
// CHECK-DAG: spv.Store "Function" %[[VAR1]], %[[RET1TRUE]] : f32
// CHECK-DAG: spv.Store "Function" %[[VAR2]], %[[RET2TRUE]] : f32
// CHECK: spv.Branch ^[[MERGE:.*]]
// CHECK-NEXT: [[FALSE]]:
- // CHECK: %[[RET2FALSE:.*]] = spv.constant 2.000000e+00 : f32
- // CHECK: %[[RET1FALSE:.*]] = spv.constant 3.000000e+00 : f32
+ // CHECK: %[[RET2FALSE:.*]] = spv.Constant 2.000000e+00 : f32
+ // CHECK: %[[RET1FALSE:.*]] = spv.Constant 3.000000e+00 : f32
// CHECK-DAG: spv.Store "Function" %[[VAR1]], %[[RET1FALSE]] : f32
// CHECK-DAG: spv.Store "Function" %[[VAR2]], %[[RET2FALSE]] : f32
// CHECK: spv.Branch ^[[MERGE]]
diff --git a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
index 0ab431b3ac24..24c3aebdd0fc 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
@@ -1,61 +1,61 @@
// RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
// CHECK-LABEL: @bool_constant_scalar
spv.func @bool_constant_scalar() "None" {
// CHECK: llvm.mlir.constant(true) : i1
- %0 = spv.constant true
+ %0 = spv.Constant true
// CHECK: llvm.mlir.constant(false) : i1
- %1 = spv.constant false
+ %1 = spv.Constant false
spv.Return
}
// CHECK-LABEL: @bool_constant_vector
spv.func @bool_constant_vector() "None" {
// CHECK: llvm.mlir.constant(dense<[true, false]> : vector<2xi1>) : vector<2xi1>
- %0 = spv.constant dense<[true, false]> : vector<2xi1>
+ %0 = spv.Constant dense<[true, false]> : vector<2xi1>
// CHECK: llvm.mlir.constant(dense<false> : vector<3xi1>) : vector<3xi1>
- %1 = spv.constant dense<false> : vector<3xi1>
+ %1 = spv.Constant dense<false> : vector<3xi1>
spv.Return
}
// CHECK-LABEL: @integer_constant_scalar
spv.func @integer_constant_scalar() "None" {
// CHECK: llvm.mlir.constant(0 : i8) : i8
- %0 = spv.constant 0 : i8
+ %0 = spv.Constant 0 : i8
// CHECK: llvm.mlir.constant(-5 : i64) : i64
- %1 = spv.constant -5 : si64
+ %1 = spv.Constant -5 : si64
// CHECK: llvm.mlir.constant(10 : i16) : i16
- %2 = spv.constant 10 : ui16
+ %2 = spv.Constant 10 : ui16
spv.Return
}
// CHECK-LABEL: @integer_constant_vector
spv.func @integer_constant_vector() "None" {
// CHECK: llvm.mlir.constant(dense<[2, 3]> : vector<2xi32>) : vector<2xi32>
- %0 = spv.constant dense<[2, 3]> : vector<2xi32>
+ %0 = spv.Constant dense<[2, 3]> : vector<2xi32>
// CHECK: llvm.mlir.constant(dense<-4> : vector<2xi32>) : vector<2xi32>
- %1 = spv.constant dense<-4> : vector<2xsi32>
+ %1 = spv.Constant dense<-4> : vector<2xsi32>
// CHECK: llvm.mlir.constant(dense<[2, 3, 4]> : vector<3xi32>) : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xui32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xui32>
spv.Return
}
// CHECK-LABEL: @float_constant_scalar
spv.func @float_constant_scalar() "None" {
// CHECK: llvm.mlir.constant(5.000000e+00 : f16) : f16
- %0 = spv.constant 5.000000e+00 : f16
+ %0 = spv.Constant 5.000000e+00 : f16
// CHECK: llvm.mlir.constant(5.000000e+00 : f64) : f64
- %1 = spv.constant 5.000000e+00 : f64
+ %1 = spv.Constant 5.000000e+00 : f64
spv.Return
}
// CHECK-LABEL: @float_constant_vector
spv.func @float_constant_vector() "None" {
// CHECK: llvm.mlir.constant(dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>) : vector<2xf32>
- %0 = spv.constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>
+ %0 = spv.Constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>
spv.Return
}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
index 4211b47d6e66..0a2489f04c68 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
@@ -14,8 +14,8 @@ spv.module Logical GLSL450 {
}
spv.func @branch_with_arguments() -> () "None" {
- %0 = spv.constant 0 : i32
- %1 = spv.constant true
+ %0 = spv.Constant 0 : i32
+ %1 = spv.Constant true
// CHECK: llvm.br ^bb1(%{{.*}}, %{{.*}} : i32, i1)
spv.Branch ^label(%0, %1: i32, i1)
// CHECK: ^bb1(%{{.*}}: i32, %{{.*}}: i1)
@@ -33,7 +33,7 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @cond_branch_without_arguments() -> () "None" {
// CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
- %cond = spv.constant true
+ %cond = spv.Constant true
// CHECK: lvm.cond_br %[[COND]], ^bb1, ^bb2
spv.BranchConditional %cond, ^true, ^false
// CHECK: ^bb1:
@@ -46,10 +46,10 @@ spv.module Logical GLSL450 {
spv.func @cond_branch_with_arguments_nested() -> () "None" {
// CHECK: %[[COND1:.*]] = llvm.mlir.constant(true) : i1
- %cond = spv.constant true
- %0 = spv.constant 0 : i32
+ %cond = spv.Constant true
+ %0 = spv.Constant 0 : i32
// CHECK: %[[COND2:.*]] = llvm.mlir.constant(false) : i1
- %false = spv.constant false
+ %false = spv.Constant false
// CHECK: llvm.cond_br %[[COND1]], ^bb1(%{{.*}}, %[[COND2]] : i32, i1), ^bb2
spv.BranchConditional %cond, ^outer_true(%0, %false: i32, i1), ^outer_false
// CHECK: ^bb1(%{{.*}}: i32, %[[COND:.*]]: i1):
@@ -103,7 +103,7 @@ spv.module Logical GLSL450 {
spv.loop {
spv.Branch ^header
^header:
- %cond = spv.constant true
+ %cond = spv.Constant true
spv.BranchConditional %cond, ^body, ^merge
^body:
// Do nothing
@@ -133,7 +133,7 @@ spv.module Logical GLSL450 {
}
spv.func @selection_with_merge_block_only() -> () "None" {
- %cond = spv.constant true
+ %cond = spv.Constant true
// CHECK: llvm.return
spv.selection {
spv.BranchConditional %cond, ^merge, ^merge
@@ -145,7 +145,7 @@ spv.module Logical GLSL450 {
spv.func @selection_with_true_block_only() -> () "None" {
// CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
- %cond = spv.constant true
+ %cond = spv.Constant true
// CHECK: llvm.cond_br %[[COND]], ^bb1, ^bb2
spv.selection {
spv.BranchConditional %cond, ^true, ^merge
@@ -165,7 +165,7 @@ spv.module Logical GLSL450 {
spv.func @selection_with_both_true_and_false_block() -> () "None" {
// CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
- %cond = spv.constant true
+ %cond = spv.Constant true
// CHECK: llvm.cond_br %[[COND]], ^bb1, ^bb2
spv.selection {
spv.BranchConditional %cond, ^true, ^false
@@ -189,7 +189,7 @@ spv.module Logical GLSL450 {
spv.func @selection_with_early_return(%arg0: i1) -> i32 "None" {
// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
- %0 = spv.constant 0 : i32
+ %0 = spv.Constant 0 : i32
// CHECK: llvm.cond_br %{{.*}}, ^bb1(%[[ZERO]] : i32), ^bb2
spv.selection {
spv.BranchConditional %arg0, ^true(%0 : i32), ^merge
@@ -203,7 +203,7 @@ spv.module Logical GLSL450 {
spv.mlir.merge
}
// CHECK: ^bb3:
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : i32
}
}
diff --git a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
index 6aab710220e7..6eebe4c15ebf 100644
--- a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
+++ b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
@@ -7,7 +7,7 @@
// CHECK-LABEL: @access_chain
spv.func @access_chain() "None" {
// CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i32) : i32
- %0 = spv.constant 1: i32
+ %0 = spv.Constant 1: i32
%1 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
// CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
// CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %[[ONE]], %[[ONE]]] : (!llvm.ptr<struct<packed (f32, array<4 x f32>)>>, i32, i32, i32) -> !llvm.ptr<f32>
@@ -176,7 +176,7 @@ spv.func @variable_scalar_with_initialization() "None" {
// CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr<i64>
// CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr<i64>
- %c = spv.constant 0 : i64
+ %c = spv.Constant 0 : i64
%0 = spv.Variable init(%c) : !spv.ptr<i64, Function>
spv.Return
}
@@ -195,7 +195,7 @@ spv.func @variable_vector_with_initialization() "None" {
// CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
// CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x vector<3xi1> : (i32) -> !llvm.ptr<vector<3xi1>>
// CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr<vector<3xi1>>
- %c = spv.constant dense<false> : vector<3xi1>
+ %c = spv.Constant dense<false> : vector<3xi1>
%0 = spv.Variable init(%c) : !spv.ptr<vector<3xi1>, Function>
spv.Return
}
diff --git a/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir b/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
index 6bb6d78d56af..7729ec2df080 100644
--- a/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
+++ b/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
@@ -410,53 +410,53 @@ module attributes {
// CHECK-LABEL: @constant
func @constant() {
- // CHECK: spv.constant true
+ // CHECK: spv.Constant true
%0 = constant true
- // CHECK: spv.constant 42 : i32
+ // CHECK: spv.Constant 42 : i32
%1 = constant 42 : i32
- // CHECK: spv.constant 5.000000e-01 : f32
+ // CHECK: spv.Constant 5.000000e-01 : f32
%2 = constant 0.5 : f32
- // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
+ // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
%3 = constant dense<[2, 3]> : vector<2xi32>
- // CHECK: spv.constant 1 : i32
+ // CHECK: spv.Constant 1 : i32
%4 = constant 1 : index
- // CHECK: spv.constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+ // CHECK: spv.Constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
%5 = constant dense<1> : tensor<2x3xi32>
- // CHECK: spv.constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
+ // CHECK: spv.Constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
%6 = constant dense<1.0> : tensor<2x3xf32>
- // CHECK: spv.constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
+ // CHECK: spv.Constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
%7 = constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32>
- // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+ // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
%8 = constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32>
- // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+ // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
%9 = constant dense<[[1, 2], [3, 4], [5, 6]]> : tensor<3x2xi32>
- // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+ // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
%10 = constant dense<[1, 2, 3, 4, 5, 6]> : tensor<6xi32>
return
}
// CHECK-LABEL: @constant_16bit
func @constant_16bit() {
- // CHECK: spv.constant 4 : i16
+ // CHECK: spv.Constant 4 : i16
%0 = constant 4 : i16
- // CHECK: spv.constant 5.000000e+00 : f16
+ // CHECK: spv.Constant 5.000000e+00 : f16
%1 = constant 5.0 : f16
- // CHECK: spv.constant dense<[2, 3]> : vector<2xi16>
+ // CHECK: spv.Constant dense<[2, 3]> : vector<2xi16>
%2 = constant dense<[2, 3]> : vector<2xi16>
- // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2>
+ // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2>
%3 = constant dense<4.0> : tensor<5xf16>
return
}
// CHECK-LABEL: @constant_64bit
func @constant_64bit() {
- // CHECK: spv.constant 4 : i64
+ // CHECK: spv.Constant 4 : i64
%0 = constant 4 : i64
- // CHECK: spv.constant 5.000000e+00 : f64
+ // CHECK: spv.Constant 5.000000e+00 : f64
%1 = constant 5.0 : f64
- // CHECK: spv.constant dense<[2, 3]> : vector<2xi64>
+ // CHECK: spv.Constant dense<[2, 3]> : vector<2xi64>
%2 = constant dense<[2, 3]> : vector<2xi64>
- // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8>
+ // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8>
%3 = constant dense<4.0> : tensor<5xf64>
return
}
@@ -472,58 +472,58 @@ module attributes {
// CHECK-LABEL: @constant_16bit
func @constant_16bit() {
- // CHECK: spv.constant 4 : i32
+ // CHECK: spv.Constant 4 : i32
%0 = constant 4 : i16
- // CHECK: spv.constant 5.000000e+00 : f32
+ // CHECK: spv.Constant 5.000000e+00 : f32
%1 = constant 5.0 : f16
- // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
+ // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
%2 = constant dense<[2, 3]> : vector<2xi16>
- // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
+ // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
%3 = constant dense<4.0> : tensor<5xf16>
- // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
+ // CHECK: spv.Constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
%4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16>
return
}
// CHECK-LABEL: @constant_64bit
func @constant_64bit() {
- // CHECK: spv.constant 4 : i32
+ // CHECK: spv.Constant 4 : i32
%0 = constant 4 : i64
- // CHECK: spv.constant 5.000000e+00 : f32
+ // CHECK: spv.Constant 5.000000e+00 : f32
%1 = constant 5.0 : f64
- // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
+ // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
%2 = constant dense<[2, 3]> : vector<2xi64>
- // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
+ // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
%3 = constant dense<4.0> : tensor<5xf64>
- // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
+ // CHECK: spv.Constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
%4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16>
return
}
// CHECK-LABEL: @corner_cases
func @corner_cases() {
- // CHECK: %{{.*}} = spv.constant -1 : i32
+ // CHECK: %{{.*}} = spv.Constant -1 : i32
%0 = constant 4294967295 : i64 // 2^32 - 1
- // CHECK: %{{.*}} = spv.constant 2147483647 : i32
+ // CHECK: %{{.*}} = spv.Constant 2147483647 : i32
%1 = constant 2147483647 : i64 // 2^31 - 1
- // CHECK: %{{.*}} = spv.constant -2147483648 : i32
+ // CHECK: %{{.*}} = spv.Constant -2147483648 : i32
%2 = constant 2147483648 : i64 // 2^31
- // CHECK: %{{.*}} = spv.constant -2147483648 : i32
+ // CHECK: %{{.*}} = spv.Constant -2147483648 : i32
%3 = constant -2147483648 : i64 // -2^31
- // CHECK: %{{.*}} = spv.constant -1 : i32
+ // CHECK: %{{.*}} = spv.Constant -1 : i32
%5 = constant -1 : i64
- // CHECK: %{{.*}} = spv.constant -2 : i32
+ // CHECK: %{{.*}} = spv.Constant -2 : i32
%6 = constant -2 : i64
- // CHECK: %{{.*}} = spv.constant -1 : i32
+ // CHECK: %{{.*}} = spv.Constant -1 : i32
%7 = constant -1 : index
- // CHECK: %{{.*}} = spv.constant -2 : i32
+ // CHECK: %{{.*}} = spv.Constant -2 : i32
%8 = constant -2 : index
- // CHECK: spv.constant false
+ // CHECK: spv.Constant false
%9 = constant false
- // CHECK: spv.constant true
+ // CHECK: spv.Constant true
%10 = constant true
return
@@ -639,8 +639,8 @@ func @uitofp_i32_f32(%arg0 : i32) -> f32 {
// CHECK-LABEL: @uitofp_i1_f32
func @uitofp_i1_f32(%arg0 : i1) -> f32 {
- // CHECK: %[[ZERO:.+]] = spv.constant 0.000000e+00 : f32
- // CHECK: %[[ONE:.+]] = spv.constant 1.000000e+00 : f32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0.000000e+00 : f32
+ // CHECK: %[[ONE:.+]] = spv.Constant 1.000000e+00 : f32
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, f32
%0 = std.uitofp %arg0 : i1 to f32
return %0 : f32
@@ -648,8 +648,8 @@ func @uitofp_i1_f32(%arg0 : i1) -> f32 {
// CHECK-LABEL: @uitofp_i1_f64
func @uitofp_i1_f64(%arg0 : i1) -> f64 {
- // CHECK: %[[ZERO:.+]] = spv.constant 0.000000e+00 : f64
- // CHECK: %[[ONE:.+]] = spv.constant 1.000000e+00 : f64
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0.000000e+00 : f64
+ // CHECK: %[[ONE:.+]] = spv.Constant 1.000000e+00 : f64
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, f64
%0 = std.uitofp %arg0 : i1 to f64
return %0 : f64
@@ -657,8 +657,8 @@ func @uitofp_i1_f64(%arg0 : i1) -> f64 {
// CHECK-LABEL: @uitofp_vec_i1_f32
func @uitofp_vec_i1_f32(%arg0 : vector<4xi1>) -> vector<4xf32> {
- // CHECK: %[[ZERO:.+]] = spv.constant dense<0.000000e+00> : vector<4xf32>
- // CHECK: %[[ONE:.+]] = spv.constant dense<1.000000e+00> : vector<4xf32>
+ // CHECK: %[[ZERO:.+]] = spv.Constant dense<0.000000e+00> : vector<4xf32>
+ // CHECK: %[[ONE:.+]] = spv.Constant dense<1.000000e+00> : vector<4xf32>
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xf32>
%0 = std.uitofp %arg0 : vector<4xi1> to vector<4xf32>
return %0 : vector<4xf32>
@@ -666,11 +666,11 @@ func @uitofp_vec_i1_f32(%arg0 : vector<4xi1>) -> vector<4xf32> {
// CHECK-LABEL: @uitofp_vec_i1_f64
spv.func @uitofp_vec_i1_f64(%arg0: vector<4xi1>) -> vector<4xf64> "None" {
- // CHECK: %[[ZERO:.+]] = spv.constant dense<0.000000e+00> : vector<4xf64>
- // CHECK: %[[ONE:.+]] = spv.constant dense<1.000000e+00> : vector<4xf64>
+ // CHECK: %[[ZERO:.+]] = spv.Constant dense<0.000000e+00> : vector<4xf64>
+ // CHECK: %[[ONE:.+]] = spv.Constant dense<1.000000e+00> : vector<4xf64>
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xf64>
- %0 = spv.constant dense<0.000000e+00> : vector<4xf64>
- %1 = spv.constant dense<1.000000e+00> : vector<4xf64>
+ %0 = spv.Constant dense<0.000000e+00> : vector<4xf64>
+ %1 = spv.Constant dense<1.000000e+00> : vector<4xf64>
%2 = spv.Select %arg0, %1, %0 : vector<4xi1>, vector<4xf64>
spv.ReturnValue %2 : vector<4xf64>
}
@@ -705,8 +705,8 @@ func @zexti2(%arg0 : i32) -> i64 {
// CHECK-LABEL: @zexti3
func @zexti3(%arg0 : i1) -> i32 {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[ONE:.+]] = spv.constant 1 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, i32
%0 = std.zexti %arg0 : i1 to i32
return %0 : i32
@@ -714,8 +714,8 @@ func @zexti3(%arg0 : i1) -> i32 {
// CHECK-LABEL: @zexti4
func @zexti4(%arg0 : vector<4xi1>) -> vector<4xi32> {
- // CHECK: %[[ZERO:.+]] = spv.constant dense<0> : vector<4xi32>
- // CHECK: %[[ONE:.+]] = spv.constant dense<1> : vector<4xi32>
+ // CHECK: %[[ZERO:.+]] = spv.Constant dense<0> : vector<4xi32>
+ // CHECK: %[[ONE:.+]] = spv.Constant dense<1> : vector<4xi32>
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xi32>
%0 = std.zexti %arg0 : vector<4xi1> to vector<4xi32>
return %0 : vector<4xi32>
@@ -723,8 +723,8 @@ func @zexti4(%arg0 : vector<4xi1>) -> vector<4xi32> {
// CHECK-LABEL: @zexti5
func @zexti5(%arg0 : vector<4xi1>) -> vector<4xi64> {
- // CHECK: %[[ZERO:.+]] = spv.constant dense<0> : vector<4xi64>
- // CHECK: %[[ONE:.+]] = spv.constant dense<1> : vector<4xi64>
+ // CHECK: %[[ZERO:.+]] = spv.Constant dense<0> : vector<4xi64>
+ // CHECK: %[[ONE:.+]] = spv.Constant dense<1> : vector<4xi64>
// CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xi64>
%0 = std.zexti %arg0 : vector<4xi1> to vector<4xi64>
return %0 : vector<4xi64>
@@ -746,11 +746,11 @@ func @trunci2(%arg0: i32) -> i16 {
// CHECK-LABEL: @trunc_to_i1
func @trunc_to_i1(%arg0: i32) -> i1 {
- // CHECK: %[[MASK:.*]] = spv.constant 1 : i32
+ // CHECK: %[[MASK:.*]] = spv.Constant 1 : i32
// CHECK: %[[MASKED_SRC:.*]] = spv.BitwiseAnd %{{.*}}, %[[MASK]] : i32
// CHECK: %[[IS_ONE:.*]] = spv.IEqual %[[MASKED_SRC]], %[[MASK]] : i32
- // CHECK-DAG: %[[TRUE:.*]] = spv.constant true
- // CHECK-DAG: %[[FALSE:.*]] = spv.constant false
+ // CHECK-DAG: %[[TRUE:.*]] = spv.Constant true
+ // CHECK-DAG: %[[FALSE:.*]] = spv.Constant false
// CHECK: spv.Select %[[IS_ONE]], %[[TRUE]], %[[FALSE]] : i1, i1
%0 = std.trunci %arg0 : i32 to i1
return %0 : i1
@@ -758,11 +758,11 @@ func @trunc_to_i1(%arg0: i32) -> i1 {
// CHECK-LABEL: @trunc_to_veci1
func @trunc_to_veci1(%arg0: vector<4xi32>) -> vector<4xi1> {
- // CHECK: %[[MASK:.*]] = spv.constant dense<1> : vector<4xi32>
+ // CHECK: %[[MASK:.*]] = spv.Constant dense<1> : vector<4xi32>
// CHECK: %[[MASKED_SRC:.*]] = spv.BitwiseAnd %{{.*}}, %[[MASK]] : vector<4xi32>
// CHECK: %[[IS_ONE:.*]] = spv.IEqual %[[MASKED_SRC]], %[[MASK]] : vector<4xi32>
- // CHECK-DAG: %[[TRUE:.*]] = spv.constant dense<true> : vector<4xi1>
- // CHECK-DAG: %[[FALSE:.*]] = spv.constant dense<false> : vector<4xi1>
+ // CHECK-DAG: %[[TRUE:.*]] = spv.Constant dense<true> : vector<4xi1>
+ // CHECK-DAG: %[[FALSE:.*]] = spv.Constant dense<false> : vector<4xi1>
// CHECK: spv.Select %[[IS_ONE]], %[[TRUE]], %[[FALSE]] : vector<4xi1>, vector<4xi1>
%0 = std.trunci %arg0 : vector<4xi32> to vector<4xi1>
return %0 : vector<4xi1>
@@ -871,13 +871,13 @@ func @select(%arg0 : i32, %arg1 : i32) {
// CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>,
// CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>)
func @load_store_zero_rank_float(%arg0: memref<f32>, %arg1: memref<f32>) {
- // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
+ // CHECK: [[ZERO1:%.*]] = spv.Constant 0 : i32
// CHECK: spv.AccessChain [[ARG0]][
// CHECK-SAME: [[ZERO1]], [[ZERO1]]
// CHECK-SAME: ] :
// CHECK: spv.Load "StorageBuffer" %{{.*}} : f32
%0 = load %arg0[] : memref<f32>
- // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32
+ // CHECK: [[ZERO2:%.*]] = spv.Constant 0 : i32
// CHECK: spv.AccessChain [[ARG1]][
// CHECK-SAME: [[ZERO2]], [[ZERO2]]
// CHECK-SAME: ] :
@@ -890,13 +890,13 @@ func @load_store_zero_rank_float(%arg0: memref<f32>, %arg1: memref<f32>) {
// CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>,
// CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>)
func @load_store_zero_rank_int(%arg0: memref<i32>, %arg1: memref<i32>) {
- // CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
+ // CHECK: [[ZERO1:%.*]] = spv.Constant 0 : i32
// CHECK: spv.AccessChain [[ARG0]][
// CHECK-SAME: [[ZERO1]], [[ZERO1]]
// CHECK-SAME: ] :
// CHECK: spv.Load "StorageBuffer" %{{.*}} : i32
%0 = load %arg0[] : memref<i32>
- // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32
+ // CHECK: [[ZERO2:%.*]] = spv.Constant 0 : i32
// CHECK: spv.AccessChain [[ARG1]][
// CHECK-SAME: [[ZERO2]], [[ZERO2]]
// CHECK-SAME: ] :
@@ -919,19 +919,19 @@ module attributes {
// CHECK-LABEL: @load_i8
func @load_i8(%arg0: memref<i8>) {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[FOUR1:.+]] = spv.constant 4 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[FOUR1:.+]] = spv.Constant 4 : i32
// CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[ZERO]], %[[FOUR1]] : i32
// CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]
// CHECK: %[[LOAD:.+]] = spv.Load "StorageBuffer" %[[PTR]]
- // CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
- // CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+ // CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
+ // CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
// CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR2]] : i32
// CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
// CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32
- // CHECK: %[[MASK:.+]] = spv.constant 255 : i32
+ // CHECK: %[[MASK:.+]] = spv.Constant 255 : i32
// CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32
- // CHECK: %[[T2:.+]] = spv.constant 24 : i32
+ // CHECK: %[[T2:.+]] = spv.Constant 24 : i32
// CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32
// CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32
%0 = load %arg0[] : memref<i8>
@@ -941,23 +941,23 @@ func @load_i8(%arg0: memref<i8>) {
// CHECK-LABEL: @load_i16
// CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32)
func @load_i16(%arg0: memref<10xi16>, %index : index) {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[OFFSET:.+]] = spv.constant 0 : i32
- // CHECK: %[[ONE:.+]] = spv.constant 1 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[OFFSET:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32
// CHECK: %[[UPDATE:.+]] = spv.IMul %[[ONE]], %[[ARG1]] : i32
// CHECK: %[[FLAT_IDX:.+]] = spv.IAdd %[[OFFSET]], %[[UPDATE]] : i32
- // CHECK: %[[TWO1:.+]] = spv.constant 2 : i32
+ // CHECK: %[[TWO1:.+]] = spv.Constant 2 : i32
// CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[FLAT_IDX]], %[[TWO1]] : i32
// CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]
// CHECK: %[[LOAD:.+]] = spv.Load "StorageBuffer" %[[PTR]]
- // CHECK: %[[TWO2:.+]] = spv.constant 2 : i32
- // CHECK: %[[SIXTEEN:.+]] = spv.constant 16 : i32
+ // CHECK: %[[TWO2:.+]] = spv.Constant 2 : i32
+ // CHECK: %[[SIXTEEN:.+]] = spv.Constant 16 : i32
// CHECK: %[[IDX:.+]] = spv.UMod %[[FLAT_IDX]], %[[TWO2]] : i32
// CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[SIXTEEN]] : i32
// CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32
- // CHECK: %[[MASK:.+]] = spv.constant 65535 : i32
+ // CHECK: %[[MASK:.+]] = spv.Constant 65535 : i32
// CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32
- // CHECK: %[[T2:.+]] = spv.constant 16 : i32
+ // CHECK: %[[T2:.+]] = spv.Constant 16 : i32
// CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32
// CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32
%0 = load %arg0[%index] : memref<10xi16>
@@ -985,17 +985,17 @@ func @load_f32(%arg0: memref<f32>) {
// CHECK-LABEL: @store_i8
// CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32)
func @store_i8(%arg0: memref<i8>, %value: i8) {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[FOUR:.+]] = spv.constant 4 : i32
- // CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[FOUR:.+]] = spv.Constant 4 : i32
+ // CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
// CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR]] : i32
// CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
- // CHECK: %[[MASK1:.+]] = spv.constant 255 : i32
+ // CHECK: %[[MASK1:.+]] = spv.Constant 255 : i32
// CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32
// CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32
// CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG1]], %[[MASK1]] : i32
// CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32
- // CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
+ // CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
// CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[ZERO]], %[[FOUR2]] : i32
// CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]]
// CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]]
@@ -1007,21 +1007,21 @@ func @store_i8(%arg0: memref<i8>, %value: i8) {
// CHECK-LABEL: @store_i16
// CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32, %[[ARG2:.+]]: i32)
func @store_i16(%arg0: memref<10xi16>, %index: index, %value: i16) {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[OFFSET:.+]] = spv.constant 0 : i32
- // CHECK: %[[ONE:.+]] = spv.constant 1 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[OFFSET:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32
// CHECK: %[[UPDATE:.+]] = spv.IMul %[[ONE]], %[[ARG1]] : i32
// CHECK: %[[FLAT_IDX:.+]] = spv.IAdd %[[OFFSET]], %[[UPDATE]] : i32
- // CHECK: %[[TWO:.+]] = spv.constant 2 : i32
- // CHECK: %[[SIXTEEN:.+]] = spv.constant 16 : i32
+ // CHECK: %[[TWO:.+]] = spv.Constant 2 : i32
+ // CHECK: %[[SIXTEEN:.+]] = spv.Constant 16 : i32
// CHECK: %[[IDX:.+]] = spv.UMod %[[FLAT_IDX]], %[[TWO]] : i32
// CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[SIXTEEN]] : i32
- // CHECK: %[[MASK1:.+]] = spv.constant 65535 : i32
+ // CHECK: %[[MASK1:.+]] = spv.Constant 65535 : i32
// CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32
// CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32
// CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG2]], %[[MASK1]] : i32
// CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32
- // CHECK: %[[TWO2:.+]] = spv.constant 2 : i32
+ // CHECK: %[[TWO2:.+]] = spv.Constant 2 : i32
// CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[FLAT_IDX]], %[[TWO2]] : i32
// CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]]
// CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]]
@@ -1062,19 +1062,19 @@ module attributes {
// CHECK-LABEL: @load_i8
func @load_i8(%arg0: memref<i8>) {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[FOUR1:.+]] = spv.constant 4 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[FOUR1:.+]] = spv.Constant 4 : i32
// CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[ZERO]], %[[FOUR1]] : i32
// CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]
// CHECK: %[[LOAD:.+]] = spv.Load "StorageBuffer" %[[PTR]]
- // CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
- // CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+ // CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
+ // CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
// CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR2]] : i32
// CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
// CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32
- // CHECK: %[[MASK:.+]] = spv.constant 255 : i32
+ // CHECK: %[[MASK:.+]] = spv.Constant 255 : i32
// CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32
- // CHECK: %[[T2:.+]] = spv.constant 24 : i32
+ // CHECK: %[[T2:.+]] = spv.Constant 24 : i32
// CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32
// CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32
%0 = load %arg0[] : memref<i8>
@@ -1093,17 +1093,17 @@ func @load_i16(%arg0: memref<i16>) {
// CHECK-LABEL: @store_i8
// CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32)
func @store_i8(%arg0: memref<i8>, %value: i8) {
- // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
- // CHECK: %[[FOUR:.+]] = spv.constant 4 : i32
- // CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+ // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+ // CHECK: %[[FOUR:.+]] = spv.Constant 4 : i32
+ // CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
// CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR]] : i32
// CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
- // CHECK: %[[MASK1:.+]] = spv.constant 255 : i32
+ // CHECK: %[[MASK1:.+]] = spv.Constant 255 : i32
// CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32
// CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32
// CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG1]], %[[MASK1]] : i32
// CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32
- // CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
+ // CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
// CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[ZERO]], %[[FOUR2]] : i32
// CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]]
// CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]]
diff --git a/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir b/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir
index 08b7fd26e1db..05837dffb191 100644
--- a/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir
@@ -113,7 +113,7 @@ func @composite_extract_no_ssa_operand() -> () {
// -----
func @composite_extract_invalid_index_type_1() -> () {
- %0 = spv.constant 10 : i32
+ %0 = spv.Constant 10 : i32
%1 = spv.Variable : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
%2 = spv.Load "Function" %1 ["Volatile"] : !spv.array<4x!spv.array<4xf32>>
// expected-error @+1 {{expected non-function type}}
diff --git a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
index 61d27c5416a7..cc6a0b2562f0 100644
--- a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
@@ -14,7 +14,7 @@ func @branch() -> () {
// -----
func @branch_argument() -> () {
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
// CHECK: spv.Branch ^bb1(%{{.*}}, %{{.*}} : i32, i32)
spv.Branch ^next(%zero, %zero: i32, i32)
^next(%arg0: i32, %arg1: i32):
@@ -31,7 +31,7 @@ func @missing_accessor() -> () {
// -----
func @wrong_accessor_count() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// expected-error @+1 {{requires 1 successor but found 2}}
"spv.Branch"()[^one, ^two] : () -> ()
^one:
@@ -47,7 +47,7 @@ func @wrong_accessor_count() -> () {
//===----------------------------------------------------------------------===//
func @cond_branch() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// CHECK: spv.BranchConditional %{{.*}}, ^bb1, ^bb2
spv.BranchConditional %true, ^one, ^two
// CHECK: ^bb1
@@ -61,8 +61,8 @@ func @cond_branch() -> () {
// -----
func @cond_branch_argument() -> () {
- %true = spv.constant true
- %zero = spv.constant 0 : i32
+ %true = spv.Constant true
+ %zero = spv.Constant 0 : i32
// CHECK: spv.BranchConditional %{{.*}}, ^bb1(%{{.*}}, %{{.*}} : i32, i32), ^bb2
spv.BranchConditional %true, ^true1(%zero, %zero: i32, i32), ^false1
^true1(%arg0: i32, %arg1: i32):
@@ -79,7 +79,7 @@ func @cond_branch_argument() -> () {
// -----
func @cond_branch_with_weights() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// CHECK: spv.BranchConditional %{{.*}} [5, 10]
spv.BranchConditional %true [5, 10], ^one, ^two
^one:
@@ -103,7 +103,7 @@ func @missing_condition() -> () {
func @wrong_condition_type() -> () {
// expected-note @+1 {{prior use here}}
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
// expected-error @+1 {{use of value '%zero' expects
diff erent type than prior uses: 'i1' vs 'i32'}}
spv.BranchConditional %zero, ^one, ^two
^one:
@@ -115,7 +115,7 @@ func @wrong_condition_type() -> () {
// -----
func @wrong_accessor_count() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// expected-error @+1 {{requires 2 successors but found 1}}
"spv.BranchConditional"(%true)[^one] {operand_segment_sizes = dense<[1, 0, 0]>: vector<3xi32>} : (i1) -> ()
^one:
@@ -127,7 +127,7 @@ func @wrong_accessor_count() -> () {
// -----
func @wrong_number_of_weights() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// expected-error @+1 {{must have exactly two branch weights}}
"spv.BranchConditional"(%true)[^one, ^two] {branch_weights = [1 : i32, 2 : i32, 3 : i32],
operand_segment_sizes = dense<[1, 0, 0]>: vector<3xi32>} : (i1) -> ()
@@ -140,7 +140,7 @@ func @wrong_number_of_weights() -> () {
// -----
func @weights_cannot_both_be_zero() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// expected-error @+1 {{branch weights cannot both be zero}}
spv.BranchConditional %true [0, 0], ^one, ^two
^one:
@@ -232,7 +232,7 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
- %0 = spv.constant 2.0 : f32
+ %0 = spv.Constant 2.0 : f32
// expected-error @+1 {{operand type mismatch: expected operand type 'i32', but provided 'f32' for operand number 1}}
spv.FunctionCall @f_type_mismatch(%arg0, %0) : (i32, f32) -> ()
spv.Return
@@ -243,7 +243,7 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 "None" {
- %cst = spv.constant 0: i32
+ %cst = spv.Constant 0: i32
// expected-error @+1 {{result type mismatch: expected 'i32', but provided 'f32'}}
%0 = spv.FunctionCall @f_type_mismatch(%arg0, %arg0) : (i32, i32) -> f32
spv.ReturnValue %cst: i32
@@ -268,8 +268,8 @@ spv.module Logical GLSL450 {
// for (int i = 0; i < count; ++i) {}
func @loop(%count : i32) -> () {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
// CHECK: spv.loop {
@@ -438,8 +438,8 @@ func @merge() -> () {
// -----
func @only_allowed_in_last_block(%cond : i1) -> () {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.selection {
@@ -460,7 +460,7 @@ func @only_allowed_in_last_block(%cond : i1) -> () {
// -----
func @only_allowed_in_last_block() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
spv.loop {
spv.Branch ^header
^header:
@@ -548,7 +548,7 @@ spv.module Logical GLSL450 {
spv.mlir.merge
}
- %zero = spv.constant 0: i32
+ %zero = spv.Constant 0: i32
spv.ReturnValue %zero: i32
}
}
@@ -560,7 +560,7 @@ spv.module Logical GLSL450 {
//===----------------------------------------------------------------------===//
func @ret_val() -> (i32) {
- %0 = spv.constant 42 : i32
+ %0 = spv.Constant 42 : i32
// CHECK: spv.ReturnValue %{{.*}} : i32
spv.ReturnValue %0 : i32
}
@@ -570,13 +570,13 @@ func @in_selection(%cond : i1) -> (i32) {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
^then:
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
// CHECK: spv.ReturnValue
spv.ReturnValue %zero : i32
^merge:
spv.mlir.merge
}
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : i32
}
@@ -587,7 +587,7 @@ func @in_loop(%cond : i1) -> (i32) {
^header:
spv.BranchConditional %cond, ^body, ^merge
^body:
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
// CHECK: spv.ReturnValue
spv.ReturnValue %zero : i32
^continue:
@@ -595,7 +595,7 @@ func @in_loop(%cond : i1) -> (i32) {
^merge:
spv.mlir.merge
}
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : i32
}
@@ -608,7 +608,7 @@ func @in_other_func_like_op(%arg: i32) -> i32 {
// -----
"foo.function"() ({
- %0 = spv.constant true
+ %0 = spv.Constant true
// expected-error @+1 {{op must appear in a function-like op's block}}
spv.ReturnValue %0 : i1
}) : () -> ()
@@ -617,7 +617,7 @@ func @in_other_func_like_op(%arg: i32) -> i32 {
spv.module Logical GLSL450 {
spv.func @value_count_mismatch() -> () "None" {
- %0 = spv.constant 42 : i32
+ %0 = spv.Constant 42 : i32
// expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
spv.ReturnValue %0 : i32
}
@@ -627,7 +627,7 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @value_type_mismatch() -> (f32) "None" {
- %0 = spv.constant 42 : i32
+ %0 = spv.Constant 42 : i32
// expected-error @+1 {{return value's type ('i32') mismatch with function's result type ('f32')}}
spv.ReturnValue %0 : i32
}
@@ -640,7 +640,7 @@ spv.module Logical GLSL450 {
spv.selection {
spv.BranchConditional %cond, ^then, ^merge
^then:
- %cst = spv.constant 0: i32
+ %cst = spv.Constant 0: i32
// expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
spv.ReturnValue %cst: i32
^merge:
@@ -658,8 +658,8 @@ spv.module Logical GLSL450 {
//===----------------------------------------------------------------------===//
func @selection(%cond: i1) -> () {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
// CHECK: spv.selection {
@@ -685,9 +685,9 @@ func @selection(%cond: i1) -> () {
// -----
func @selection(%cond: i1) -> () {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
- %two = spv.constant 2: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
+ %two = spv.Constant 2: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
// CHECK: spv.selection {
diff --git a/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir b/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir
index 7cd0b631b04a..7de8a36183c7 100644
--- a/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir
@@ -103,7 +103,7 @@ spv.func @cooperative_matrix_fdiv(%a : !spv.coopmatrix<8x16xf32, Subgroup>, %b :
// CHECK-LABEL: @cooperative_matrix_access_chain
spv.func @cooperative_matrix_access_chain(%a : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>) -> !spv.ptr<f32, Function> "None" {
- %0 = spv.constant 0: i32
+ %0 = spv.Constant 0: i32
// CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
%1 = spv.AccessChain %a[%0] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
spv.ReturnValue %1 : !spv.ptr<f32, Function>
diff --git a/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir b/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir
index b2c34b85f194..25d91416d2ef 100644
--- a/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir
@@ -178,24 +178,24 @@ func @logicalUnary(%arg0 : i32)
//===----------------------------------------------------------------------===//
func @select_op_bool(%arg0: i1) -> () {
- %0 = spv.constant true
- %1 = spv.constant false
+ %0 = spv.Constant true
+ %1 = spv.Constant false
// CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i1
%2 = spv.Select %arg0, %0, %1 : i1, i1
return
}
func @select_op_int(%arg0: i1) -> () {
- %0 = spv.constant 2 : i32
- %1 = spv.constant 3 : i32
+ %0 = spv.Constant 2 : i32
+ %1 = spv.Constant 3 : i32
// CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i32
%2 = spv.Select %arg0, %0, %1 : i1, i32
return
}
func @select_op_float(%arg0: i1) -> () {
- %0 = spv.constant 2.0 : f32
- %1 = spv.constant 3.0 : f32
+ %0 = spv.Constant 2.0 : f32
+ %1 = spv.Constant 3.0 : f32
// CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, f32
%2 = spv.Select %arg0, %0, %1 : i1, f32
return
@@ -210,16 +210,16 @@ func @select_op_ptr(%arg0: i1) -> () {
}
func @select_op_vec(%arg0: i1) -> () {
- %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
- %1 = spv.constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
+ %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+ %1 = spv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
// CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, vector<3xf32>
%2 = spv.Select %arg0, %0, %1 : i1, vector<3xf32>
return
}
func @select_op_vec_condn_vec(%arg0: vector<3xi1>) -> () {
- %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
- %1 = spv.constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
+ %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+ %1 = spv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
// CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi1>, vector<3xf32>
%2 = spv.Select %arg0, %0, %1 : vector<3xi1>, vector<3xf32>
return
@@ -228,8 +228,8 @@ func @select_op_vec_condn_vec(%arg0: vector<3xi1>) -> () {
// -----
func @select_op(%arg0: i1) -> () {
- %0 = spv.constant 2 : i32
- %1 = spv.constant 3 : i32
+ %0 = spv.Constant 2 : i32
+ %1 = spv.Constant 3 : i32
// expected-error @+2 {{expected ','}}
%2 = spv.Select %arg0, %0, %1 : i1
return
@@ -238,8 +238,8 @@ func @select_op(%arg0: i1) -> () {
// -----
func @select_op(%arg1: vector<3xi1>) -> () {
- %0 = spv.constant 2 : i32
- %1 = spv.constant 3 : i32
+ %0 = spv.Constant 2 : i32
+ %1 = spv.Constant 3 : i32
// expected-error @+1 {{result expected to be of vector type when condition is of vector type}}
%2 = spv.Select %arg1, %0, %1 : vector<3xi1>, i32
return
@@ -248,8 +248,8 @@ func @select_op(%arg1: vector<3xi1>) -> () {
// -----
func @select_op(%arg1: vector<4xi1>) -> () {
- %0 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32>
+ %0 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32>
// expected-error @+1 {{result should have the same number of elements as the condition when condition is of vector type}}
%2 = spv.Select %arg1, %0, %1 : vector<4xi1>, vector<3xi32>
return
@@ -258,8 +258,8 @@ func @select_op(%arg1: vector<4xi1>) -> () {
// -----
func @select_op(%arg1: vector<4xi1>) -> () {
- %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
- %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32>
+ %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+ %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32>
// expected-error @+1 {{all of {true_value, false_value, result} have same type}}
%2 = "spv.Select"(%arg1, %0, %1) : (vector<4xi1>, vector<3xf32>, vector<3xi32>) -> vector<3xi32>
return
@@ -268,8 +268,8 @@ func @select_op(%arg1: vector<4xi1>) -> () {
// -----
func @select_op(%arg1: vector<4xi1>) -> () {
- %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
- %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32>
+ %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+ %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32>
// expected-error @+1 {{all of {true_value, false_value, result} have same type}}
%2 = "spv.Select"(%arg1, %1, %0) : (vector<4xi1>, vector<3xi32>, vector<3xf32>) -> vector<3xi32>
return
diff --git a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
index 5e82cf91e7af..ba76516a21e4 100644
--- a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
@@ -5,7 +5,7 @@
//===----------------------------------------------------------------------===//
func @access_chain_struct() -> () {
- %0 = spv.constant 1: i32
+ %0 = spv.Constant 1: i32
%1 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
// CHECK: spv.AccessChain {{.*}}[{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Function>
%2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
@@ -46,7 +46,7 @@ func @access_chain_rtarray(%arg0 : i32) -> () {
// -----
func @access_chain_non_composite() -> () {
- %0 = spv.constant 1: i32
+ %0 = spv.Constant 1: i32
%1 = spv.Variable : !spv.ptr<f32, Function>
// expected-error @+1 {{cannot extract from non-composite type 'f32' with index 0}}
%2 = spv.AccessChain %1[%0] : !spv.ptr<f32, Function>, i32
@@ -112,7 +112,7 @@ func @access_chain_invalid_index_1(%index0 : i32) -> () {
func @access_chain_invalid_index_2(%index0 : i32) -> () {
%0 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
- // expected-error @+1 {{index must be an integer spv.constant to access element of spv.struct}}
+ // expected-error @+1 {{index must be an integer spv.Constant to access element of spv.struct}}
%1 = spv.AccessChain %0[%index0, %index0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
return
}
@@ -122,7 +122,7 @@ func @access_chain_invalid_index_2(%index0 : i32) -> () {
func @access_chain_invalid_constant_type_1() -> () {
%0 = std.constant 1: i32
%1 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
- // expected-error @+1 {{index must be an integer spv.constant to access element of spv.struct, but provided std.constant}}
+ // expected-error @+1 {{index must be an integer spv.Constant to access element of spv.struct, but provided std.constant}}
%2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
return
}
@@ -130,7 +130,7 @@ func @access_chain_invalid_constant_type_1() -> () {
// -----
func @access_chain_out_of_bounds() -> () {
- %index0 = "spv.constant"() { value = 12: i32} : () -> i32
+ %index0 = "spv.Constant"() { value = 12: i32} : () -> i32
%0 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
// expected-error @+1 {{'spv.AccessChain' op index 12 out of bounds for '!spv.struct<(f32, !spv.array<4 x f32>)>'}}
%1 = spv.AccessChain %0[%index0, %index0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
@@ -487,7 +487,7 @@ func @variable(%arg0: f32) -> () {
// -----
func @variable_init_normal_constant() -> () {
- %0 = spv.constant 4.0 : f32
+ %0 = spv.Constant 4.0 : f32
// CHECK: spv.Variable init(%0) : !spv.ptr<f32, Function>
%1 = spv.Variable init(%0) : !spv.ptr<f32, Function>
return
@@ -529,7 +529,7 @@ func @variable_bind() -> () {
// -----
func @variable_init_bind() -> () {
- %0 = spv.constant 4.0 : f32
+ %0 = spv.Constant 4.0 : f32
// expected-error @+1 {{cannot have 'binding' attribute (only allowed in spv.globalVariable)}}
%1 = spv.Variable init(%0) {binding = 5 : i32} : !spv.ptr<f32, Function>
return
diff --git a/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir b/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
index 61febd37460a..602d5d2c75e9 100644
--- a/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
@@ -33,7 +33,7 @@ func @group_non_uniform_ballot(%predicate: i1) -> vector<4xsi32> {
//===----------------------------------------------------------------------===//
func @group_non_uniform_broadcast_scalar(%value: f32) -> f32 {
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
// CHECK: spv.GroupNonUniformBroadcast Workgroup %{{.*}}, %{{.*}} : f32, i32
%0 = spv.GroupNonUniformBroadcast Workgroup %value, %one : f32, i32
return %0: f32
@@ -42,7 +42,7 @@ func @group_non_uniform_broadcast_scalar(%value: f32) -> f32 {
// -----
func @group_non_uniform_broadcast_vector(%value: vector<4xf32>) -> vector<4xf32> {
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
// CHECK: spv.GroupNonUniformBroadcast Subgroup %{{.*}}, %{{.*}} : vector<4xf32>, i32
%0 = spv.GroupNonUniformBroadcast Subgroup %value, %one : vector<4xf32>, i32
return %0: vector<4xf32>
@@ -51,7 +51,7 @@ func @group_non_uniform_broadcast_vector(%value: vector<4xf32>) -> vector<4xf32>
// -----
func @group_non_uniform_broadcast_negative_scope(%value: f32, %localid: i32 ) -> f32 {
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
// expected-error @+1 {{execution scope must be 'Workgroup' or 'Subgroup'}}
%0 = spv.GroupNonUniformBroadcast Device %value, %one : f32, i32
return %0: f32
@@ -101,7 +101,7 @@ func @group_non_uniform_fadd_reduce(%val: f32) -> f32 {
// CHECK-LABEL: @group_non_uniform_fadd_clustered_reduce
func @group_non_uniform_fadd_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> {
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
// CHECK: %{{.+}} = spv.GroupNonUniformFAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>
%0 = spv.GroupNonUniformFAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xf32>
return %0: vector<2xf32>
@@ -120,7 +120,7 @@ func @group_non_uniform_fmul_reduce(%val: f32) -> f32 {
// CHECK-LABEL: @group_non_uniform_fmul_clustered_reduce
func @group_non_uniform_fmul_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> {
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
// CHECK: %{{.+}} = spv.GroupNonUniformFMul "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>
%0 = spv.GroupNonUniformFMul "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xf32>
return %0: vector<2xf32>
@@ -167,7 +167,7 @@ func @group_non_uniform_iadd_reduce(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce
func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
// CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
%0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
return %0: vector<2xi32>
@@ -200,7 +200,7 @@ func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>, %size: i32) -
// -----
func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
- %five = spv.constant 5 : i32
+ %five = spv.Constant 5 : i32
// expected-error @+1 {{cluster size operand must be a power of two}}
%0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%five) : vector<2xi32>
return %0: vector<2xi32>
@@ -221,7 +221,7 @@ func @group_non_uniform_imul_reduce(%val: i32) -> i32 {
// CHECK-LABEL: @group_non_uniform_imul_clustered_reduce
func @group_non_uniform_imul_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
// CHECK: %{{.+}} = spv.GroupNonUniformIMul "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
%0 = spv.GroupNonUniformIMul "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
return %0: vector<2xi32>
diff --git a/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir b/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir
index 9b8e66bab567..dbabaed48cba 100644
--- a/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir
+++ b/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir
@@ -7,7 +7,7 @@
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
+ %0 = spv.Constant 1: i32
// CHECK: [[VAR1:%.*]] = spv.mlir.addressof @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
// CHECK-NEXT: spv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
%1 = spv.mlir.addressof @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
@@ -49,29 +49,29 @@ spv.module Logical GLSL450 {
// -----
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
func @const() -> () {
- // CHECK: spv.constant true
- // CHECK: spv.constant 42 : i32
- // CHECK: spv.constant 5.000000e-01 : f32
- // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
- // CHECK: spv.constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>>
- // CHECK: spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
- // CHECK: spv.constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
- // CHECK: spv.constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
- // CHECK: spv.constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
-
- %0 = spv.constant true
- %1 = spv.constant 42 : i32
- %2 = spv.constant 0.5 : f32
- %3 = spv.constant dense<[2, 3]> : vector<2xi32>
- %4 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
- %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
- %6 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
- %7 = spv.constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
- %8 = spv.constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+ // CHECK: spv.Constant true
+ // CHECK: spv.Constant 42 : i32
+ // CHECK: spv.Constant 5.000000e-01 : f32
+ // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
+ // CHECK: spv.Constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>>
+ // CHECK: spv.Constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+ // CHECK: spv.Constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+ // CHECK: spv.Constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+ // CHECK: spv.Constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+
+ %0 = spv.Constant true
+ %1 = spv.Constant 42 : i32
+ %2 = spv.Constant 0.5 : f32
+ %3 = spv.Constant dense<[2, 3]> : vector<2xi32>
+ %4 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
+ %5 = spv.Constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+ %6 = spv.Constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+ %7 = spv.Constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+ %8 = spv.Constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
return
}
@@ -79,7 +79,7 @@ func @const() -> () {
func @unaccepted_std_attr() -> () {
// expected-error @+1 {{cannot have value of type 'none'}}
- %0 = spv.constant unit : none
+ %0 = spv.Constant unit : none
return
}
@@ -87,7 +87,7 @@ func @unaccepted_std_attr() -> () {
func @array_constant() -> () {
// expected-error @+1 {{has array element whose type ('vector<2xi32>') does not match the result element type ('vector<2xf32>')}}
- %0 = spv.constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spv.array<2xvector<2xf32>>
+ %0 = spv.Constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spv.array<2xvector<2xf32>>
return
}
@@ -95,7 +95,7 @@ func @array_constant() -> () {
func @array_constant() -> () {
// expected-error @+1 {{must have spv.array result type for array value}}
- %0 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.rtarray<vector<2xf32>>
+ %0 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.rtarray<vector<2xf32>>
return
}
@@ -103,7 +103,7 @@ func @array_constant() -> () {
func @non_nested_array_constant() -> () {
// expected-error @+1 {{only support nested array result type}}
- %0 = spv.constant dense<3.0> : tensor<2x2xf32> : !spv.array<2xvector<2xf32>>
+ %0 = spv.Constant dense<3.0> : tensor<2x2xf32> : !spv.array<2xvector<2xf32>>
return
}
@@ -111,21 +111,21 @@ func @non_nested_array_constant() -> () {
func @value_result_type_mismatch() -> () {
// expected-error @+1 {{must have spv.array result type for array value}}
- %0 = "spv.constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>)
+ %0 = "spv.Constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>)
}
// -----
func @value_result_type_mismatch() -> () {
// expected-error @+1 {{result element type ('i32') does not match value element type ('f32')}}
- %0 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>>
+ %0 = spv.Constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>>
}
// -----
func @value_result_num_elements_mismatch() -> () {
// expected-error @+1 {{result number of elements (6) does not match value number of elements (4)}}
- %0 = spv.constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>>
+ %0 = spv.Constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>>
return
}
@@ -308,7 +308,7 @@ spv.module Logical GLSL450 {
// TODO: Fix test case after initialization with normal constant is addressed
// spv.module Logical GLSL450 {
-// %0 = spv.constant 4.0 : f32
+// %0 = spv.Constant 4.0 : f32
// // CHECK1: spv.Variable init(%0) : !spv.ptr<f32, Private>
// spv.globalVariable @var1 init(%0) : !spv.ptr<f32, Private>
// }
@@ -337,7 +337,7 @@ spv.module Logical GLSL450 {
// TODO: Fix test case after initialization with constant is addressed
// spv.module Logical GLSL450 {
-// %0 = spv.constant 4.0 : f32
+// %0 = spv.Constant 4.0 : f32
// // CHECK1: spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
// spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
// }
@@ -480,7 +480,7 @@ spv.module Logical GLSL450 {
// expected-error at +2 {{expects regions to end with 'spv.mlir.endmodule'}}
// expected-note at +1 {{in custom textual format, the absence of terminator implies 'spv.mlir.endmodule'}}
"spv.module"() ({
- %0 = spv.constant true
+ %0 = spv.Constant true
}) {addressing_model = 0 : i32, memory_model = 1 : i32} : () -> ()
// -----
@@ -561,7 +561,7 @@ spv.module Logical GLSL450 {
spv.func @compute() -> f32 "None" {
// CHECK: spv.mlir.referenceof @sc3 : f32
%0 = spv.mlir.referenceof @sc3 : f32
- %1 = spv.constant 6.0 : f32
+ %1 = spv.Constant 6.0 : f32
%2 = spv.FAdd %0, %1 : f32
spv.ReturnValue %2 : f32
}
@@ -801,10 +801,10 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @foo() -> i32 "None" {
- // CHECK: [[LHS:%.*]] = spv.constant
- %0 = spv.constant 1: i32
- // CHECK: [[RHS:%.*]] = spv.constant
- %1 = spv.constant 1: i32
+ // CHECK: [[LHS:%.*]] = spv.Constant
+ %0 = spv.Constant 1: i32
+ // CHECK: [[RHS:%.*]] = spv.Constant
+ %1 = spv.Constant 1: i32
// CHECK: spv.SpecConstantOperation wraps "spv.IAdd"([[LHS]], [[RHS]]) : (i32, i32) -> i32
%2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32
@@ -831,7 +831,7 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @foo() -> i32 "None" {
- %0 = spv.constant 1: i32
+ %0 = spv.Constant 1: i32
// expected-error @+1 {{op expects parent op 'spv.SpecConstantOperation'}}
spv.mlir.yield %0 : i32
}
diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
index d9c7406406bc..e0c70f7df519 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
@@ -18,7 +18,7 @@ spv.module Logical GLSL450 {
attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
// CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]
// CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]]
- // CHECK: [[CONST0:%.*]] = spv.constant 0 : i32
+ // CHECK: [[CONST0:%.*]] = spv.Constant 0 : i32
// CHECK: [[ARG0PTR:%.*]] = spv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]]
// CHECK: [[ARG0:%.*]] = spv.Load "StorageBuffer" [[ARG0PTR]]
// CHECK: spv.Return
diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
index 7db0f888bc80..c2a56d2aa106 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
@@ -40,19 +40,19 @@ spv.module Logical GLSL450 {
{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.mlir.addressof [[VAR6]]
- // CHECK: [[CONST6:%.*]] = spv.constant 0 : i32
+ // CHECK: [[CONST6:%.*]] = spv.Constant 0 : i32
// CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
// CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]]
// CHECK: [[ADDRESSARG5:%.*]] = spv.mlir.addressof [[VAR5]]
- // CHECK: [[CONST5:%.*]] = spv.constant 0 : i32
+ // CHECK: [[CONST5:%.*]] = spv.Constant 0 : i32
// CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]
// CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG5PTR]]
// CHECK: [[ADDRESSARG4:%.*]] = spv.mlir.addressof [[VAR4]]
- // CHECK: [[CONST4:%.*]] = spv.constant 0 : i32
+ // CHECK: [[CONST4:%.*]] = spv.Constant 0 : i32
// CHECK: [[ARG4PTR:%.*]] = spv.AccessChain [[ADDRESSARG4]]{{\[}}[[CONST4]]
// CHECK: [[ARG4:%.*]] = spv.Load "StorageBuffer" [[ARG4PTR]]
// CHECK: [[ADDRESSARG3:%.*]] = spv.mlir.addressof [[VAR3]]
- // CHECK: [[CONST3:%.*]] = spv.constant 0 : i32
+ // CHECK: [[CONST3:%.*]] = spv.Constant 0 : i32
// CHECK: [[ARG3PTR:%.*]] = spv.AccessChain [[ADDRESSARG3]]{{\[}}[[CONST3]]
// CHECK: [[ARG3:%.*]] = spv.Load "StorageBuffer" [[ARG3PTR]]
// CHECK: [[ADDRESSARG2:%.*]] = spv.mlir.addressof [[VAR2]]
@@ -102,7 +102,7 @@ spv.module Logical GLSL450 {
// CHECK: spv.IAdd [[ARG4]]
%37 = spv.IAdd %arg4, %11 : i32
// CHECK: spv.AccessChain [[ARG0]]
- %c0 = spv.constant 0 : i32
+ %c0 = spv.Constant 0 : i32
%38 = spv.AccessChain %arg0[%c0, %36, %37] : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32>>)>, StorageBuffer>, i32, i32, i32
%39 = spv.Load "StorageBuffer" %38 : f32
// CHECK: spv.AccessChain [[ARG1]]
diff --git a/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir b/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
index cb955b4512ee..66c1ddc694d3 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
@@ -5,11 +5,11 @@
//===----------------------------------------------------------------------===//
func @combine_full_access_chain() -> f32 {
- // CHECK: %[[INDEX:.*]] = spv.constant 0
+ // CHECK: %[[INDEX:.*]] = spv.Constant 0
// CHECK-NEXT: %[[VAR:.*]] = spv.Variable
// CHECK-NEXT: %[[PTR:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]], %[[INDEX]]]
// CHECK-NEXT: spv.Load "Function" %[[PTR]]
- %c0 = spv.constant 0: i32
+ %c0 = spv.Constant 0: i32
%0 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
%1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>, i32
%2 = spv.AccessChain %1[%c0, %c0] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, i32, i32
@@ -20,13 +20,13 @@ func @combine_full_access_chain() -> f32 {
// -----
func @combine_access_chain_multi_use() -> !spv.array<4xf32> {
- // CHECK: %[[INDEX:.*]] = spv.constant 0
+ // CHECK: %[[INDEX:.*]] = spv.Constant 0
// CHECK-NEXT: %[[VAR:.*]] = spv.Variable
// CHECK-NEXT: %[[PTR_0:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]]]
// CHECK-NEXT: %[[PTR_1:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]], %[[INDEX]]]
// CHECK-NEXT: spv.Load "Function" %[[PTR_0]]
// CHECK-NEXT: spv.Load "Function" %[[PTR_1]]
- %c0 = spv.constant 0: i32
+ %c0 = spv.Constant 0: i32
%0 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
%1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>, i32
%2 = spv.AccessChain %1[%c0] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, i32
@@ -39,14 +39,14 @@ func @combine_access_chain_multi_use() -> !spv.array<4xf32> {
// -----
func @dont_combine_access_chain_without_common_base() -> !spv.array<4xi32> {
- // CHECK: %[[INDEX:.*]] = spv.constant 1
+ // CHECK: %[[INDEX:.*]] = spv.Constant 1
// CHECK-NEXT: %[[VAR_0:.*]] = spv.Variable
// CHECK-NEXT: %[[VAR_1:.*]] = spv.Variable
// CHECK-NEXT: %[[VAR_0_PTR:.*]] = spv.AccessChain %[[VAR_0]][%[[INDEX]]]
// CHECK-NEXT: %[[VAR_1_PTR:.*]] = spv.AccessChain %[[VAR_1]][%[[INDEX]]]
// CHECK-NEXT: spv.Load "Function" %[[VAR_0_PTR]]
// CHECK-NEXT: spv.Load "Function" %[[VAR_1_PTR]]
- %c1 = spv.constant 1: i32
+ %c1 = spv.Constant 1: i32
%0 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
%1 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
%2 = spv.AccessChain %0[%c1] : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>, i32
@@ -92,10 +92,10 @@ func @convert_bitcast_multi_use(%arg0 : vector<2xf32>, %arg1 : !spv.ptr<i64, Uni
// CHECK-LABEL: extract_vector
func @extract_vector() -> (i32, i32, i32) {
- // CHECK: spv.constant 42 : i32
- // CHECK: spv.constant -33 : i32
- // CHECK: spv.constant 6 : i32
- %0 = spv.constant dense<[42, -33, 6]> : vector<3xi32>
+ // CHECK: spv.Constant 42 : i32
+ // CHECK: spv.Constant -33 : i32
+ // CHECK: spv.Constant 6 : i32
+ %0 = spv.Constant dense<[42, -33, 6]> : vector<3xi32>
%1 = spv.CompositeExtract %0[0 : i32] : vector<3xi32>
%2 = spv.CompositeExtract %0[1 : i32] : vector<3xi32>
%3 = spv.CompositeExtract %0[2 : i32] : vector<3xi32>
@@ -106,9 +106,9 @@ func @extract_vector() -> (i32, i32, i32) {
// CHECK-LABEL: extract_array_final
func @extract_array_final() -> (i32, i32) {
- // CHECK: spv.constant 4 : i32
- // CHECK: spv.constant -5 : i32
- %0 = spv.constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+ // CHECK: spv.Constant 4 : i32
+ // CHECK: spv.Constant -5 : i32
+ %0 = spv.Constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
%1 = spv.CompositeExtract %0[0 : i32, 0 : i32] : !spv.array<1 x vector<2 x i32>>
%2 = spv.CompositeExtract %0[0 : i32, 1 : i32] : !spv.array<1 x vector<2 x i32>>
return %1, %2 : i32, i32
@@ -118,8 +118,8 @@ func @extract_array_final() -> (i32, i32) {
// CHECK-LABEL: extract_array_interm
func @extract_array_interm() -> (vector<2xi32>) {
- // CHECK: spv.constant dense<[4, -5]> : vector<2xi32>
- %0 = spv.constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+ // CHECK: spv.Constant dense<[4, -5]> : vector<2xi32>
+ %0 = spv.Constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
%1 = spv.CompositeExtract %0[0 : i32] : !spv.array<1 x vector<2 x i32>>
return %1 : vector<2xi32>
}
@@ -138,15 +138,15 @@ func @extract_from_not_constant() -> i32 {
// -----
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
// TODO: test constants in
diff erent blocks
func @deduplicate_scalar_constant() -> (i32, i32) {
- // CHECK: %[[CST:.*]] = spv.constant 42 : i32
- %0 = spv.constant 42 : i32
- %1 = spv.constant 42 : i32
+ // CHECK: %[[CST:.*]] = spv.Constant 42 : i32
+ %0 = spv.Constant 42 : i32
+ %1 = spv.Constant 42 : i32
// CHECK-NEXT: return %[[CST]], %[[CST]]
return %0, %1 : i32, i32
}
@@ -154,9 +154,9 @@ func @deduplicate_scalar_constant() -> (i32, i32) {
// -----
func @deduplicate_vector_constant() -> (vector<3xi32>, vector<3xi32>) {
- // CHECK: %[[CST:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %0 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[CST:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %0 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
// CHECK-NEXT: return %[[CST]], %[[CST]]
return %0, %1 : vector<3xi32>, vector<3xi32>
}
@@ -164,9 +164,9 @@ func @deduplicate_vector_constant() -> (vector<3xi32>, vector<3xi32>) {
// -----
func @deduplicate_composite_constant() -> (!spv.array<1 x vector<2xi32>>, !spv.array<1 x vector<2xi32>>) {
- // CHECK: %[[CST:.*]] = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
- %0 = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
- %1 = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+ // CHECK: %[[CST:.*]] = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+ %0 = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+ %1 = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
// CHECK-NEXT: return %[[CST]], %[[CST]]
return %0, %1 : !spv.array<1 x vector<2xi32>>, !spv.array<1 x vector<2xi32>>
}
@@ -180,7 +180,7 @@ func @deduplicate_composite_constant() -> (!spv.array<1 x vector<2xi32>>, !spv.a
// CHECK-LABEL: @iadd_zero
// CHECK-SAME: (%[[ARG:.*]]: i32)
func @iadd_zero(%arg0: i32) -> (i32, i32) {
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
%0 = spv.IAdd %arg0, %zero : i32
%1 = spv.IAdd %zero, %arg0 : i32
// CHECK: return %[[ARG]], %[[ARG]]
@@ -189,12 +189,12 @@ func @iadd_zero(%arg0: i32) -> (i32, i32) {
// CHECK-LABEL: @const_fold_scalar_iadd_normal
func @const_fold_scalar_iadd_normal() -> (i32, i32, i32) {
- %c5 = spv.constant 5 : i32
- %cn8 = spv.constant -8 : i32
+ %c5 = spv.Constant 5 : i32
+ %cn8 = spv.Constant -8 : i32
- // CHECK: spv.constant 10
- // CHECK: spv.constant -16
- // CHECK: spv.constant -3
+ // CHECK: spv.Constant 10
+ // CHECK: spv.Constant -16
+ // CHECK: spv.Constant -3
%0 = spv.IAdd %c5, %c5 : i32
%1 = spv.IAdd %cn8, %cn8 : i32
%2 = spv.IAdd %c5, %cn8 : i32
@@ -203,34 +203,34 @@ func @const_fold_scalar_iadd_normal() -> (i32, i32, i32) {
// CHECK-LABEL: @const_fold_scalar_iadd_flow
func @const_fold_scalar_iadd_flow() -> (i32, i32, i32, i32) {
- %c1 = spv.constant 1 : i32
- %c2 = spv.constant 2 : i32
- %c3 = spv.constant 4294967295 : i32 // 2^32 - 1: 0xffff ffff
- %c4 = spv.constant -2147483648 : i32 // -2^31 : 0x8000 0000
- %c5 = spv.constant -1 : i32 // : 0xffff ffff
- %c6 = spv.constant -2 : i32 // : 0xffff fffe
+ %c1 = spv.Constant 1 : i32
+ %c2 = spv.Constant 2 : i32
+ %c3 = spv.Constant 4294967295 : i32 // 2^32 - 1: 0xffff ffff
+ %c4 = spv.Constant -2147483648 : i32 // -2^31 : 0x8000 0000
+ %c5 = spv.Constant -1 : i32 // : 0xffff ffff
+ %c6 = spv.Constant -2 : i32 // : 0xffff fffe
// 0x0000 0001 + 0xffff ffff = 0x1 0000 0000 -> 0x0000 0000
- // CHECK: spv.constant 0
+ // CHECK: spv.Constant 0
%0 = spv.IAdd %c1, %c3 : i32
// 0x0000 0002 + 0xffff ffff = 0x1 0000 0001 -> 0x0000 0001
- // CHECK: spv.constant 1
+ // CHECK: spv.Constant 1
%1 = spv.IAdd %c2, %c3 : i32
// 0x8000 0000 + 0xffff ffff = 0x1 7fff ffff -> 0x7fff ffff
- // CHECK: spv.constant 2147483647
+ // CHECK: spv.Constant 2147483647
%2 = spv.IAdd %c4, %c5 : i32
// 0x8000 0000 + 0xffff fffe = 0x1 7fff fffe -> 0x7fff fffe
- // CHECK: spv.constant 2147483646
+ // CHECK: spv.Constant 2147483646
%3 = spv.IAdd %c4, %c6 : i32
return %0, %1, %2, %3: i32, i32, i32, i32
}
// CHECK-LABEL: @const_fold_vector_iadd
func @const_fold_vector_iadd() -> vector<3xi32> {
- %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32>
- %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32>
+ %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32>
+ %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32>
- // CHECK: spv.constant dense<[39, -70, 155]>
+ // CHECK: spv.Constant dense<[39, -70, 155]>
%0 = spv.IAdd %vc1, %vc2 : vector<3xi32>
return %0: vector<3xi32>
}
@@ -244,9 +244,9 @@ func @const_fold_vector_iadd() -> vector<3xi32> {
// CHECK-LABEL: @imul_zero_one
// CHECK-SAME: (%[[ARG:.*]]: i32)
func @imul_zero_one(%arg0: i32) -> (i32, i32) {
- // CHECK: %[[ZERO:.*]] = spv.constant 0
- %zero = spv.constant 0 : i32
- %one = spv.constant 1: i32
+ // CHECK: %[[ZERO:.*]] = spv.Constant 0
+ %zero = spv.Constant 0 : i32
+ %one = spv.Constant 1: i32
%0 = spv.IMul %arg0, %zero : i32
%1 = spv.IMul %one, %arg0 : i32
// CHECK: return %[[ZERO]], %[[ARG]]
@@ -255,13 +255,13 @@ func @imul_zero_one(%arg0: i32) -> (i32, i32) {
// CHECK-LABEL: @const_fold_scalar_imul_normal
func @const_fold_scalar_imul_normal() -> (i32, i32, i32) {
- %c5 = spv.constant 5 : i32
- %cn8 = spv.constant -8 : i32
- %c7 = spv.constant 7 : i32
+ %c5 = spv.Constant 5 : i32
+ %cn8 = spv.Constant -8 : i32
+ %c7 = spv.Constant 7 : i32
- // CHECK: spv.constant 35
- // CHECK: spv.constant -40
- // CHECK: spv.constant -56
+ // CHECK: spv.Constant 35
+ // CHECK: spv.Constant -40
+ // CHECK: spv.Constant -56
%0 = spv.IMul %c7, %c5 : i32
%1 = spv.IMul %c5, %cn8 : i32
%2 = spv.IMul %cn8, %c7 : i32
@@ -270,18 +270,18 @@ func @const_fold_scalar_imul_normal() -> (i32, i32, i32) {
// CHECK-LABEL: @const_fold_scalar_imul_flow
func @const_fold_scalar_imul_flow() -> (i32, i32, i32) {
- %c1 = spv.constant 2 : i32
- %c2 = spv.constant 4 : i32
- %c3 = spv.constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff
- %c4 = spv.constant 2147483647 : i32 // 2^31 - 1 : 0x7fff ffff
+ %c1 = spv.Constant 2 : i32
+ %c2 = spv.Constant 4 : i32
+ %c3 = spv.Constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff
+ %c4 = spv.Constant 2147483647 : i32 // 2^31 - 1 : 0x7fff ffff
// (0xffff ffff << 1) = 0x1 ffff fffe -> 0xffff fffe
- // CHECK: %[[CST2:.*]] = spv.constant -2
+ // CHECK: %[[CST2:.*]] = spv.Constant -2
%0 = spv.IMul %c1, %c3 : i32
// (0x7fff ffff << 1) = 0x0 ffff fffe -> 0xffff fffe
%1 = spv.IMul %c1, %c4 : i32
// (0x7fff ffff << 2) = 0x1 ffff fffc -> 0xffff fffc
- // CHECK: %[[CST4:.*]] = spv.constant -4
+ // CHECK: %[[CST4:.*]] = spv.Constant -4
%2 = spv.IMul %c4, %c2 : i32
// CHECK: return %[[CST2]], %[[CST2]], %[[CST4]]
return %0, %1, %2: i32, i32, i32
@@ -290,10 +290,10 @@ func @const_fold_scalar_imul_flow() -> (i32, i32, i32) {
// CHECK-LABEL: @const_fold_vector_imul
func @const_fold_vector_imul() -> vector<3xi32> {
- %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32>
- %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32>
+ %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32>
+ %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32>
- // CHECK: spv.constant dense<[-126, 825, 3556]>
+ // CHECK: spv.Constant dense<[-126, 825, 3556]>
%0 = spv.IMul %vc1, %vc2 : vector<3xi32>
return %0: vector<3xi32>
}
@@ -306,20 +306,20 @@ func @const_fold_vector_imul() -> vector<3xi32> {
// CHECK-LABEL: @isub_x_x
func @isub_x_x(%arg0: i32) -> i32 {
- // CHECK: spv.constant 0
+ // CHECK: spv.Constant 0
%0 = spv.ISub %arg0, %arg0: i32
return %0: i32
}
// CHECK-LABEL: @const_fold_scalar_isub_normal
func @const_fold_scalar_isub_normal() -> (i32, i32, i32) {
- %c5 = spv.constant 5 : i32
- %cn8 = spv.constant -8 : i32
- %c7 = spv.constant 7 : i32
+ %c5 = spv.Constant 5 : i32
+ %cn8 = spv.Constant -8 : i32
+ %c7 = spv.Constant 7 : i32
- // CHECK: spv.constant 2
- // CHECK: spv.constant 13
- // CHECK: spv.constant -15
+ // CHECK: spv.Constant 2
+ // CHECK: spv.Constant 13
+ // CHECK: spv.Constant -15
%0 = spv.ISub %c7, %c5 : i32
%1 = spv.ISub %c5, %cn8 : i32
%2 = spv.ISub %cn8, %c7 : i32
@@ -328,34 +328,34 @@ func @const_fold_scalar_isub_normal() -> (i32, i32, i32) {
// CHECK-LABEL: @const_fold_scalar_isub_flow
func @const_fold_scalar_isub_flow() -> (i32, i32, i32, i32) {
- %c1 = spv.constant 0 : i32
- %c2 = spv.constant 1 : i32
- %c3 = spv.constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff
- %c4 = spv.constant 2147483647 : i32 // 2^31 : 0x7fff ffff
- %c5 = spv.constant -1 : i32 // : 0xffff ffff
- %c6 = spv.constant -2 : i32 // : 0xffff fffe
+ %c1 = spv.Constant 0 : i32
+ %c2 = spv.Constant 1 : i32
+ %c3 = spv.Constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff
+ %c4 = spv.Constant 2147483647 : i32 // 2^31 : 0x7fff ffff
+ %c5 = spv.Constant -1 : i32 // : 0xffff ffff
+ %c6 = spv.Constant -2 : i32 // : 0xffff fffe
// 0x0000 0000 - 0xffff ffff -> 0x0000 0000 + 0x0000 0001 = 0x0000 0001
- // CHECK: spv.constant 1
+ // CHECK: spv.Constant 1
%0 = spv.ISub %c1, %c3 : i32
// 0x0000 0001 - 0xffff ffff -> 0x0000 0001 + 0x0000 0001 = 0x0000 0002
- // CHECK: spv.constant 2
+ // CHECK: spv.Constant 2
%1 = spv.ISub %c2, %c3 : i32
// 0xffff ffff - 0x7fff ffff -> 0xffff ffff + 0x8000 0001 = 0x1 8000 0000
- // CHECK: spv.constant -2147483648
+ // CHECK: spv.Constant -2147483648
%2 = spv.ISub %c5, %c4 : i32
// 0xffff fffe - 0x7fff ffff -> 0xffff fffe + 0x8000 0001 = 0x1 7fff ffff
- // CHECK: spv.constant 2147483647
+ // CHECK: spv.Constant 2147483647
%3 = spv.ISub %c6, %c4 : i32
return %0, %1, %2, %3: i32, i32, i32, i32
}
// CHECK-LABEL: @const_fold_vector_isub
func @const_fold_vector_isub() -> vector<3xi32> {
- %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32>
- %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32>
+ %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32>
+ %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32>
- // CHECK: spv.constant dense<[45, -40, 99]>
+ // CHECK: spv.Constant dense<[45, -40, 99]>
%0 = spv.ISub %vc1, %vc2 : vector<3xi32>
return %0: vector<3xi32>
}
@@ -369,9 +369,9 @@ func @const_fold_vector_isub() -> vector<3xi32> {
// CHECK-LABEL: @convert_logical_and_true_false_scalar
// CHECK-SAME: %[[ARG:.+]]: i1
func @convert_logical_and_true_false_scalar(%arg: i1) -> (i1, i1) {
- %true = spv.constant true
- // CHECK: %[[FALSE:.+]] = spv.constant false
- %false = spv.constant false
+ %true = spv.Constant true
+ // CHECK: %[[FALSE:.+]] = spv.Constant false
+ %false = spv.Constant false
%0 = spv.LogicalAnd %true, %arg: i1
%1 = spv.LogicalAnd %arg, %false: i1
// CHECK: return %[[ARG]], %[[FALSE]]
@@ -381,9 +381,9 @@ func @convert_logical_and_true_false_scalar(%arg: i1) -> (i1, i1) {
// CHECK-LABEL: @convert_logical_and_true_false_vector
// CHECK-SAME: %[[ARG:.+]]: vector<3xi1>
func @convert_logical_and_true_false_vector(%arg: vector<3xi1>) -> (vector<3xi1>, vector<3xi1>) {
- %true = spv.constant dense<true> : vector<3xi1>
- // CHECK: %[[FALSE:.+]] = spv.constant dense<false>
- %false = spv.constant dense<false> : vector<3xi1>
+ %true = spv.Constant dense<true> : vector<3xi1>
+ // CHECK: %[[FALSE:.+]] = spv.Constant dense<false>
+ %false = spv.Constant dense<false> : vector<3xi1>
%0 = spv.LogicalAnd %true, %arg: vector<3xi1>
%1 = spv.LogicalAnd %arg, %false: vector<3xi1>
// CHECK: return %[[ARG]], %[[FALSE]]
@@ -456,9 +456,9 @@ func @convert_logical_not_to_logical_equal(%arg0: vector<3xi1>, %arg1: vector<3x
// CHECK-LABEL: @convert_logical_or_true_false_scalar
// CHECK-SAME: %[[ARG:.+]]: i1
func @convert_logical_or_true_false_scalar(%arg: i1) -> (i1, i1) {
- // CHECK: %[[TRUE:.+]] = spv.constant true
- %true = spv.constant true
- %false = spv.constant false
+ // CHECK: %[[TRUE:.+]] = spv.Constant true
+ %true = spv.Constant true
+ %false = spv.Constant false
%0 = spv.LogicalOr %true, %arg: i1
%1 = spv.LogicalOr %arg, %false: i1
// CHECK: return %[[TRUE]], %[[ARG]]
@@ -468,9 +468,9 @@ func @convert_logical_or_true_false_scalar(%arg: i1) -> (i1, i1) {
// CHECK-LABEL: @convert_logical_or_true_false_vector
// CHECK-SAME: %[[ARG:.+]]: vector<3xi1>
func @convert_logical_or_true_false_vector(%arg: vector<3xi1>) -> (vector<3xi1>, vector<3xi1>) {
- // CHECK: %[[TRUE:.+]] = spv.constant dense<true>
- %true = spv.constant dense<true> : vector<3xi1>
- %false = spv.constant dense<false> : vector<3xi1>
+ // CHECK: %[[TRUE:.+]] = spv.Constant dense<true>
+ %true = spv.Constant dense<true> : vector<3xi1>
+ %false = spv.Constant dense<false> : vector<3xi1>
%0 = spv.LogicalOr %true, %arg: vector<3xi1>
%1 = spv.LogicalOr %arg, %false: vector<3xi1>
// CHECK: return %[[TRUE]], %[[ARG]]
@@ -484,11 +484,11 @@ func @convert_logical_or_true_false_vector(%arg: vector<3xi1>) -> (vector<3xi1>,
//===----------------------------------------------------------------------===//
func @canonicalize_selection_op_scalar_type(%cond: i1) -> () {
- %0 = spv.constant 0: i32
- // CHECK: %[[TRUE_VALUE:.*]] = spv.constant 1 : i32
- %1 = spv.constant 1: i32
- // CHECK: %[[FALSE_VALUE:.*]] = spv.constant 2 : i32
- %2 = spv.constant 2: i32
+ %0 = spv.Constant 0: i32
+ // CHECK: %[[TRUE_VALUE:.*]] = spv.Constant 1 : i32
+ %1 = spv.Constant 1: i32
+ // CHECK: %[[FALSE_VALUE:.*]] = spv.Constant 2 : i32
+ %2 = spv.Constant 2: i32
// CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<i32, Function>
%3 = spv.Variable init(%0) : !spv.ptr<i32, Function>
@@ -515,11 +515,11 @@ func @canonicalize_selection_op_scalar_type(%cond: i1) -> () {
// -----
func @canonicalize_selection_op_vector_type(%cond: i1) -> () {
- %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
- // CHECK: %[[TRUE_VALUE:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- // CHECK: %[[FALSE_VALUE:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+ %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+ // CHECK: %[[TRUE_VALUE:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[FALSE_VALUE:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
// CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
%3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
@@ -547,11 +547,11 @@ func @canonicalize_selection_op_vector_type(%cond: i1) -> () {
// Store to a
diff erent variables.
func @cannot_canonicalize_selection_op_0(%cond: i1) -> () {
- %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+ %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
// CHECK: %[[DST_VAR_0:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
%3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
// CHECK: %[[DST_VAR_1:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
@@ -584,11 +584,11 @@ func @cannot_canonicalize_selection_op_0(%cond: i1) -> () {
// A conditional block consists of more than 2 operations.
func @cannot_canonicalize_selection_op_1(%cond: i1) -> () {
- %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+ %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
// CHECK: %[[DST_VAR_0:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
%3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
// CHECK: %[[DST_VAR_1:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
@@ -620,11 +620,11 @@ func @cannot_canonicalize_selection_op_1(%cond: i1) -> () {
// A control-flow goes into `^then` block from `^else` block.
func @cannot_canonicalize_selection_op_2(%cond: i1) -> () {
- %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+ %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
// CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
%3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
@@ -652,11 +652,11 @@ func @cannot_canonicalize_selection_op_2(%cond: i1) -> () {
// `spv.Return` as a block terminator.
func @cannot_canonicalize_selection_op_3(%cond: i1) -> () {
- %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+ %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
// CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
%3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
@@ -684,11 +684,11 @@ func @cannot_canonicalize_selection_op_3(%cond: i1) -> () {
// Different memory access attributes.
func @cannot_canonicalize_selection_op_4(%cond: i1) -> () {
- %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
- // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+ %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+ // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
// CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
%3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
diff --git a/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir b/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir
index 90e9b85b9035..8f31ba716f85 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir
@@ -2,10 +2,10 @@
// CHECK: func @clamp_fordlessthan(%[[INPUT:.*]]: f32)
func @clamp_fordlessthan(%input: f32) -> f32 {
- // CHECK: %[[MIN:.*]] = spv.constant
- %min = spv.constant 0.5 : f32
- // CHECK: %[[MAX:.*]] = spv.constant
- %max = spv.constant 1.0 : f32
+ // CHECK: %[[MIN:.*]] = spv.Constant
+ %min = spv.Constant 0.5 : f32
+ // CHECK: %[[MAX:.*]] = spv.Constant
+ %max = spv.Constant 1.0 : f32
// CHECK: [[RES:%.*]] = spv.GLSL.FClamp %[[INPUT]], %[[MIN]], %[[MAX]]
%0 = spv.FOrdLessThan %min, %input : f32
@@ -21,10 +21,10 @@ func @clamp_fordlessthan(%input: f32) -> f32 {
// CHECK: func @clamp_fordlessthanequal(%[[INPUT:.*]]: f32)
func @clamp_fordlessthanequal(%input: f32) -> f32 {
- // CHECK: %[[MIN:.*]] = spv.constant
- %min = spv.constant 0.5 : f32
- // CHECK: %[[MAX:.*]] = spv.constant
- %max = spv.constant 1.0 : f32
+ // CHECK: %[[MIN:.*]] = spv.Constant
+ %min = spv.Constant 0.5 : f32
+ // CHECK: %[[MAX:.*]] = spv.Constant
+ %max = spv.Constant 1.0 : f32
// CHECK: [[RES:%.*]] = spv.GLSL.FClamp %[[INPUT]], %[[MIN]], %[[MAX]]
%0 = spv.FOrdLessThanEqual %min, %input : f32
@@ -40,10 +40,10 @@ func @clamp_fordlessthanequal(%input: f32) -> f32 {
// CHECK: func @clamp_slessthan(%[[INPUT:.*]]: si32)
func @clamp_slessthan(%input: si32) -> si32 {
- // CHECK: %[[MIN:.*]] = spv.constant
- %min = spv.constant 0 : si32
- // CHECK: %[[MAX:.*]] = spv.constant
- %max = spv.constant 10 : si32
+ // CHECK: %[[MIN:.*]] = spv.Constant
+ %min = spv.Constant 0 : si32
+ // CHECK: %[[MAX:.*]] = spv.Constant
+ %max = spv.Constant 10 : si32
// CHECK: [[RES:%.*]] = spv.GLSL.SClamp %[[INPUT]], %[[MIN]], %[[MAX]]
%0 = spv.SLessThan %min, %input : si32
@@ -59,10 +59,10 @@ func @clamp_slessthan(%input: si32) -> si32 {
// CHECK: func @clamp_slessthanequal(%[[INPUT:.*]]: si32)
func @clamp_slessthanequal(%input: si32) -> si32 {
- // CHECK: %[[MIN:.*]] = spv.constant
- %min = spv.constant 0 : si32
- // CHECK: %[[MAX:.*]] = spv.constant
- %max = spv.constant 10 : si32
+ // CHECK: %[[MIN:.*]] = spv.Constant
+ %min = spv.Constant 0 : si32
+ // CHECK: %[[MAX:.*]] = spv.Constant
+ %max = spv.Constant 10 : si32
// CHECK: [[RES:%.*]] = spv.GLSL.SClamp %[[INPUT]], %[[MIN]], %[[MAX]]
%0 = spv.SLessThanEqual %min, %input : si32
@@ -78,10 +78,10 @@ func @clamp_slessthanequal(%input: si32) -> si32 {
// CHECK: func @clamp_ulessthan(%[[INPUT:.*]]: i32)
func @clamp_ulessthan(%input: i32) -> i32 {
- // CHECK: %[[MIN:.*]] = spv.constant
- %min = spv.constant 0 : i32
- // CHECK: %[[MAX:.*]] = spv.constant
- %max = spv.constant 10 : i32
+ // CHECK: %[[MIN:.*]] = spv.Constant
+ %min = spv.Constant 0 : i32
+ // CHECK: %[[MAX:.*]] = spv.Constant
+ %max = spv.Constant 10 : i32
// CHECK: [[RES:%.*]] = spv.GLSL.UClamp %[[INPUT]], %[[MIN]], %[[MAX]]
%0 = spv.ULessThan %min, %input : i32
@@ -97,10 +97,10 @@ func @clamp_ulessthan(%input: i32) -> i32 {
// CHECK: func @clamp_ulessthanequal(%[[INPUT:.*]]: i32)
func @clamp_ulessthanequal(%input: i32) -> i32 {
- // CHECK: %[[MIN:.*]] = spv.constant
- %min = spv.constant 0 : i32
- // CHECK: %[[MAX:.*]] = spv.constant
- %max = spv.constant 10 : i32
+ // CHECK: %[[MIN:.*]] = spv.Constant
+ %min = spv.Constant 0 : i32
+ // CHECK: %[[MAX:.*]] = spv.Constant
+ %max = spv.Constant 10 : i32
// CHECK: [[RES:%.*]] = spv.GLSL.UClamp %[[INPUT]], %[[MIN]], %[[MAX]]
%0 = spv.ULessThanEqual %min, %input : i32
diff --git a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
index 983fc2611223..7b663591da56 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
@@ -17,13 +17,13 @@ spv.module Logical GLSL450 {
spv.module Logical GLSL450 {
spv.func @callee() -> i32 "None" {
- %0 = spv.constant 42 : i32
+ %0 = spv.Constant 42 : i32
spv.ReturnValue %0 : i32
}
// CHECK-LABEL: @calling_single_block_retval_func
spv.func @calling_single_block_retval_func() -> i32 "None" {
- // CHECK-NEXT: %[[CST:.*]] = spv.constant 42
+ // CHECK-NEXT: %[[CST:.*]] = spv.Constant 42
%0 = spv.FunctionCall @callee() : () -> (i32)
// CHECK-NEXT: spv.ReturnValue %[[CST]]
spv.ReturnValue %0 : i32
@@ -36,12 +36,12 @@ 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.mlir.addressof @data : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
- %1 = spv.constant 0: i32
+ %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
+ %3 = spv.Constant 42: i32
spv.Store "StorageBuffer" %2, %3 : i32
spv.Return
}
@@ -49,11 +49,11 @@ spv.module Logical GLSL450 {
// CHECK-LABEL: @calling_multi_block_ret_func
spv.func @calling_multi_block_ret_func() "None" {
// CHECK-NEXT: spv.mlir.addressof
- // CHECK-NEXT: spv.constant 0
+ // CHECK-NEXT: spv.Constant 0
// CHECK-NEXT: spv.AccessChain
// CHECK-NEXT: spv.Branch ^bb1
// CHECK-NEXT: ^bb1:
- // CHECK-NEXT: spv.constant
+ // CHECK-NEXT: spv.Constant
// CHECK-NEXT: spv.Store
// CHECK-NEXT: spv.Branch ^bb2
spv.FunctionCall @callee() : () -> ()
@@ -81,7 +81,7 @@ spv.module Logical GLSL450 {
// CHECK-LABEL: @calling_selection_ret_func
spv.func @calling_selection_ret_func() "None" {
- %0 = spv.constant true
+ %0 = spv.Constant true
// CHECK: spv.FunctionCall
spv.FunctionCall @callee(%0) : (i1) -> ()
spv.Return
@@ -104,8 +104,8 @@ spv.module Logical GLSL450 {
// CHECK-LABEL: @calling_selection_no_ret_func
spv.func @calling_selection_no_ret_func() "None" {
- // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true
- %0 = spv.constant true
+ // CHECK-NEXT: %[[TRUE:.*]] = spv.Constant true
+ %0 = spv.Constant true
// CHECK-NEXT: spv.selection
// CHECK-NEXT: spv.BranchConditional %[[TRUE]], ^bb1, ^bb2
// CHECK-NEXT: ^bb1:
@@ -137,7 +137,7 @@ spv.module Logical GLSL450 {
// CHECK-LABEL: @calling_loop_ret_func
spv.func @calling_loop_ret_func() "None" {
- %0 = spv.constant true
+ %0 = spv.Constant true
// CHECK: spv.FunctionCall
spv.FunctionCall @callee(%0) : (i1) -> ()
spv.Return
@@ -164,8 +164,8 @@ spv.module Logical GLSL450 {
// CHECK-LABEL: @calling_loop_no_ret_func
spv.func @calling_loop_no_ret_func() "None" {
- // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true
- %0 = spv.constant true
+ // CHECK-NEXT: %[[TRUE:.*]] = spv.Constant true
+ %0 = spv.Constant true
// CHECK-NEXT: spv.loop
// CHECK-NEXT: spv.Branch ^bb1
// CHECK-NEXT: ^bb1:
@@ -189,7 +189,7 @@ spv.module Logical GLSL450 {
// CHECK: @inline_into_selection_region
spv.func @inline_into_selection_region() "None" {
- %1 = spv.constant 0 : i32
+ %1 = spv.Constant 0 : i32
// CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv.mlir.addressof @arg_0
// CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv.mlir.addressof @arg_1
// CHECK-DAG: [[LOADPTR:%.*]] = spv.AccessChain [[ADDRESS_ARG0]]
diff --git a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
index 2a7eb653c603..f21f97cc7563 100644
--- a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
+++ b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
@@ -20,7 +20,7 @@ spv.module Logical GLSL450 {
spv.globalVariable @var5 bind(1,3) : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, StorageBuffer>
spv.func @kernel() -> () "None" {
- %c0 = spv.constant 0 : i32
+ %c0 = spv.Constant 0 : i32
// CHECK: {{%.*}} = spv.mlir.addressof @var0 : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
%0 = spv.mlir.addressof @var0 : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
// CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
diff --git a/mlir/test/Target/SPIRV/constant.mlir b/mlir/test/Target/SPIRV/constant.mlir
index 520669262e12..e9955ce707f3 100644
--- a/mlir/test/Target/SPIRV/constant.mlir
+++ b/mlir/test/Target/SPIRV/constant.mlir
@@ -3,10 +3,10 @@
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @bool_const
spv.func @bool_const() -> () "None" {
- // CHECK: spv.constant true
- %0 = spv.constant true
- // CHECK: spv.constant false
- %1 = spv.constant false
+ // CHECK: spv.Constant true
+ %0 = spv.Constant true
+ // CHECK: spv.Constant false
+ %1 = spv.Constant false
%2 = spv.Variable init(%0): !spv.ptr<i1, Function>
%3 = spv.Variable init(%1): !spv.ptr<i1, Function>
@@ -15,12 +15,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @i32_const
spv.func @i32_const() -> () "None" {
- // CHECK: spv.constant 0 : i32
- %0 = spv.constant 0 : i32
- // CHECK: spv.constant 10 : i32
- %1 = spv.constant 10 : i32
- // CHECK: spv.constant -5 : i32
- %2 = spv.constant -5 : i32
+ // CHECK: spv.Constant 0 : i32
+ %0 = spv.Constant 0 : i32
+ // CHECK: spv.Constant 10 : i32
+ %1 = spv.Constant 10 : i32
+ // CHECK: spv.Constant -5 : i32
+ %2 = spv.Constant -5 : i32
%3 = spv.IAdd %0, %1 : i32
%4 = spv.IAdd %2, %3 : i32
@@ -29,12 +29,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @si32_const
spv.func @si32_const() -> () "None" {
- // CHECK: spv.constant 0 : si32
- %0 = spv.constant 0 : si32
- // CHECK: spv.constant 10 : si32
- %1 = spv.constant 10 : si32
- // CHECK: spv.constant -5 : si32
- %2 = spv.constant -5 : si32
+ // CHECK: spv.Constant 0 : si32
+ %0 = spv.Constant 0 : si32
+ // CHECK: spv.Constant 10 : si32
+ %1 = spv.Constant 10 : si32
+ // CHECK: spv.Constant -5 : si32
+ %2 = spv.Constant -5 : si32
%3 = spv.IAdd %0, %1 : si32
%4 = spv.IAdd %2, %3 : si32
@@ -46,12 +46,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// because they all use 1 as the signedness bit. So we always treat them
// as signless integers.
spv.func @ui32_const() -> () "None" {
- // CHECK: spv.constant 0 : i32
- %0 = spv.constant 0 : ui32
- // CHECK: spv.constant 10 : i32
- %1 = spv.constant 10 : ui32
- // CHECK: spv.constant -5 : i32
- %2 = spv.constant 4294967291 : ui32
+ // CHECK: spv.Constant 0 : i32
+ %0 = spv.Constant 0 : ui32
+ // CHECK: spv.Constant 10 : i32
+ %1 = spv.Constant 10 : ui32
+ // CHECK: spv.Constant -5 : i32
+ %2 = spv.Constant 4294967291 : ui32
%3 = spv.IAdd %0, %1 : ui32
%4 = spv.IAdd %2, %3 : ui32
@@ -60,14 +60,14 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @i64_const
spv.func @i64_const() -> () "None" {
- // CHECK: spv.constant 4294967296 : i64
- %0 = spv.constant 4294967296 : i64 // 2^32
- // CHECK: spv.constant -4294967296 : i64
- %1 = spv.constant -4294967296 : i64 // -2^32
- // CHECK: spv.constant 9223372036854775807 : i64
- %2 = spv.constant 9223372036854775807 : i64 // 2^63 - 1
- // CHECK: spv.constant -9223372036854775808 : i64
- %3 = spv.constant -9223372036854775808 : i64 // -2^63
+ // CHECK: spv.Constant 4294967296 : i64
+ %0 = spv.Constant 4294967296 : i64 // 2^32
+ // CHECK: spv.Constant -4294967296 : i64
+ %1 = spv.Constant -4294967296 : i64 // -2^32
+ // CHECK: spv.Constant 9223372036854775807 : i64
+ %2 = spv.Constant 9223372036854775807 : i64 // 2^63 - 1
+ // CHECK: spv.Constant -9223372036854775808 : i64
+ %3 = spv.Constant -9223372036854775808 : i64 // -2^63
%4 = spv.IAdd %0, %1 : i64
%5 = spv.IAdd %2, %3 : i64
@@ -76,10 +76,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @i16_const
spv.func @i16_const() -> () "None" {
- // CHECK: spv.constant -32768 : i16
- %0 = spv.constant -32768 : i16 // -2^15
- // CHECK: spv.constant 32767 : i16
- %1 = spv.constant 32767 : i16 // 2^15 - 1
+ // CHECK: spv.Constant -32768 : i16
+ %0 = spv.Constant -32768 : i16 // -2^15
+ // CHECK: spv.Constant 32767 : i16
+ %1 = spv.Constant 32767 : i16 // 2^15 - 1
%2 = spv.IAdd %0, %1 : i16
spv.Return
@@ -87,18 +87,18 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @float_const
spv.func @float_const() -> () "None" {
- // CHECK: spv.constant 0.000000e+00 : f32
- %0 = spv.constant 0. : f32
- // CHECK: spv.constant 1.000000e+00 : f32
- %1 = spv.constant 1. : f32
- // CHECK: spv.constant -0.000000e+00 : f32
- %2 = spv.constant -0. : f32
- // CHECK: spv.constant -1.000000e+00 : f32
- %3 = spv.constant -1. : f32
- // CHECK: spv.constant 7.500000e-01 : f32
- %4 = spv.constant 0.75 : f32
- // CHECK: spv.constant -2.500000e-01 : f32
- %5 = spv.constant -0.25 : f32
+ // CHECK: spv.Constant 0.000000e+00 : f32
+ %0 = spv.Constant 0. : f32
+ // CHECK: spv.Constant 1.000000e+00 : f32
+ %1 = spv.Constant 1. : f32
+ // CHECK: spv.Constant -0.000000e+00 : f32
+ %2 = spv.Constant -0. : f32
+ // CHECK: spv.Constant -1.000000e+00 : f32
+ %3 = spv.Constant -1. : f32
+ // CHECK: spv.Constant 7.500000e-01 : f32
+ %4 = spv.Constant 0.75 : f32
+ // CHECK: spv.Constant -2.500000e-01 : f32
+ %5 = spv.Constant -0.25 : f32
%6 = spv.FAdd %0, %1 : f32
%7 = spv.FAdd %2, %3 : f32
@@ -109,10 +109,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @double_const
spv.func @double_const() -> () "None" {
// TODO: test range boundary values
- // CHECK: spv.constant 1.024000e+03 : f64
- %0 = spv.constant 1024. : f64
- // CHECK: spv.constant -1.024000e+03 : f64
- %1 = spv.constant -1024. : f64
+ // CHECK: spv.Constant 1.024000e+03 : f64
+ %0 = spv.Constant 1024. : f64
+ // CHECK: spv.Constant -1.024000e+03 : f64
+ %1 = spv.Constant -1024. : f64
%2 = spv.FAdd %0, %1 : f64
spv.Return
@@ -120,10 +120,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @half_const
spv.func @half_const() -> () "None" {
- // CHECK: spv.constant 5.120000e+02 : f16
- %0 = spv.constant 512. : f16
- // CHECK: spv.constant -5.120000e+02 : f16
- %1 = spv.constant -512. : f16
+ // CHECK: spv.Constant 5.120000e+02 : f16
+ %0 = spv.Constant 512. : f16
+ // CHECK: spv.Constant -5.120000e+02 : f16
+ %1 = spv.Constant -512. : f16
%2 = spv.FAdd %0, %1 : f16
spv.Return
@@ -131,12 +131,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @bool_vector_const
spv.func @bool_vector_const() -> () "None" {
- // CHECK: spv.constant dense<false> : vector<2xi1>
- %0 = spv.constant dense<false> : vector<2xi1>
- // CHECK: spv.constant dense<[true, true, true]> : vector<3xi1>
- %1 = spv.constant dense<true> : vector<3xi1>
- // CHECK: spv.constant dense<[false, true]> : vector<2xi1>
- %2 = spv.constant dense<[false, true]> : vector<2xi1>
+ // CHECK: spv.Constant dense<false> : vector<2xi1>
+ %0 = spv.Constant dense<false> : vector<2xi1>
+ // CHECK: spv.Constant dense<[true, true, true]> : vector<3xi1>
+ %1 = spv.Constant dense<true> : vector<3xi1>
+ // CHECK: spv.Constant dense<[false, true]> : vector<2xi1>
+ %2 = spv.Constant dense<[false, true]> : vector<2xi1>
%3 = spv.Variable init(%0): !spv.ptr<vector<2xi1>, Function>
%4 = spv.Variable init(%1): !spv.ptr<vector<3xi1>, Function>
@@ -146,12 +146,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @int_vector_const
spv.func @int_vector_const() -> () "None" {
- // CHECK: spv.constant dense<0> : vector<3xi32>
- %0 = spv.constant dense<0> : vector<3xi32>
- // CHECK: spv.constant dense<1> : vector<3xi32>
- %1 = spv.constant dense<1> : vector<3xi32>
- // CHECK: spv.constant dense<[2, -3, 4]> : vector<3xi32>
- %2 = spv.constant dense<[2, -3, 4]> : vector<3xi32>
+ // CHECK: spv.Constant dense<0> : vector<3xi32>
+ %0 = spv.Constant dense<0> : vector<3xi32>
+ // CHECK: spv.Constant dense<1> : vector<3xi32>
+ %1 = spv.Constant dense<1> : vector<3xi32>
+ // CHECK: spv.Constant dense<[2, -3, 4]> : vector<3xi32>
+ %2 = spv.Constant dense<[2, -3, 4]> : vector<3xi32>
%3 = spv.IAdd %0, %1 : vector<3xi32>
%4 = spv.IAdd %2, %3 : vector<3xi32>
@@ -160,12 +160,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @fp_vector_const
spv.func @fp_vector_const() -> () "None" {
- // CHECK: spv.constant dense<0.000000e+00> : vector<4xf32>
- %0 = spv.constant dense<0.> : vector<4xf32>
- // CHECK: spv.constant dense<-1.500000e+01> : vector<4xf32>
- %1 = spv.constant dense<-15.> : vector<4xf32>
- // CHECK: spv.constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32>
- %2 = spv.constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32>
+ // CHECK: spv.Constant dense<0.000000e+00> : vector<4xf32>
+ %0 = spv.Constant dense<0.> : vector<4xf32>
+ // CHECK: spv.Constant dense<-1.500000e+01> : vector<4xf32>
+ %1 = spv.Constant dense<-15.> : vector<4xf32>
+ // CHECK: spv.Constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32>
+ %2 = spv.Constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32>
%3 = spv.FAdd %0, %1 : vector<4xf32>
%4 = spv.FAdd %2, %3 : vector<4xf32>
@@ -174,51 +174,51 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @ui64_array_const
spv.func @ui64_array_const() -> (!spv.array<3xui64>) "None" {
- // CHECK: spv.constant [5, 6, 7] : !spv.array<3 x i64>
- %0 = spv.constant [5 : ui64, 6 : ui64, 7 : ui64] : !spv.array<3 x ui64>
+ // CHECK: spv.Constant [5, 6, 7] : !spv.array<3 x i64>
+ %0 = spv.Constant [5 : ui64, 6 : ui64, 7 : ui64] : !spv.array<3 x ui64>
spv.ReturnValue %0: !spv.array<3xui64>
}
// CHECK-LABEL: @si32_array_const
spv.func @si32_array_const() -> (!spv.array<3xsi32>) "None" {
- // CHECK: spv.constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
- %0 = spv.constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
+ // CHECK: spv.Constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
+ %0 = spv.Constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
spv.ReturnValue %0 : !spv.array<3xsi32>
}
// CHECK-LABEL: @float_array_const
spv.func @float_array_const() -> (!spv.array<2 x vector<2xf32>>) "None" {
- // CHECK: spv.constant [dense<3.000000e+00> : vector<2xf32>, dense<[4.000000e+00, 5.000000e+00]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
- %0 = spv.constant [dense<3.0> : vector<2xf32>, dense<[4., 5.]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
+ // CHECK: spv.Constant [dense<3.000000e+00> : vector<2xf32>, dense<[4.000000e+00, 5.000000e+00]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
+ %0 = spv.Constant [dense<3.0> : vector<2xf32>, dense<[4., 5.]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
spv.ReturnValue %0 : !spv.array<2 x vector<2xf32>>
}
// CHECK-LABEL: @ignore_not_used_const
spv.func @ignore_not_used_const() -> () "None" {
- %0 = spv.constant false
+ %0 = spv.Constant false
// CHECK-NEXT: spv.Return
spv.Return
}
// CHECK-LABEL: @materialize_const_at_each_use
spv.func @materialize_const_at_each_use() -> (i32) "None" {
- // CHECK: %[[USE1:.*]] = spv.constant 42 : i32
- // CHECK: %[[USE2:.*]] = spv.constant 42 : i32
+ // CHECK: %[[USE1:.*]] = spv.Constant 42 : i32
+ // CHECK: %[[USE2:.*]] = spv.Constant 42 : i32
// CHECK: spv.IAdd %[[USE1]], %[[USE2]]
- %0 = spv.constant 42 : i32
+ %0 = spv.Constant 42 : i32
%1 = spv.IAdd %0, %0 : i32
spv.ReturnValue %1 : i32
}
// CHECK-LABEL: @const_variable
spv.func @const_variable(%arg0 : i32, %arg1 : i32) -> () "None" {
- // CHECK: %[[CONST:.*]] = spv.constant 5 : i32
+ // CHECK: %[[CONST:.*]] = spv.Constant 5 : i32
// CHECK: spv.Variable init(%[[CONST]]) : !spv.ptr<i32, Function>
// CHECK: spv.IAdd %arg0, %arg1
%0 = spv.IAdd %arg0, %arg1 : i32
- %1 = spv.constant 5 : i32
+ %1 = spv.Constant 5 : i32
%2 = spv.Variable init(%1) : !spv.ptr<i32, Function>
%3 = spv.Load "Function" %2 : i32
%4 = spv.IAdd %0, %3 : i32
@@ -227,15 +227,15 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @multi_dimensions_const
spv.func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" {
- // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
- %0 = spv.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+ // CHECK: spv.Constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+ %0 = spv.Constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
}
// CHECK-LABEL: @multi_dimensions_splat_const
spv.func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" {
- // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
- %0 = spv.constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+ // CHECK: spv.Constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+ %0 = spv.Constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
}
}
diff --git a/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir b/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir
index 3f8c4bff4738..bf7deebad6ef 100644
--- a/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir
+++ b/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir
@@ -94,7 +94,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [CooperativeMatrixNV], [SPV_N
// CHECK-LABEL: @cooperative_matrix_access_chain
spv.func @cooperative_matrix_access_chain(%a : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>) -> !spv.ptr<f32, Function> "None" {
- %0 = spv.constant 0: i32
+ %0 = spv.Constant 0: i32
// CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
%1 = spv.AccessChain %a[%0] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
spv.ReturnValue %1 : !spv.ptr<f32, Function>
diff --git a/mlir/test/Target/SPIRV/debug.mlir b/mlir/test/Target/SPIRV/debug.mlir
index 1b4abefd4dbe..88fb919156c8 100644
--- a/mlir/test/Target/SPIRV/debug.mlir
+++ b/mlir/test/Target/SPIRV/debug.mlir
@@ -44,7 +44,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
}
spv.func @local_var() "None" {
- %zero = spv.constant 0: i32
+ %zero = spv.Constant 0: i32
// CHECK: loc({{".*debug.mlir"}}:49:12)
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.Return
@@ -68,8 +68,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
}
spv.func @loop(%count : i32) -> () "None" {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%ivar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
%jvar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.loop {
@@ -121,9 +121,9 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
}
spv.func @selection(%cond: i1) -> () "None" {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
- %two = spv.constant 2: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
+ %two = spv.Constant 2: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.selection {
// CHECK: loc({{".*debug.mlir"}}:128:5)
diff --git a/mlir/test/Target/SPIRV/function-call.mlir b/mlir/test/Target/SPIRV/function-call.mlir
index 04d375af2079..92be451fc6a5 100644
--- a/mlir/test/Target/SPIRV/function-call.mlir
+++ b/mlir/test/Target/SPIRV/function-call.mlir
@@ -3,7 +3,7 @@
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
+ %0 = spv.Constant 16 : i32
%1 = spv.mlir.addressof @var1 : !spv.ptr<!spv.array<4xf32>, Input>
// CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}) : (i32) -> i32
%3 = spv.FunctionCall @f_0(%0) : (i32) -> i32
@@ -24,7 +24,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
}
spv.func @f_loop_with_function_call(%count : i32) -> () "None" {
- %zero = spv.constant 0: i32
+ %zero = spv.Constant 0: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
spv.loop {
spv.Branch ^header
@@ -44,7 +44,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.Return
}
spv.func @f_inc(%arg0 : !spv.ptr<i32, Function>) -> () "None" {
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
%0 = spv.Load "Function" %arg0 : i32
%1 = spv.IAdd %0, %one : i32
spv.Store "Function" %arg0, %1 : i32
diff --git a/mlir/test/Target/SPIRV/global-variable.mlir b/mlir/test/Target/SPIRV/global-variable.mlir
index 1cdba5876932..16c1061c1b2c 100644
--- a/mlir/test/Target/SPIRV/global-variable.mlir
+++ b/mlir/test/Target/SPIRV/global-variable.mlir
@@ -28,7 +28,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @foo() "None" {
// CHECK: %[[ADDR:.*]] = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
%0 = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
- %1 = spv.constant 0: i32
+ %1 = spv.Constant 0: i32
// CHECK: spv.AccessChain %[[ADDR]]
%2 = spv.AccessChain %0[%1] : !spv.ptr<vector<3xi32>, Input>, i32
spv.Return
diff --git a/mlir/test/Target/SPIRV/logical-ops.mlir b/mlir/test/Target/SPIRV/logical-ops.mlir
index 7be965db6783..6a5d532d7978 100644
--- a/mlir/test/Target/SPIRV/logical-ops.mlir
+++ b/mlir/test/Target/SPIRV/logical-ops.mlir
@@ -93,16 +93,16 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
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
+ %0 = spv.Constant 4.0 : f32
+ %1 = spv.Constant 5.0 : f32
%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>
+ %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>
+ %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
diff --git a/mlir/test/Target/SPIRV/loop.mlir b/mlir/test/Target/SPIRV/loop.mlir
index d54754e86494..a18a420f8850 100644
--- a/mlir/test/Target/SPIRV/loop.mlir
+++ b/mlir/test/Target/SPIRV/loop.mlir
@@ -5,8 +5,8 @@
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// for (int i = 0; i < count; ++i) {}
spv.func @loop(%count : i32) -> () "None" {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
// CHECK: spv.Branch ^bb1
@@ -35,7 +35,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
^continue:
// CHECK-NEXT: spv.Load
%val1 = spv.Load "Function" %var : i32
-// CHECK-NEXT: spv.constant 1
+// CHECK-NEXT: spv.Constant 1
// CHECK-NEXT: spv.IAdd
%add = spv.IAdd %val1, %one : i32
// CHECK-NEXT: spv.Store
@@ -64,13 +64,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
spv.func @loop_kernel() "None" {
%0 = spv.mlir.addressof @GV1 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
- %1 = spv.constant 0 : i32
+ %1 = spv.Constant 0 : i32
%2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
%3 = spv.mlir.addressof @GV2 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
%5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
- %6 = spv.constant 4 : i32
- %7 = spv.constant 42 : i32
- %8 = spv.constant 2 : 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 {
@@ -112,8 +112,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// for (int j = 0; j < count; ++j) { }
// }
spv.func @loop(%count : i32) -> () "None" {
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
%ivar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
%jvar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
@@ -135,7 +135,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-NEXT: ^bb2:
^body:
-// CHECK-NEXT: spv.constant 0
+// CHECK-NEXT: spv.Constant 0
// CHECK-NEXT: spv.Store
spv.Store "Function" %jvar, %zero : i32
// CHECK-NEXT: spv.Branch ^bb3
@@ -164,7 +164,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
^continue:
// CHECK-NEXT: spv.Load
%jval1 = spv.Load "Function" %jvar : i32
-// CHECK-NEXT: spv.constant 1
+// CHECK-NEXT: spv.Constant 1
// CHECK-NEXT: spv.IAdd
%add = spv.IAdd %jval1, %one : i32
// CHECK-NEXT: spv.Store
@@ -185,7 +185,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
^continue:
// CHECK-NEXT: spv.Load
%ival1 = spv.Load "Function" %ivar : i32
-// CHECK-NEXT: spv.constant 1
+// CHECK-NEXT: spv.Constant 1
// CHECK-NEXT: spv.IAdd
%add = spv.IAdd %ival1, %one : i32
// CHECK-NEXT: spv.Store
diff --git a/mlir/test/Target/SPIRV/memory-ops.mlir b/mlir/test/Target/SPIRV/memory-ops.mlir
index 8bf021707b12..8730044d769e 100644
--- a/mlir/test/Target/SPIRV/memory-ops.mlir
+++ b/mlir/test/Target/SPIRV/memory-ops.mlir
@@ -30,13 +30,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @load_store_zero_rank_float(%arg0: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>) "None" {
// CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>
// CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : f32
- %0 = spv.constant 0 : i32
+ %0 = spv.Constant 0 : i32
%1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
%2 = spv.Load "StorageBuffer" %1 : f32
// CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>
// CHECK-NEXT: spv.Store "StorageBuffer" [[STORE_PTR]], [[VAL]] : f32
- %3 = spv.constant 0 : i32
+ %3 = spv.Constant 0 : i32
%4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
spv.Store "StorageBuffer" %4, %2 : f32
spv.Return
@@ -45,13 +45,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @load_store_zero_rank_int(%arg0: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>) "None" {
// CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>
// CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : i32
- %0 = spv.constant 0 : i32
+ %0 = spv.Constant 0 : i32
%1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
%2 = spv.Load "StorageBuffer" %1 : i32
// CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>
// CHECK-NEXT: spv.Store "StorageBuffer" [[STORE_PTR]], [[VAL]] : i32
- %3 = spv.constant 0 : i32
+ %3 = spv.Constant 0 : i32
%4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
spv.Store "StorageBuffer" %4, %2 : i32
spv.Return
diff --git a/mlir/test/Target/SPIRV/non-uniform-ops.mlir b/mlir/test/Target/SPIRV/non-uniform-ops.mlir
index 5cb035f02282..e0d576f1b52c 100644
--- a/mlir/test/Target/SPIRV/non-uniform-ops.mlir
+++ b/mlir/test/Target/SPIRV/non-uniform-ops.mlir
@@ -10,7 +10,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @group_non_uniform_broadcast
spv.func @group_non_uniform_broadcast(%value: f32) -> f32 "None" {
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
// CHECK: spv.GroupNonUniformBroadcast Subgroup %{{.*}}, %{{.*}} : f32, i32
%0 = spv.GroupNonUniformBroadcast Subgroup %value, %one : f32, i32
spv.ReturnValue %0: f32
@@ -60,7 +60,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce
spv.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> "None" {
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
// CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
%0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
spv.ReturnValue %0: vector<2xi32>
diff --git a/mlir/test/Target/SPIRV/phi.mlir b/mlir/test/Target/SPIRV/phi.mlir
index bb94ecc1bacc..ba320b317135 100644
--- a/mlir/test/Target/SPIRV/phi.mlir
+++ b/mlir/test/Target/SPIRV/phi.mlir
@@ -4,8 +4,8 @@
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
-// CHECK: %[[CST:.*]] = spv.constant 0
- %zero = spv.constant 0 : i32
+// CHECK: %[[CST:.*]] = spv.Constant 0
+ %zero = spv.Constant 0 : i32
// CHECK-NEXT: spv.Branch ^bb1(%[[CST]] : i32)
spv.Branch ^bb1(%zero : i32)
// CHECK-NEXT: ^bb1(%{{.*}}: i32):
@@ -25,10 +25,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
-// CHECK: %[[ZERO:.*]] = spv.constant 0
- %zero = spv.constant 0 : i32
-// CHECK-NEXT: %[[ONE:.*]] = spv.constant 1
- %one = spv.constant 1.0 : f32
+// CHECK: %[[ZERO:.*]] = spv.Constant 0
+ %zero = spv.Constant 0 : i32
+// CHECK-NEXT: %[[ONE:.*]] = spv.Constant 1
+ %one = spv.Constant 1.0 : f32
// CHECK-NEXT: spv.Branch ^bb1(%[[ZERO]], %[[ONE]] : i32, f32)
spv.Branch ^bb1(%zero, %one : i32, f32)
@@ -49,8 +49,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @foo() -> () "None" {
-// CHECK: %[[CST0:.*]] = spv.constant 0
- %zero = spv.constant 0 : i32
+// CHECK: %[[CST0:.*]] = spv.Constant 0
+ %zero = spv.Constant 0 : i32
// CHECK-NEXT: spv.Branch ^bb1(%[[CST0]] : i32)
spv.Branch ^bb1(%zero : i32)
@@ -58,7 +58,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
^bb1(%arg0: i32):
// CHECK-NEXT: %[[ADD:.*]] = spv.IAdd %[[ARG]], %[[ARG]] : i32
%0 = spv.IAdd %arg0, %arg0 : i32
-// CHECK-NEXT: %[[CST1:.*]] = spv.constant 0
+// CHECK-NEXT: %[[CST1:.*]] = spv.Constant 0
// CHECK-NEXT: spv.Branch ^bb2(%[[CST1]], %[[ADD]] : i32, i32)
spv.Branch ^bb2(%zero, %0 : i32, i32)
@@ -83,8 +83,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.Branch ^bb1
// CHECK-NEXT: ^bb1:
-// CHECK-NEXT: %[[ZERO:.*]] = spv.constant 0
-// CHECK-NEXT: %[[ONE:.*]] = spv.constant 1
+// CHECK-NEXT: %[[ZERO:.*]] = spv.Constant 0
+// CHECK-NEXT: %[[ONE:.*]] = spv.Constant 1
// CHECK-NEXT: spv.Branch ^bb2(%[[ZERO]], %[[ONE]] : i32, f32)
// CHECK-NEXT: ^bb2(%{{.*}}: i32, %{{.*}}: f32):
@@ -94,8 +94,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// This block is reordered to follow domination order.
^bb1:
- %zero = spv.constant 0 : i32
- %one = spv.constant 1.0 : f32
+ %zero = spv.Constant 0 : i32
+ %one = spv.Constant 1.0 : f32
spv.Branch ^bb2(%zero, %one : i32, f32)
}
@@ -115,21 +115,21 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK: spv.selection
spv.selection {
- %true = spv.constant true
+ %true = spv.Constant true
// CHECK: spv.BranchConditional %{{.*}}, ^bb1, ^bb2
spv.BranchConditional %true, ^true, ^false
// CHECK-NEXT: ^bb1:
^true:
-// CHECK-NEXT: %[[ZERO:.*]] = spv.constant 0
- %zero = spv.constant 0 : i32
+// CHECK-NEXT: %[[ZERO:.*]] = spv.Constant 0
+ %zero = spv.Constant 0 : i32
// CHECK-NEXT: spv.Branch ^bb3(%[[ZERO]] : i32)
spv.Branch ^phi(%zero: i32)
// CHECK-NEXT: ^bb2:
^false:
-// CHECK-NEXT: %[[ONE:.*]] = spv.constant 1
- %one = spv.constant 1 : i32
+// CHECK-NEXT: %[[ONE:.*]] = spv.Constant 1
+ %one = spv.Constant 1 : i32
// CHECK-NEXT: spv.Branch ^bb3(%[[ONE]] : i32)
spv.Branch ^phi(%one: i32)
@@ -162,9 +162,9 @@ 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
+ %3 = spv.Constant 12 : i32
+ %4 = spv.Constant 32 : i32
+ %5 = spv.Constant 4 : i32
%6 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
%7 = spv.Load "Input" %6 : vector<3xi32>
%8 = spv.CompositeExtract %7[0 : i32] : vector<3xi32>
@@ -243,12 +243,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @fmul_kernel() "None" {
- %cst4 = spv.constant 4 : i32
+ %cst4 = spv.Constant 4 : i32
- %val1 = spv.constant 43 : i32
- %val2 = spv.constant 44 : i32
+ %val1 = spv.Constant 43 : i32
+ %val2 = spv.Constant 44 : i32
-// CHECK: spv.constant 43
+// CHECK: spv.Constant 43
// CHECK-NEXT: spv.Branch ^[[BB1:.+]](%{{.+}} : i32)
// CHECK-NEXT: ^[[BB1]](%{{.+}}: i32):
// CHECK-NEXT: spv.loop
@@ -264,7 +264,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.mlir.merge
}
-// CHECK: spv.constant 44
+// CHECK: spv.Constant 44
// CHECK-NEXT: spv.Branch ^[[BB2:.+]](%{{.+}} : i32)
// CHECK-NEXT: ^[[BB2]](%{{.+}}: i32):
// CHECK-NEXT: spv.loop
diff --git a/mlir/test/Target/SPIRV/selection.mlir b/mlir/test/Target/SPIRV/selection.mlir
index e05f4a135144..31dfcbcc8eee 100644
--- a/mlir/test/Target/SPIRV/selection.mlir
+++ b/mlir/test/Target/SPIRV/selection.mlir
@@ -6,13 +6,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @selection(%cond: i1) -> () "None" {
// CHECK: spv.Branch ^bb1
// CHECK-NEXT: ^bb1:
- %zero = spv.constant 0: i32
- %one = spv.constant 1: i32
- %two = spv.constant 2: i32
+ %zero = spv.Constant 0: i32
+ %one = spv.Constant 1: i32
+ %two = spv.Constant 2: i32
%var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
// CHECK-NEXT: spv.selection control(Flatten)
-// CHECK-NEXT: spv.constant 0
+// CHECK-NEXT: spv.Constant 0
// CHECK-NEXT: spv.Variable
spv.selection control(Flatten) {
// CHECK-NEXT: spv.BranchConditional %{{.*}} [5, 10], ^bb1, ^bb2
@@ -20,7 +20,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-NEXT: ^bb1:
^then:
-// CHECK-NEXT: spv.constant 1
+// CHECK-NEXT: spv.Constant 1
// CHECK-NEXT: spv.Store
spv.Store "Function" %var, %one : i32
// CHECK-NEXT: spv.Branch ^bb3
@@ -28,7 +28,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK-NEXT: ^bb2:
^else:
-// CHECK-NEXT: spv.constant 2
+// CHECK-NEXT: spv.Constant 2
// CHECK-NEXT: spv.Store
spv.Store "Function" %var, %two : i32
// CHECK-NEXT: spv.Branch ^bb3
@@ -67,7 +67,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
// CHECK: ^bb1:
^then:
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
spv.ReturnValue %zero : i32
// CHECK: ^bb2:
@@ -76,7 +76,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.mlir.merge
}
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : i32
}
diff --git a/mlir/test/Target/SPIRV/spec-constant.mlir b/mlir/test/Target/SPIRV/spec-constant.mlir
index 70b6e32b0a07..a69834bd1f2b 100644
--- a/mlir/test/Target/SPIRV/spec-constant.mlir
+++ b/mlir/test/Target/SPIRV/spec-constant.mlir
@@ -94,17 +94,17 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
spv.func @use_composite() -> (i32) "None" {
// CHECK: [[USE1:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32
- // CHECK: [[USE2:%.*]] = spv.constant 0 : i32
+ // CHECK: [[USE2:%.*]] = spv.Constant 0 : i32
// CHECK: [[RES1:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE1]], [[USE2]]) : (i32, i32) -> i32
// CHECK: [[USE3:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32
- // CHECK: [[USE4:%.*]] = spv.constant 0 : i32
+ // CHECK: [[USE4:%.*]] = spv.Constant 0 : i32
// CHECK: [[RES2:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE3]], [[USE4]]) : (i32, i32) -> i32
%0 = spv.mlir.referenceof @sc_i32_1 : i32
- %1 = spv.constant 0 : i32
+ %1 = spv.Constant 0 : i32
%2 = spv.SpecConstantOperation wraps "spv.ISub"(%0, %1) : (i32, i32) -> i32
// CHECK: [[RES3:%.*]] = spv.SpecConstantOperation wraps "spv.IMul"([[RES1]], [[RES2]]) : (i32, i32) -> i32
diff --git a/mlir/test/Target/SPIRV/undef.mlir b/mlir/test/Target/SPIRV/undef.mlir
index 6a287b2d2665..efea01001fc6 100644
--- a/mlir/test/Target/SPIRV/undef.mlir
+++ b/mlir/test/Target/SPIRV/undef.mlir
@@ -15,7 +15,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
%6 = spv.CompositeExtract %5[1 : i32, 2 : i32] : !spv.array<4x!spv.array<4xi32>>
// CHECK: {{%.*}} = spv.undef : !spv.ptr<!spv.struct<(f32)>, StorageBuffer>
%7 = spv.undef : !spv.ptr<!spv.struct<(f32)>, StorageBuffer>
- %8 = spv.constant 0 : i32
+ %8 = spv.Constant 0 : i32
%9 = spv.AccessChain %7[%8] : !spv.ptr<!spv.struct<(f32)>, StorageBuffer>, i32
spv.Return
}
More information about the Mlir-commits
mailing list