diff --git a/mlir/docs/Dialects/SPIR-V.md b/mlir/docs/Dialects/SPIR-V.md --- a/mlir/docs/Dialects/SPIR-V.md +++ b/mlir/docs/Dialects/SPIR-V.md @@ -88,7 +88,7 @@ * 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 @@ -166,7 +166,7 @@ #### 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 @@ -475,7 +475,7 @@ 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 %2 = "spv.IAdd"(%0, %0) : (i32, i32) -> i32 ``` @@ -581,9 +581,9 @@ ```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 spv.selection { @@ -669,8 +669,8 @@ ```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 spv.loop { @@ -732,15 +732,15 @@ %var = spv.Variable : !spv.ptr spv.selection { - %true = spv.constant true + %true = spv.Constant true spv.BranchConditional %true, ^true, ^false ^true: - %zero = spv.constant 0 : i32 + %zero = spv.Constant 0 : i32 spv.Branch ^phi(%zero: i32) ^false: - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 spv.Branch ^phi(%one: i32) ^phi(%arg: i32): @@ -964,7 +964,7 @@ instructions. * Types are serialized into `OpType*` instructions in the SPIR-V binary module section for types, constants, and global variables. -* `spv.constant`s are unified and placed in the SPIR-V binary module section +* `spv.Constant`s are unified and placed in the SPIR-V binary module section for types, constants, and global variables. * Attributes on ops, if not part of the op's binary encoding, are emitted as `OpDecorate*` instructions in the SPIR-V binary module section for @@ -980,7 +980,7 @@ capabilities, extended instruction sets, etc.) will be placed as attributes on `spv.module`. * `OpType*` instructions will be converted into proper `mlir::Type`s. -* `OpConstant*` instructions are materialized as `spv.constant` at each use +* `OpConstant*` instructions are materialized as `spv.Constant` at each use site. * `OpVariable` instructions will be converted to `spv.globalVariable` ops if in module-level; otherwise they will be converted into `spv.Variable` ops. diff --git a/mlir/docs/SPIRVToLLVMDialectConversion.md b/mlir/docs/SPIRVToLLVMDialectConversion.md --- a/mlir/docs/SPIRVToLLVMDialectConversion.md +++ b/mlir/docs/SPIRVToLLVMDialectConversion.md @@ -434,7 +434,7 @@ ```mlir // Access the 1st element of the array -%i = spv.constant 1: i32 +%i = spv.Constant 1: i32 %var = spv.Variable : !spv.ptr>, Function> %el = spv.AccessChain %var[%i, %i] : !spv.ptr>, Function>, i32, i32 @@ -529,7 +529,7 @@ `spv.Variable` is modelled as `llvm.alloca` op. If initialized, an additional store instruction is used. Note that there is no initialization for arrays and structs since constants of these types are not supported in LLVM dialect (TODO). -Also, at the moment initialization is only possible via `spv.constant`. +Also, at the moment initialization is only possible via `spv.Constant`. ```mlir // Conversion of VariableOp without initialization @@ -538,7 +538,7 @@ // 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 => %res = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr llvm.store %c, %res : !llvm.ptr ``` @@ -584,14 +584,14 @@ %res1 = spv.ShiftRightArithmetic %0, %1 : i32, i16 => %res1 = llvm.ashr %0, %ext: i32 ``` -### `spv.constant` +### `spv.Constant` -At the moment `spv.constant` conversion supports scalar and vector constants +At the moment `spv.Constant` conversion supports scalar and vector constants **only**. #### Mapping -`spv.constant` is mapped to `llvm.mlir.constant`. This is a straightforward +`spv.Constant` is mapped to `llvm.mlir.constant`. This is a straightforward conversion pattern with a special case when the argument is signed or unsigned. #### Special case @@ -608,10 +608,10 @@ ```mlir // %0 = llvm.mlir.constant(0 : i8) : i8 -%0 = spv.constant 0 : i8 +%0 = spv.Constant 0 : i8 // %1 = llvm.mlir.constant(dense<[2, 3, 4]> : vector<3xi32>) : vector<3xi32> -%1 = spv.constant dense<[2, 3, 4]> : vector<3xui32> +%1 = spv.Constant dense<[2, 3, 4]> : vector<3xui32> ``` ### Not implemented ops @@ -672,7 +672,7 @@ ```mlir // Conversion of selection -%cond = spv.constant true %cond = llvm.mlir.constant(true) : i1 +%cond = spv.Constant true %cond = llvm.mlir.constant(true) : i1 spv.selection { spv.BranchConditional %cond, ^true, ^false llvm.cond_br %cond, ^true, ^false @@ -693,7 +693,7 @@ ```mlir // Conversion of loop -%cond = spv.constant true %cond = llvm.mlir.constant(true) : i1 +%cond = spv.Constant true %cond = llvm.mlir.constant(true) : i1 spv.loop { spv.Branch ^header llvm.br ^header diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td --- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td @@ -58,7 +58,7 @@ #### Example: ```mlir - %0 = "spv.constant"() { value = 1: i32} : () -> i32 + %0 = "spv.Constant"() { value = 1: i32} : () -> i32 %1 = spv.Variable : !spv.ptr>, Function> %2 = spv.AccessChain %1[%0] : !spv.ptr>, Function> %3 = spv.Load "Function" %2 ["Volatile"] : !spv.array<4xf32> @@ -276,7 +276,7 @@ #### Example: ```mlir - %0 = spv.constant ... + %0 = spv.Constant ... %1 = spv.Variable : !spv.ptr %2 = spv.Variable init(%0): !spv.ptr diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td --- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td @@ -254,7 +254,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : f32 %vector = ... : vector<4xf32> %0 = spv.GroupNonUniformFAdd "Workgroup" "Reduce" %scalar : f32 @@ -314,7 +314,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : f32 %vector = ... : vector<4xf32> %0 = spv.GroupNonUniformFMax "Workgroup" "Reduce" %scalar : f32 @@ -374,7 +374,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : f32 %vector = ... : vector<4xf32> %0 = spv.GroupNonUniformFMin "Workgroup" "Reduce" %scalar : f32 @@ -431,7 +431,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : f32 %vector = ... : vector<4xf32> %0 = spv.GroupNonUniformFMul "Workgroup" "Reduce" %scalar : f32 @@ -486,7 +486,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : i32 %vector = ... : vector<4xi32> %0 = spv.GroupNonUniformIAdd "Workgroup" "Reduce" %scalar : i32 @@ -541,7 +541,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : i32 %vector = ... : vector<4xi32> %0 = spv.GroupNonUniformIMul "Workgroup" "Reduce" %scalar : i32 @@ -598,7 +598,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : i32 %vector = ... : vector<4xi32> %0 = spv.GroupNonUniformSMax "Workgroup" "Reduce" %scalar : i32 @@ -655,7 +655,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : i32 %vector = ... : vector<4xi32> %0 = spv.GroupNonUniformSMin "Workgroup" "Reduce" %scalar : i32 @@ -713,7 +713,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : i32 %vector = ... : vector<4xi32> %0 = spv.GroupNonUniformUMax "Workgroup" "Reduce" %scalar : i32 @@ -771,7 +771,7 @@ #### Example: ```mlir - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 %scalar = ... : i32 %vector = ... : vector<4xi32> %0 = spv.GroupNonUniformUMin "Workgroup" "Reduce" %scalar : i32 diff --git a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td --- a/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td +++ b/mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td @@ -67,7 +67,7 @@ // ----- -def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> { +def SPV_ConstantOp : SPV_Op<"Constant", [ConstantLike, NoSideEffect]> { let summary = "The op that declares a SPIR-V normal constant"; let description = [{ @@ -81,7 +81,7 @@ * ... Having such a plethora of constant instructions renders IR transformations - more tedious. Therefore, we use a single `spv.constant` op to represent + more tedious. Therefore, we use a single `spv.Constant` op to represent them all. Note that conversion between those SPIR-V constant instructions and this op is purely mechanical; so it can be scoped to the binary (de)serialization process. @@ -89,16 +89,16 @@ ``` - spv-constant-op ::= ssa-id `=` `spv.constant` attribute-value + spv.Constant-op ::= ssa-id `=` `spv.Constant` attribute-value (`:` spirv-type)? ``` #### Example: ```mlir - %0 = spv.constant true - %1 = spv.constant dense<[2, 3]> : vector<2xf32> - %2 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>> + %0 = spv.Constant true + %1 = spv.Constant dense<[2, 3]> : vector<2xf32> + %2 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>> ``` TODO: support constant structs @@ -572,7 +572,7 @@ * `OpSpecConstantTrue` and `OpSpecConstantFalse` for boolean constants * `OpSpecConstant` for scalar constants - Similar as `spv.constant`, this op represents all of the above cases. + Similar as `spv.Constant`, this op represents all of the above cases. `OpSpecConstantComposite` and `OpSpecConstantOp` are modelled with separate ops. @@ -731,8 +731,8 @@ #### Example: ```mlir - %0 = spv.constant 1: i32 - %1 = spv.constant 1: i32 + %0 = spv.Constant 1: i32 + %1 = spv.Constant 1: i32 %2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32 ``` diff --git a/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp b/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp --- a/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp +++ b/mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp @@ -356,7 +356,7 @@ } }; -/// Converts composite std.constant operation to spv.constant. +/// Converts composite std.constant operation to spv.Constant. class ConstantCompositeOpPattern final : public OpConversionPattern { public: @@ -367,7 +367,7 @@ 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 { public: using OpConversionPattern::OpConversionPattern; diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp --- a/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp @@ -138,11 +138,11 @@ } //===----------------------------------------------------------------------===// -// spv.constant +// spv.Constant //===----------------------------------------------------------------------===// OpFoldResult spirv::ConstantOp::fold(ArrayRef operands) { - assert(operands.empty() && "spv.constant has no operands"); + assert(operands.empty() && "spv.Constant has no operands"); return value(); } diff --git a/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp b/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp --- a/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp @@ -941,7 +941,7 @@ Operation *op = indexSSA.getDefiningOp(); if (!op) { emitError(baseLoc, "'spv.AccessChain' op index must be an " - "integer spv.constant to access " + "integer spv.Constant to access " "element of spv.struct"); return nullptr; } @@ -950,7 +950,7 @@ // integer literals of other bitwidths. if (failed(extractValueFromConstOp(op, index))) { emitError(baseLoc, - "'spv.AccessChain' index must be an integer spv.constant to " + "'spv.AccessChain' index must be an integer spv.Constant to " "access element of spv.struct, but provided ") << op->getName(); return nullptr; @@ -1483,7 +1483,7 @@ } //===----------------------------------------------------------------------===// -// spv.constant +// spv.Constant //===----------------------------------------------------------------------===// static ParseResult parseConstantOp(OpAsmParser &parser, OperationState &state) { diff --git a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp --- a/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp +++ b/mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp @@ -38,7 +38,7 @@ 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(unknownLoc, constInfo->second, constInfo->first); } diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -88,7 +88,7 @@ // Note that this ignores the workgroup size specification in gpu.launch. // We may want to define gpu.workgroup_size and convert it to the entry // point ABI we want here. - // CHECK: spv.constant 32 : i32 + // CHECK: spv.Constant 32 : i32 %0 = "gpu.block_dim"() {dimension = "x"} : () -> index gpu.return } @@ -110,7 +110,7 @@ gpu.func @builtin_workgroup_size_y() kernel attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { // The constant value is obtained from the spv.entry_point_abi. - // CHECK: spv.constant 4 : i32 + // CHECK: spv.Constant 4 : i32 %0 = "gpu.block_dim"() {dimension = "y"} : () -> index gpu.return } @@ -132,7 +132,7 @@ gpu.func @builtin_workgroup_size_z() kernel attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} { // The constant value is obtained from the spv.entry_point_abi. - // CHECK: spv.constant 1 : i32 + // CHECK: spv.Constant 1 : i32 %0 = "gpu.block_dim"() {dimension = "z"} : () -> index gpu.return } diff --git a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir --- a/mlir/test/Conversion/GPUToSPIRV/load-store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load-store.mlir @@ -59,12 +59,12 @@ %12 = addi %arg3, %0 : index // CHECK: %[[INDEX2:.*]] = spv.IAdd %[[ARG4]], %[[LOCALINVOCATIONIDX]] %13 = addi %arg4, %3 : index - // CHECK: %[[ZERO:.*]] = spv.constant 0 : i32 - // CHECK: %[[OFFSET1_0:.*]] = spv.constant 0 : i32 - // CHECK: %[[STRIDE1_1:.*]] = spv.constant 4 : i32 + // CHECK: %[[ZERO:.*]] = spv.Constant 0 : i32 + // CHECK: %[[OFFSET1_0:.*]] = spv.Constant 0 : i32 + // CHECK: %[[STRIDE1_1:.*]] = spv.Constant 4 : i32 // CHECK: %[[UPDATE1_1:.*]] = spv.IMul %[[STRIDE1_1]], %[[INDEX1]] : i32 // CHECK: %[[OFFSET1_1:.*]] = spv.IAdd %[[OFFSET1_0]], %[[UPDATE1_1]] : i32 - // CHECK: %[[STRIDE1_2:.*]] = spv.constant 1 : i32 + // CHECK: %[[STRIDE1_2:.*]] = spv.Constant 1 : i32 // CHECK: %[[UPDATE1_2:.*]] = spv.IMul %[[STRIDE1_2]], %[[INDEX2]] : i32 // CHECK: %[[OFFSET1_2:.*]] = spv.IAdd %[[OFFSET1_1]], %[[UPDATE1_2]] : i32 // CHECK: %[[PTR1:.*]] = spv.AccessChain %[[ARG0]]{{\[}}%[[ZERO]], %[[OFFSET1_2]]{{\]}} diff --git a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir --- a/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir +++ b/mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir @@ -9,7 +9,7 @@ spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr [0])>, StorageBuffer> spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} { %0 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr [0])>, StorageBuffer> - %2 = spv.constant 0 : i32 + %2 = spv.Constant 0 : i32 %3 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr [0])>, StorageBuffer> %4 = spv.AccessChain %0[%2, %2] : !spv.ptr [0])>, StorageBuffer>, i32, i32 %5 = spv.Load "StorageBuffer" %4 : f32 diff --git a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir --- a/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir +++ b/mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir @@ -23,7 +23,7 @@ // CHECK: @single_workgroup_reduction // CHECK-SAME: (%[[INPUT:.+]]: !spv.ptr{{.+}}, %[[OUTPUT:.+]]: !spv.ptr{{.+}}) -// CHECK: %[[ZERO:.+]] = spv.constant 0 : i32 +// CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32 // CHECK: %[[ID:.+]] = spv.Load "Input" %{{.+}} : vector<3xi32> // CHECK: %[[X:.+]] = spv.CompositeExtract %[[ID]][0 : i32] diff --git a/mlir/test/Conversion/SCFToSPIRV/for.mlir b/mlir/test/Conversion/SCFToSPIRV/for.mlir --- a/mlir/test/Conversion/SCFToSPIRV/for.mlir +++ b/mlir/test/Conversion/SCFToSPIRV/for.mlir @@ -6,11 +6,11 @@ } { func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) { - // CHECK: %[[LB:.*]] = spv.constant 4 : i32 + // CHECK: %[[LB:.*]] = spv.Constant 4 : i32 %lb = constant 4 : index - // CHECK: %[[UB:.*]] = spv.constant 42 : i32 + // CHECK: %[[UB:.*]] = spv.Constant 42 : i32 %ub = constant 42 : index - // CHECK: %[[STEP:.*]] = spv.constant 2 : i32 + // CHECK: %[[STEP:.*]] = spv.Constant 2 : i32 %step = constant 2 : index // CHECK: spv.loop { // CHECK-NEXT: spv.Branch ^[[HEADER:.*]](%[[LB]] : i32) @@ -18,15 +18,15 @@ // CHECK: %[[CMP:.*]] = spv.SLessThan %[[INDVAR]], %[[UB]] : i32 // CHECK: spv.BranchConditional %[[CMP]], ^[[BODY:.*]], ^[[MERGE:.*]] // CHECK: ^[[BODY]]: - // CHECK: %[[ZERO1:.*]] = spv.constant 0 : i32 - // CHECK: %[[OFFSET1:.*]] = spv.constant 0 : i32 - // CHECK: %[[STRIDE1:.*]] = spv.constant 1 : i32 + // CHECK: %[[ZERO1:.*]] = spv.Constant 0 : i32 + // CHECK: %[[OFFSET1:.*]] = spv.Constant 0 : i32 + // CHECK: %[[STRIDE1:.*]] = spv.Constant 1 : i32 // CHECK: %[[UPDATE1:.*]] = spv.IMul %[[STRIDE1]], %[[INDVAR]] : i32 // CHECK: %[[INDEX1:.*]] = spv.IAdd %[[OFFSET1]], %[[UPDATE1]] : i32 // CHECK: spv.AccessChain {{%.*}}{{\[}}%[[ZERO1]], %[[INDEX1]]{{\]}} - // CHECK: %[[ZERO2:.*]] = spv.constant 0 : i32 - // CHECK: %[[OFFSET2:.*]] = spv.constant 0 : i32 - // CHECK: %[[STRIDE2:.*]] = spv.constant 1 : i32 + // CHECK: %[[ZERO2:.*]] = spv.Constant 0 : i32 + // CHECK: %[[OFFSET2:.*]] = spv.Constant 0 : i32 + // CHECK: %[[STRIDE2:.*]] = spv.Constant 1 : i32 // CHECK: %[[UPDATE2:.*]] = spv.IMul %[[STRIDE2]], %[[INDVAR]] : i32 // CHECK: %[[INDEX2:.*]] = spv.IAdd %[[OFFSET2]], %[[UPDATE2]] : i32 // CHECK: spv.AccessChain {{%.*}}[%[[ZERO2]], %[[INDEX2]]] @@ -44,15 +44,15 @@ // 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 // CHECK: %[[VAR2:.*]] = spv.Variable : !spv.ptr diff --git a/mlir/test/Conversion/SCFToSPIRV/if.mlir b/mlir/test/Conversion/SCFToSPIRV/if.mlir --- a/mlir/test/Conversion/SCFToSPIRV/if.mlir +++ b/mlir/test/Conversion/SCFToSPIRV/if.mlir @@ -86,14 +86,14 @@ // CHECK: spv.selection { // CHECK-NEXT: spv.BranchConditional {{%.*}}, [[TRUE:\^.*]], [[FALSE:\^.*]] // CHECK-NEXT: [[TRUE]]: - // CHECK: %[[RET1TRUE:.*]] = spv.constant 0.000000e+00 : f32 - // CHECK: %[[RET2TRUE:.*]] = spv.constant 1.000000e+00 : f32 + // CHECK: %[[RET1TRUE:.*]] = spv.Constant 0.000000e+00 : f32 + // CHECK: %[[RET2TRUE:.*]] = spv.Constant 1.000000e+00 : f32 // CHECK-DAG: spv.Store "Function" %[[VAR1]], %[[RET1TRUE]] : f32 // CHECK-DAG: spv.Store "Function" %[[VAR2]], %[[RET2TRUE]] : f32 // CHECK: spv.Branch ^[[MERGE:.*]] // CHECK-NEXT: [[FALSE]]: - // CHECK: %[[RET2FALSE:.*]] = spv.constant 2.000000e+00 : f32 - // CHECK: %[[RET1FALSE:.*]] = spv.constant 3.000000e+00 : f32 + // CHECK: %[[RET2FALSE:.*]] = spv.Constant 2.000000e+00 : f32 + // CHECK: %[[RET1FALSE:.*]] = spv.Constant 3.000000e+00 : f32 // CHECK-DAG: spv.Store "Function" %[[VAR1]], %[[RET1FALSE]] : f32 // CHECK-DAG: spv.Store "Function" %[[VAR2]], %[[RET2FALSE]] : f32 // CHECK: spv.Branch ^[[MERGE]] diff --git a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir @@ -1,61 +1,61 @@ // RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s //===----------------------------------------------------------------------===// -// spv.constant +// spv.Constant //===----------------------------------------------------------------------===// // CHECK-LABEL: @bool_constant_scalar spv.func @bool_constant_scalar() "None" { // CHECK: llvm.mlir.constant(true) : i1 - %0 = spv.constant true + %0 = spv.Constant true // CHECK: llvm.mlir.constant(false) : i1 - %1 = spv.constant false + %1 = spv.Constant false spv.Return } // CHECK-LABEL: @bool_constant_vector spv.func @bool_constant_vector() "None" { // CHECK: llvm.mlir.constant(dense<[true, false]> : vector<2xi1>) : vector<2xi1> - %0 = spv.constant dense<[true, false]> : vector<2xi1> + %0 = spv.Constant dense<[true, false]> : vector<2xi1> // CHECK: llvm.mlir.constant(dense : vector<3xi1>) : vector<3xi1> - %1 = spv.constant dense : vector<3xi1> + %1 = spv.Constant dense : vector<3xi1> spv.Return } // CHECK-LABEL: @integer_constant_scalar spv.func @integer_constant_scalar() "None" { // CHECK: llvm.mlir.constant(0 : i8) : i8 - %0 = spv.constant 0 : i8 + %0 = spv.Constant 0 : i8 // CHECK: llvm.mlir.constant(-5 : i64) : i64 - %1 = spv.constant -5 : si64 + %1 = spv.Constant -5 : si64 // CHECK: llvm.mlir.constant(10 : i16) : i16 - %2 = spv.constant 10 : ui16 + %2 = spv.Constant 10 : ui16 spv.Return } // CHECK-LABEL: @integer_constant_vector spv.func @integer_constant_vector() "None" { // CHECK: llvm.mlir.constant(dense<[2, 3]> : vector<2xi32>) : vector<2xi32> - %0 = spv.constant dense<[2, 3]> : vector<2xi32> + %0 = spv.Constant dense<[2, 3]> : vector<2xi32> // CHECK: llvm.mlir.constant(dense<-4> : vector<2xi32>) : vector<2xi32> - %1 = spv.constant dense<-4> : vector<2xsi32> + %1 = spv.Constant dense<-4> : vector<2xsi32> // CHECK: llvm.mlir.constant(dense<[2, 3, 4]> : vector<3xi32>) : vector<3xi32> - %2 = spv.constant dense<[2, 3, 4]> : vector<3xui32> + %2 = spv.Constant dense<[2, 3, 4]> : vector<3xui32> spv.Return } // CHECK-LABEL: @float_constant_scalar spv.func @float_constant_scalar() "None" { // CHECK: llvm.mlir.constant(5.000000e+00 : f16) : f16 - %0 = spv.constant 5.000000e+00 : f16 + %0 = spv.Constant 5.000000e+00 : f16 // CHECK: llvm.mlir.constant(5.000000e+00 : f64) : f64 - %1 = spv.constant 5.000000e+00 : f64 + %1 = spv.Constant 5.000000e+00 : f64 spv.Return } // CHECK-LABEL: @float_constant_vector spv.func @float_constant_vector() "None" { // CHECK: llvm.mlir.constant(dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>) : vector<2xf32> - %0 = spv.constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32> + %0 = spv.Constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32> spv.Return } diff --git a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir @@ -14,8 +14,8 @@ } spv.func @branch_with_arguments() -> () "None" { - %0 = spv.constant 0 : i32 - %1 = spv.constant true + %0 = spv.Constant 0 : i32 + %1 = spv.Constant true // CHECK: llvm.br ^bb1(%{{.*}}, %{{.*}} : i32, i1) spv.Branch ^label(%0, %1: i32, i1) // CHECK: ^bb1(%{{.*}}: i32, %{{.*}}: i1) @@ -33,7 +33,7 @@ spv.module Logical GLSL450 { spv.func @cond_branch_without_arguments() -> () "None" { // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1 - %cond = spv.constant true + %cond = spv.Constant true // CHECK: lvm.cond_br %[[COND]], ^bb1, ^bb2 spv.BranchConditional %cond, ^true, ^false // CHECK: ^bb1: @@ -46,10 +46,10 @@ spv.func @cond_branch_with_arguments_nested() -> () "None" { // CHECK: %[[COND1:.*]] = llvm.mlir.constant(true) : i1 - %cond = spv.constant true - %0 = spv.constant 0 : i32 + %cond = spv.Constant true + %0 = spv.Constant 0 : i32 // CHECK: %[[COND2:.*]] = llvm.mlir.constant(false) : i1 - %false = spv.constant false + %false = spv.Constant false // CHECK: llvm.cond_br %[[COND1]], ^bb1(%{{.*}}, %[[COND2]] : i32, i1), ^bb2 spv.BranchConditional %cond, ^outer_true(%0, %false: i32, i1), ^outer_false // CHECK: ^bb1(%{{.*}}: i32, %[[COND:.*]]: i1): @@ -103,7 +103,7 @@ spv.loop { spv.Branch ^header ^header: - %cond = spv.constant true + %cond = spv.Constant true spv.BranchConditional %cond, ^body, ^merge ^body: // Do nothing @@ -133,7 +133,7 @@ } spv.func @selection_with_merge_block_only() -> () "None" { - %cond = spv.constant true + %cond = spv.Constant true // CHECK: llvm.return spv.selection { spv.BranchConditional %cond, ^merge, ^merge @@ -145,7 +145,7 @@ spv.func @selection_with_true_block_only() -> () "None" { // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1 - %cond = spv.constant true + %cond = spv.Constant true // CHECK: llvm.cond_br %[[COND]], ^bb1, ^bb2 spv.selection { spv.BranchConditional %cond, ^true, ^merge @@ -165,7 +165,7 @@ spv.func @selection_with_both_true_and_false_block() -> () "None" { // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1 - %cond = spv.constant true + %cond = spv.Constant true // CHECK: llvm.cond_br %[[COND]], ^bb1, ^bb2 spv.selection { spv.BranchConditional %cond, ^true, ^false @@ -189,7 +189,7 @@ spv.func @selection_with_early_return(%arg0: i1) -> i32 "None" { // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32 - %0 = spv.constant 0 : i32 + %0 = spv.Constant 0 : i32 // CHECK: llvm.cond_br %{{.*}}, ^bb1(%[[ZERO]] : i32), ^bb2 spv.selection { spv.BranchConditional %arg0, ^true(%0 : i32), ^merge @@ -203,7 +203,7 @@ spv.mlir.merge } // CHECK: ^bb3: - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 spv.ReturnValue %one : i32 } } diff --git a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir --- a/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir +++ b/mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir @@ -7,7 +7,7 @@ // CHECK-LABEL: @access_chain spv.func @access_chain() "None" { // CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i32) : i32 - %0 = spv.constant 1: i32 + %0 = spv.Constant 1: i32 %1 = spv.Variable : !spv.ptr)>, Function> // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32 // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %[[ONE]], %[[ONE]]] : (!llvm.ptr)>>, i32, i32, i32) -> !llvm.ptr @@ -176,7 +176,7 @@ // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32 // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr - %c = spv.constant 0 : i64 + %c = spv.Constant 0 : i64 %0 = spv.Variable init(%c) : !spv.ptr spv.Return } @@ -195,7 +195,7 @@ // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32 // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x vector<3xi1> : (i32) -> !llvm.ptr> // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr> - %c = spv.constant dense : vector<3xi1> + %c = spv.Constant dense : vector<3xi1> %0 = spv.Variable init(%c) : !spv.ptr, Function> spv.Return } diff --git a/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir b/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir --- a/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir +++ b/mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir @@ -410,53 +410,53 @@ // CHECK-LABEL: @constant func @constant() { - // CHECK: spv.constant true + // CHECK: spv.Constant true %0 = constant true - // CHECK: spv.constant 42 : i32 + // CHECK: spv.Constant 42 : i32 %1 = constant 42 : i32 - // CHECK: spv.constant 5.000000e-01 : f32 + // CHECK: spv.Constant 5.000000e-01 : f32 %2 = constant 0.5 : f32 - // CHECK: spv.constant dense<[2, 3]> : vector<2xi32> + // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32> %3 = constant dense<[2, 3]> : vector<2xi32> - // CHECK: spv.constant 1 : i32 + // CHECK: spv.Constant 1 : i32 %4 = constant 1 : index - // CHECK: spv.constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4> + // CHECK: spv.Constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4> %5 = constant dense<1> : tensor<2x3xi32> - // CHECK: spv.constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4> + // CHECK: spv.Constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4> %6 = constant dense<1.0> : tensor<2x3xf32> - // CHECK: spv.constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4> + // CHECK: spv.Constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4> %7 = constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> - // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4> + // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4> %8 = constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> - // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4> + // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4> %9 = constant dense<[[1, 2], [3, 4], [5, 6]]> : tensor<3x2xi32> - // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4> + // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4> %10 = constant dense<[1, 2, 3, 4, 5, 6]> : tensor<6xi32> return } // CHECK-LABEL: @constant_16bit func @constant_16bit() { - // CHECK: spv.constant 4 : i16 + // CHECK: spv.Constant 4 : i16 %0 = constant 4 : i16 - // CHECK: spv.constant 5.000000e+00 : f16 + // CHECK: spv.Constant 5.000000e+00 : f16 %1 = constant 5.0 : f16 - // CHECK: spv.constant dense<[2, 3]> : vector<2xi16> + // CHECK: spv.Constant dense<[2, 3]> : vector<2xi16> %2 = constant dense<[2, 3]> : vector<2xi16> - // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2> + // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2> %3 = constant dense<4.0> : tensor<5xf16> return } // CHECK-LABEL: @constant_64bit func @constant_64bit() { - // CHECK: spv.constant 4 : i64 + // CHECK: spv.Constant 4 : i64 %0 = constant 4 : i64 - // CHECK: spv.constant 5.000000e+00 : f64 + // CHECK: spv.Constant 5.000000e+00 : f64 %1 = constant 5.0 : f64 - // CHECK: spv.constant dense<[2, 3]> : vector<2xi64> + // CHECK: spv.Constant dense<[2, 3]> : vector<2xi64> %2 = constant dense<[2, 3]> : vector<2xi64> - // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8> + // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8> %3 = constant dense<4.0> : tensor<5xf64> return } @@ -472,58 +472,58 @@ // CHECK-LABEL: @constant_16bit func @constant_16bit() { - // CHECK: spv.constant 4 : i32 + // CHECK: spv.Constant 4 : i32 %0 = constant 4 : i16 - // CHECK: spv.constant 5.000000e+00 : f32 + // CHECK: spv.Constant 5.000000e+00 : f32 %1 = constant 5.0 : f16 - // CHECK: spv.constant dense<[2, 3]> : vector<2xi32> + // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32> %2 = constant dense<[2, 3]> : vector<2xi16> - // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4> + // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4> %3 = constant dense<4.0> : tensor<5xf16> - // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4> + // CHECK: spv.Constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4> %4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16> return } // CHECK-LABEL: @constant_64bit func @constant_64bit() { - // CHECK: spv.constant 4 : i32 + // CHECK: spv.Constant 4 : i32 %0 = constant 4 : i64 - // CHECK: spv.constant 5.000000e+00 : f32 + // CHECK: spv.Constant 5.000000e+00 : f32 %1 = constant 5.0 : f64 - // CHECK: spv.constant dense<[2, 3]> : vector<2xi32> + // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32> %2 = constant dense<[2, 3]> : vector<2xi64> - // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4> + // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4> %3 = constant dense<4.0> : tensor<5xf64> - // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4> + // CHECK: spv.Constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4> %4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16> return } // CHECK-LABEL: @corner_cases func @corner_cases() { - // CHECK: %{{.*}} = spv.constant -1 : i32 + // CHECK: %{{.*}} = spv.Constant -1 : i32 %0 = constant 4294967295 : i64 // 2^32 - 1 - // CHECK: %{{.*}} = spv.constant 2147483647 : i32 + // CHECK: %{{.*}} = spv.Constant 2147483647 : i32 %1 = constant 2147483647 : i64 // 2^31 - 1 - // CHECK: %{{.*}} = spv.constant -2147483648 : i32 + // CHECK: %{{.*}} = spv.Constant -2147483648 : i32 %2 = constant 2147483648 : i64 // 2^31 - // CHECK: %{{.*}} = spv.constant -2147483648 : i32 + // CHECK: %{{.*}} = spv.Constant -2147483648 : i32 %3 = constant -2147483648 : i64 // -2^31 - // CHECK: %{{.*}} = spv.constant -1 : i32 + // CHECK: %{{.*}} = spv.Constant -1 : i32 %5 = constant -1 : i64 - // CHECK: %{{.*}} = spv.constant -2 : i32 + // CHECK: %{{.*}} = spv.Constant -2 : i32 %6 = constant -2 : i64 - // CHECK: %{{.*}} = spv.constant -1 : i32 + // CHECK: %{{.*}} = spv.Constant -1 : i32 %7 = constant -1 : index - // CHECK: %{{.*}} = spv.constant -2 : i32 + // CHECK: %{{.*}} = spv.Constant -2 : i32 %8 = constant -2 : index - // CHECK: spv.constant false + // CHECK: spv.Constant false %9 = constant false - // CHECK: spv.constant true + // CHECK: spv.Constant true %10 = constant true return @@ -639,8 +639,8 @@ // CHECK-LABEL: @uitofp_i1_f32 func @uitofp_i1_f32(%arg0 : i1) -> f32 { - // CHECK: %[[ZERO:.+]] = spv.constant 0.000000e+00 : f32 - // CHECK: %[[ONE:.+]] = spv.constant 1.000000e+00 : f32 + // CHECK: %[[ZERO:.+]] = spv.Constant 0.000000e+00 : f32 + // CHECK: %[[ONE:.+]] = spv.Constant 1.000000e+00 : f32 // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, f32 %0 = std.uitofp %arg0 : i1 to f32 return %0 : f32 @@ -648,8 +648,8 @@ // CHECK-LABEL: @uitofp_i1_f64 func @uitofp_i1_f64(%arg0 : i1) -> f64 { - // CHECK: %[[ZERO:.+]] = spv.constant 0.000000e+00 : f64 - // CHECK: %[[ONE:.+]] = spv.constant 1.000000e+00 : f64 + // CHECK: %[[ZERO:.+]] = spv.Constant 0.000000e+00 : f64 + // CHECK: %[[ONE:.+]] = spv.Constant 1.000000e+00 : f64 // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, f64 %0 = std.uitofp %arg0 : i1 to f64 return %0 : f64 @@ -657,8 +657,8 @@ // CHECK-LABEL: @uitofp_vec_i1_f32 func @uitofp_vec_i1_f32(%arg0 : vector<4xi1>) -> vector<4xf32> { - // CHECK: %[[ZERO:.+]] = spv.constant dense<0.000000e+00> : vector<4xf32> - // CHECK: %[[ONE:.+]] = spv.constant dense<1.000000e+00> : vector<4xf32> + // CHECK: %[[ZERO:.+]] = spv.Constant dense<0.000000e+00> : vector<4xf32> + // CHECK: %[[ONE:.+]] = spv.Constant dense<1.000000e+00> : vector<4xf32> // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xf32> %0 = std.uitofp %arg0 : vector<4xi1> to vector<4xf32> return %0 : vector<4xf32> @@ -666,11 +666,11 @@ // CHECK-LABEL: @uitofp_vec_i1_f64 spv.func @uitofp_vec_i1_f64(%arg0: vector<4xi1>) -> vector<4xf64> "None" { - // CHECK: %[[ZERO:.+]] = spv.constant dense<0.000000e+00> : vector<4xf64> - // CHECK: %[[ONE:.+]] = spv.constant dense<1.000000e+00> : vector<4xf64> + // CHECK: %[[ZERO:.+]] = spv.Constant dense<0.000000e+00> : vector<4xf64> + // CHECK: %[[ONE:.+]] = spv.Constant dense<1.000000e+00> : vector<4xf64> // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xf64> - %0 = spv.constant dense<0.000000e+00> : vector<4xf64> - %1 = spv.constant dense<1.000000e+00> : vector<4xf64> + %0 = spv.Constant dense<0.000000e+00> : vector<4xf64> + %1 = spv.Constant dense<1.000000e+00> : vector<4xf64> %2 = spv.Select %arg0, %1, %0 : vector<4xi1>, vector<4xf64> spv.ReturnValue %2 : vector<4xf64> } @@ -705,8 +705,8 @@ // CHECK-LABEL: @zexti3 func @zexti3(%arg0 : i1) -> i32 { - // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32 - // CHECK: %[[ONE:.+]] = spv.constant 1 : i32 + // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32 + // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32 // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, i32 %0 = std.zexti %arg0 : i1 to i32 return %0 : i32 @@ -714,8 +714,8 @@ // CHECK-LABEL: @zexti4 func @zexti4(%arg0 : vector<4xi1>) -> vector<4xi32> { - // CHECK: %[[ZERO:.+]] = spv.constant dense<0> : vector<4xi32> - // CHECK: %[[ONE:.+]] = spv.constant dense<1> : vector<4xi32> + // CHECK: %[[ZERO:.+]] = spv.Constant dense<0> : vector<4xi32> + // CHECK: %[[ONE:.+]] = spv.Constant dense<1> : vector<4xi32> // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xi32> %0 = std.zexti %arg0 : vector<4xi1> to vector<4xi32> return %0 : vector<4xi32> @@ -723,8 +723,8 @@ // CHECK-LABEL: @zexti5 func @zexti5(%arg0 : vector<4xi1>) -> vector<4xi64> { - // CHECK: %[[ZERO:.+]] = spv.constant dense<0> : vector<4xi64> - // CHECK: %[[ONE:.+]] = spv.constant dense<1> : vector<4xi64> + // CHECK: %[[ZERO:.+]] = spv.Constant dense<0> : vector<4xi64> + // CHECK: %[[ONE:.+]] = spv.Constant dense<1> : vector<4xi64> // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xi64> %0 = std.zexti %arg0 : vector<4xi1> to vector<4xi64> return %0 : vector<4xi64> @@ -746,11 +746,11 @@ // CHECK-LABEL: @trunc_to_i1 func @trunc_to_i1(%arg0: i32) -> i1 { - // CHECK: %[[MASK:.*]] = spv.constant 1 : i32 + // CHECK: %[[MASK:.*]] = spv.Constant 1 : i32 // CHECK: %[[MASKED_SRC:.*]] = spv.BitwiseAnd %{{.*}}, %[[MASK]] : i32 // CHECK: %[[IS_ONE:.*]] = spv.IEqual %[[MASKED_SRC]], %[[MASK]] : i32 - // CHECK-DAG: %[[TRUE:.*]] = spv.constant true - // CHECK-DAG: %[[FALSE:.*]] = spv.constant false + // CHECK-DAG: %[[TRUE:.*]] = spv.Constant true + // CHECK-DAG: %[[FALSE:.*]] = spv.Constant false // CHECK: spv.Select %[[IS_ONE]], %[[TRUE]], %[[FALSE]] : i1, i1 %0 = std.trunci %arg0 : i32 to i1 return %0 : i1 @@ -758,11 +758,11 @@ // 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 : vector<4xi1> - // CHECK-DAG: %[[FALSE:.*]] = spv.constant dense : vector<4xi1> + // CHECK-DAG: %[[TRUE:.*]] = spv.Constant dense : vector<4xi1> + // CHECK-DAG: %[[FALSE:.*]] = spv.Constant dense : vector<4xi1> // CHECK: spv.Select %[[IS_ONE]], %[[TRUE]], %[[FALSE]] : vector<4xi1>, vector<4xi1> %0 = std.trunci %arg0 : vector<4xi32> to vector<4xi1> return %0 : vector<4xi1> @@ -871,13 +871,13 @@ // CHECK: [[ARG0:%.*]]: !spv.ptr [0])>, StorageBuffer>, // CHECK: [[ARG1:%.*]]: !spv.ptr [0])>, StorageBuffer>) func @load_store_zero_rank_float(%arg0: memref, %arg1: memref) { - // 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 - // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 + // CHECK: [[ZERO2:%.*]] = spv.Constant 0 : i32 // CHECK: spv.AccessChain [[ARG1]][ // CHECK-SAME: [[ZERO2]], [[ZERO2]] // CHECK-SAME: ] : @@ -890,13 +890,13 @@ // CHECK: [[ARG0:%.*]]: !spv.ptr [0])>, StorageBuffer>, // CHECK: [[ARG1:%.*]]: !spv.ptr [0])>, StorageBuffer>) func @load_store_zero_rank_int(%arg0: memref, %arg1: memref) { - // 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 - // CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32 + // CHECK: [[ZERO2:%.*]] = spv.Constant 0 : i32 // CHECK: spv.AccessChain [[ARG1]][ // CHECK-SAME: [[ZERO2]], [[ZERO2]] // CHECK-SAME: ] : @@ -919,19 +919,19 @@ // CHECK-LABEL: @load_i8 func @load_i8(%arg0: memref) { - // 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 @@ -941,23 +941,23 @@ // CHECK-LABEL: @load_i16 // CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32) func @load_i16(%arg0: memref<10xi16>, %index : index) { - // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32 - // CHECK: %[[OFFSET:.+]] = spv.constant 0 : i32 - // CHECK: %[[ONE:.+]] = spv.constant 1 : i32 + // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32 + // CHECK: %[[OFFSET:.+]] = spv.Constant 0 : i32 + // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32 // CHECK: %[[UPDATE:.+]] = spv.IMul %[[ONE]], %[[ARG1]] : i32 // CHECK: %[[FLAT_IDX:.+]] = spv.IAdd %[[OFFSET]], %[[UPDATE]] : i32 - // CHECK: %[[TWO1:.+]] = spv.constant 2 : i32 + // CHECK: %[[TWO1:.+]] = spv.Constant 2 : i32 // CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[FLAT_IDX]], %[[TWO1]] : i32 // CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]] // CHECK: %[[LOAD:.+]] = spv.Load "StorageBuffer" %[[PTR]] - // CHECK: %[[TWO2:.+]] = spv.constant 2 : i32 - // CHECK: %[[SIXTEEN:.+]] = spv.constant 16 : i32 + // CHECK: %[[TWO2:.+]] = spv.Constant 2 : i32 + // CHECK: %[[SIXTEEN:.+]] = spv.Constant 16 : i32 // CHECK: %[[IDX:.+]] = spv.UMod %[[FLAT_IDX]], %[[TWO2]] : i32 // CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[SIXTEEN]] : i32 // CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32 - // CHECK: %[[MASK:.+]] = spv.constant 65535 : i32 + // CHECK: %[[MASK:.+]] = spv.Constant 65535 : i32 // CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32 - // CHECK: %[[T2:.+]] = spv.constant 16 : i32 + // CHECK: %[[T2:.+]] = spv.Constant 16 : i32 // CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32 // CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32 %0 = load %arg0[%index] : memref<10xi16> @@ -985,17 +985,17 @@ // CHECK-LABEL: @store_i8 // CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32) func @store_i8(%arg0: memref, %value: i8) { - // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32 - // CHECK: %[[FOUR:.+]] = spv.constant 4 : i32 - // CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32 + // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32 + // CHECK: %[[FOUR:.+]] = spv.Constant 4 : i32 + // CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32 // CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR]] : i32 // CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32 - // CHECK: %[[MASK1:.+]] = spv.constant 255 : i32 + // CHECK: %[[MASK1:.+]] = spv.Constant 255 : i32 // CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32 // CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32 // CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG1]], %[[MASK1]] : i32 // CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32 - // CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32 + // CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32 // CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[ZERO]], %[[FOUR2]] : i32 // CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]] // CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]] @@ -1007,21 +1007,21 @@ // CHECK-LABEL: @store_i16 // CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32, %[[ARG2:.+]]: i32) func @store_i16(%arg0: memref<10xi16>, %index: index, %value: i16) { - // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32 - // CHECK: %[[OFFSET:.+]] = spv.constant 0 : i32 - // CHECK: %[[ONE:.+]] = spv.constant 1 : i32 + // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32 + // CHECK: %[[OFFSET:.+]] = spv.Constant 0 : i32 + // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32 // CHECK: %[[UPDATE:.+]] = spv.IMul %[[ONE]], %[[ARG1]] : i32 // CHECK: %[[FLAT_IDX:.+]] = spv.IAdd %[[OFFSET]], %[[UPDATE]] : i32 - // CHECK: %[[TWO:.+]] = spv.constant 2 : i32 - // CHECK: %[[SIXTEEN:.+]] = spv.constant 16 : i32 + // CHECK: %[[TWO:.+]] = spv.Constant 2 : i32 + // CHECK: %[[SIXTEEN:.+]] = spv.Constant 16 : i32 // CHECK: %[[IDX:.+]] = spv.UMod %[[FLAT_IDX]], %[[TWO]] : i32 // CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[SIXTEEN]] : i32 - // CHECK: %[[MASK1:.+]] = spv.constant 65535 : i32 + // CHECK: %[[MASK1:.+]] = spv.Constant 65535 : i32 // CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32 // CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32 // CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG2]], %[[MASK1]] : i32 // CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32 - // CHECK: %[[TWO2:.+]] = spv.constant 2 : i32 + // CHECK: %[[TWO2:.+]] = spv.Constant 2 : i32 // CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[FLAT_IDX]], %[[TWO2]] : i32 // CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]] // CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]] @@ -1062,19 +1062,19 @@ // CHECK-LABEL: @load_i8 func @load_i8(%arg0: memref) { - // 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 @@ -1093,17 +1093,17 @@ // CHECK-LABEL: @store_i8 // CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32) func @store_i8(%arg0: memref, %value: i8) { - // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32 - // CHECK: %[[FOUR:.+]] = spv.constant 4 : i32 - // CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32 + // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32 + // CHECK: %[[FOUR:.+]] = spv.Constant 4 : i32 + // CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32 // CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR]] : i32 // CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32 - // CHECK: %[[MASK1:.+]] = spv.constant 255 : i32 + // CHECK: %[[MASK1:.+]] = spv.Constant 255 : i32 // CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32 // CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32 // CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG1]], %[[MASK1]] : i32 // CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32 - // CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32 + // CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32 // CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[ZERO]], %[[FOUR2]] : i32 // CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]] // CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]] diff --git a/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir b/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/composite-ops.mlir @@ -113,7 +113,7 @@ // ----- func @composite_extract_invalid_index_type_1() -> () { - %0 = spv.constant 10 : i32 + %0 = spv.Constant 10 : i32 %1 = spv.Variable : !spv.ptr>, Function> %2 = spv.Load "Function" %1 ["Volatile"] : !spv.array<4x!spv.array<4xf32>> // expected-error @+1 {{expected non-function type}} diff --git a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir @@ -14,7 +14,7 @@ // ----- func @branch_argument() -> () { - %zero = spv.constant 0 : i32 + %zero = spv.Constant 0 : i32 // CHECK: spv.Branch ^bb1(%{{.*}}, %{{.*}} : i32, i32) spv.Branch ^next(%zero, %zero: i32, i32) ^next(%arg0: i32, %arg1: i32): @@ -31,7 +31,7 @@ // ----- func @wrong_accessor_count() -> () { - %true = spv.constant true + %true = spv.Constant true // expected-error @+1 {{requires 1 successor but found 2}} "spv.Branch"()[^one, ^two] : () -> () ^one: @@ -47,7 +47,7 @@ //===----------------------------------------------------------------------===// func @cond_branch() -> () { - %true = spv.constant true + %true = spv.Constant true // CHECK: spv.BranchConditional %{{.*}}, ^bb1, ^bb2 spv.BranchConditional %true, ^one, ^two // CHECK: ^bb1 @@ -61,8 +61,8 @@ // ----- func @cond_branch_argument() -> () { - %true = spv.constant true - %zero = spv.constant 0 : i32 + %true = spv.Constant true + %zero = spv.Constant 0 : i32 // CHECK: spv.BranchConditional %{{.*}}, ^bb1(%{{.*}}, %{{.*}} : i32, i32), ^bb2 spv.BranchConditional %true, ^true1(%zero, %zero: i32, i32), ^false1 ^true1(%arg0: i32, %arg1: i32): @@ -79,7 +79,7 @@ // ----- func @cond_branch_with_weights() -> () { - %true = spv.constant true + %true = spv.Constant true // CHECK: spv.BranchConditional %{{.*}} [5, 10] spv.BranchConditional %true [5, 10], ^one, ^two ^one: @@ -103,7 +103,7 @@ func @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: @@ -115,7 +115,7 @@ // ----- func @wrong_accessor_count() -> () { - %true = spv.constant true + %true = spv.Constant true // expected-error @+1 {{requires 2 successors but found 1}} "spv.BranchConditional"(%true)[^one] {operand_segment_sizes = dense<[1, 0, 0]>: vector<3xi32>} : (i1) -> () ^one: @@ -127,7 +127,7 @@ // ----- func @wrong_number_of_weights() -> () { - %true = spv.constant true + %true = spv.Constant true // expected-error @+1 {{must have exactly two branch weights}} "spv.BranchConditional"(%true)[^one, ^two] {branch_weights = [1 : i32, 2 : i32, 3 : i32], operand_segment_sizes = dense<[1, 0, 0]>: vector<3xi32>} : (i1) -> () @@ -140,7 +140,7 @@ // ----- func @weights_cannot_both_be_zero() -> () { - %true = spv.constant true + %true = spv.Constant true // expected-error @+1 {{branch weights cannot both be zero}} spv.BranchConditional %true [0, 0], ^one, ^two ^one: @@ -232,7 +232,7 @@ spv.module Logical GLSL450 { spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" { - %0 = spv.constant 2.0 : f32 + %0 = spv.Constant 2.0 : f32 // expected-error @+1 {{operand type mismatch: expected operand type 'i32', but provided 'f32' for operand number 1}} spv.FunctionCall @f_type_mismatch(%arg0, %0) : (i32, f32) -> () spv.Return @@ -243,7 +243,7 @@ spv.module Logical GLSL450 { spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 "None" { - %cst = spv.constant 0: i32 + %cst = spv.Constant 0: i32 // expected-error @+1 {{result type mismatch: expected 'i32', but provided 'f32'}} %0 = spv.FunctionCall @f_type_mismatch(%arg0, %arg0) : (i32, i32) -> f32 spv.ReturnValue %cst: i32 @@ -268,8 +268,8 @@ // 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 // CHECK: spv.loop { @@ -438,8 +438,8 @@ // ----- 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 spv.selection { @@ -460,7 +460,7 @@ // ----- func @only_allowed_in_last_block() -> () { - %true = spv.constant true + %true = spv.Constant true spv.loop { spv.Branch ^header ^header: @@ -548,7 +548,7 @@ spv.mlir.merge } - %zero = spv.constant 0: i32 + %zero = spv.Constant 0: i32 spv.ReturnValue %zero: i32 } } @@ -560,7 +560,7 @@ //===----------------------------------------------------------------------===// func @ret_val() -> (i32) { - %0 = spv.constant 42 : i32 + %0 = spv.Constant 42 : i32 // CHECK: spv.ReturnValue %{{.*}} : i32 spv.ReturnValue %0 : i32 } @@ -570,13 +570,13 @@ spv.selection { spv.BranchConditional %cond, ^then, ^merge ^then: - %zero = spv.constant 0 : i32 + %zero = spv.Constant 0 : i32 // CHECK: spv.ReturnValue spv.ReturnValue %zero : i32 ^merge: spv.mlir.merge } - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 spv.ReturnValue %one : i32 } @@ -587,7 +587,7 @@ ^header: spv.BranchConditional %cond, ^body, ^merge ^body: - %zero = spv.constant 0 : i32 + %zero = spv.Constant 0 : i32 // CHECK: spv.ReturnValue spv.ReturnValue %zero : i32 ^continue: @@ -595,7 +595,7 @@ ^merge: spv.mlir.merge } - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 spv.ReturnValue %one : i32 } @@ -608,7 +608,7 @@ // ----- "foo.function"() ({ - %0 = spv.constant true + %0 = spv.Constant true // expected-error @+1 {{op must appear in a function-like op's block}} spv.ReturnValue %0 : i1 }) : () -> () @@ -617,7 +617,7 @@ spv.module Logical GLSL450 { spv.func @value_count_mismatch() -> () "None" { - %0 = spv.constant 42 : i32 + %0 = spv.Constant 42 : i32 // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}} spv.ReturnValue %0 : i32 } @@ -627,7 +627,7 @@ spv.module Logical GLSL450 { spv.func @value_type_mismatch() -> (f32) "None" { - %0 = spv.constant 42 : i32 + %0 = spv.Constant 42 : i32 // expected-error @+1 {{return value's type ('i32') mismatch with function's result type ('f32')}} spv.ReturnValue %0 : i32 } @@ -640,7 +640,7 @@ spv.selection { spv.BranchConditional %cond, ^then, ^merge ^then: - %cst = spv.constant 0: i32 + %cst = spv.Constant 0: i32 // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}} spv.ReturnValue %cst: i32 ^merge: @@ -658,8 +658,8 @@ //===----------------------------------------------------------------------===// 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 // CHECK: spv.selection { @@ -685,9 +685,9 @@ // ----- 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 // CHECK: spv.selection { diff --git a/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir b/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir @@ -103,7 +103,7 @@ // CHECK-LABEL: @cooperative_matrix_access_chain spv.func @cooperative_matrix_access_chain(%a : !spv.ptr, Function>) -> !spv.ptr "None" { - %0 = spv.constant 0: i32 + %0 = spv.Constant 0: i32 // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr, Function>, i32 %1 = spv.AccessChain %a[%0] : !spv.ptr, Function>, i32 spv.ReturnValue %1 : !spv.ptr diff --git a/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir b/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/logical-ops.mlir @@ -178,24 +178,24 @@ //===----------------------------------------------------------------------===// func @select_op_bool(%arg0: i1) -> () { - %0 = spv.constant true - %1 = spv.constant false + %0 = spv.Constant true + %1 = spv.Constant false // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i1 %2 = spv.Select %arg0, %0, %1 : i1, i1 return } func @select_op_int(%arg0: i1) -> () { - %0 = spv.constant 2 : i32 - %1 = spv.constant 3 : i32 + %0 = spv.Constant 2 : i32 + %1 = spv.Constant 3 : i32 // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i32 %2 = spv.Select %arg0, %0, %1 : i1, i32 return } func @select_op_float(%arg0: i1) -> () { - %0 = spv.constant 2.0 : f32 - %1 = spv.constant 3.0 : f32 + %0 = spv.Constant 2.0 : f32 + %1 = spv.Constant 3.0 : f32 // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, f32 %2 = spv.Select %arg0, %0, %1 : i1, f32 return @@ -210,16 +210,16 @@ } func @select_op_vec(%arg0: i1) -> () { - %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> - %1 = spv.constant dense<[5.0, 6.0, 7.0]> : vector<3xf32> + %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> + %1 = spv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32> // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, vector<3xf32> %2 = spv.Select %arg0, %0, %1 : i1, vector<3xf32> return } func @select_op_vec_condn_vec(%arg0: vector<3xi1>) -> () { - %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> - %1 = spv.constant dense<[5.0, 6.0, 7.0]> : vector<3xf32> + %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> + %1 = spv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32> // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi1>, vector<3xf32> %2 = spv.Select %arg0, %0, %1 : vector<3xi1>, vector<3xf32> return @@ -228,8 +228,8 @@ // ----- func @select_op(%arg0: i1) -> () { - %0 = spv.constant 2 : i32 - %1 = spv.constant 3 : i32 + %0 = spv.Constant 2 : i32 + %1 = spv.Constant 3 : i32 // expected-error @+2 {{expected ','}} %2 = spv.Select %arg0, %0, %1 : i1 return @@ -238,8 +238,8 @@ // ----- func @select_op(%arg1: vector<3xi1>) -> () { - %0 = spv.constant 2 : i32 - %1 = spv.constant 3 : i32 + %0 = spv.Constant 2 : i32 + %1 = spv.Constant 3 : i32 // expected-error @+1 {{result expected to be of vector type when condition is of vector type}} %2 = spv.Select %arg1, %0, %1 : vector<3xi1>, i32 return @@ -248,8 +248,8 @@ // ----- func @select_op(%arg1: vector<4xi1>) -> () { - %0 = spv.constant dense<[2, 3, 4]> : vector<3xi32> - %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32> + %0 = spv.Constant dense<[2, 3, 4]> : vector<3xi32> + %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32> // expected-error @+1 {{result should have the same number of elements as the condition when condition is of vector type}} %2 = spv.Select %arg1, %0, %1 : vector<4xi1>, vector<3xi32> return @@ -258,8 +258,8 @@ // ----- func @select_op(%arg1: vector<4xi1>) -> () { - %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> - %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32> + %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> + %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32> // expected-error @+1 {{all of {true_value, false_value, result} have same type}} %2 = "spv.Select"(%arg1, %0, %1) : (vector<4xi1>, vector<3xf32>, vector<3xi32>) -> vector<3xi32> return @@ -268,8 +268,8 @@ // ----- func @select_op(%arg1: vector<4xi1>) -> () { - %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> - %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32> + %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32> + %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32> // expected-error @+1 {{all of {true_value, false_value, result} have same type}} %2 = "spv.Select"(%arg1, %1, %0) : (vector<4xi1>, vector<3xi32>, vector<3xf32>) -> vector<3xi32> return diff --git a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/memory-ops.mlir @@ -5,7 +5,7 @@ //===----------------------------------------------------------------------===// func @access_chain_struct() -> () { - %0 = spv.constant 1: i32 + %0 = spv.Constant 1: i32 %1 = spv.Variable : !spv.ptr)>, Function> // CHECK: spv.AccessChain {{.*}}[{{.*}}, {{.*}}] : !spv.ptr)>, Function> %2 = spv.AccessChain %1[%0, %0] : !spv.ptr)>, Function>, i32, i32 @@ -46,7 +46,7 @@ // ----- func @access_chain_non_composite() -> () { - %0 = spv.constant 1: i32 + %0 = spv.Constant 1: i32 %1 = spv.Variable : !spv.ptr // expected-error @+1 {{cannot extract from non-composite type 'f32' with index 0}} %2 = spv.AccessChain %1[%0] : !spv.ptr, i32 @@ -112,7 +112,7 @@ func @access_chain_invalid_index_2(%index0 : i32) -> () { %0 = spv.Variable : !spv.ptr)>, 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)>, Function>, i32, i32 return } @@ -122,7 +122,7 @@ func @access_chain_invalid_constant_type_1() -> () { %0 = std.constant 1: i32 %1 = spv.Variable : !spv.ptr)>, 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)>, Function>, i32, i32 return } @@ -130,7 +130,7 @@ // ----- 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)>, 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)>, Function>, i32, i32 @@ -487,7 +487,7 @@ // ----- func @variable_init_normal_constant() -> () { - %0 = spv.constant 4.0 : f32 + %0 = spv.Constant 4.0 : f32 // CHECK: spv.Variable init(%0) : !spv.ptr %1 = spv.Variable init(%0) : !spv.ptr return @@ -529,7 +529,7 @@ // ----- 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 return diff --git a/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir b/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir @@ -33,7 +33,7 @@ //===----------------------------------------------------------------------===// func @group_non_uniform_broadcast_scalar(%value: f32) -> f32 { - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 // CHECK: spv.GroupNonUniformBroadcast Workgroup %{{.*}}, %{{.*}} : f32, i32 %0 = spv.GroupNonUniformBroadcast Workgroup %value, %one : f32, i32 return %0: f32 @@ -42,7 +42,7 @@ // ----- func @group_non_uniform_broadcast_vector(%value: vector<4xf32>) -> vector<4xf32> { - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 // CHECK: spv.GroupNonUniformBroadcast Subgroup %{{.*}}, %{{.*}} : vector<4xf32>, i32 %0 = spv.GroupNonUniformBroadcast Subgroup %value, %one : vector<4xf32>, i32 return %0: vector<4xf32> @@ -51,7 +51,7 @@ // ----- func @group_non_uniform_broadcast_negative_scope(%value: f32, %localid: i32 ) -> f32 { - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 // expected-error @+1 {{execution scope must be 'Workgroup' or 'Subgroup'}} %0 = spv.GroupNonUniformBroadcast Device %value, %one : f32, i32 return %0: f32 @@ -101,7 +101,7 @@ // CHECK-LABEL: @group_non_uniform_fadd_clustered_reduce func @group_non_uniform_fadd_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> { - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 // CHECK: %{{.+}} = spv.GroupNonUniformFAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xf32> %0 = spv.GroupNonUniformFAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xf32> return %0: vector<2xf32> @@ -120,7 +120,7 @@ // CHECK-LABEL: @group_non_uniform_fmul_clustered_reduce func @group_non_uniform_fmul_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> { - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 // CHECK: %{{.+}} = spv.GroupNonUniformFMul "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xf32> %0 = spv.GroupNonUniformFMul "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xf32> return %0: vector<2xf32> @@ -167,7 +167,7 @@ // CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> { - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 // CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32> %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32> return %0: vector<2xi32> @@ -200,7 +200,7 @@ // ----- func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> { - %five = spv.constant 5 : i32 + %five = spv.Constant 5 : i32 // expected-error @+1 {{cluster size operand must be a power of two}} %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%five) : vector<2xi32> return %0: vector<2xi32> @@ -221,7 +221,7 @@ // CHECK-LABEL: @group_non_uniform_imul_clustered_reduce func @group_non_uniform_imul_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> { - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 // CHECK: %{{.+}} = spv.GroupNonUniformIMul "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32> %0 = spv.GroupNonUniformIMul "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32> return %0: vector<2xi32> diff --git a/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir b/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir --- a/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir +++ b/mlir/test/Dialect/SPIRV/IR/structure-ops.mlir @@ -7,7 +7,7 @@ spv.module Logical GLSL450 { spv.globalVariable @var1 : !spv.ptr)>, Input> spv.func @access_chain() -> () "None" { - %0 = spv.constant 1: i32 + %0 = spv.Constant 1: i32 // CHECK: [[VAR1:%.*]] = spv.mlir.addressof @var1 : !spv.ptr)>, Input> // CHECK-NEXT: spv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spv.ptr)>, Input> %1 = spv.mlir.addressof @var1 : !spv.ptr)>, Input> @@ -49,29 +49,29 @@ // ----- //===----------------------------------------------------------------------===// -// spv.constant +// spv.Constant //===----------------------------------------------------------------------===// func @const() -> () { - // CHECK: spv.constant true - // CHECK: spv.constant 42 : i32 - // CHECK: spv.constant 5.000000e-01 : f32 - // CHECK: spv.constant dense<[2, 3]> : vector<2xi32> - // CHECK: spv.constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>> - // CHECK: spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> - // CHECK: spv.constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> - // CHECK: spv.constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> - // CHECK: spv.constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> - - %0 = spv.constant true - %1 = spv.constant 42 : i32 - %2 = spv.constant 0.5 : f32 - %3 = spv.constant dense<[2, 3]> : vector<2xi32> - %4 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>> - %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> - %6 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> - %7 = spv.constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> - %8 = spv.constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> + // CHECK: spv.Constant true + // CHECK: spv.Constant 42 : i32 + // CHECK: spv.Constant 5.000000e-01 : f32 + // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32> + // CHECK: spv.Constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>> + // CHECK: spv.Constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> + // CHECK: spv.Constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> + // CHECK: spv.Constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> + // CHECK: spv.Constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> + + %0 = spv.Constant true + %1 = spv.Constant 42 : i32 + %2 = spv.Constant 0.5 : f32 + %3 = spv.Constant dense<[2, 3]> : vector<2xi32> + %4 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>> + %5 = spv.Constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> + %6 = spv.Constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> + %7 = spv.Constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>> + %8 = spv.Constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>> return } @@ -79,7 +79,7 @@ func @unaccepted_std_attr() -> () { // expected-error @+1 {{cannot have value of type 'none'}} - %0 = spv.constant unit : none + %0 = spv.Constant unit : none return } @@ -87,7 +87,7 @@ func @array_constant() -> () { // expected-error @+1 {{has array element whose type ('vector<2xi32>') does not match the result element type ('vector<2xf32>')}} - %0 = spv.constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spv.array<2xvector<2xf32>> + %0 = spv.Constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spv.array<2xvector<2xf32>> return } @@ -95,7 +95,7 @@ func @array_constant() -> () { // expected-error @+1 {{must have spv.array result type for array value}} - %0 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.rtarray> + %0 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.rtarray> return } @@ -103,7 +103,7 @@ func @non_nested_array_constant() -> () { // expected-error @+1 {{only support nested array result type}} - %0 = spv.constant dense<3.0> : tensor<2x2xf32> : !spv.array<2xvector<2xf32>> + %0 = spv.Constant dense<3.0> : tensor<2x2xf32> : !spv.array<2xvector<2xf32>> return } @@ -111,21 +111,21 @@ func @value_result_type_mismatch() -> () { // expected-error @+1 {{must have spv.array result type for array value}} - %0 = "spv.constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>) + %0 = "spv.Constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>) } // ----- func @value_result_type_mismatch() -> () { // expected-error @+1 {{result element type ('i32') does not match value element type ('f32')}} - %0 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>> + %0 = spv.Constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>> } // ----- func @value_result_num_elements_mismatch() -> () { // expected-error @+1 {{result number of elements (6) does not match value number of elements (4)}} - %0 = spv.constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>> + %0 = spv.Constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>> return } @@ -308,7 +308,7 @@ // 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 // spv.globalVariable @var1 init(%0) : !spv.ptr // } @@ -337,7 +337,7 @@ // 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 // spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr // } @@ -480,7 +480,7 @@ // 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} : () -> () // ----- @@ -561,7 +561,7 @@ spv.func @compute() -> f32 "None" { // CHECK: spv.mlir.referenceof @sc3 : f32 %0 = spv.mlir.referenceof @sc3 : f32 - %1 = spv.constant 6.0 : f32 + %1 = spv.Constant 6.0 : f32 %2 = spv.FAdd %0, %1 : f32 spv.ReturnValue %2 : f32 } @@ -801,10 +801,10 @@ spv.module Logical GLSL450 { spv.func @foo() -> i32 "None" { - // CHECK: [[LHS:%.*]] = spv.constant - %0 = spv.constant 1: i32 - // CHECK: [[RHS:%.*]] = spv.constant - %1 = spv.constant 1: i32 + // CHECK: [[LHS:%.*]] = spv.Constant + %0 = spv.Constant 1: i32 + // CHECK: [[RHS:%.*]] = spv.Constant + %1 = spv.Constant 1: i32 // CHECK: spv.SpecConstantOperation wraps "spv.IAdd"([[LHS]], [[RHS]]) : (i32, i32) -> i32 %2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32 @@ -831,7 +831,7 @@ spv.module Logical GLSL450 { spv.func @foo() -> i32 "None" { - %0 = spv.constant 1: i32 + %0 = spv.Constant 1: i32 // expected-error @+1 {{op expects parent op 'spv.SpecConstantOperation'}} spv.mlir.yield %0 : i32 } diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir @@ -18,7 +18,7 @@ attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]] // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]] - // CHECK: [[CONST0:%.*]] = spv.constant 0 : i32 + // CHECK: [[CONST0:%.*]] = spv.Constant 0 : i32 // CHECK: [[ARG0PTR:%.*]] = spv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]] // CHECK: [[ARG0:%.*]] = spv.Load "StorageBuffer" [[ARG0PTR]] // CHECK: spv.Return diff --git a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir @@ -40,19 +40,19 @@ {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}) "None" attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} { // CHECK: [[ADDRESSARG6:%.*]] = spv.mlir.addressof [[VAR6]] - // CHECK: [[CONST6:%.*]] = spv.constant 0 : i32 + // CHECK: [[CONST6:%.*]] = spv.Constant 0 : i32 // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]] // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]] // CHECK: [[ADDRESSARG5:%.*]] = spv.mlir.addressof [[VAR5]] - // CHECK: [[CONST5:%.*]] = spv.constant 0 : i32 + // CHECK: [[CONST5:%.*]] = spv.Constant 0 : i32 // CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]] // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG5PTR]] // CHECK: [[ADDRESSARG4:%.*]] = spv.mlir.addressof [[VAR4]] - // CHECK: [[CONST4:%.*]] = spv.constant 0 : i32 + // CHECK: [[CONST4:%.*]] = spv.Constant 0 : i32 // CHECK: [[ARG4PTR:%.*]] = spv.AccessChain [[ADDRESSARG4]]{{\[}}[[CONST4]] // CHECK: [[ARG4:%.*]] = spv.Load "StorageBuffer" [[ARG4PTR]] // CHECK: [[ADDRESSARG3:%.*]] = spv.mlir.addressof [[VAR3]] - // CHECK: [[CONST3:%.*]] = spv.constant 0 : i32 + // CHECK: [[CONST3:%.*]] = spv.Constant 0 : i32 // CHECK: [[ARG3PTR:%.*]] = spv.AccessChain [[ADDRESSARG3]]{{\[}}[[CONST3]] // CHECK: [[ARG3:%.*]] = spv.Load "StorageBuffer" [[ARG3PTR]] // CHECK: [[ADDRESSARG2:%.*]] = spv.mlir.addressof [[VAR2]] @@ -102,7 +102,7 @@ // 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>)>, StorageBuffer>, i32, i32, i32 %39 = spv.Load "StorageBuffer" %38 : f32 // CHECK: spv.AccessChain [[ARG1]] diff --git a/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir b/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir @@ -5,11 +5,11 @@ //===----------------------------------------------------------------------===// func @combine_full_access_chain() -> f32 { - // CHECK: %[[INDEX:.*]] = spv.constant 0 + // CHECK: %[[INDEX:.*]] = spv.Constant 0 // CHECK-NEXT: %[[VAR:.*]] = spv.Variable // CHECK-NEXT: %[[PTR:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]], %[[INDEX]]] // CHECK-NEXT: spv.Load "Function" %[[PTR]] - %c0 = spv.constant 0: i32 + %c0 = spv.Constant 0: i32 %0 = spv.Variable : !spv.ptr>, !spv.array<4xi32>)>, Function> %1 = spv.AccessChain %0[%c0] : !spv.ptr>, !spv.array<4xi32>)>, Function>, i32 %2 = spv.AccessChain %1[%c0, %c0] : !spv.ptr>, Function>, i32, i32 @@ -20,13 +20,13 @@ // ----- 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.array<4xi32>)>, Function> %1 = spv.AccessChain %0[%c0] : !spv.ptr>, !spv.array<4xi32>)>, Function>, i32 %2 = spv.AccessChain %1[%c0] : !spv.ptr>, Function>, i32 @@ -39,14 +39,14 @@ // ----- 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.array<4xi32>)>, Function> %1 = spv.Variable : !spv.ptr>, !spv.array<4xi32>)>, Function> %2 = spv.AccessChain %0[%c1] : !spv.ptr>, !spv.array<4xi32>)>, Function>, i32 @@ -92,10 +92,10 @@ // CHECK-LABEL: extract_vector func @extract_vector() -> (i32, i32, i32) { - // CHECK: spv.constant 42 : i32 - // CHECK: spv.constant -33 : i32 - // CHECK: spv.constant 6 : i32 - %0 = spv.constant dense<[42, -33, 6]> : vector<3xi32> + // CHECK: spv.Constant 42 : i32 + // CHECK: spv.Constant -33 : i32 + // CHECK: spv.Constant 6 : i32 + %0 = spv.Constant dense<[42, -33, 6]> : vector<3xi32> %1 = spv.CompositeExtract %0[0 : i32] : vector<3xi32> %2 = spv.CompositeExtract %0[1 : i32] : vector<3xi32> %3 = spv.CompositeExtract %0[2 : i32] : vector<3xi32> @@ -106,9 +106,9 @@ // CHECK-LABEL: extract_array_final func @extract_array_final() -> (i32, i32) { - // CHECK: spv.constant 4 : i32 - // CHECK: spv.constant -5 : i32 - %0 = spv.constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> + // CHECK: spv.Constant 4 : i32 + // CHECK: spv.Constant -5 : i32 + %0 = spv.Constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> %1 = spv.CompositeExtract %0[0 : i32, 0 : i32] : !spv.array<1 x vector<2 x i32>> %2 = spv.CompositeExtract %0[0 : i32, 1 : i32] : !spv.array<1 x vector<2 x i32>> return %1, %2 : i32, i32 @@ -118,8 +118,8 @@ // CHECK-LABEL: extract_array_interm func @extract_array_interm() -> (vector<2xi32>) { - // CHECK: spv.constant dense<[4, -5]> : vector<2xi32> - %0 = spv.constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> + // CHECK: spv.Constant dense<[4, -5]> : vector<2xi32> + %0 = spv.Constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<1 x vector<2 x i32>> return %1 : vector<2xi32> } @@ -138,15 +138,15 @@ // ----- //===----------------------------------------------------------------------===// -// 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 } @@ -154,9 +154,9 @@ // ----- func @deduplicate_vector_constant() -> (vector<3xi32>, vector<3xi32>) { - // CHECK: %[[CST:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32> - %0 = spv.constant dense<[1, 2, 3]> : vector<3xi32> - %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32> + // CHECK: %[[CST:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32> + %0 = spv.Constant dense<[1, 2, 3]> : vector<3xi32> + %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32> // CHECK-NEXT: return %[[CST]], %[[CST]] return %0, %1 : vector<3xi32>, vector<3xi32> } @@ -164,9 +164,9 @@ // ----- func @deduplicate_composite_constant() -> (!spv.array<1 x vector<2xi32>>, !spv.array<1 x vector<2xi32>>) { - // CHECK: %[[CST:.*]] = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> - %0 = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> - %1 = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> + // CHECK: %[[CST:.*]] = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> + %0 = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> + %1 = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>> // CHECK-NEXT: return %[[CST]], %[[CST]] return %0, %1 : !spv.array<1 x vector<2xi32>>, !spv.array<1 x vector<2xi32>> } @@ -180,7 +180,7 @@ // CHECK-LABEL: @iadd_zero // CHECK-SAME: (%[[ARG:.*]]: i32) func @iadd_zero(%arg0: i32) -> (i32, i32) { - %zero = spv.constant 0 : i32 + %zero = spv.Constant 0 : i32 %0 = spv.IAdd %arg0, %zero : i32 %1 = spv.IAdd %zero, %arg0 : i32 // CHECK: return %[[ARG]], %[[ARG]] @@ -189,12 +189,12 @@ // CHECK-LABEL: @const_fold_scalar_iadd_normal func @const_fold_scalar_iadd_normal() -> (i32, i32, i32) { - %c5 = spv.constant 5 : i32 - %cn8 = spv.constant -8 : i32 + %c5 = spv.Constant 5 : i32 + %cn8 = spv.Constant -8 : i32 - // CHECK: spv.constant 10 - // CHECK: spv.constant -16 - // CHECK: spv.constant -3 + // CHECK: spv.Constant 10 + // CHECK: spv.Constant -16 + // CHECK: spv.Constant -3 %0 = spv.IAdd %c5, %c5 : i32 %1 = spv.IAdd %cn8, %cn8 : i32 %2 = spv.IAdd %c5, %cn8 : i32 @@ -203,34 +203,34 @@ // CHECK-LABEL: @const_fold_scalar_iadd_flow func @const_fold_scalar_iadd_flow() -> (i32, i32, i32, i32) { - %c1 = spv.constant 1 : i32 - %c2 = spv.constant 2 : i32 - %c3 = spv.constant 4294967295 : i32 // 2^32 - 1: 0xffff ffff - %c4 = spv.constant -2147483648 : i32 // -2^31 : 0x8000 0000 - %c5 = spv.constant -1 : i32 // : 0xffff ffff - %c6 = spv.constant -2 : i32 // : 0xffff fffe + %c1 = spv.Constant 1 : i32 + %c2 = spv.Constant 2 : i32 + %c3 = spv.Constant 4294967295 : i32 // 2^32 - 1: 0xffff ffff + %c4 = spv.Constant -2147483648 : i32 // -2^31 : 0x8000 0000 + %c5 = spv.Constant -1 : i32 // : 0xffff ffff + %c6 = spv.Constant -2 : i32 // : 0xffff fffe // 0x0000 0001 + 0xffff ffff = 0x1 0000 0000 -> 0x0000 0000 - // CHECK: spv.constant 0 + // CHECK: spv.Constant 0 %0 = spv.IAdd %c1, %c3 : i32 // 0x0000 0002 + 0xffff ffff = 0x1 0000 0001 -> 0x0000 0001 - // CHECK: spv.constant 1 + // CHECK: spv.Constant 1 %1 = spv.IAdd %c2, %c3 : i32 // 0x8000 0000 + 0xffff ffff = 0x1 7fff ffff -> 0x7fff ffff - // CHECK: spv.constant 2147483647 + // CHECK: spv.Constant 2147483647 %2 = spv.IAdd %c4, %c5 : i32 // 0x8000 0000 + 0xffff fffe = 0x1 7fff fffe -> 0x7fff fffe - // CHECK: spv.constant 2147483646 + // CHECK: spv.Constant 2147483646 %3 = spv.IAdd %c4, %c6 : i32 return %0, %1, %2, %3: i32, i32, i32, i32 } // CHECK-LABEL: @const_fold_vector_iadd func @const_fold_vector_iadd() -> vector<3xi32> { - %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32> - %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32> + %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32> + %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32> - // CHECK: spv.constant dense<[39, -70, 155]> + // CHECK: spv.Constant dense<[39, -70, 155]> %0 = spv.IAdd %vc1, %vc2 : vector<3xi32> return %0: vector<3xi32> } @@ -244,9 +244,9 @@ // CHECK-LABEL: @imul_zero_one // CHECK-SAME: (%[[ARG:.*]]: i32) func @imul_zero_one(%arg0: i32) -> (i32, i32) { - // CHECK: %[[ZERO:.*]] = spv.constant 0 - %zero = spv.constant 0 : i32 - %one = spv.constant 1: i32 + // CHECK: %[[ZERO:.*]] = spv.Constant 0 + %zero = spv.Constant 0 : i32 + %one = spv.Constant 1: i32 %0 = spv.IMul %arg0, %zero : i32 %1 = spv.IMul %one, %arg0 : i32 // CHECK: return %[[ZERO]], %[[ARG]] @@ -255,13 +255,13 @@ // CHECK-LABEL: @const_fold_scalar_imul_normal func @const_fold_scalar_imul_normal() -> (i32, i32, i32) { - %c5 = spv.constant 5 : i32 - %cn8 = spv.constant -8 : i32 - %c7 = spv.constant 7 : i32 + %c5 = spv.Constant 5 : i32 + %cn8 = spv.Constant -8 : i32 + %c7 = spv.Constant 7 : i32 - // CHECK: spv.constant 35 - // CHECK: spv.constant -40 - // CHECK: spv.constant -56 + // CHECK: spv.Constant 35 + // CHECK: spv.Constant -40 + // CHECK: spv.Constant -56 %0 = spv.IMul %c7, %c5 : i32 %1 = spv.IMul %c5, %cn8 : i32 %2 = spv.IMul %cn8, %c7 : i32 @@ -270,18 +270,18 @@ // CHECK-LABEL: @const_fold_scalar_imul_flow func @const_fold_scalar_imul_flow() -> (i32, i32, i32) { - %c1 = spv.constant 2 : i32 - %c2 = spv.constant 4 : i32 - %c3 = spv.constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff - %c4 = spv.constant 2147483647 : i32 // 2^31 - 1 : 0x7fff ffff + %c1 = spv.Constant 2 : i32 + %c2 = spv.Constant 4 : i32 + %c3 = spv.Constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff + %c4 = spv.Constant 2147483647 : i32 // 2^31 - 1 : 0x7fff ffff // (0xffff ffff << 1) = 0x1 ffff fffe -> 0xffff fffe - // CHECK: %[[CST2:.*]] = spv.constant -2 + // CHECK: %[[CST2:.*]] = spv.Constant -2 %0 = spv.IMul %c1, %c3 : i32 // (0x7fff ffff << 1) = 0x0 ffff fffe -> 0xffff fffe %1 = spv.IMul %c1, %c4 : i32 // (0x7fff ffff << 2) = 0x1 ffff fffc -> 0xffff fffc - // CHECK: %[[CST4:.*]] = spv.constant -4 + // CHECK: %[[CST4:.*]] = spv.Constant -4 %2 = spv.IMul %c4, %c2 : i32 // CHECK: return %[[CST2]], %[[CST2]], %[[CST4]] return %0, %1, %2: i32, i32, i32 @@ -290,10 +290,10 @@ // CHECK-LABEL: @const_fold_vector_imul func @const_fold_vector_imul() -> vector<3xi32> { - %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32> - %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32> + %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32> + %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32> - // CHECK: spv.constant dense<[-126, 825, 3556]> + // CHECK: spv.Constant dense<[-126, 825, 3556]> %0 = spv.IMul %vc1, %vc2 : vector<3xi32> return %0: vector<3xi32> } @@ -306,20 +306,20 @@ // CHECK-LABEL: @isub_x_x func @isub_x_x(%arg0: i32) -> i32 { - // CHECK: spv.constant 0 + // CHECK: spv.Constant 0 %0 = spv.ISub %arg0, %arg0: i32 return %0: i32 } // CHECK-LABEL: @const_fold_scalar_isub_normal func @const_fold_scalar_isub_normal() -> (i32, i32, i32) { - %c5 = spv.constant 5 : i32 - %cn8 = spv.constant -8 : i32 - %c7 = spv.constant 7 : i32 + %c5 = spv.Constant 5 : i32 + %cn8 = spv.Constant -8 : i32 + %c7 = spv.Constant 7 : i32 - // CHECK: spv.constant 2 - // CHECK: spv.constant 13 - // CHECK: spv.constant -15 + // CHECK: spv.Constant 2 + // CHECK: spv.Constant 13 + // CHECK: spv.Constant -15 %0 = spv.ISub %c7, %c5 : i32 %1 = spv.ISub %c5, %cn8 : i32 %2 = spv.ISub %cn8, %c7 : i32 @@ -328,34 +328,34 @@ // CHECK-LABEL: @const_fold_scalar_isub_flow func @const_fold_scalar_isub_flow() -> (i32, i32, i32, i32) { - %c1 = spv.constant 0 : i32 - %c2 = spv.constant 1 : i32 - %c3 = spv.constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff - %c4 = spv.constant 2147483647 : i32 // 2^31 : 0x7fff ffff - %c5 = spv.constant -1 : i32 // : 0xffff ffff - %c6 = spv.constant -2 : i32 // : 0xffff fffe + %c1 = spv.Constant 0 : i32 + %c2 = spv.Constant 1 : i32 + %c3 = spv.Constant 4294967295 : i32 // 2^32 - 1 : 0xffff ffff + %c4 = spv.Constant 2147483647 : i32 // 2^31 : 0x7fff ffff + %c5 = spv.Constant -1 : i32 // : 0xffff ffff + %c6 = spv.Constant -2 : i32 // : 0xffff fffe // 0x0000 0000 - 0xffff ffff -> 0x0000 0000 + 0x0000 0001 = 0x0000 0001 - // CHECK: spv.constant 1 + // CHECK: spv.Constant 1 %0 = spv.ISub %c1, %c3 : i32 // 0x0000 0001 - 0xffff ffff -> 0x0000 0001 + 0x0000 0001 = 0x0000 0002 - // CHECK: spv.constant 2 + // CHECK: spv.Constant 2 %1 = spv.ISub %c2, %c3 : i32 // 0xffff ffff - 0x7fff ffff -> 0xffff ffff + 0x8000 0001 = 0x1 8000 0000 - // CHECK: spv.constant -2147483648 + // CHECK: spv.Constant -2147483648 %2 = spv.ISub %c5, %c4 : i32 // 0xffff fffe - 0x7fff ffff -> 0xffff fffe + 0x8000 0001 = 0x1 7fff ffff - // CHECK: spv.constant 2147483647 + // CHECK: spv.Constant 2147483647 %3 = spv.ISub %c6, %c4 : i32 return %0, %1, %2, %3: i32, i32, i32, i32 } // CHECK-LABEL: @const_fold_vector_isub func @const_fold_vector_isub() -> vector<3xi32> { - %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32> - %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32> + %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32> + %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32> - // CHECK: spv.constant dense<[45, -40, 99]> + // CHECK: spv.Constant dense<[45, -40, 99]> %0 = spv.ISub %vc1, %vc2 : vector<3xi32> return %0: vector<3xi32> } @@ -369,9 +369,9 @@ // CHECK-LABEL: @convert_logical_and_true_false_scalar // CHECK-SAME: %[[ARG:.+]]: i1 func @convert_logical_and_true_false_scalar(%arg: i1) -> (i1, i1) { - %true = spv.constant true - // CHECK: %[[FALSE:.+]] = spv.constant false - %false = spv.constant false + %true = spv.Constant true + // CHECK: %[[FALSE:.+]] = spv.Constant false + %false = spv.Constant false %0 = spv.LogicalAnd %true, %arg: i1 %1 = spv.LogicalAnd %arg, %false: i1 // CHECK: return %[[ARG]], %[[FALSE]] @@ -381,9 +381,9 @@ // 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 : vector<3xi1> - // CHECK: %[[FALSE:.+]] = spv.constant dense - %false = spv.constant dense : vector<3xi1> + %true = spv.Constant dense : vector<3xi1> + // CHECK: %[[FALSE:.+]] = spv.Constant dense + %false = spv.Constant dense : vector<3xi1> %0 = spv.LogicalAnd %true, %arg: vector<3xi1> %1 = spv.LogicalAnd %arg, %false: vector<3xi1> // CHECK: return %[[ARG]], %[[FALSE]] @@ -456,9 +456,9 @@ // CHECK-LABEL: @convert_logical_or_true_false_scalar // CHECK-SAME: %[[ARG:.+]]: i1 func @convert_logical_or_true_false_scalar(%arg: i1) -> (i1, i1) { - // CHECK: %[[TRUE:.+]] = spv.constant true - %true = spv.constant true - %false = spv.constant false + // CHECK: %[[TRUE:.+]] = spv.Constant true + %true = spv.Constant true + %false = spv.Constant false %0 = spv.LogicalOr %true, %arg: i1 %1 = spv.LogicalOr %arg, %false: i1 // CHECK: return %[[TRUE]], %[[ARG]] @@ -468,9 +468,9 @@ // 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 = spv.constant dense : vector<3xi1> - %false = spv.constant dense : vector<3xi1> + // CHECK: %[[TRUE:.+]] = spv.Constant dense + %true = spv.Constant dense : vector<3xi1> + %false = spv.Constant dense : vector<3xi1> %0 = spv.LogicalOr %true, %arg: vector<3xi1> %1 = spv.LogicalOr %arg, %false: vector<3xi1> // CHECK: return %[[TRUE]], %[[ARG]] @@ -484,11 +484,11 @@ //===----------------------------------------------------------------------===// func @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 %3 = spv.Variable init(%0) : !spv.ptr @@ -515,11 +515,11 @@ // ----- 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, Function> %3 = spv.Variable init(%0) : !spv.ptr, Function> @@ -547,11 +547,11 @@ // 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, Function> %3 = spv.Variable init(%0) : !spv.ptr, Function> // CHECK: %[[DST_VAR_1:.*]] = spv.Variable init({{%.*}}) : !spv.ptr, Function> @@ -584,11 +584,11 @@ // 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, Function> %3 = spv.Variable init(%0) : !spv.ptr, Function> // CHECK: %[[DST_VAR_1:.*]] = spv.Variable init({{%.*}}) : !spv.ptr, Function> @@ -620,11 +620,11 @@ // 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, Function> %3 = spv.Variable init(%0) : !spv.ptr, Function> @@ -652,11 +652,11 @@ // `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, Function> %3 = spv.Variable init(%0) : !spv.ptr, Function> @@ -684,11 +684,11 @@ // 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, Function> %3 = spv.Variable init(%0) : !spv.ptr, Function> diff --git a/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir b/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir @@ -2,10 +2,10 @@ // CHECK: func @clamp_fordlessthan(%[[INPUT:.*]]: f32) func @clamp_fordlessthan(%input: f32) -> f32 { - // CHECK: %[[MIN:.*]] = spv.constant - %min = spv.constant 0.5 : f32 - // CHECK: %[[MAX:.*]] = spv.constant - %max = spv.constant 1.0 : f32 + // CHECK: %[[MIN:.*]] = spv.Constant + %min = spv.Constant 0.5 : f32 + // CHECK: %[[MAX:.*]] = spv.Constant + %max = spv.Constant 1.0 : f32 // CHECK: [[RES:%.*]] = spv.GLSL.FClamp %[[INPUT]], %[[MIN]], %[[MAX]] %0 = spv.FOrdLessThan %min, %input : f32 @@ -21,10 +21,10 @@ // CHECK: func @clamp_fordlessthanequal(%[[INPUT:.*]]: f32) func @clamp_fordlessthanequal(%input: f32) -> f32 { - // CHECK: %[[MIN:.*]] = spv.constant - %min = spv.constant 0.5 : f32 - // CHECK: %[[MAX:.*]] = spv.constant - %max = spv.constant 1.0 : f32 + // CHECK: %[[MIN:.*]] = spv.Constant + %min = spv.Constant 0.5 : f32 + // CHECK: %[[MAX:.*]] = spv.Constant + %max = spv.Constant 1.0 : f32 // CHECK: [[RES:%.*]] = spv.GLSL.FClamp %[[INPUT]], %[[MIN]], %[[MAX]] %0 = spv.FOrdLessThanEqual %min, %input : f32 @@ -40,10 +40,10 @@ // CHECK: func @clamp_slessthan(%[[INPUT:.*]]: si32) func @clamp_slessthan(%input: si32) -> si32 { - // CHECK: %[[MIN:.*]] = spv.constant - %min = spv.constant 0 : si32 - // CHECK: %[[MAX:.*]] = spv.constant - %max = spv.constant 10 : si32 + // CHECK: %[[MIN:.*]] = spv.Constant + %min = spv.Constant 0 : si32 + // CHECK: %[[MAX:.*]] = spv.Constant + %max = spv.Constant 10 : si32 // CHECK: [[RES:%.*]] = spv.GLSL.SClamp %[[INPUT]], %[[MIN]], %[[MAX]] %0 = spv.SLessThan %min, %input : si32 @@ -59,10 +59,10 @@ // CHECK: func @clamp_slessthanequal(%[[INPUT:.*]]: si32) func @clamp_slessthanequal(%input: si32) -> si32 { - // CHECK: %[[MIN:.*]] = spv.constant - %min = spv.constant 0 : si32 - // CHECK: %[[MAX:.*]] = spv.constant - %max = spv.constant 10 : si32 + // CHECK: %[[MIN:.*]] = spv.Constant + %min = spv.Constant 0 : si32 + // CHECK: %[[MAX:.*]] = spv.Constant + %max = spv.Constant 10 : si32 // CHECK: [[RES:%.*]] = spv.GLSL.SClamp %[[INPUT]], %[[MIN]], %[[MAX]] %0 = spv.SLessThanEqual %min, %input : si32 @@ -78,10 +78,10 @@ // CHECK: func @clamp_ulessthan(%[[INPUT:.*]]: i32) func @clamp_ulessthan(%input: i32) -> i32 { - // CHECK: %[[MIN:.*]] = spv.constant - %min = spv.constant 0 : i32 - // CHECK: %[[MAX:.*]] = spv.constant - %max = spv.constant 10 : i32 + // CHECK: %[[MIN:.*]] = spv.Constant + %min = spv.Constant 0 : i32 + // CHECK: %[[MAX:.*]] = spv.Constant + %max = spv.Constant 10 : i32 // CHECK: [[RES:%.*]] = spv.GLSL.UClamp %[[INPUT]], %[[MIN]], %[[MAX]] %0 = spv.ULessThan %min, %input : i32 @@ -97,10 +97,10 @@ // CHECK: func @clamp_ulessthanequal(%[[INPUT:.*]]: i32) func @clamp_ulessthanequal(%input: i32) -> i32 { - // CHECK: %[[MIN:.*]] = spv.constant - %min = spv.constant 0 : i32 - // CHECK: %[[MAX:.*]] = spv.constant - %max = spv.constant 10 : i32 + // CHECK: %[[MIN:.*]] = spv.Constant + %min = spv.Constant 0 : i32 + // CHECK: %[[MAX:.*]] = spv.Constant + %max = spv.Constant 10 : i32 // CHECK: [[RES:%.*]] = spv.GLSL.UClamp %[[INPUT]], %[[MIN]], %[[MAX]] %0 = spv.ULessThanEqual %min, %input : i32 diff --git a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/inlining.mlir @@ -17,13 +17,13 @@ spv.module Logical GLSL450 { spv.func @callee() -> i32 "None" { - %0 = spv.constant 42 : i32 + %0 = spv.Constant 42 : i32 spv.ReturnValue %0 : i32 } // CHECK-LABEL: @calling_single_block_retval_func spv.func @calling_single_block_retval_func() -> i32 "None" { - // CHECK-NEXT: %[[CST:.*]] = spv.constant 42 + // CHECK-NEXT: %[[CST:.*]] = spv.Constant 42 %0 = spv.FunctionCall @callee() : () -> (i32) // CHECK-NEXT: spv.ReturnValue %[[CST]] spv.ReturnValue %0 : i32 @@ -36,12 +36,12 @@ spv.globalVariable @data bind(0, 0) : !spv.ptr [0])>, StorageBuffer> spv.func @callee() "None" { %0 = spv.mlir.addressof @data : !spv.ptr [0])>, StorageBuffer> - %1 = spv.constant 0: i32 + %1 = spv.Constant 0: i32 %2 = spv.AccessChain %0[%1, %1] : !spv.ptr [0])>, StorageBuffer>, i32, i32 spv.Branch ^next ^next: - %3 = spv.constant 42: i32 + %3 = spv.Constant 42: i32 spv.Store "StorageBuffer" %2, %3 : i32 spv.Return } @@ -49,11 +49,11 @@ // CHECK-LABEL: @calling_multi_block_ret_func spv.func @calling_multi_block_ret_func() "None" { // CHECK-NEXT: spv.mlir.addressof - // CHECK-NEXT: spv.constant 0 + // CHECK-NEXT: spv.Constant 0 // CHECK-NEXT: spv.AccessChain // CHECK-NEXT: spv.Branch ^bb1 // CHECK-NEXT: ^bb1: - // CHECK-NEXT: spv.constant + // CHECK-NEXT: spv.Constant // CHECK-NEXT: spv.Store // CHECK-NEXT: spv.Branch ^bb2 spv.FunctionCall @callee() : () -> () @@ -81,7 +81,7 @@ // CHECK-LABEL: @calling_selection_ret_func spv.func @calling_selection_ret_func() "None" { - %0 = spv.constant true + %0 = spv.Constant true // CHECK: spv.FunctionCall spv.FunctionCall @callee(%0) : (i1) -> () spv.Return @@ -104,8 +104,8 @@ // CHECK-LABEL: @calling_selection_no_ret_func spv.func @calling_selection_no_ret_func() "None" { - // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true - %0 = spv.constant true + // CHECK-NEXT: %[[TRUE:.*]] = spv.Constant true + %0 = spv.Constant true // CHECK-NEXT: spv.selection // CHECK-NEXT: spv.BranchConditional %[[TRUE]], ^bb1, ^bb2 // CHECK-NEXT: ^bb1: @@ -137,7 +137,7 @@ // CHECK-LABEL: @calling_loop_ret_func spv.func @calling_loop_ret_func() "None" { - %0 = spv.constant true + %0 = spv.Constant true // CHECK: spv.FunctionCall spv.FunctionCall @callee(%0) : (i1) -> () spv.Return @@ -164,8 +164,8 @@ // CHECK-LABEL: @calling_loop_no_ret_func spv.func @calling_loop_no_ret_func() "None" { - // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true - %0 = spv.constant true + // CHECK-NEXT: %[[TRUE:.*]] = spv.Constant true + %0 = spv.Constant true // CHECK-NEXT: spv.loop // CHECK-NEXT: spv.Branch ^bb1 // CHECK-NEXT: ^bb1: @@ -189,7 +189,7 @@ // CHECK: @inline_into_selection_region spv.func @inline_into_selection_region() "None" { - %1 = spv.constant 0 : i32 + %1 = spv.Constant 0 : i32 // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv.mlir.addressof @arg_0 // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv.mlir.addressof @arg_1 // CHECK-DAG: [[LOADPTR:%.*]] = spv.AccessChain [[ADDRESS_ARG0]] diff --git a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir --- a/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir +++ b/mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir @@ -20,7 +20,7 @@ spv.globalVariable @var5 bind(1,3) : !spv.ptr)>, StorageBuffer> spv.func @kernel() -> () "None" { - %c0 = spv.constant 0 : i32 + %c0 = spv.Constant 0 : i32 // CHECK: {{%.*}} = spv.mlir.addressof @var0 : !spv.ptr [4], f32 [12])>, Uniform> %0 = spv.mlir.addressof @var0 : !spv.ptr, f32)>, Uniform> // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr [4], f32 [12])>, Uniform> diff --git a/mlir/test/Target/SPIRV/constant.mlir b/mlir/test/Target/SPIRV/constant.mlir --- a/mlir/test/Target/SPIRV/constant.mlir +++ b/mlir/test/Target/SPIRV/constant.mlir @@ -3,10 +3,10 @@ spv.module Logical GLSL450 requires #spv.vce { // 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 %3 = spv.Variable init(%1): !spv.ptr @@ -15,12 +15,12 @@ // CHECK-LABEL: @i32_const spv.func @i32_const() -> () "None" { - // CHECK: spv.constant 0 : i32 - %0 = spv.constant 0 : i32 - // CHECK: spv.constant 10 : i32 - %1 = spv.constant 10 : i32 - // CHECK: spv.constant -5 : i32 - %2 = spv.constant -5 : i32 + // CHECK: spv.Constant 0 : i32 + %0 = spv.Constant 0 : i32 + // CHECK: spv.Constant 10 : i32 + %1 = spv.Constant 10 : i32 + // CHECK: spv.Constant -5 : i32 + %2 = spv.Constant -5 : i32 %3 = spv.IAdd %0, %1 : i32 %4 = spv.IAdd %2, %3 : i32 @@ -29,12 +29,12 @@ // CHECK-LABEL: @si32_const spv.func @si32_const() -> () "None" { - // CHECK: spv.constant 0 : si32 - %0 = spv.constant 0 : si32 - // CHECK: spv.constant 10 : si32 - %1 = spv.constant 10 : si32 - // CHECK: spv.constant -5 : si32 - %2 = spv.constant -5 : si32 + // CHECK: spv.Constant 0 : si32 + %0 = spv.Constant 0 : si32 + // CHECK: spv.Constant 10 : si32 + %1 = spv.Constant 10 : si32 + // CHECK: spv.Constant -5 : si32 + %2 = spv.Constant -5 : si32 %3 = spv.IAdd %0, %1 : si32 %4 = spv.IAdd %2, %3 : si32 @@ -46,12 +46,12 @@ // because they all use 1 as the signedness bit. So we always treat them // as signless integers. spv.func @ui32_const() -> () "None" { - // CHECK: spv.constant 0 : i32 - %0 = spv.constant 0 : ui32 - // CHECK: spv.constant 10 : i32 - %1 = spv.constant 10 : ui32 - // CHECK: spv.constant -5 : i32 - %2 = spv.constant 4294967291 : ui32 + // CHECK: spv.Constant 0 : i32 + %0 = spv.Constant 0 : ui32 + // CHECK: spv.Constant 10 : i32 + %1 = spv.Constant 10 : ui32 + // CHECK: spv.Constant -5 : i32 + %2 = spv.Constant 4294967291 : ui32 %3 = spv.IAdd %0, %1 : ui32 %4 = spv.IAdd %2, %3 : ui32 @@ -60,14 +60,14 @@ // CHECK-LABEL: @i64_const spv.func @i64_const() -> () "None" { - // CHECK: spv.constant 4294967296 : i64 - %0 = spv.constant 4294967296 : i64 // 2^32 - // CHECK: spv.constant -4294967296 : i64 - %1 = spv.constant -4294967296 : i64 // -2^32 - // CHECK: spv.constant 9223372036854775807 : i64 - %2 = spv.constant 9223372036854775807 : i64 // 2^63 - 1 - // CHECK: spv.constant -9223372036854775808 : i64 - %3 = spv.constant -9223372036854775808 : i64 // -2^63 + // CHECK: spv.Constant 4294967296 : i64 + %0 = spv.Constant 4294967296 : i64 // 2^32 + // CHECK: spv.Constant -4294967296 : i64 + %1 = spv.Constant -4294967296 : i64 // -2^32 + // CHECK: spv.Constant 9223372036854775807 : i64 + %2 = spv.Constant 9223372036854775807 : i64 // 2^63 - 1 + // CHECK: spv.Constant -9223372036854775808 : i64 + %3 = spv.Constant -9223372036854775808 : i64 // -2^63 %4 = spv.IAdd %0, %1 : i64 %5 = spv.IAdd %2, %3 : i64 @@ -76,10 +76,10 @@ // CHECK-LABEL: @i16_const spv.func @i16_const() -> () "None" { - // CHECK: spv.constant -32768 : i16 - %0 = spv.constant -32768 : i16 // -2^15 - // CHECK: spv.constant 32767 : i16 - %1 = spv.constant 32767 : i16 // 2^15 - 1 + // CHECK: spv.Constant -32768 : i16 + %0 = spv.Constant -32768 : i16 // -2^15 + // CHECK: spv.Constant 32767 : i16 + %1 = spv.Constant 32767 : i16 // 2^15 - 1 %2 = spv.IAdd %0, %1 : i16 spv.Return @@ -87,18 +87,18 @@ // CHECK-LABEL: @float_const spv.func @float_const() -> () "None" { - // CHECK: spv.constant 0.000000e+00 : f32 - %0 = spv.constant 0. : f32 - // CHECK: spv.constant 1.000000e+00 : f32 - %1 = spv.constant 1. : f32 - // CHECK: spv.constant -0.000000e+00 : f32 - %2 = spv.constant -0. : f32 - // CHECK: spv.constant -1.000000e+00 : f32 - %3 = spv.constant -1. : f32 - // CHECK: spv.constant 7.500000e-01 : f32 - %4 = spv.constant 0.75 : f32 - // CHECK: spv.constant -2.500000e-01 : f32 - %5 = spv.constant -0.25 : f32 + // CHECK: spv.Constant 0.000000e+00 : f32 + %0 = spv.Constant 0. : f32 + // CHECK: spv.Constant 1.000000e+00 : f32 + %1 = spv.Constant 1. : f32 + // CHECK: spv.Constant -0.000000e+00 : f32 + %2 = spv.Constant -0. : f32 + // CHECK: spv.Constant -1.000000e+00 : f32 + %3 = spv.Constant -1. : f32 + // CHECK: spv.Constant 7.500000e-01 : f32 + %4 = spv.Constant 0.75 : f32 + // CHECK: spv.Constant -2.500000e-01 : f32 + %5 = spv.Constant -0.25 : f32 %6 = spv.FAdd %0, %1 : f32 %7 = spv.FAdd %2, %3 : f32 @@ -109,10 +109,10 @@ // CHECK-LABEL: @double_const spv.func @double_const() -> () "None" { // TODO: test range boundary values - // CHECK: spv.constant 1.024000e+03 : f64 - %0 = spv.constant 1024. : f64 - // CHECK: spv.constant -1.024000e+03 : f64 - %1 = spv.constant -1024. : f64 + // CHECK: spv.Constant 1.024000e+03 : f64 + %0 = spv.Constant 1024. : f64 + // CHECK: spv.Constant -1.024000e+03 : f64 + %1 = spv.Constant -1024. : f64 %2 = spv.FAdd %0, %1 : f64 spv.Return @@ -120,10 +120,10 @@ // CHECK-LABEL: @half_const spv.func @half_const() -> () "None" { - // CHECK: spv.constant 5.120000e+02 : f16 - %0 = spv.constant 512. : f16 - // CHECK: spv.constant -5.120000e+02 : f16 - %1 = spv.constant -512. : f16 + // CHECK: spv.Constant 5.120000e+02 : f16 + %0 = spv.Constant 512. : f16 + // CHECK: spv.Constant -5.120000e+02 : f16 + %1 = spv.Constant -512. : f16 %2 = spv.FAdd %0, %1 : f16 spv.Return @@ -131,12 +131,12 @@ // CHECK-LABEL: @bool_vector_const spv.func @bool_vector_const() -> () "None" { - // CHECK: spv.constant dense : vector<2xi1> - %0 = spv.constant dense : vector<2xi1> - // CHECK: spv.constant dense<[true, true, true]> : vector<3xi1> - %1 = spv.constant dense : vector<3xi1> - // CHECK: spv.constant dense<[false, true]> : vector<2xi1> - %2 = spv.constant dense<[false, true]> : vector<2xi1> + // CHECK: spv.Constant dense : vector<2xi1> + %0 = spv.Constant dense : vector<2xi1> + // CHECK: spv.Constant dense<[true, true, true]> : vector<3xi1> + %1 = spv.Constant dense : 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, Function> %4 = spv.Variable init(%1): !spv.ptr, Function> @@ -146,12 +146,12 @@ // CHECK-LABEL: @int_vector_const spv.func @int_vector_const() -> () "None" { - // CHECK: spv.constant dense<0> : vector<3xi32> - %0 = spv.constant dense<0> : vector<3xi32> - // CHECK: spv.constant dense<1> : vector<3xi32> - %1 = spv.constant dense<1> : vector<3xi32> - // CHECK: spv.constant dense<[2, -3, 4]> : vector<3xi32> - %2 = spv.constant dense<[2, -3, 4]> : vector<3xi32> + // CHECK: spv.Constant dense<0> : vector<3xi32> + %0 = spv.Constant dense<0> : vector<3xi32> + // CHECK: spv.Constant dense<1> : vector<3xi32> + %1 = spv.Constant dense<1> : vector<3xi32> + // CHECK: spv.Constant dense<[2, -3, 4]> : vector<3xi32> + %2 = spv.Constant dense<[2, -3, 4]> : vector<3xi32> %3 = spv.IAdd %0, %1 : vector<3xi32> %4 = spv.IAdd %2, %3 : vector<3xi32> @@ -160,12 +160,12 @@ // CHECK-LABEL: @fp_vector_const spv.func @fp_vector_const() -> () "None" { - // CHECK: spv.constant dense<0.000000e+00> : vector<4xf32> - %0 = spv.constant dense<0.> : vector<4xf32> - // CHECK: spv.constant dense<-1.500000e+01> : vector<4xf32> - %1 = spv.constant dense<-15.> : vector<4xf32> - // CHECK: spv.constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32> - %2 = spv.constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32> + // CHECK: spv.Constant dense<0.000000e+00> : vector<4xf32> + %0 = spv.Constant dense<0.> : vector<4xf32> + // CHECK: spv.Constant dense<-1.500000e+01> : vector<4xf32> + %1 = spv.Constant dense<-15.> : vector<4xf32> + // CHECK: spv.Constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32> + %2 = spv.Constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32> %3 = spv.FAdd %0, %1 : vector<4xf32> %4 = spv.FAdd %2, %3 : vector<4xf32> @@ -174,51 +174,51 @@ // 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 // 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 %3 = spv.Load "Function" %2 : i32 %4 = spv.IAdd %0, %3 : i32 @@ -227,15 +227,15 @@ // CHECK-LABEL: @multi_dimensions_const spv.func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" { - // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> - %0 = spv.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> + // CHECK: spv.Constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> + %0 = spv.Constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> } // CHECK-LABEL: @multi_dimensions_splat_const spv.func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" { - // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> - %0 = spv.constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> + // CHECK: spv.Constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> + %0 = spv.Constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24> } } diff --git a/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir b/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir --- a/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir +++ b/mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir @@ -94,7 +94,7 @@ // CHECK-LABEL: @cooperative_matrix_access_chain spv.func @cooperative_matrix_access_chain(%a : !spv.ptr, Function>) -> !spv.ptr "None" { - %0 = spv.constant 0: i32 + %0 = spv.Constant 0: i32 // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr, Function>, i32 %1 = spv.AccessChain %a[%0] : !spv.ptr, Function>, i32 spv.ReturnValue %1 : !spv.ptr diff --git a/mlir/test/Target/SPIRV/debug.mlir b/mlir/test/Target/SPIRV/debug.mlir --- a/mlir/test/Target/SPIRV/debug.mlir +++ b/mlir/test/Target/SPIRV/debug.mlir @@ -44,7 +44,7 @@ } 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 spv.Return @@ -68,8 +68,8 @@ } 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 %jvar = spv.Variable init(%zero) : !spv.ptr spv.loop { @@ -121,9 +121,9 @@ } 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 spv.selection { // CHECK: loc({{".*debug.mlir"}}:128:5) diff --git a/mlir/test/Target/SPIRV/function-call.mlir b/mlir/test/Target/SPIRV/function-call.mlir --- a/mlir/test/Target/SPIRV/function-call.mlir +++ b/mlir/test/Target/SPIRV/function-call.mlir @@ -3,7 +3,7 @@ spv.module Logical GLSL450 requires #spv.vce { spv.globalVariable @var1 : !spv.ptr, Input> spv.func @fmain() -> i32 "None" { - %0 = spv.constant 16 : i32 + %0 = spv.Constant 16 : i32 %1 = spv.mlir.addressof @var1 : !spv.ptr, Input> // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}) : (i32) -> i32 %3 = spv.FunctionCall @f_0(%0) : (i32) -> i32 @@ -24,7 +24,7 @@ } 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 spv.loop { spv.Branch ^header @@ -44,7 +44,7 @@ spv.Return } spv.func @f_inc(%arg0 : !spv.ptr) -> () "None" { - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 %0 = spv.Load "Function" %arg0 : i32 %1 = spv.IAdd %0, %one : i32 spv.Store "Function" %arg0, %1 : i32 diff --git a/mlir/test/Target/SPIRV/global-variable.mlir b/mlir/test/Target/SPIRV/global-variable.mlir --- a/mlir/test/Target/SPIRV/global-variable.mlir +++ b/mlir/test/Target/SPIRV/global-variable.mlir @@ -28,7 +28,7 @@ spv.func @foo() "None" { // CHECK: %[[ADDR:.*]] = spv.mlir.addressof @globalInvocationID : !spv.ptr, Input> %0 = spv.mlir.addressof @globalInvocationID : !spv.ptr, Input> - %1 = spv.constant 0: i32 + %1 = spv.Constant 0: i32 // CHECK: spv.AccessChain %[[ADDR]] %2 = spv.AccessChain %0[%1] : !spv.ptr, Input>, i32 spv.Return diff --git a/mlir/test/Target/SPIRV/logical-ops.mlir b/mlir/test/Target/SPIRV/logical-ops.mlir --- a/mlir/test/Target/SPIRV/logical-ops.mlir +++ b/mlir/test/Target/SPIRV/logical-ops.mlir @@ -93,16 +93,16 @@ spv.module Logical GLSL450 requires #spv.vce { spv.specConstant @condition_scalar = true spv.func @select() -> () "None" { - %0 = spv.constant 4.0 : f32 - %1 = spv.constant 5.0 : f32 + %0 = spv.Constant 4.0 : f32 + %1 = spv.Constant 5.0 : f32 %2 = spv.mlir.referenceof @condition_scalar : i1 // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, f32 %3 = spv.Select %2, %0, %1 : i1, f32 - %4 = spv.constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32> - %5 = spv.constant dense<[6.0, 7.0, 8.0, 9.0]> : vector<4xf32> + %4 = spv.Constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32> + %5 = spv.Constant dense<[6.0, 7.0, 8.0, 9.0]> : vector<4xf32> // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, vector<4xf32> %6 = spv.Select %2, %4, %5 : i1, vector<4xf32> - %7 = spv.constant dense<[true, true, true, true]> : vector<4xi1> + %7 = spv.Constant dense<[true, true, true, true]> : vector<4xi1> // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : vector<4xi1>, vector<4xf32> %8 = spv.Select %7, %4, %5 : vector<4xi1>, vector<4xf32> spv.Return diff --git a/mlir/test/Target/SPIRV/loop.mlir b/mlir/test/Target/SPIRV/loop.mlir --- a/mlir/test/Target/SPIRV/loop.mlir +++ b/mlir/test/Target/SPIRV/loop.mlir @@ -5,8 +5,8 @@ spv.module Logical GLSL450 requires #spv.vce { // 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 // CHECK: spv.Branch ^bb1 @@ -35,7 +35,7 @@ ^continue: // CHECK-NEXT: spv.Load %val1 = spv.Load "Function" %var : i32 -// CHECK-NEXT: spv.constant 1 +// CHECK-NEXT: spv.Constant 1 // CHECK-NEXT: spv.IAdd %add = spv.IAdd %val1, %one : i32 // CHECK-NEXT: spv.Store @@ -64,13 +64,13 @@ spv.globalVariable @GV2 bind(0, 1) : !spv.ptr [0])>, StorageBuffer> spv.func @loop_kernel() "None" { %0 = spv.mlir.addressof @GV1 : !spv.ptr [0])>, StorageBuffer> - %1 = spv.constant 0 : i32 + %1 = spv.Constant 0 : i32 %2 = spv.AccessChain %0[%1] : !spv.ptr [0])>, StorageBuffer>, i32 %3 = spv.mlir.addressof @GV2 : !spv.ptr [0])>, StorageBuffer> %5 = spv.AccessChain %3[%1] : !spv.ptr [0])>, StorageBuffer>, i32 - %6 = spv.constant 4 : i32 - %7 = spv.constant 42 : i32 - %8 = spv.constant 2 : i32 + %6 = spv.Constant 4 : i32 + %7 = spv.Constant 42 : i32 + %8 = spv.Constant 2 : i32 // CHECK: spv.Branch ^bb1(%{{.*}} : i32) // CHECK-NEXT: ^bb1(%[[OUTARG:.*]]: i32): // CHECK-NEXT: spv.loop { @@ -112,8 +112,8 @@ // 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 %jvar = spv.Variable init(%zero) : !spv.ptr @@ -135,7 +135,7 @@ // CHECK-NEXT: ^bb2: ^body: -// CHECK-NEXT: spv.constant 0 +// CHECK-NEXT: spv.Constant 0 // CHECK-NEXT: spv.Store spv.Store "Function" %jvar, %zero : i32 // CHECK-NEXT: spv.Branch ^bb3 @@ -164,7 +164,7 @@ ^continue: // CHECK-NEXT: spv.Load %jval1 = spv.Load "Function" %jvar : i32 -// CHECK-NEXT: spv.constant 1 +// CHECK-NEXT: spv.Constant 1 // CHECK-NEXT: spv.IAdd %add = spv.IAdd %jval1, %one : i32 // CHECK-NEXT: spv.Store @@ -185,7 +185,7 @@ ^continue: // CHECK-NEXT: spv.Load %ival1 = spv.Load "Function" %ivar : i32 -// CHECK-NEXT: spv.constant 1 +// CHECK-NEXT: spv.Constant 1 // CHECK-NEXT: spv.IAdd %add = spv.IAdd %ival1, %one : i32 // CHECK-NEXT: spv.Store diff --git a/mlir/test/Target/SPIRV/memory-ops.mlir b/mlir/test/Target/SPIRV/memory-ops.mlir --- a/mlir/test/Target/SPIRV/memory-ops.mlir +++ b/mlir/test/Target/SPIRV/memory-ops.mlir @@ -30,13 +30,13 @@ spv.func @load_store_zero_rank_float(%arg0: !spv.ptr [0])>, StorageBuffer>, %arg1: !spv.ptr [0])>, StorageBuffer>) "None" { // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [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 [0])>, StorageBuffer>, i32, i32 %2 = spv.Load "StorageBuffer" %1 : f32 // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [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 [0])>, StorageBuffer>, i32, i32 spv.Store "StorageBuffer" %4, %2 : f32 spv.Return @@ -45,13 +45,13 @@ spv.func @load_store_zero_rank_int(%arg0: !spv.ptr [0])>, StorageBuffer>, %arg1: !spv.ptr [0])>, StorageBuffer>) "None" { // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [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 [0])>, StorageBuffer>, i32, i32 %2 = spv.Load "StorageBuffer" %1 : i32 // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [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 [0])>, StorageBuffer>, i32, i32 spv.Store "StorageBuffer" %4, %2 : i32 spv.Return diff --git a/mlir/test/Target/SPIRV/non-uniform-ops.mlir b/mlir/test/Target/SPIRV/non-uniform-ops.mlir --- a/mlir/test/Target/SPIRV/non-uniform-ops.mlir +++ b/mlir/test/Target/SPIRV/non-uniform-ops.mlir @@ -10,7 +10,7 @@ // CHECK-LABEL: @group_non_uniform_broadcast spv.func @group_non_uniform_broadcast(%value: f32) -> f32 "None" { - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 // CHECK: spv.GroupNonUniformBroadcast Subgroup %{{.*}}, %{{.*}} : f32, i32 %0 = spv.GroupNonUniformBroadcast Subgroup %value, %one : f32, i32 spv.ReturnValue %0: f32 @@ -60,7 +60,7 @@ // CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce spv.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> "None" { - %four = spv.constant 4 : i32 + %four = spv.Constant 4 : i32 // CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32> %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32> spv.ReturnValue %0: vector<2xi32> diff --git a/mlir/test/Target/SPIRV/phi.mlir b/mlir/test/Target/SPIRV/phi.mlir --- a/mlir/test/Target/SPIRV/phi.mlir +++ b/mlir/test/Target/SPIRV/phi.mlir @@ -4,8 +4,8 @@ spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { -// CHECK: %[[CST:.*]] = spv.constant 0 - %zero = spv.constant 0 : i32 +// CHECK: %[[CST:.*]] = spv.Constant 0 + %zero = spv.Constant 0 : i32 // CHECK-NEXT: spv.Branch ^bb1(%[[CST]] : i32) spv.Branch ^bb1(%zero : i32) // CHECK-NEXT: ^bb1(%{{.*}}: i32): @@ -25,10 +25,10 @@ spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { -// CHECK: %[[ZERO:.*]] = spv.constant 0 - %zero = spv.constant 0 : i32 -// CHECK-NEXT: %[[ONE:.*]] = spv.constant 1 - %one = spv.constant 1.0 : f32 +// CHECK: %[[ZERO:.*]] = spv.Constant 0 + %zero = spv.Constant 0 : i32 +// CHECK-NEXT: %[[ONE:.*]] = spv.Constant 1 + %one = spv.Constant 1.0 : f32 // CHECK-NEXT: spv.Branch ^bb1(%[[ZERO]], %[[ONE]] : i32, f32) spv.Branch ^bb1(%zero, %one : i32, f32) @@ -49,8 +49,8 @@ spv.module Logical GLSL450 requires #spv.vce { spv.func @foo() -> () "None" { -// CHECK: %[[CST0:.*]] = spv.constant 0 - %zero = spv.constant 0 : i32 +// CHECK: %[[CST0:.*]] = spv.Constant 0 + %zero = spv.Constant 0 : i32 // CHECK-NEXT: spv.Branch ^bb1(%[[CST0]] : i32) spv.Branch ^bb1(%zero : i32) @@ -58,7 +58,7 @@ ^bb1(%arg0: i32): // CHECK-NEXT: %[[ADD:.*]] = spv.IAdd %[[ARG]], %[[ARG]] : i32 %0 = spv.IAdd %arg0, %arg0 : i32 -// CHECK-NEXT: %[[CST1:.*]] = spv.constant 0 +// CHECK-NEXT: %[[CST1:.*]] = spv.Constant 0 // CHECK-NEXT: spv.Branch ^bb2(%[[CST1]], %[[ADD]] : i32, i32) spv.Branch ^bb2(%zero, %0 : i32, i32) @@ -83,8 +83,8 @@ spv.Branch ^bb1 // CHECK-NEXT: ^bb1: -// CHECK-NEXT: %[[ZERO:.*]] = spv.constant 0 -// CHECK-NEXT: %[[ONE:.*]] = spv.constant 1 +// CHECK-NEXT: %[[ZERO:.*]] = spv.Constant 0 +// CHECK-NEXT: %[[ONE:.*]] = spv.Constant 1 // CHECK-NEXT: spv.Branch ^bb2(%[[ZERO]], %[[ONE]] : i32, f32) // CHECK-NEXT: ^bb2(%{{.*}}: i32, %{{.*}}: f32): @@ -94,8 +94,8 @@ // This block is reordered to follow domination order. ^bb1: - %zero = spv.constant 0 : i32 - %one = spv.constant 1.0 : f32 + %zero = spv.Constant 0 : i32 + %one = spv.Constant 1.0 : f32 spv.Branch ^bb2(%zero, %one : i32, f32) } @@ -115,21 +115,21 @@ // CHECK: spv.selection spv.selection { - %true = spv.constant true + %true = spv.Constant true // CHECK: spv.BranchConditional %{{.*}}, ^bb1, ^bb2 spv.BranchConditional %true, ^true, ^false // CHECK-NEXT: ^bb1: ^true: -// CHECK-NEXT: %[[ZERO:.*]] = spv.constant 0 - %zero = spv.constant 0 : i32 +// CHECK-NEXT: %[[ZERO:.*]] = spv.Constant 0 + %zero = spv.Constant 0 : i32 // CHECK-NEXT: spv.Branch ^bb3(%[[ZERO]] : i32) spv.Branch ^phi(%zero: i32) // CHECK-NEXT: ^bb2: ^false: -// CHECK-NEXT: %[[ONE:.*]] = spv.constant 1 - %one = spv.constant 1 : i32 +// CHECK-NEXT: %[[ONE:.*]] = spv.Constant 1 + %one = spv.Constant 1 : i32 // CHECK-NEXT: spv.Branch ^bb3(%[[ONE]] : i32) spv.Branch ^phi(%one: i32) @@ -162,9 +162,9 @@ spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr, Input> spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr, 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, Input> %7 = spv.Load "Input" %6 : vector<3xi32> %8 = spv.CompositeExtract %7[0 : i32] : vector<3xi32> @@ -243,12 +243,12 @@ spv.module Logical GLSL450 requires #spv.vce { spv.func @fmul_kernel() "None" { - %cst4 = spv.constant 4 : i32 + %cst4 = spv.Constant 4 : i32 - %val1 = spv.constant 43 : i32 - %val2 = spv.constant 44 : i32 + %val1 = spv.Constant 43 : i32 + %val2 = spv.Constant 44 : i32 -// CHECK: spv.constant 43 +// CHECK: spv.Constant 43 // CHECK-NEXT: spv.Branch ^[[BB1:.+]](%{{.+}} : i32) // CHECK-NEXT: ^[[BB1]](%{{.+}}: i32): // CHECK-NEXT: spv.loop @@ -264,7 +264,7 @@ spv.mlir.merge } -// CHECK: spv.constant 44 +// CHECK: spv.Constant 44 // CHECK-NEXT: spv.Branch ^[[BB2:.+]](%{{.+}} : i32) // CHECK-NEXT: ^[[BB2]](%{{.+}}: i32): // CHECK-NEXT: spv.loop diff --git a/mlir/test/Target/SPIRV/selection.mlir b/mlir/test/Target/SPIRV/selection.mlir --- a/mlir/test/Target/SPIRV/selection.mlir +++ b/mlir/test/Target/SPIRV/selection.mlir @@ -6,13 +6,13 @@ 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 // CHECK-NEXT: spv.selection control(Flatten) -// CHECK-NEXT: spv.constant 0 +// CHECK-NEXT: spv.Constant 0 // CHECK-NEXT: spv.Variable spv.selection control(Flatten) { // CHECK-NEXT: spv.BranchConditional %{{.*}} [5, 10], ^bb1, ^bb2 @@ -20,7 +20,7 @@ // CHECK-NEXT: ^bb1: ^then: -// CHECK-NEXT: spv.constant 1 +// CHECK-NEXT: spv.Constant 1 // CHECK-NEXT: spv.Store spv.Store "Function" %var, %one : i32 // CHECK-NEXT: spv.Branch ^bb3 @@ -28,7 +28,7 @@ // CHECK-NEXT: ^bb2: ^else: -// CHECK-NEXT: spv.constant 2 +// CHECK-NEXT: spv.Constant 2 // CHECK-NEXT: spv.Store spv.Store "Function" %var, %two : i32 // CHECK-NEXT: spv.Branch ^bb3 @@ -67,7 +67,7 @@ // CHECK: ^bb1: ^then: - %zero = spv.constant 0 : i32 + %zero = spv.Constant 0 : i32 spv.ReturnValue %zero : i32 // CHECK: ^bb2: @@ -76,7 +76,7 @@ spv.mlir.merge } - %one = spv.constant 1 : i32 + %one = spv.Constant 1 : i32 spv.ReturnValue %one : i32 } diff --git a/mlir/test/Target/SPIRV/spec-constant.mlir b/mlir/test/Target/SPIRV/spec-constant.mlir --- a/mlir/test/Target/SPIRV/spec-constant.mlir +++ b/mlir/test/Target/SPIRV/spec-constant.mlir @@ -94,17 +94,17 @@ spv.func @use_composite() -> (i32) "None" { // CHECK: [[USE1:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32 - // CHECK: [[USE2:%.*]] = spv.constant 0 : i32 + // CHECK: [[USE2:%.*]] = spv.Constant 0 : i32 // CHECK: [[RES1:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE1]], [[USE2]]) : (i32, i32) -> i32 // CHECK: [[USE3:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32 - // CHECK: [[USE4:%.*]] = spv.constant 0 : i32 + // CHECK: [[USE4:%.*]] = spv.Constant 0 : i32 // CHECK: [[RES2:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE3]], [[USE4]]) : (i32, i32) -> i32 %0 = spv.mlir.referenceof @sc_i32_1 : i32 - %1 = spv.constant 0 : i32 + %1 = spv.Constant 0 : i32 %2 = spv.SpecConstantOperation wraps "spv.ISub"(%0, %1) : (i32, i32) -> i32 // CHECK: [[RES3:%.*]] = spv.SpecConstantOperation wraps "spv.IMul"([[RES1]], [[RES2]]) : (i32, i32) -> i32 diff --git a/mlir/test/Target/SPIRV/undef.mlir b/mlir/test/Target/SPIRV/undef.mlir --- a/mlir/test/Target/SPIRV/undef.mlir +++ b/mlir/test/Target/SPIRV/undef.mlir @@ -15,7 +15,7 @@ %6 = spv.CompositeExtract %5[1 : i32, 2 : i32] : !spv.array<4x!spv.array<4xi32>> // CHECK: {{%.*}} = spv.undef : !spv.ptr, StorageBuffer> %7 = spv.undef : !spv.ptr, StorageBuffer> - %8 = spv.constant 0 : i32 + %8 = spv.Constant 0 : i32 %9 = spv.AccessChain %7[%8] : !spv.ptr, StorageBuffer>, i32 spv.Return }