* Ops with `snake_case` names are those that have different 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
#### Unify and localize constants
* Various normal constant instructions are represented by the same
- `spv.constant` op. Those instructions are just for constants of different
+ `spv.Constant` op. Those instructions are just for constants of different
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
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
```
```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 {
```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 {
%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):
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
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.
```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
`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
// 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>
```
%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
```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
```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
```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
#### 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>
#### Example:
```mlir
- %0 = spv.constant ...
+ %0 = spv.Constant ...
%1 = spv.Variable : !spv.ptr<f32, Function>
%2 = spv.Variable init(%0): !spv.ptr<f32, Function>
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFAdd "Workgroup" "Reduce" %scalar : f32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFMax "Workgroup" "Reduce" %scalar : f32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFMin "Workgroup" "Reduce" %scalar : f32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : f32
%vector = ... : vector<4xf32>
%0 = spv.GroupNonUniformFMul "Workgroup" "Reduce" %scalar : f32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformIAdd "Workgroup" "Reduce" %scalar : i32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformIMul "Workgroup" "Reduce" %scalar : i32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformSMax "Workgroup" "Reduce" %scalar : i32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformSMin "Workgroup" "Reduce" %scalar : i32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformUMax "Workgroup" "Reduce" %scalar : i32
#### Example:
```mlir
- %four = spv.constant 4 : i32
+ %four = spv.Constant 4 : i32
%scalar = ... : i32
%vector = ... : vector<4xi32>
%0 = spv.GroupNonUniformUMin "Workgroup" "Reduce" %scalar : i32
// -----
-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 = [{
* ...
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.
<!-- 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
* `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.
#### 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
```
}
};
-/// Converts composite std.constant operation to spv.constant.
+/// Converts composite std.constant operation to spv.Constant.
class ConstantCompositeOpPattern final
: public OpConversionPattern<ConstantOp> {
public:
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;
}
//===----------------------------------------------------------------------===//
-// 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();
}
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;
}
// 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;
}
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
static ParseResult parseConstantOp(OpAsmParser &parser, OperationState &state) {
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);
}
// 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
}
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
}
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
}
%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]]{{\]}}
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
// 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]
} {
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)
// 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]]]
// 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>
// 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]]
// 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
}
}
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)
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:
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):
spv.loop {
spv.Branch ^header
^header:
- %cond = spv.constant true
+ %cond = spv.Constant true
spv.BranchConditional %cond, ^body, ^merge
^body:
// Do nothing
}
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
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
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
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
spv.mlir.merge
}
// CHECK: ^bb3:
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : i32
}
}
// 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>
// 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
}
// 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
}
// 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
}
// 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
// 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
// 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
// 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>
// 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>
}
// 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
// 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>
// 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>
// 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
// 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>
// 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: ] :
// 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: ] :
// 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>
// 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>
// 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]]
// 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]]
// 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>
// 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]]
// -----
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}}
// -----
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):
// -----
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:
//===----------------------------------------------------------------------===//
func @cond_branch() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
// CHECK: spv.BranchConditional %{{.*}}, ^bb1, ^bb2
spv.BranchConditional %true, ^one, ^two
// CHECK: ^bb1
// -----
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):
// -----
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:
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 different type than prior uses: 'i1' vs 'i32'}}
spv.BranchConditional %zero, ^one, ^two
^one:
// -----
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:
// -----
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) -> ()
// -----
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:
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
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
// 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 {
// -----
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 {
// -----
func @only_allowed_in_last_block() -> () {
- %true = spv.constant true
+ %true = spv.Constant true
spv.loop {
spv.Branch ^header
^header:
spv.mlir.merge
}
- %zero = spv.constant 0: i32
+ %zero = spv.Constant 0: i32
spv.ReturnValue %zero: i32
}
}
//===----------------------------------------------------------------------===//
func @ret_val() -> (i32) {
- %0 = spv.constant 42 : i32
+ %0 = spv.Constant 42 : i32
// CHECK: spv.ReturnValue %{{.*}} : i32
spv.ReturnValue %0 : 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
}
^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:
^merge:
spv.mlir.merge
}
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : 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
}) : () -> ()
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
}
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
}
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:
//===----------------------------------------------------------------------===//
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 {
// -----
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 {
// 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>
//===----------------------------------------------------------------------===//
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
}
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
// -----
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
// -----
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
// -----
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
// -----
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
// -----
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
//===----------------------------------------------------------------------===//
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
// -----
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
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
}
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
}
// -----
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
// -----
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
// -----
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
//===----------------------------------------------------------------------===//
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
// -----
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>
// -----
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
// 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>
// 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>
// 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>
// -----
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>
// 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>
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>
// -----
//===----------------------------------------------------------------------===//
-// 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
}
func @unaccepted_std_attr() -> () {
// expected-error @+1 {{cannot have value of type 'none'}}
- %0 = spv.constant unit : none
+ %0 = spv.Constant unit : none
return
}
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
}
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
}
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
}
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
}
// 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>
// }
// 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>
// }
// expected-error@+2 {{expects regions to end with 'spv.mlir.endmodule'}}
// expected-note@+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} : () -> ()
// -----
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
}
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
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
}
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
{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]]
// 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]]
//===----------------------------------------------------------------------===//
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
// -----
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
// -----
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
// 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>
// 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
// 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>
}
// -----
//===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
//===----------------------------------------------------------------------===//
// TODO: test constants in different 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
}
// -----
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>
}
// -----
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>>
}
// 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]]
// 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
// 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>
}
// 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]]
// 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
// 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
// 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>
}
// 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
// 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>
}
// 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]]
// 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]]
// 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]]
// 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]]
//===----------------------------------------------------------------------===//
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>
// -----
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>
// Store to a different 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>
// 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>
// 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>
// `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>
// 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>
// 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
// 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
// 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
// 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
// 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
// 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
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
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
}
// 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() : () -> ()
// 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
// 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:
// 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
// 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:
// 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]]
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>
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>
// 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
// 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
// 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
// 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
// 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
// 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
// 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
// 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
// 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>
// 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>
// 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>
// 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
// 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>
}
}
// 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>
}
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
}
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 {
}
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)
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
}
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
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
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
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
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
^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
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 {
// 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>
// 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
^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
^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
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
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
// 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
// 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>
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):
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)
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)
^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)
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):
// 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)
}
// 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)
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>
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
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
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
// 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
// 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
// CHECK: ^bb1:
^then:
- %zero = spv.constant 0 : i32
+ %zero = spv.Constant 0 : i32
spv.ReturnValue %zero : i32
// CHECK: ^bb2:
spv.mlir.merge
}
- %one = spv.constant 1 : i32
+ %one = spv.Constant 1 : i32
spv.ReturnValue %one : i32
}
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
%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
}
projects / platform / upstream / llvm.git / commitdiff