[MLIR][SPIRV] Rename `spv.constant` to `spv.Constant`.
authorKareemErgawy-TomTom <kareem.ergawy@gmail.com>
Thu, 4 Mar 2021 21:15:46 +0000 (16:15 -0500)
committerLei Zhang <antiagainst@google.com>
Thu, 4 Mar 2021 21:15:56 +0000 (16:15 -0500)
To unify the naming scheme across all ops in the SPIR-V dialect, we are
moving from `spv.camelCase` to `spv.CamelCase` everywhere.

Reviewed By: antiagainst

Differential Revision: https://reviews.llvm.org/D97917

45 files changed:
mlir/docs/Dialects/SPIR-V.md
mlir/docs/SPIRVToLLVMDialectConversion.md
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVMemoryOps.td
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVNonUniformOps.td
mlir/include/mlir/Dialect/SPIRV/IR/SPIRVStructureOps.td
mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp
mlir/lib/Dialect/SPIRV/IR/SPIRVCanonicalization.cpp
mlir/lib/Dialect/SPIRV/IR/SPIRVOps.cpp
mlir/lib/Target/SPIRV/Deserialization/DeserializeOps.cpp
mlir/test/Conversion/GPUToSPIRV/builtins.mlir
mlir/test/Conversion/GPUToSPIRV/load-store.mlir
mlir/test/Conversion/GPUToVulkan/lower-gpu-launch-vulkan-launch.mlir
mlir/test/Conversion/LinalgToSPIRV/linalg-to-spirv.mlir
mlir/test/Conversion/SCFToSPIRV/for.mlir
mlir/test/Conversion/SCFToSPIRV/if.mlir
mlir/test/Conversion/SPIRVToLLVM/constant-op-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/control-flow-ops-to-llvm.mlir
mlir/test/Conversion/SPIRVToLLVM/memory-ops-to-llvm.mlir
mlir/test/Conversion/StandardToSPIRV/std-ops-to-spirv.mlir
mlir/test/Dialect/SPIRV/IR/composite-ops.mlir
mlir/test/Dialect/SPIRV/IR/control-flow-ops.mlir
mlir/test/Dialect/SPIRV/IR/cooperative-matrix-ops.mlir
mlir/test/Dialect/SPIRV/IR/logical-ops.mlir
mlir/test/Dialect/SPIRV/IR/memory-ops.mlir
mlir/test/Dialect/SPIRV/IR/non-uniform-ops.mlir
mlir/test/Dialect/SPIRV/IR/structure-ops.mlir
mlir/test/Dialect/SPIRV/Transforms/abi-interface.mlir
mlir/test/Dialect/SPIRV/Transforms/abi-load-store.mlir
mlir/test/Dialect/SPIRV/Transforms/canonicalize.mlir
mlir/test/Dialect/SPIRV/Transforms/glsl_canonicalize.mlir
mlir/test/Dialect/SPIRV/Transforms/inlining.mlir
mlir/test/Dialect/SPIRV/Transforms/layout-decoration.mlir
mlir/test/Target/SPIRV/constant.mlir
mlir/test/Target/SPIRV/cooperative-matrix-ops.mlir
mlir/test/Target/SPIRV/debug.mlir
mlir/test/Target/SPIRV/function-call.mlir
mlir/test/Target/SPIRV/global-variable.mlir
mlir/test/Target/SPIRV/logical-ops.mlir
mlir/test/Target/SPIRV/loop.mlir
mlir/test/Target/SPIRV/memory-ops.mlir
mlir/test/Target/SPIRV/non-uniform-ops.mlir
mlir/test/Target/SPIRV/phi.mlir
mlir/test/Target/SPIRV/selection.mlir
mlir/test/Target/SPIRV/spec-constant.mlir
mlir/test/Target/SPIRV/undef.mlir

index 295dad8..2d9e668 100644 (file)
@@ -88,7 +88,7 @@ The SPIR-V dialect adopts the following conventions for IR:
 *   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 @@ instructions are represented in the SPIR-V dialect:
 #### Unify and localize constants
 
 *   Various normal constant instructions are represented by the same
-    `spv.constant` op. Those instructions are just for constants of 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 @@ For example,
 can be represented in the dialect as
 
 ```mlir
-%0 = "spv.constant"() { value = 42 : i32 } : () -> i32
+%0 = "spv.Constant"() { value = 42 : i32 } : () -> i32
 %1 = "spv.Variable"(%0) { storage_class = "Function" } : (i32) -> !spv.ptr<i32, Function>
 %2 = "spv.IAdd"(%0, %0) : (i32, i32) -> i32
 ```
@@ -581,9 +581,9 @@ It will be represented as
 
 ```mlir
 func @selection(%cond: i1) -> () {
-  %zero = spv.constant 0: i32
-  %one = spv.constant 1: i32
-  %two = spv.constant 2: i32
+  %zero = spv.Constant 0: i32
+  %one = spv.Constant 1: i32
+  %two = spv.Constant 2: i32
   %x = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
   spv.selection {
@@ -669,8 +669,8 @@ It will be represented as
 
 ```mlir
 func @loop(%count : i32) -> () {
-  %zero = spv.constant 0: i32
-  %one = spv.constant 1: i32
+  %zero = spv.Constant 0: i32
+  %one = spv.Constant 1: i32
   %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
   spv.loop {
@@ -732,15 +732,15 @@ func @foo() -> () {
   %var = spv.Variable : !spv.ptr<i32, Function>
 
   spv.selection {
-    %true = spv.constant true
+    %true = spv.Constant true
     spv.BranchConditional %true, ^true, ^false
 
   ^true:
-    %zero = spv.constant 0 : i32
+    %zero = spv.Constant 0 : i32
     spv.Branch ^phi(%zero: i32)
 
   ^false:
-    %one = spv.constant 1 : i32
+    %one = spv.Constant 1 : i32
     spv.Branch ^phi(%one: i32)
 
   ^phi(%arg: i32):
@@ -964,7 +964,7 @@ the representational differences between SPIR-V dialect and binary format:
     instructions.
 *   Types are serialized into `OpType*` instructions in the SPIR-V binary module
     section for types, constants, and global variables.
-*   `spv.constant`s are unified and placed in the SPIR-V binary module section
+*   `spv.Constant`s are unified and placed in the SPIR-V binary module section
     for types, constants, and global variables.
 *   Attributes on ops, if not part of the op's binary encoding, are emitted as
     `OpDecorate*` instructions in the SPIR-V binary module section for
@@ -980,7 +980,7 @@ Similarly, a few transformations are performed during deserialization:
     capabilities, extended instruction sets, etc.) will be placed as attributes
     on `spv.module`.
 *   `OpType*` instructions will be converted into proper `mlir::Type`s.
-*   `OpConstant*` instructions are materialized as `spv.constant` at each use
+*   `OpConstant*` instructions are materialized as `spv.Constant` at each use
     site.
 *   `OpVariable` instructions will be converted to `spv.globalVariable` ops if
     in module-level; otherwise they will be converted into `spv.Variable` ops.
index f16bd46..4623be9 100644 (file)
@@ -434,7 +434,7 @@ order to go through the pointer.
 
 ```mlir
 // Access the 1st element of the array
-%i   = spv.constant 1: i32
+%i   = spv.Constant 1: i32
 %var = spv.Variable : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>
 %el  = spv.AccessChain %var[%i, %i] : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>, i32, i32
 
@@ -529,7 +529,7 @@ also a function-level variable.
 `spv.Variable` is modelled as `llvm.alloca` op. If initialized, an additional
 store instruction is used. Note that there is no initialization for arrays and
 structs since constants of these types are not supported in LLVM dialect (TODO).
-Also, at the moment initialization is only possible via `spv.constant`.
+Also, at the moment initialization is only possible via `spv.Constant`.
 
 ```mlir
 // Conversion of VariableOp without initialization
@@ -538,7 +538,7 @@ Also, at the moment initialization is only possible via `spv.constant`.
 
 // Conversion of VariableOp with initialization
                                                                %c    = llvm.mlir.constant(0 : i64) : i64
-%c   = spv.constant 0 : i64                                    %size = llvm.mlir.constant(1 : i32) : i32
+%c   = spv.Constant 0 : i64                                    %size = llvm.mlir.constant(1 : i32) : i32
 %res = spv.Variable init(%c) : !spv.ptr<i64, Function>    =>   %res  = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr<i64>
                                                                llvm.store %c, %res : !llvm.ptr<i64>
 ```
@@ -584,14 +584,14 @@ bitwidth. This leads to the following conversions:
 %res1 = spv.ShiftRightArithmetic %0, %1 : i32, i16  =>  %res1 = llvm.ashr %0, %ext: i32
 ```
 
-### `spv.constant`
+### `spv.Constant`
 
-At the moment `spv.constant` conversion supports scalar and vector constants
+At the moment `spv.Constant` conversion supports scalar and vector constants
 **only**.
 
 #### Mapping
 
-`spv.constant` is mapped to `llvm.mlir.constant`. This is a straightforward
+`spv.Constant` is mapped to `llvm.mlir.constant`. This is a straightforward
 conversion pattern with a special case when the argument is signed or unsigned.
 
 #### Special case
@@ -608,10 +608,10 @@ cover all possible corner cases.
 
 ```mlir
 // %0 = llvm.mlir.constant(0 : i8) : i8
-%0 = spv.constant  0 : i8
+%0 = spv.Constant  0 : i8
 
 // %1 = llvm.mlir.constant(dense<[2, 3, 4]> : vector<3xi32>) : vector<3xi32>
-%1 = spv.constant dense<[2, 3, 4]> : vector<3xui32>
+%1 = spv.Constant dense<[2, 3, 4]> : vector<3xui32>
 ```
 
 ### Not implemented ops
@@ -672,7 +672,7 @@ blocks being reachable. Moreover, selection and loop control attributes (such as
 
 ```mlir
 // Conversion of selection
-%cond = spv.constant true                               %cond = llvm.mlir.constant(true) : i1
+%cond = spv.Constant true                               %cond = llvm.mlir.constant(true) : i1
 spv.selection {
   spv.BranchConditional %cond, ^true, ^false            llvm.cond_br %cond, ^true, ^false
 
@@ -693,7 +693,7 @@ spv.selection {
 
 ```mlir
 // Conversion of loop
-%cond = spv.constant true                               %cond = llvm.mlir.constant(true) : i1
+%cond = spv.Constant true                               %cond = llvm.mlir.constant(true) : i1
 spv.loop {
   spv.Branch ^header                                    llvm.br ^header
 
index 9f514a5..0cffdb5 100644 (file)
@@ -58,7 +58,7 @@ def SPV_AccessChainOp : SPV_Op<"AccessChain", [NoSideEffect]> {
     #### Example:
 
     ```mlir
-    %0 = "spv.constant"() { value = 1: i32} : () -> i32
+    %0 = "spv.Constant"() { value = 1: i32} : () -> i32
     %1 = spv.Variable : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>
     %2 = spv.AccessChain %1[%0] : !spv.ptr<!spv.struct<f32, !spv.array<4xf32>>, Function>
     %3 = spv.Load "Function" %2 ["Volatile"] : !spv.array<4xf32>
@@ -276,7 +276,7 @@ def SPV_VariableOp : SPV_Op<"Variable", []> {
     #### Example:
 
     ```mlir
-    %0 = spv.constant ...
+    %0 = spv.Constant ...
 
     %1 = spv.Variable : !spv.ptr<f32, Function>
     %2 = spv.Variable init(%0): !spv.ptr<f32, Function>
index e7a9f44..050c492 100644 (file)
@@ -254,7 +254,7 @@ def SPV_GroupNonUniformFAddOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : f32
     %vector = ... : vector<4xf32>
     %0 = spv.GroupNonUniformFAdd "Workgroup" "Reduce" %scalar : f32
@@ -314,7 +314,7 @@ def SPV_GroupNonUniformFMaxOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : f32
     %vector = ... : vector<4xf32>
     %0 = spv.GroupNonUniformFMax "Workgroup" "Reduce" %scalar : f32
@@ -374,7 +374,7 @@ def SPV_GroupNonUniformFMinOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : f32
     %vector = ... : vector<4xf32>
     %0 = spv.GroupNonUniformFMin "Workgroup" "Reduce" %scalar : f32
@@ -431,7 +431,7 @@ def SPV_GroupNonUniformFMulOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : f32
     %vector = ... : vector<4xf32>
     %0 = spv.GroupNonUniformFMul "Workgroup" "Reduce" %scalar : f32
@@ -486,7 +486,7 @@ def SPV_GroupNonUniformIAddOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : i32
     %vector = ... : vector<4xi32>
     %0 = spv.GroupNonUniformIAdd "Workgroup" "Reduce" %scalar : i32
@@ -541,7 +541,7 @@ def SPV_GroupNonUniformIMulOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : i32
     %vector = ... : vector<4xi32>
     %0 = spv.GroupNonUniformIMul "Workgroup" "Reduce" %scalar : i32
@@ -598,7 +598,7 @@ def SPV_GroupNonUniformSMaxOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : i32
     %vector = ... : vector<4xi32>
     %0 = spv.GroupNonUniformSMax "Workgroup" "Reduce" %scalar : i32
@@ -655,7 +655,7 @@ def SPV_GroupNonUniformSMinOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : i32
     %vector = ... : vector<4xi32>
     %0 = spv.GroupNonUniformSMin "Workgroup" "Reduce" %scalar : i32
@@ -713,7 +713,7 @@ def SPV_GroupNonUniformUMaxOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : i32
     %vector = ... : vector<4xi32>
     %0 = spv.GroupNonUniformUMax "Workgroup" "Reduce" %scalar : i32
@@ -771,7 +771,7 @@ def SPV_GroupNonUniformUMinOp : SPV_GroupNonUniformArithmeticOp<"GroupNonUniform
     #### Example:
 
     ```mlir
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     %scalar = ... : i32
     %vector = ... : vector<4xi32>
     %0 = spv.GroupNonUniformUMin "Workgroup" "Reduce" %scalar : i32
index bcc4acd..188b4d7 100644 (file)
@@ -67,7 +67,7 @@ def SPV_AddressOfOp : SPV_Op<"mlir.addressof", [InFunctionScope, NoSideEffect]>
 
 // -----
 
-def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> {
+def SPV_ConstantOp : SPV_Op<"Constant", [ConstantLike, NoSideEffect]> {
   let summary = "The op that declares a SPIR-V normal constant";
 
   let description = [{
@@ -81,7 +81,7 @@ def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> {
     * ...
 
     Having such a plethora of constant instructions renders IR transformations
-    more tedious. Therefore, we use a single `spv.constant` op to represent
+    more tedious. Therefore, we use a single `spv.Constant` op to represent
     them all. Note that conversion between those SPIR-V constant instructions
     and this op is purely mechanical; so it can be scoped to the binary
     (de)serialization process.
@@ -89,16 +89,16 @@ def SPV_ConstantOp : SPV_Op<"constant", [ConstantLike, NoSideEffect]> {
     <!-- End of AutoGen section -->
 
     ```
-    spv-constant-op ::= ssa-id `=` `spv.constant` attribute-value
+    spv.Constant-op ::= ssa-id `=` `spv.Constant` attribute-value
                         (`:` spirv-type)?
     ```
 
     #### Example:
 
     ```mlir
-    %0 = spv.constant true
-    %1 = spv.constant dense<[2, 3]> : vector<2xf32>
-    %2 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
+    %0 = spv.Constant true
+    %1 = spv.Constant dense<[2, 3]> : vector<2xf32>
+    %2 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
     ```
 
     TODO: support constant structs
@@ -572,7 +572,7 @@ def SPV_SpecConstantOp : SPV_Op<"SpecConstant", [InModuleScope, Symbol]> {
     * `OpSpecConstantTrue` and `OpSpecConstantFalse` for boolean constants
     * `OpSpecConstant` for scalar constants
 
-    Similar as `spv.constant`, this op represents all of the above cases.
+    Similar as `spv.Constant`, this op represents all of the above cases.
     `OpSpecConstantComposite` and `OpSpecConstantOp` are modelled with
     separate ops.
 
@@ -731,8 +731,8 @@ def SPV_SpecConstantOperationOp : SPV_Op<"SpecConstantOperation", [
 
     #### Example:
     ```mlir
-    %0 = spv.constant 1: i32
-    %1 = spv.constant 1: i32
+    %0 = spv.Constant 1: i32
+    %1 = spv.Constant 1: i32
 
     %2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32
     ```
index e07934b..6b115e6 100644 (file)
@@ -356,7 +356,7 @@ public:
   }
 };
 
-/// Converts composite std.constant operation to spv.constant.
+/// Converts composite std.constant operation to spv.Constant.
 class ConstantCompositeOpPattern final
     : public OpConversionPattern<ConstantOp> {
 public:
@@ -367,7 +367,7 @@ public:
                   ConversionPatternRewriter &rewriter) const override;
 };
 
-/// Converts scalar std.constant operation to spv.constant.
+/// Converts scalar std.constant operation to spv.Constant.
 class ConstantScalarOpPattern final : public OpConversionPattern<ConstantOp> {
 public:
   using OpConversionPattern<ConstantOp>::OpConversionPattern;
index 7f268ca..8c47a89 100644 (file)
@@ -138,11 +138,11 @@ OpFoldResult spirv::CompositeExtractOp::fold(ArrayRef<Attribute> operands) {
 }
 
 //===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
 //===----------------------------------------------------------------------===//
 
 OpFoldResult spirv::ConstantOp::fold(ArrayRef<Attribute> operands) {
-  assert(operands.empty() && "spv.constant has no operands");
+  assert(operands.empty() && "spv.Constant has no operands");
   return value();
 }
 
index cc5e204..07fa245 100644 (file)
@@ -941,7 +941,7 @@ static Type getElementPtrType(Type type, ValueRange indices, Location baseLoc) {
       Operation *op = indexSSA.getDefiningOp();
       if (!op) {
         emitError(baseLoc, "'spv.AccessChain' op index must be an "
-                           "integer spv.constant to access "
+                           "integer spv.Constant to access "
                            "element of spv.struct");
         return nullptr;
       }
@@ -950,7 +950,7 @@ static Type getElementPtrType(Type type, ValueRange indices, Location baseLoc) {
       // integer literals of other bitwidths.
       if (failed(extractValueFromConstOp(op, index))) {
         emitError(baseLoc,
-                  "'spv.AccessChain' index must be an integer spv.constant to "
+                  "'spv.AccessChain' index must be an integer spv.Constant to "
                   "access element of spv.struct, but provided ")
             << op->getName();
         return nullptr;
@@ -1483,7 +1483,7 @@ static void print(spirv::CompositeInsertOp compositeInsertOp,
 }
 
 //===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
 //===----------------------------------------------------------------------===//
 
 static ParseResult parseConstantOp(OpAsmParser &parser, OperationState &state) {
index 7424f08..06e7f81 100644 (file)
@@ -38,7 +38,7 @@ static inline spirv::Opcode extractOpcode(uint32_t word) {
 
 Value spirv::Deserializer::getValue(uint32_t id) {
   if (auto constInfo = getConstant(id)) {
-    // Materialize a `spv.constant` op at every use site.
+    // Materialize a `spv.Constant` op at every use site.
     return opBuilder.create<spirv::ConstantOp>(unknownLoc, constInfo->second,
                                                constInfo->first);
   }
index 34fc4af..05adf87 100644 (file)
@@ -88,7 +88,7 @@ module attributes {gpu.container_module} {
       // Note that this ignores the workgroup size specification in gpu.launch.
       // We may want to define gpu.workgroup_size and convert it to the entry
       // point ABI we want here.
-      // CHECK: spv.constant 32 : i32
+      // CHECK: spv.Constant 32 : i32
       %0 = "gpu.block_dim"() {dimension = "x"} : () -> index
       gpu.return
     }
@@ -110,7 +110,7 @@ module attributes {gpu.container_module} {
     gpu.func @builtin_workgroup_size_y() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
-      // CHECK: spv.constant 4 : i32
+      // CHECK: spv.Constant 4 : i32
       %0 = "gpu.block_dim"() {dimension = "y"} : () -> index
       gpu.return
     }
@@ -132,7 +132,7 @@ module attributes {gpu.container_module} {
     gpu.func @builtin_workgroup_size_z() kernel
       attributes {spv.entry_point_abi = {local_size = dense<[32, 4, 1]>: vector<3xi32>}} {
       // The constant value is obtained from the spv.entry_point_abi.
-      // CHECK: spv.constant 1 : i32
+      // CHECK: spv.Constant 1 : i32
       %0 = "gpu.block_dim"() {dimension = "z"} : () -> index
       gpu.return
     }
index 40ce2d4..de4af83 100644 (file)
@@ -59,12 +59,12 @@ module attributes {
       %12 = addi %arg3, %0 : index
       // CHECK: %[[INDEX2:.*]] = spv.IAdd %[[ARG4]], %[[LOCALINVOCATIONIDX]]
       %13 = addi %arg4, %3 : index
-      // CHECK: %[[ZERO:.*]] = spv.constant 0 : i32
-      // CHECK: %[[OFFSET1_0:.*]] = spv.constant 0 : i32
-      // CHECK: %[[STRIDE1_1:.*]] = spv.constant 4 : i32
+      // CHECK: %[[ZERO:.*]] = spv.Constant 0 : i32
+      // CHECK: %[[OFFSET1_0:.*]] = spv.Constant 0 : i32
+      // CHECK: %[[STRIDE1_1:.*]] = spv.Constant 4 : i32
       // CHECK: %[[UPDATE1_1:.*]] = spv.IMul %[[STRIDE1_1]], %[[INDEX1]] : i32
       // CHECK: %[[OFFSET1_1:.*]] = spv.IAdd %[[OFFSET1_0]], %[[UPDATE1_1]] : i32
-      // CHECK: %[[STRIDE1_2:.*]] = spv.constant 1 : i32
+      // CHECK: %[[STRIDE1_2:.*]] = spv.Constant 1 : i32
       // CHECK: %[[UPDATE1_2:.*]] = spv.IMul %[[STRIDE1_2]], %[[INDEX2]] : i32
       // CHECK: %[[OFFSET1_2:.*]] = spv.IAdd %[[OFFSET1_1]], %[[UPDATE1_2]] : i32
       // CHECK: %[[PTR1:.*]] = spv.AccessChain %[[ARG0]]{{\[}}%[[ZERO]], %[[OFFSET1_2]]{{\]}}
index d99d28e..9b64d49 100644 (file)
@@ -9,7 +9,7 @@ module attributes {gpu.container_module} {
     spv.globalVariable @kernel_arg_0 bind(0, 0) : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
     spv.func @kernel() "None" attributes {workgroup_attributions = 0 : i64} {
       %0 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
-      %2 = spv.constant 0 : i32
+      %2 = spv.Constant 0 : i32
       %3 = spv.mlir.addressof @kernel_arg_0 : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>
       %4 = spv.AccessChain %0[%2, %2] : !spv.ptr<!spv.struct<(!spv.array<12 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
       %5 = spv.Load "StorageBuffer" %4 : f32
index e81c4cd..7a23e53 100644 (file)
@@ -23,7 +23,7 @@ module attributes {
 // CHECK:      @single_workgroup_reduction
 // CHECK-SAME: (%[[INPUT:.+]]: !spv.ptr{{.+}}, %[[OUTPUT:.+]]: !spv.ptr{{.+}})
 
-// CHECK:        %[[ZERO:.+]] = spv.constant 0 : i32
+// CHECK:        %[[ZERO:.+]] = spv.Constant 0 : i32
 // CHECK:        %[[ID:.+]] = spv.Load "Input" %{{.+}} : vector<3xi32>
 // CHECK:        %[[X:.+]] = spv.CompositeExtract %[[ID]][0 : i32]
 
index 3e4545b..c0aa5ef 100644 (file)
@@ -6,11 +6,11 @@ module attributes {
 } {
 
 func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
-  // CHECK: %[[LB:.*]] = spv.constant 4 : i32
+  // CHECK: %[[LB:.*]] = spv.Constant 4 : i32
   %lb = constant 4 : index
-  // CHECK: %[[UB:.*]] = spv.constant 42 : i32
+  // CHECK: %[[UB:.*]] = spv.Constant 42 : i32
   %ub = constant 42 : index
-  // CHECK: %[[STEP:.*]] = spv.constant 2 : i32
+  // CHECK: %[[STEP:.*]] = spv.Constant 2 : i32
   %step = constant 2 : index
   // CHECK:      spv.loop {
   // CHECK-NEXT:   spv.Branch ^[[HEADER:.*]](%[[LB]] : i32)
@@ -18,15 +18,15 @@ func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
   // CHECK:        %[[CMP:.*]] = spv.SLessThan %[[INDVAR]], %[[UB]] : i32
   // CHECK:        spv.BranchConditional %[[CMP]], ^[[BODY:.*]], ^[[MERGE:.*]]
   // CHECK:      ^[[BODY]]:
-  // CHECK:        %[[ZERO1:.*]] = spv.constant 0 : i32
-  // CHECK:        %[[OFFSET1:.*]] = spv.constant 0 : i32
-  // CHECK:        %[[STRIDE1:.*]] = spv.constant 1 : i32
+  // CHECK:        %[[ZERO1:.*]] = spv.Constant 0 : i32
+  // CHECK:        %[[OFFSET1:.*]] = spv.Constant 0 : i32
+  // CHECK:        %[[STRIDE1:.*]] = spv.Constant 1 : i32
   // CHECK:        %[[UPDATE1:.*]] = spv.IMul %[[STRIDE1]], %[[INDVAR]] : i32
   // CHECK:        %[[INDEX1:.*]] = spv.IAdd %[[OFFSET1]], %[[UPDATE1]] : i32
   // CHECK:        spv.AccessChain {{%.*}}{{\[}}%[[ZERO1]], %[[INDEX1]]{{\]}}
-  // CHECK:        %[[ZERO2:.*]] = spv.constant 0 : i32
-  // CHECK:        %[[OFFSET2:.*]] = spv.constant 0 : i32
-  // CHECK:        %[[STRIDE2:.*]] = spv.constant 1 : i32
+  // CHECK:        %[[ZERO2:.*]] = spv.Constant 0 : i32
+  // CHECK:        %[[OFFSET2:.*]] = spv.Constant 0 : i32
+  // CHECK:        %[[STRIDE2:.*]] = spv.Constant 1 : i32
   // CHECK:        %[[UPDATE2:.*]] = spv.IMul %[[STRIDE2]], %[[INDVAR]] : i32
   // CHECK:        %[[INDEX2:.*]] = spv.IAdd %[[OFFSET2]], %[[UPDATE2]] : i32
   // CHECK:        spv.AccessChain {{%.*}}[%[[ZERO2]], %[[INDEX2]]]
@@ -44,15 +44,15 @@ func @loop_kernel(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
 
 // CHECK-LABEL: @loop_yield
 func @loop_yield(%arg2 : memref<10xf32>, %arg3 : memref<10xf32>) {
-  // CHECK: %[[LB:.*]] = spv.constant 4 : i32
+  // CHECK: %[[LB:.*]] = spv.Constant 4 : i32
   %lb = constant 4 : index
-  // CHECK: %[[UB:.*]] = spv.constant 42 : i32
+  // CHECK: %[[UB:.*]] = spv.Constant 42 : i32
   %ub = constant 42 : index
-  // CHECK: %[[STEP:.*]] = spv.constant 2 : i32
+  // CHECK: %[[STEP:.*]] = spv.Constant 2 : i32
   %step = constant 2 : index
-  // CHECK: %[[INITVAR1:.*]] = spv.constant 0.000000e+00 : f32
+  // CHECK: %[[INITVAR1:.*]] = spv.Constant 0.000000e+00 : f32
   %s0 = constant 0.0 : f32
-  // CHECK: %[[INITVAR2:.*]] = spv.constant 1.000000e+00 : f32
+  // CHECK: %[[INITVAR2:.*]] = spv.Constant 1.000000e+00 : f32
   %s1 = constant 1.0 : f32
   // CHECK: %[[VAR1:.*]] = spv.Variable : !spv.ptr<f32, Function>
   // CHECK: %[[VAR2:.*]] = spv.Variable : !spv.ptr<f32, Function>
index d7c048f..0db2dfc 100644 (file)
@@ -86,14 +86,14 @@ func @simple_if_yield(%arg2 : memref<10xf32>, %arg3 : i1) {
   // CHECK:       spv.selection {
   // CHECK-NEXT:    spv.BranchConditional {{%.*}}, [[TRUE:\^.*]], [[FALSE:\^.*]]
   // CHECK-NEXT:  [[TRUE]]:
-  // CHECK:         %[[RET1TRUE:.*]] = spv.constant 0.000000e+00 : f32
-  // CHECK:         %[[RET2TRUE:.*]] = spv.constant 1.000000e+00 : f32
+  // CHECK:         %[[RET1TRUE:.*]] = spv.Constant 0.000000e+00 : f32
+  // CHECK:         %[[RET2TRUE:.*]] = spv.Constant 1.000000e+00 : f32
   // CHECK-DAG:     spv.Store "Function" %[[VAR1]], %[[RET1TRUE]] : f32
   // CHECK-DAG:     spv.Store "Function" %[[VAR2]], %[[RET2TRUE]] : f32
   // CHECK:         spv.Branch ^[[MERGE:.*]]
   // CHECK-NEXT:  [[FALSE]]:
-  // CHECK:         %[[RET2FALSE:.*]] = spv.constant 2.000000e+00 : f32
-  // CHECK:         %[[RET1FALSE:.*]] = spv.constant 3.000000e+00 : f32
+  // CHECK:         %[[RET2FALSE:.*]] = spv.Constant 2.000000e+00 : f32
+  // CHECK:         %[[RET1FALSE:.*]] = spv.Constant 3.000000e+00 : f32
   // CHECK-DAG:     spv.Store "Function" %[[VAR1]], %[[RET1FALSE]] : f32
   // CHECK-DAG:     spv.Store "Function" %[[VAR2]], %[[RET2FALSE]] : f32
   // CHECK:         spv.Branch ^[[MERGE]]
index 0ab431b..24c3aeb 100644 (file)
@@ -1,61 +1,61 @@
 // RUN: mlir-opt -convert-spirv-to-llvm %s | FileCheck %s
 
 //===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
 //===----------------------------------------------------------------------===//
 
 // CHECK-LABEL: @bool_constant_scalar
 spv.func @bool_constant_scalar() "None" {
   // CHECK: llvm.mlir.constant(true) : i1
-  %0 = spv.constant true
+  %0 = spv.Constant true
   // CHECK: llvm.mlir.constant(false) : i1
-  %1 = spv.constant false
+  %1 = spv.Constant false
   spv.Return
 }
 
 // CHECK-LABEL: @bool_constant_vector
 spv.func @bool_constant_vector() "None" {
   // CHECK: llvm.mlir.constant(dense<[true, false]> : vector<2xi1>) : vector<2xi1>
-  %0 = spv.constant dense<[true, false]> : vector<2xi1>
+  %0 = spv.Constant dense<[true, false]> : vector<2xi1>
   // CHECK: llvm.mlir.constant(dense<false> : vector<3xi1>) : vector<3xi1>
-  %1 = spv.constant dense<false> : vector<3xi1>
+  %1 = spv.Constant dense<false> : vector<3xi1>
   spv.Return
 }
 
 // CHECK-LABEL: @integer_constant_scalar
 spv.func @integer_constant_scalar() "None" {
   // CHECK: llvm.mlir.constant(0 : i8) : i8
-  %0 = spv.constant  0 : i8
+  %0 = spv.Constant  0 : i8
   // CHECK: llvm.mlir.constant(-5 : i64) : i64
-  %1 = spv.constant -5 : si64
+  %1 = spv.Constant -5 : si64
   // CHECK: llvm.mlir.constant(10 : i16) : i16
-  %2 = spv.constant  10 : ui16
+  %2 = spv.Constant  10 : ui16
   spv.Return
 }
 
 // CHECK-LABEL: @integer_constant_vector
 spv.func @integer_constant_vector() "None" {
   // CHECK: llvm.mlir.constant(dense<[2, 3]> : vector<2xi32>) : vector<2xi32>
-  %0 = spv.constant dense<[2, 3]> : vector<2xi32>
+  %0 = spv.Constant dense<[2, 3]> : vector<2xi32>
   // CHECK: llvm.mlir.constant(dense<-4> : vector<2xi32>) : vector<2xi32>
-  %1 = spv.constant dense<-4> : vector<2xsi32>
+  %1 = spv.Constant dense<-4> : vector<2xsi32>
   // CHECK: llvm.mlir.constant(dense<[2, 3, 4]> : vector<3xi32>) : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xui32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xui32>
   spv.Return
 }
 
 // CHECK-LABEL: @float_constant_scalar
 spv.func @float_constant_scalar() "None" {
   // CHECK: llvm.mlir.constant(5.000000e+00 : f16) : f16
-  %0 = spv.constant 5.000000e+00 : f16
+  %0 = spv.Constant 5.000000e+00 : f16
   // CHECK: llvm.mlir.constant(5.000000e+00 : f64) : f64
-  %1 = spv.constant 5.000000e+00 : f64
+  %1 = spv.Constant 5.000000e+00 : f64
   spv.Return
 }
 
 // CHECK-LABEL: @float_constant_vector
 spv.func @float_constant_vector() "None" {
   // CHECK: llvm.mlir.constant(dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>) : vector<2xf32>
-  %0 = spv.constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>
+  %0 = spv.Constant dense<[2.000000e+00, 3.000000e+00]> : vector<2xf32>
   spv.Return
 }
index 4211b47..0a2489f 100644 (file)
@@ -14,8 +14,8 @@ spv.module Logical GLSL450 {
   }
 
   spv.func @branch_with_arguments() -> () "None" {
-    %0 = spv.constant 0 : i32
-    %1 = spv.constant true
+    %0 = spv.Constant 0 : i32
+    %1 = spv.Constant true
     // CHECK: llvm.br ^bb1(%{{.*}}, %{{.*}} : i32, i1)
     spv.Branch ^label(%0, %1: i32, i1)
   // CHECK: ^bb1(%{{.*}}: i32, %{{.*}}: i1)
@@ -33,7 +33,7 @@ spv.module Logical GLSL450 {
 spv.module Logical GLSL450 {
   spv.func @cond_branch_without_arguments() -> () "None" {
     // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
-    %cond = spv.constant true
+    %cond = spv.Constant true
     // CHECK: lvm.cond_br %[[COND]], ^bb1, ^bb2
     spv.BranchConditional %cond, ^true, ^false
     // CHECK: ^bb1:
@@ -46,10 +46,10 @@ spv.module Logical GLSL450 {
 
   spv.func @cond_branch_with_arguments_nested() -> () "None" {
     // CHECK: %[[COND1:.*]] = llvm.mlir.constant(true) : i1
-    %cond = spv.constant true
-    %0 = spv.constant 0 : i32
+    %cond = spv.Constant true
+    %0 = spv.Constant 0 : i32
     // CHECK: %[[COND2:.*]] = llvm.mlir.constant(false) : i1
-    %false = spv.constant false
+    %false = spv.Constant false
     // CHECK: llvm.cond_br %[[COND1]], ^bb1(%{{.*}}, %[[COND2]] : i32, i1), ^bb2
     spv.BranchConditional %cond, ^outer_true(%0, %false: i32, i1), ^outer_false
   // CHECK: ^bb1(%{{.*}}: i32, %[[COND:.*]]: i1):
@@ -103,7 +103,7 @@ spv.module Logical GLSL450 {
     spv.loop {
       spv.Branch ^header
     ^header:
-      %cond = spv.constant true
+      %cond = spv.Constant true
       spv.BranchConditional %cond, ^body, ^merge
     ^body:
       // Do nothing
@@ -133,7 +133,7 @@ spv.module Logical GLSL450 {
   }
 
   spv.func @selection_with_merge_block_only() -> () "None" {
-    %cond = spv.constant true
+    %cond = spv.Constant true
     // CHECK: llvm.return
     spv.selection {
       spv.BranchConditional %cond, ^merge, ^merge
@@ -145,7 +145,7 @@ spv.module Logical GLSL450 {
 
   spv.func @selection_with_true_block_only() -> () "None" {
     // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
-    %cond = spv.constant true
+    %cond = spv.Constant true
     // CHECK: llvm.cond_br %[[COND]], ^bb1, ^bb2
     spv.selection {
       spv.BranchConditional %cond, ^true, ^merge
@@ -165,7 +165,7 @@ spv.module Logical GLSL450 {
 
   spv.func @selection_with_both_true_and_false_block() -> () "None" {
     // CHECK: %[[COND:.*]] = llvm.mlir.constant(true) : i1
-    %cond = spv.constant true
+    %cond = spv.Constant true
     // CHECK: llvm.cond_br %[[COND]], ^bb1, ^bb2
     spv.selection {
       spv.BranchConditional %cond, ^true, ^false
@@ -189,7 +189,7 @@ spv.module Logical GLSL450 {
 
   spv.func @selection_with_early_return(%arg0: i1) -> i32 "None" {
     // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
-    %0 = spv.constant 0 : i32
+    %0 = spv.Constant 0 : i32
     // CHECK: llvm.cond_br %{{.*}}, ^bb1(%[[ZERO]] : i32), ^bb2
     spv.selection {
       spv.BranchConditional %arg0, ^true(%0 : i32), ^merge
@@ -203,7 +203,7 @@ spv.module Logical GLSL450 {
       spv.mlir.merge
     }
     // CHECK: ^bb3:
-    %one = spv.constant 1 : i32
+    %one = spv.Constant 1 : i32
     spv.ReturnValue %one : i32
   }
 }
index 6aab710..6eebe4c 100644 (file)
@@ -7,7 +7,7 @@
 // CHECK-LABEL: @access_chain
 spv.func @access_chain() "None" {
   // CHECK: %[[ONE:.*]] = llvm.mlir.constant(1 : i32) : i32
-  %0 = spv.constant 1: i32
+  %0 = spv.Constant 1: i32
   %1 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
   // CHECK: %[[ZERO:.*]] = llvm.mlir.constant(0 : i32) : i32
   // CHECK: llvm.getelementptr %{{.*}}[%[[ZERO]], %[[ONE]], %[[ONE]]] : (!llvm.ptr<struct<packed (f32, array<4 x f32>)>>, i32, i32, i32) -> !llvm.ptr<f32>
@@ -176,7 +176,7 @@ spv.func @variable_scalar_with_initialization() "None" {
   // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
   // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x i64 : (i32) -> !llvm.ptr<i64>
   // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr<i64>
-  %c = spv.constant 0 : i64
+  %c = spv.Constant 0 : i64
   %0 = spv.Variable init(%c) : !spv.ptr<i64, Function>
   spv.Return
 }
@@ -195,7 +195,7 @@ spv.func @variable_vector_with_initialization() "None" {
   // CHECK: %[[SIZE:.*]] = llvm.mlir.constant(1 : i32) : i32
   // CHECK: %[[ALLOCATED:.*]] = llvm.alloca %[[SIZE]] x vector<3xi1> : (i32) -> !llvm.ptr<vector<3xi1>>
   // CHECK: llvm.store %[[VALUE]], %[[ALLOCATED]] : !llvm.ptr<vector<3xi1>>
-  %c = spv.constant dense<false> : vector<3xi1>
+  %c = spv.Constant dense<false> : vector<3xi1>
   %0 = spv.Variable init(%c) : !spv.ptr<vector<3xi1>, Function>
   spv.Return
 }
index 6bb6d78..7729ec2 100644 (file)
@@ -410,53 +410,53 @@ module attributes {
 
 // CHECK-LABEL: @constant
 func @constant() {
-  // CHECK: spv.constant true
+  // CHECK: spv.Constant true
   %0 = constant true
-  // CHECK: spv.constant 42 : i32
+  // CHECK: spv.Constant 42 : i32
   %1 = constant 42 : i32
-  // CHECK: spv.constant 5.000000e-01 : f32
+  // CHECK: spv.Constant 5.000000e-01 : f32
   %2 = constant 0.5 : f32
-  // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
+  // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
   %3 = constant dense<[2, 3]> : vector<2xi32>
-  // CHECK: spv.constant 1 : i32
+  // CHECK: spv.Constant 1 : i32
   %4 = constant 1 : index
-  // CHECK: spv.constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+  // CHECK: spv.Constant dense<1> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %5 = constant dense<1> : tensor<2x3xi32>
-  // CHECK: spv.constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
+  // CHECK: spv.Constant dense<1.000000e+00> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
   %6 = constant dense<1.0> : tensor<2x3xf32>
-  // CHECK: spv.constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
+  // CHECK: spv.Constant dense<{{\[}}1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00, 5.000000e+00, 6.000000e+00]> : tensor<6xf32> : !spv.array<6 x f32, stride=4>
   %7 = constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32>
-  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+  // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %8 = constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32>
-  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+  // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %9 =  constant dense<[[1, 2], [3, 4], [5, 6]]> : tensor<3x2xi32>
-  // CHECK: spv.constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
+  // CHECK: spv.Constant dense<{{\[}}1, 2, 3, 4, 5, 6]> : tensor<6xi32> : !spv.array<6 x i32, stride=4>
   %10 =  constant dense<[1, 2, 3, 4, 5, 6]> : tensor<6xi32>
   return
 }
 
 // CHECK-LABEL: @constant_16bit
 func @constant_16bit() {
-  // CHECK: spv.constant 4 : i16
+  // CHECK: spv.Constant 4 : i16
   %0 = constant 4 : i16
-  // CHECK: spv.constant 5.000000e+00 : f16
+  // CHECK: spv.Constant 5.000000e+00 : f16
   %1 = constant 5.0 : f16
-  // CHECK: spv.constant dense<[2, 3]> : vector<2xi16>
+  // CHECK: spv.Constant dense<[2, 3]> : vector<2xi16>
   %2 = constant dense<[2, 3]> : vector<2xi16>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2>
+  // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf16> : !spv.array<5 x f16, stride=2>
   %3 = constant dense<4.0> : tensor<5xf16>
   return
 }
 
 // CHECK-LABEL: @constant_64bit
 func @constant_64bit() {
-  // CHECK: spv.constant 4 : i64
+  // CHECK: spv.Constant 4 : i64
   %0 = constant 4 : i64
-  // CHECK: spv.constant 5.000000e+00 : f64
+  // CHECK: spv.Constant 5.000000e+00 : f64
   %1 = constant 5.0 : f64
-  // CHECK: spv.constant dense<[2, 3]> : vector<2xi64>
+  // CHECK: spv.Constant dense<[2, 3]> : vector<2xi64>
   %2 = constant dense<[2, 3]> : vector<2xi64>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8>
+  // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf64> : !spv.array<5 x f64, stride=8>
   %3 = constant dense<4.0> : tensor<5xf64>
   return
 }
@@ -472,58 +472,58 @@ module attributes {
 
 // CHECK-LABEL: @constant_16bit
 func @constant_16bit() {
-  // CHECK: spv.constant 4 : i32
+  // CHECK: spv.Constant 4 : i32
   %0 = constant 4 : i16
-  // CHECK: spv.constant 5.000000e+00 : f32
+  // CHECK: spv.Constant 5.000000e+00 : f32
   %1 = constant 5.0 : f16
-  // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
+  // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
   %2 = constant dense<[2, 3]> : vector<2xi16>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
+  // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
   %3 = constant dense<4.0> : tensor<5xf16>
-  // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
+  // CHECK: spv.Constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
   %4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16>
   return
 }
 
 // CHECK-LABEL: @constant_64bit
 func @constant_64bit() {
-  // CHECK: spv.constant 4 : i32
+  // CHECK: spv.Constant 4 : i32
   %0 = constant 4 : i64
-  // CHECK: spv.constant 5.000000e+00 : f32
+  // CHECK: spv.Constant 5.000000e+00 : f32
   %1 = constant 5.0 : f64
-  // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
+  // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
   %2 = constant dense<[2, 3]> : vector<2xi64>
-  // CHECK: spv.constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
+  // CHECK: spv.Constant dense<4.000000e+00> : tensor<5xf32> : !spv.array<5 x f32, stride=4>
   %3 = constant dense<4.0> : tensor<5xf64>
-  // CHECK: spv.constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
+  // CHECK: spv.Constant dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf32> : !spv.array<4 x f32, stride=4>
   %4 = constant dense<[[1.0, 2.0], [3.0, 4.0]]> : tensor<2x2xf16>
   return
 }
 
 // CHECK-LABEL: @corner_cases
 func @corner_cases() {
-  // CHECK: %{{.*}} = spv.constant -1 : i32
+  // CHECK: %{{.*}} = spv.Constant -1 : i32
   %0 = constant 4294967295  : i64 // 2^32 - 1
-  // CHECK: %{{.*}} = spv.constant 2147483647 : i32
+  // CHECK: %{{.*}} = spv.Constant 2147483647 : i32
   %1 = constant 2147483647  : i64 // 2^31 - 1
-  // CHECK: %{{.*}} = spv.constant -2147483648 : i32
+  // CHECK: %{{.*}} = spv.Constant -2147483648 : i32
   %2 = constant 2147483648  : i64 // 2^31
-  // CHECK: %{{.*}} = spv.constant -2147483648 : i32
+  // CHECK: %{{.*}} = spv.Constant -2147483648 : i32
   %3 = constant -2147483648 : i64 // -2^31
 
-  // CHECK: %{{.*}} = spv.constant -1 : i32
+  // CHECK: %{{.*}} = spv.Constant -1 : i32
   %5 = constant -1 : i64
-  // CHECK: %{{.*}} = spv.constant -2 : i32
+  // CHECK: %{{.*}} = spv.Constant -2 : i32
   %6 = constant -2 : i64
-  // CHECK: %{{.*}} = spv.constant -1 : i32
+  // CHECK: %{{.*}} = spv.Constant -1 : i32
   %7 = constant -1 : index
-  // CHECK: %{{.*}} = spv.constant -2 : i32
+  // CHECK: %{{.*}} = spv.Constant -2 : i32
   %8 = constant -2 : index
 
 
-  // CHECK: spv.constant false
+  // CHECK: spv.Constant false
   %9 = constant false
-  // CHECK: spv.constant true
+  // CHECK: spv.Constant true
   %10 = constant true
 
   return
@@ -639,8 +639,8 @@ func @uitofp_i32_f32(%arg0 : i32) -> f32 {
 
 // CHECK-LABEL: @uitofp_i1_f32
 func @uitofp_i1_f32(%arg0 : i1) -> f32 {
-  // CHECK: %[[ZERO:.+]] = spv.constant 0.000000e+00 : f32
-  // CHECK: %[[ONE:.+]] = spv.constant 1.000000e+00 : f32
+  // CHECK: %[[ZERO:.+]] = spv.Constant 0.000000e+00 : f32
+  // CHECK: %[[ONE:.+]] = spv.Constant 1.000000e+00 : f32
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, f32
   %0 = std.uitofp %arg0 : i1 to f32
   return %0 : f32
@@ -648,8 +648,8 @@ func @uitofp_i1_f32(%arg0 : i1) -> f32 {
 
 // CHECK-LABEL: @uitofp_i1_f64
 func @uitofp_i1_f64(%arg0 : i1) -> f64 {
-  // CHECK: %[[ZERO:.+]] = spv.constant 0.000000e+00 : f64
-  // CHECK: %[[ONE:.+]] = spv.constant 1.000000e+00 : f64
+  // CHECK: %[[ZERO:.+]] = spv.Constant 0.000000e+00 : f64
+  // CHECK: %[[ONE:.+]] = spv.Constant 1.000000e+00 : f64
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, f64
   %0 = std.uitofp %arg0 : i1 to f64
   return %0 : f64
@@ -657,8 +657,8 @@ func @uitofp_i1_f64(%arg0 : i1) -> f64 {
 
 // CHECK-LABEL: @uitofp_vec_i1_f32
 func @uitofp_vec_i1_f32(%arg0 : vector<4xi1>) -> vector<4xf32> {
-  // CHECK: %[[ZERO:.+]] = spv.constant dense<0.000000e+00> : vector<4xf32>
-  // CHECK: %[[ONE:.+]] = spv.constant dense<1.000000e+00> : vector<4xf32>
+  // CHECK: %[[ZERO:.+]] = spv.Constant dense<0.000000e+00> : vector<4xf32>
+  // CHECK: %[[ONE:.+]] = spv.Constant dense<1.000000e+00> : vector<4xf32>
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xf32>
   %0 = std.uitofp %arg0 : vector<4xi1> to vector<4xf32>
   return %0 : vector<4xf32>
@@ -666,11 +666,11 @@ func @uitofp_vec_i1_f32(%arg0 : vector<4xi1>) -> vector<4xf32> {
 
 // CHECK-LABEL: @uitofp_vec_i1_f64
 spv.func @uitofp_vec_i1_f64(%arg0: vector<4xi1>) -> vector<4xf64> "None" {
-  // CHECK: %[[ZERO:.+]] = spv.constant dense<0.000000e+00> : vector<4xf64>
-  // CHECK: %[[ONE:.+]] = spv.constant dense<1.000000e+00> : vector<4xf64>
+  // CHECK: %[[ZERO:.+]] = spv.Constant dense<0.000000e+00> : vector<4xf64>
+  // CHECK: %[[ONE:.+]] = spv.Constant dense<1.000000e+00> : vector<4xf64>
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xf64>
-  %0 = spv.constant dense<0.000000e+00> : vector<4xf64>
-  %1 = spv.constant dense<1.000000e+00> : vector<4xf64>
+  %0 = spv.Constant dense<0.000000e+00> : vector<4xf64>
+  %1 = spv.Constant dense<1.000000e+00> : vector<4xf64>
   %2 = spv.Select %arg0, %1, %0 : vector<4xi1>, vector<4xf64>
   spv.ReturnValue %2 : vector<4xf64>
 }
@@ -705,8 +705,8 @@ func @zexti2(%arg0 : i32) -> i64 {
 
 // CHECK-LABEL: @zexti3
 func @zexti3(%arg0 : i1) -> i32 {
-  // CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  // CHECK: %[[ONE:.+]] = spv.constant 1 : i32
+  // CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  // CHECK: %[[ONE:.+]] = spv.Constant 1 : i32
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : i1, i32
   %0 = std.zexti %arg0 : i1 to i32
   return %0 : i32
@@ -714,8 +714,8 @@ func @zexti3(%arg0 : i1) -> i32 {
 
 // CHECK-LABEL: @zexti4
 func @zexti4(%arg0 : vector<4xi1>) -> vector<4xi32> {
-  // CHECK: %[[ZERO:.+]] = spv.constant dense<0> : vector<4xi32>
-  // CHECK: %[[ONE:.+]] = spv.constant dense<1> : vector<4xi32>
+  // CHECK: %[[ZERO:.+]] = spv.Constant dense<0> : vector<4xi32>
+  // CHECK: %[[ONE:.+]] = spv.Constant dense<1> : vector<4xi32>
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xi32>
   %0 = std.zexti %arg0 : vector<4xi1> to vector<4xi32>
   return %0 : vector<4xi32>
@@ -723,8 +723,8 @@ func @zexti4(%arg0 : vector<4xi1>) -> vector<4xi32> {
 
 // CHECK-LABEL: @zexti5
 func @zexti5(%arg0 : vector<4xi1>) -> vector<4xi64> {
-  // CHECK: %[[ZERO:.+]] = spv.constant dense<0> : vector<4xi64>
-  // CHECK: %[[ONE:.+]] = spv.constant dense<1> : vector<4xi64>
+  // CHECK: %[[ZERO:.+]] = spv.Constant dense<0> : vector<4xi64>
+  // CHECK: %[[ONE:.+]] = spv.Constant dense<1> : vector<4xi64>
   // CHECK: spv.Select %{{.*}}, %[[ONE]], %[[ZERO]] : vector<4xi1>, vector<4xi64>
   %0 = std.zexti %arg0 : vector<4xi1> to vector<4xi64>
   return %0 : vector<4xi64>
@@ -746,11 +746,11 @@ func @trunci2(%arg0: i32) -> i16 {
 
 // CHECK-LABEL: @trunc_to_i1
 func @trunc_to_i1(%arg0: i32) -> i1 {
-  // CHECK: %[[MASK:.*]] = spv.constant 1 : i32
+  // CHECK: %[[MASK:.*]] = spv.Constant 1 : i32
   // CHECK: %[[MASKED_SRC:.*]] = spv.BitwiseAnd %{{.*}}, %[[MASK]] : i32
   // CHECK: %[[IS_ONE:.*]] = spv.IEqual %[[MASKED_SRC]], %[[MASK]] : i32
-  // CHECK-DAG: %[[TRUE:.*]] = spv.constant true
-  // CHECK-DAG: %[[FALSE:.*]] = spv.constant false
+  // CHECK-DAG: %[[TRUE:.*]] = spv.Constant true
+  // CHECK-DAG: %[[FALSE:.*]] = spv.Constant false
   // CHECK: spv.Select %[[IS_ONE]], %[[TRUE]], %[[FALSE]] : i1, i1
   %0 = std.trunci %arg0 : i32 to i1
   return %0 : i1
@@ -758,11 +758,11 @@ func @trunc_to_i1(%arg0: i32) -> i1 {
 
 // CHECK-LABEL: @trunc_to_veci1
 func @trunc_to_veci1(%arg0: vector<4xi32>) -> vector<4xi1> {
-  // CHECK: %[[MASK:.*]] = spv.constant dense<1> : vector<4xi32>
+  // CHECK: %[[MASK:.*]] = spv.Constant dense<1> : vector<4xi32>
   // CHECK: %[[MASKED_SRC:.*]] = spv.BitwiseAnd %{{.*}}, %[[MASK]] : vector<4xi32>
   // CHECK: %[[IS_ONE:.*]] = spv.IEqual %[[MASKED_SRC]], %[[MASK]] : vector<4xi32>
-  // CHECK-DAG: %[[TRUE:.*]] = spv.constant dense<true> : vector<4xi1>
-  // CHECK-DAG: %[[FALSE:.*]] = spv.constant dense<false> : vector<4xi1>
+  // CHECK-DAG: %[[TRUE:.*]] = spv.Constant dense<true> : vector<4xi1>
+  // CHECK-DAG: %[[FALSE:.*]] = spv.Constant dense<false> : vector<4xi1>
   // CHECK: spv.Select %[[IS_ONE]], %[[TRUE]], %[[FALSE]] : vector<4xi1>, vector<4xi1>
   %0 = std.trunci %arg0 : vector<4xi32> to vector<4xi1>
   return %0 : vector<4xi1>
@@ -871,13 +871,13 @@ func @select(%arg0 : i32, %arg1 : i32) {
 // CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>,
 // CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>)
 func @load_store_zero_rank_float(%arg0: memref<f32>, %arg1: memref<f32>) {
-  //      CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
+  //      CHECK: [[ZERO1:%.*]] = spv.Constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG0]][
   // CHECK-SAME: [[ZERO1]], [[ZERO1]]
   // CHECK-SAME: ] :
   //      CHECK: spv.Load "StorageBuffer" %{{.*}} : f32
   %0 = load %arg0[] : memref<f32>
-  //      CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32
+  //      CHECK: [[ZERO2:%.*]] = spv.Constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG1]][
   // CHECK-SAME: [[ZERO2]], [[ZERO2]]
   // CHECK-SAME: ] :
@@ -890,13 +890,13 @@ func @load_store_zero_rank_float(%arg0: memref<f32>, %arg1: memref<f32>) {
 // CHECK: [[ARG0:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>,
 // CHECK: [[ARG1:%.*]]: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>)
 func @load_store_zero_rank_int(%arg0: memref<i32>, %arg1: memref<i32>) {
-  //      CHECK: [[ZERO1:%.*]] = spv.constant 0 : i32
+  //      CHECK: [[ZERO1:%.*]] = spv.Constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG0]][
   // CHECK-SAME: [[ZERO1]], [[ZERO1]]
   // CHECK-SAME: ] :
   //      CHECK: spv.Load "StorageBuffer" %{{.*}} : i32
   %0 = load %arg0[] : memref<i32>
-  //      CHECK: [[ZERO2:%.*]] = spv.constant 0 : i32
+  //      CHECK: [[ZERO2:%.*]] = spv.Constant 0 : i32
   //      CHECK: spv.AccessChain [[ARG1]][
   // CHECK-SAME: [[ZERO2]], [[ZERO2]]
   // CHECK-SAME: ] :
@@ -919,19 +919,19 @@ module attributes {
 
 // CHECK-LABEL: @load_i8
 func @load_i8(%arg0: memref<i8>) {
-  //     CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[FOUR1:.+]] = spv.constant 4 : i32
+  //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[FOUR1:.+]] = spv.Constant 4 : i32
   //     CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[ZERO]], %[[FOUR1]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]
   //     CHECK: %[[LOAD:.+]] = spv.Load  "StorageBuffer" %[[PTR]]
-  //     CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
-  //     CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+  //     CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
+  //     CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
   //     CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR2]] : i32
   //     CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
   //     CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32
-  //     CHECK: %[[MASK:.+]] = spv.constant 255 : i32
+  //     CHECK: %[[MASK:.+]] = spv.Constant 255 : i32
   //     CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32
-  //     CHECK: %[[T2:.+]] = spv.constant 24 : i32
+  //     CHECK: %[[T2:.+]] = spv.Constant 24 : i32
   //     CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32
   //     CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32
   %0 = load %arg0[] : memref<i8>
@@ -941,23 +941,23 @@ func @load_i8(%arg0: memref<i8>) {
 // CHECK-LABEL: @load_i16
 //       CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32)
 func @load_i16(%arg0: memref<10xi16>, %index : index) {
-  //     CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[OFFSET:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[ONE:.+]] = spv.constant 1 : i32
+  //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[OFFSET:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[ONE:.+]] = spv.Constant 1 : i32
   //     CHECK: %[[UPDATE:.+]] = spv.IMul %[[ONE]], %[[ARG1]] : i32
   //     CHECK: %[[FLAT_IDX:.+]] = spv.IAdd %[[OFFSET]], %[[UPDATE]] : i32
-  //     CHECK: %[[TWO1:.+]] = spv.constant 2 : i32
+  //     CHECK: %[[TWO1:.+]] = spv.Constant 2 : i32
   //     CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[FLAT_IDX]], %[[TWO1]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]
   //     CHECK: %[[LOAD:.+]] = spv.Load  "StorageBuffer" %[[PTR]]
-  //     CHECK: %[[TWO2:.+]] = spv.constant 2 : i32
-  //     CHECK: %[[SIXTEEN:.+]] = spv.constant 16 : i32
+  //     CHECK: %[[TWO2:.+]] = spv.Constant 2 : i32
+  //     CHECK: %[[SIXTEEN:.+]] = spv.Constant 16 : i32
   //     CHECK: %[[IDX:.+]] = spv.UMod %[[FLAT_IDX]], %[[TWO2]] : i32
   //     CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[SIXTEEN]] : i32
   //     CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32
-  //     CHECK: %[[MASK:.+]] = spv.constant 65535 : i32
+  //     CHECK: %[[MASK:.+]] = spv.Constant 65535 : i32
   //     CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32
-  //     CHECK: %[[T2:.+]] = spv.constant 16 : i32
+  //     CHECK: %[[T2:.+]] = spv.Constant 16 : i32
   //     CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32
   //     CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32
   %0 = load %arg0[%index] : memref<10xi16>
@@ -985,17 +985,17 @@ func @load_f32(%arg0: memref<f32>) {
 // CHECK-LABEL: @store_i8
 //       CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32)
 func @store_i8(%arg0: memref<i8>, %value: i8) {
-  //     CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[FOUR:.+]] = spv.constant 4 : i32
-  //     CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+  //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[FOUR:.+]] = spv.Constant 4 : i32
+  //     CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
   //     CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR]] : i32
   //     CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
-  //     CHECK: %[[MASK1:.+]] = spv.constant 255 : i32
+  //     CHECK: %[[MASK1:.+]] = spv.Constant 255 : i32
   //     CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32
   //     CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32
   //     CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG1]], %[[MASK1]] : i32
   //     CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32
-  //     CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
+  //     CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
   //     CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[ZERO]], %[[FOUR2]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]]
   //     CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]]
@@ -1007,21 +1007,21 @@ func @store_i8(%arg0: memref<i8>, %value: i8) {
 // CHECK-LABEL: @store_i16
 //       CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32, %[[ARG2:.+]]: i32)
 func @store_i16(%arg0: memref<10xi16>, %index: index, %value: i16) {
-  //     CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[OFFSET:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[ONE:.+]] = spv.constant 1 : i32
+  //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[OFFSET:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[ONE:.+]] = spv.Constant 1 : i32
   //     CHECK: %[[UPDATE:.+]] = spv.IMul %[[ONE]], %[[ARG1]] : i32
   //     CHECK: %[[FLAT_IDX:.+]] = spv.IAdd %[[OFFSET]], %[[UPDATE]] : i32
-  //     CHECK: %[[TWO:.+]] = spv.constant 2 : i32
-  //     CHECK: %[[SIXTEEN:.+]] = spv.constant 16 : i32
+  //     CHECK: %[[TWO:.+]] = spv.Constant 2 : i32
+  //     CHECK: %[[SIXTEEN:.+]] = spv.Constant 16 : i32
   //     CHECK: %[[IDX:.+]] = spv.UMod %[[FLAT_IDX]], %[[TWO]] : i32
   //     CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[SIXTEEN]] : i32
-  //     CHECK: %[[MASK1:.+]] = spv.constant 65535 : i32
+  //     CHECK: %[[MASK1:.+]] = spv.Constant 65535 : i32
   //     CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32
   //     CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32
   //     CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG2]], %[[MASK1]] : i32
   //     CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32
-  //     CHECK: %[[TWO2:.+]] = spv.constant 2 : i32
+  //     CHECK: %[[TWO2:.+]] = spv.Constant 2 : i32
   //     CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[FLAT_IDX]], %[[TWO2]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]]
   //     CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]]
@@ -1062,19 +1062,19 @@ module attributes {
 
 // CHECK-LABEL: @load_i8
 func @load_i8(%arg0: memref<i8>) {
-  //     CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[FOUR1:.+]] = spv.constant 4 : i32
+  //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[FOUR1:.+]] = spv.Constant 4 : i32
   //     CHECK: %[[QUOTIENT:.+]] = spv.SDiv %[[ZERO]], %[[FOUR1]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %{{.+}}[%[[ZERO]], %[[QUOTIENT]]]
   //     CHECK: %[[LOAD:.+]] = spv.Load  "StorageBuffer" %[[PTR]]
-  //     CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
-  //     CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+  //     CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
+  //     CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
   //     CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR2]] : i32
   //     CHECK: %[[BITS:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
   //     CHECK: %[[VALUE:.+]] = spv.ShiftRightArithmetic %[[LOAD]], %[[BITS]] : i32, i32
-  //     CHECK: %[[MASK:.+]] = spv.constant 255 : i32
+  //     CHECK: %[[MASK:.+]] = spv.Constant 255 : i32
   //     CHECK: %[[T1:.+]] = spv.BitwiseAnd %[[VALUE]], %[[MASK]] : i32
-  //     CHECK: %[[T2:.+]] = spv.constant 24 : i32
+  //     CHECK: %[[T2:.+]] = spv.Constant 24 : i32
   //     CHECK: %[[T3:.+]] = spv.ShiftLeftLogical %[[T1]], %[[T2]] : i32, i32
   //     CHECK: spv.ShiftRightArithmetic %[[T3]], %[[T2]] : i32, i32
   %0 = load %arg0[] : memref<i8>
@@ -1093,17 +1093,17 @@ func @load_i16(%arg0: memref<i16>) {
 // CHECK-LABEL: @store_i8
 //       CHECK: (%[[ARG0:.+]]: {{.*}}, %[[ARG1:.+]]: i32)
 func @store_i8(%arg0: memref<i8>, %value: i8) {
-  //     CHECK: %[[ZERO:.+]] = spv.constant 0 : i32
-  //     CHECK: %[[FOUR:.+]] = spv.constant 4 : i32
-  //     CHECK: %[[EIGHT:.+]] = spv.constant 8 : i32
+  //     CHECK: %[[ZERO:.+]] = spv.Constant 0 : i32
+  //     CHECK: %[[FOUR:.+]] = spv.Constant 4 : i32
+  //     CHECK: %[[EIGHT:.+]] = spv.Constant 8 : i32
   //     CHECK: %[[IDX:.+]] = spv.UMod %[[ZERO]], %[[FOUR]] : i32
   //     CHECK: %[[OFFSET:.+]] = spv.IMul %[[IDX]], %[[EIGHT]] : i32
-  //     CHECK: %[[MASK1:.+]] = spv.constant 255 : i32
+  //     CHECK: %[[MASK1:.+]] = spv.Constant 255 : i32
   //     CHECK: %[[TMP1:.+]] = spv.ShiftLeftLogical %[[MASK1]], %[[OFFSET]] : i32, i32
   //     CHECK: %[[MASK:.+]] = spv.Not %[[TMP1]] : i32
   //     CHECK: %[[CLAMPED_VAL:.+]] = spv.BitwiseAnd %[[ARG1]], %[[MASK1]] : i32
   //     CHECK: %[[STORE_VAL:.+]] = spv.ShiftLeftLogical %[[CLAMPED_VAL]], %[[OFFSET]] : i32, i32
-  //     CHECK: %[[FOUR2:.+]] = spv.constant 4 : i32
+  //     CHECK: %[[FOUR2:.+]] = spv.Constant 4 : i32
   //     CHECK: %[[ACCESS_IDX:.+]] = spv.SDiv %[[ZERO]], %[[FOUR2]] : i32
   //     CHECK: %[[PTR:.+]] = spv.AccessChain %[[ARG0]][%[[ZERO]], %[[ACCESS_IDX]]]
   //     CHECK: spv.AtomicAnd "Device" "AcquireRelease" %[[PTR]], %[[MASK]]
index 08b7fd2..05837df 100644 (file)
@@ -113,7 +113,7 @@ func @composite_extract_no_ssa_operand() -> () {
 // -----
 
 func @composite_extract_invalid_index_type_1() -> () {
-  %0 = spv.constant 10 : i32
+  %0 = spv.Constant 10 : i32
   %1 = spv.Variable : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>
   %2 = spv.Load "Function" %1 ["Volatile"] : !spv.array<4x!spv.array<4xf32>>
   // expected-error @+1 {{expected non-function type}}
index 61d27c5..cc6a0b2 100644 (file)
@@ -14,7 +14,7 @@ func @branch() -> () {
 // -----
 
 func @branch_argument() -> () {
-  %zero = spv.constant 0 : i32
+  %zero = spv.Constant 0 : i32
   // CHECK: spv.Branch ^bb1(%{{.*}}, %{{.*}} : i32, i32)
   spv.Branch ^next(%zero, %zero: i32, i32)
 ^next(%arg0: i32, %arg1: i32):
@@ -31,7 +31,7 @@ func @missing_accessor() -> () {
 // -----
 
 func @wrong_accessor_count() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   // expected-error @+1 {{requires 1 successor but found 2}}
   "spv.Branch"()[^one, ^two] : () -> ()
 ^one:
@@ -47,7 +47,7 @@ func @wrong_accessor_count() -> () {
 //===----------------------------------------------------------------------===//
 
 func @cond_branch() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   // CHECK: spv.BranchConditional %{{.*}}, ^bb1, ^bb2
   spv.BranchConditional %true, ^one, ^two
 // CHECK: ^bb1
@@ -61,8 +61,8 @@ func @cond_branch() -> () {
 // -----
 
 func @cond_branch_argument() -> () {
-  %true = spv.constant true
-  %zero = spv.constant 0 : i32
+  %true = spv.Constant true
+  %zero = spv.Constant 0 : i32
   // CHECK: spv.BranchConditional %{{.*}}, ^bb1(%{{.*}}, %{{.*}} : i32, i32), ^bb2
   spv.BranchConditional %true, ^true1(%zero, %zero: i32, i32), ^false1
 ^true1(%arg0: i32, %arg1: i32):
@@ -79,7 +79,7 @@ func @cond_branch_argument() -> () {
 // -----
 
 func @cond_branch_with_weights() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   // CHECK: spv.BranchConditional %{{.*}} [5, 10]
   spv.BranchConditional %true [5, 10], ^one, ^two
 ^one:
@@ -103,7 +103,7 @@ func @missing_condition() -> () {
 
 func @wrong_condition_type() -> () {
   // expected-note @+1 {{prior use here}}
-  %zero = spv.constant 0 : i32
+  %zero = spv.Constant 0 : i32
   // expected-error @+1 {{use of value '%zero' expects different type than prior uses: 'i1' vs 'i32'}}
   spv.BranchConditional %zero, ^one, ^two
 ^one:
@@ -115,7 +115,7 @@ func @wrong_condition_type() -> () {
 // -----
 
 func @wrong_accessor_count() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   // expected-error @+1 {{requires 2 successors but found 1}}
   "spv.BranchConditional"(%true)[^one] {operand_segment_sizes = dense<[1, 0, 0]>: vector<3xi32>} : (i1) -> ()
 ^one:
@@ -127,7 +127,7 @@ func @wrong_accessor_count() -> () {
 // -----
 
 func @wrong_number_of_weights() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   // expected-error @+1 {{must have exactly two branch weights}}
   "spv.BranchConditional"(%true)[^one, ^two] {branch_weights = [1 : i32, 2 : i32, 3 : i32],
                                               operand_segment_sizes = dense<[1, 0, 0]>: vector<3xi32>} : (i1) -> ()
@@ -140,7 +140,7 @@ func @wrong_number_of_weights() -> () {
 // -----
 
 func @weights_cannot_both_be_zero() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   // expected-error @+1 {{branch weights cannot both be zero}}
   spv.BranchConditional %true [0, 0], ^one, ^two
 ^one:
@@ -232,7 +232,7 @@ spv.module Logical GLSL450 {
 
 spv.module Logical GLSL450 {
   spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> () "None" {
-    %0 = spv.constant 2.0 : f32
+    %0 = spv.Constant 2.0 : f32
     // expected-error @+1 {{operand type mismatch: expected operand type 'i32', but provided 'f32' for operand number 1}}
     spv.FunctionCall @f_type_mismatch(%arg0, %0) : (i32, f32) -> ()
     spv.Return
@@ -243,7 +243,7 @@ spv.module Logical GLSL450 {
 
 spv.module Logical GLSL450 {
   spv.func @f_type_mismatch(%arg0 : i32, %arg1 : i32) -> i32 "None" {
-    %cst = spv.constant 0: i32
+    %cst = spv.Constant 0: i32
     // expected-error @+1 {{result type mismatch: expected 'i32', but provided 'f32'}}
     %0 = spv.FunctionCall @f_type_mismatch(%arg0, %arg0) : (i32, i32) -> f32
     spv.ReturnValue %cst: i32
@@ -268,8 +268,8 @@ spv.module Logical GLSL450 {
 
 // for (int i = 0; i < count; ++i) {}
 func @loop(%count : i32) -> () {
-  %zero = spv.constant 0: i32
-  %one = spv.constant 1: i32
+  %zero = spv.Constant 0: i32
+  %one = spv.Constant 1: i32
   %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
   // CHECK: spv.loop {
@@ -438,8 +438,8 @@ func @merge() -> () {
 // -----
 
 func @only_allowed_in_last_block(%cond : i1) -> () {
-  %zero = spv.constant 0: i32
-  %one = spv.constant 1: i32
+  %zero = spv.Constant 0: i32
+  %one = spv.Constant 1: i32
   %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
   spv.selection {
@@ -460,7 +460,7 @@ func @only_allowed_in_last_block(%cond : i1) -> () {
 // -----
 
 func @only_allowed_in_last_block() -> () {
-  %true = spv.constant true
+  %true = spv.Constant true
   spv.loop {
     spv.Branch ^header
   ^header:
@@ -548,7 +548,7 @@ spv.module Logical GLSL450 {
       spv.mlir.merge
     }
 
-    %zero = spv.constant 0: i32
+    %zero = spv.Constant 0: i32
     spv.ReturnValue %zero: i32
   }
 }
@@ -560,7 +560,7 @@ spv.module Logical GLSL450 {
 //===----------------------------------------------------------------------===//
 
 func @ret_val() -> (i32) {
-  %0 = spv.constant 42 : i32
+  %0 = spv.Constant 42 : i32
   // CHECK: spv.ReturnValue %{{.*}} : i32
   spv.ReturnValue %0 : i32
 }
@@ -570,13 +570,13 @@ func @in_selection(%cond : i1) -> (i32) {
   spv.selection {
     spv.BranchConditional %cond, ^then, ^merge
   ^then:
-    %zero = spv.constant 0 : i32
+    %zero = spv.Constant 0 : i32
     // CHECK: spv.ReturnValue
     spv.ReturnValue %zero : i32
   ^merge:
     spv.mlir.merge
   }
-  %one = spv.constant 1 : i32
+  %one = spv.Constant 1 : i32
   spv.ReturnValue %one : i32
 }
 
@@ -587,7 +587,7 @@ func @in_loop(%cond : i1) -> (i32) {
   ^header:
     spv.BranchConditional %cond, ^body, ^merge
   ^body:
-    %zero = spv.constant 0 : i32
+    %zero = spv.Constant 0 : i32
     // CHECK: spv.ReturnValue
     spv.ReturnValue %zero : i32
   ^continue:
@@ -595,7 +595,7 @@ func @in_loop(%cond : i1) -> (i32) {
   ^merge:
     spv.mlir.merge
   }
-  %one = spv.constant 1 : i32
+  %one = spv.Constant 1 : i32
   spv.ReturnValue %one : i32
 }
 
@@ -608,7 +608,7 @@ func @in_other_func_like_op(%arg: i32) -> i32 {
 // -----
 
 "foo.function"() ({
-  %0 = spv.constant true
+  %0 = spv.Constant true
   // expected-error @+1 {{op must appear in a function-like op's block}}
   spv.ReturnValue %0 : i1
 })  : () -> ()
@@ -617,7 +617,7 @@ func @in_other_func_like_op(%arg: i32) -> i32 {
 
 spv.module Logical GLSL450 {
   spv.func @value_count_mismatch() -> () "None" {
-    %0 = spv.constant 42 : i32
+    %0 = spv.Constant 42 : i32
     // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
     spv.ReturnValue %0 : i32
   }
@@ -627,7 +627,7 @@ spv.module Logical GLSL450 {
 
 spv.module Logical GLSL450 {
   spv.func @value_type_mismatch() -> (f32) "None" {
-    %0 = spv.constant 42 : i32
+    %0 = spv.Constant 42 : i32
     // expected-error @+1 {{return value's type ('i32') mismatch with function's result type ('f32')}}
     spv.ReturnValue %0 : i32
   }
@@ -640,7 +640,7 @@ spv.module Logical GLSL450 {
     spv.selection {
       spv.BranchConditional %cond, ^then, ^merge
     ^then:
-      %cst = spv.constant 0: i32
+      %cst = spv.Constant 0: i32
       // expected-error @+1 {{op returns 1 value but enclosing function requires 0 results}}
       spv.ReturnValue %cst: i32
     ^merge:
@@ -658,8 +658,8 @@ spv.module Logical GLSL450 {
 //===----------------------------------------------------------------------===//
 
 func @selection(%cond: i1) -> () {
-  %zero = spv.constant 0: i32
-  %one = spv.constant 1: i32
+  %zero = spv.Constant 0: i32
+  %one = spv.Constant 1: i32
   %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
   // CHECK: spv.selection {
@@ -685,9 +685,9 @@ func @selection(%cond: i1) -> () {
 // -----
 
 func @selection(%cond: i1) -> () {
-  %zero = spv.constant 0: i32
-  %one = spv.constant 1: i32
-  %two = spv.constant 2: i32
+  %zero = spv.Constant 0: i32
+  %one = spv.Constant 1: i32
+  %two = spv.Constant 2: i32
   %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
   // CHECK: spv.selection {
index 7cd0b63..7de8a36 100644 (file)
@@ -103,7 +103,7 @@ spv.func @cooperative_matrix_fdiv(%a : !spv.coopmatrix<8x16xf32, Subgroup>, %b :
 
 // CHECK-LABEL: @cooperative_matrix_access_chain
 spv.func @cooperative_matrix_access_chain(%a : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>) -> !spv.ptr<f32, Function> "None" {
-  %0 = spv.constant 0: i32
+  %0 = spv.Constant 0: i32
   // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
   %1 = spv.AccessChain %a[%0] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
   spv.ReturnValue %1 : !spv.ptr<f32, Function>
index b2c34b8..25d9141 100644 (file)
@@ -178,24 +178,24 @@ func @logicalUnary(%arg0 : i32)
 //===----------------------------------------------------------------------===//
 
 func @select_op_bool(%arg0: i1) -> () {
-  %0 = spv.constant true
-  %1 = spv.constant false
+  %0 = spv.Constant true
+  %1 = spv.Constant false
   // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i1
   %2 = spv.Select %arg0, %0, %1 : i1, i1
   return
 }
 
 func @select_op_int(%arg0: i1) -> () {
-  %0 = spv.constant 2 : i32
-  %1 = spv.constant 3 : i32
+  %0 = spv.Constant 2 : i32
+  %1 = spv.Constant 3 : i32
   // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, i32
   %2 = spv.Select %arg0, %0, %1 : i1, i32
   return
 }
 
 func @select_op_float(%arg0: i1) -> () {
-  %0 = spv.constant 2.0 : f32
-  %1 = spv.constant 3.0 : f32
+  %0 = spv.Constant 2.0 : f32
+  %1 = spv.Constant 3.0 : f32
   // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, f32
   %2 = spv.Select %arg0, %0, %1 : i1, f32
   return
@@ -210,16 +210,16 @@ func @select_op_ptr(%arg0: i1) -> () {
 }
 
 func @select_op_vec(%arg0: i1) -> () {
-  %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
-  %1 = spv.constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
+  %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+  %1 = spv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
   // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : i1, vector<3xf32>
   %2 = spv.Select %arg0, %0, %1 : i1, vector<3xf32>
   return
 }
 
 func @select_op_vec_condn_vec(%arg0: vector<3xi1>) -> () {
-  %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
-  %1 = spv.constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
+  %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+  %1 = spv.Constant dense<[5.0, 6.0, 7.0]> : vector<3xf32>
   // CHECK : spv.Select {{%.*}}, {{%.*}}, {{%.*}} : vector<3xi1>, vector<3xf32>
   %2 = spv.Select %arg0, %0, %1 : vector<3xi1>, vector<3xf32>
   return
@@ -228,8 +228,8 @@ func @select_op_vec_condn_vec(%arg0: vector<3xi1>) -> () {
 // -----
 
 func @select_op(%arg0: i1) -> () {
-  %0 = spv.constant 2 : i32
-  %1 = spv.constant 3 : i32
+  %0 = spv.Constant 2 : i32
+  %1 = spv.Constant 3 : i32
   // expected-error @+2 {{expected ','}}
   %2 = spv.Select %arg0, %0, %1 : i1
   return
@@ -238,8 +238,8 @@ func @select_op(%arg0: i1) -> () {
 // -----
 
 func @select_op(%arg1: vector<3xi1>) -> () {
-  %0 = spv.constant 2 : i32
-  %1 = spv.constant 3 : i32
+  %0 = spv.Constant 2 : i32
+  %1 = spv.Constant 3 : i32
   // expected-error @+1 {{result expected to be of vector type when condition is of vector type}}
   %2 = spv.Select %arg1, %0, %1 : vector<3xi1>, i32
   return
@@ -248,8 +248,8 @@ func @select_op(%arg1: vector<3xi1>) -> () {
 // -----
 
 func @select_op(%arg1: vector<4xi1>) -> () {
-  %0 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32>
+  %0 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32>
   // expected-error @+1 {{result should have the same number of elements as the condition when condition is of vector type}}
   %2 = spv.Select %arg1, %0, %1 : vector<4xi1>, vector<3xi32>
   return
@@ -258,8 +258,8 @@ func @select_op(%arg1: vector<4xi1>) -> () {
 // -----
 
 func @select_op(%arg1: vector<4xi1>) -> () {
-  %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
-  %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32>
+  %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+  %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32>
   // expected-error @+1 {{all of {true_value, false_value, result} have same type}}
   %2 = "spv.Select"(%arg1, %0, %1) : (vector<4xi1>, vector<3xf32>, vector<3xi32>) -> vector<3xi32>
   return
@@ -268,8 +268,8 @@ func @select_op(%arg1: vector<4xi1>) -> () {
 // -----
 
 func @select_op(%arg1: vector<4xi1>) -> () {
-  %0 = spv.constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
-  %1 = spv.constant dense<[5, 6, 7]> : vector<3xi32>
+  %0 = spv.Constant dense<[2.0, 3.0, 4.0]> : vector<3xf32>
+  %1 = spv.Constant dense<[5, 6, 7]> : vector<3xi32>
   // expected-error @+1 {{all of {true_value, false_value, result} have same type}}
   %2 = "spv.Select"(%arg1, %1, %0) : (vector<4xi1>, vector<3xi32>, vector<3xf32>) -> vector<3xi32>
   return
index 5e82cf9..ba76516 100644 (file)
@@ -5,7 +5,7 @@
 //===----------------------------------------------------------------------===//
 
 func @access_chain_struct() -> () {
-  %0 = spv.constant 1: i32
+  %0 = spv.Constant 1: i32
   %1 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
   // CHECK: spv.AccessChain {{.*}}[{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Function>
   %2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
@@ -46,7 +46,7 @@ func @access_chain_rtarray(%arg0 : i32) -> () {
 // -----
 
 func @access_chain_non_composite() -> () {
-  %0 = spv.constant 1: i32
+  %0 = spv.Constant 1: i32
   %1 = spv.Variable : !spv.ptr<f32, Function>
   // expected-error @+1 {{cannot extract from non-composite type 'f32' with index 0}}
   %2 = spv.AccessChain %1[%0] : !spv.ptr<f32, Function>, i32
@@ -112,7 +112,7 @@ func @access_chain_invalid_index_1(%index0 : i32) -> () {
 
 func @access_chain_invalid_index_2(%index0 : i32) -> () {
   %0 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
-  // expected-error @+1 {{index must be an integer spv.constant to access element of spv.struct}}
+  // expected-error @+1 {{index must be an integer spv.Constant to access element of spv.struct}}
   %1 = spv.AccessChain %0[%index0, %index0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
   return
 }
@@ -122,7 +122,7 @@ func @access_chain_invalid_index_2(%index0 : i32) -> () {
 func @access_chain_invalid_constant_type_1() -> () {
   %0 = std.constant 1: i32
   %1 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
-  // expected-error @+1 {{index must be an integer spv.constant to access element of spv.struct, but provided std.constant}}
+  // expected-error @+1 {{index must be an integer spv.Constant to access element of spv.struct, but provided std.constant}}
   %2 = spv.AccessChain %1[%0, %0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
   return
 }
@@ -130,7 +130,7 @@ func @access_chain_invalid_constant_type_1() -> () {
 // -----
 
 func @access_chain_out_of_bounds() -> () {
-  %index0 = "spv.constant"() { value = 12: i32} : () -> i32
+  %index0 = "spv.Constant"() { value = 12: i32} : () -> i32
   %0 = spv.Variable : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>
   // expected-error @+1 {{'spv.AccessChain' op index 12 out of bounds for '!spv.struct<(f32, !spv.array<4 x f32>)>'}}
   %1 = spv.AccessChain %0[%index0, %index0] : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Function>, i32, i32
@@ -487,7 +487,7 @@ func @variable(%arg0: f32) -> () {
 // -----
 
 func @variable_init_normal_constant() -> () {
-  %0 = spv.constant 4.0 : f32
+  %0 = spv.Constant 4.0 : f32
   // CHECK: spv.Variable init(%0) : !spv.ptr<f32, Function>
   %1 = spv.Variable init(%0) : !spv.ptr<f32, Function>
   return
@@ -529,7 +529,7 @@ func @variable_bind() -> () {
 // -----
 
 func @variable_init_bind() -> () {
-  %0 = spv.constant 4.0 : f32
+  %0 = spv.Constant 4.0 : f32
   // expected-error @+1 {{cannot have 'binding' attribute (only allowed in spv.globalVariable)}}
   %1 = spv.Variable init(%0) {binding = 5 : i32} : !spv.ptr<f32, Function>
   return
index 61febd3..602d5d2 100644 (file)
@@ -33,7 +33,7 @@ func @group_non_uniform_ballot(%predicate: i1) -> vector<4xsi32> {
 //===----------------------------------------------------------------------===//
 
 func @group_non_uniform_broadcast_scalar(%value: f32) -> f32 {
-  %one = spv.constant 1 : i32
+  %one = spv.Constant 1 : i32
   // CHECK: spv.GroupNonUniformBroadcast Workgroup %{{.*}}, %{{.*}} : f32, i32
   %0 = spv.GroupNonUniformBroadcast Workgroup %value, %one : f32, i32
   return %0: f32
@@ -42,7 +42,7 @@ func @group_non_uniform_broadcast_scalar(%value: f32) -> f32 {
 // -----
 
 func @group_non_uniform_broadcast_vector(%value: vector<4xf32>) -> vector<4xf32> {
-  %one = spv.constant 1 : i32
+  %one = spv.Constant 1 : i32
   // CHECK: spv.GroupNonUniformBroadcast Subgroup %{{.*}}, %{{.*}} : vector<4xf32>, i32
   %0 = spv.GroupNonUniformBroadcast Subgroup %value, %one : vector<4xf32>, i32
   return %0: vector<4xf32>
@@ -51,7 +51,7 @@ func @group_non_uniform_broadcast_vector(%value: vector<4xf32>) -> vector<4xf32>
 // -----
 
 func @group_non_uniform_broadcast_negative_scope(%value: f32, %localid: i32 ) -> f32 {
-  %one = spv.constant 1 : i32
+  %one = spv.Constant 1 : i32
   // expected-error @+1 {{execution scope must be 'Workgroup' or 'Subgroup'}} 
   %0 = spv.GroupNonUniformBroadcast Device %value, %one : f32, i32
   return %0: f32
@@ -101,7 +101,7 @@ func @group_non_uniform_fadd_reduce(%val: f32) -> f32 {
 
 // CHECK-LABEL: @group_non_uniform_fadd_clustered_reduce
 func @group_non_uniform_fadd_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> {
-  %four = spv.constant 4 : i32
+  %four = spv.Constant 4 : i32
   // CHECK: %{{.+}} = spv.GroupNonUniformFAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>
   %0 = spv.GroupNonUniformFAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xf32>
   return %0: vector<2xf32>
@@ -120,7 +120,7 @@ func @group_non_uniform_fmul_reduce(%val: f32) -> f32 {
 
 // CHECK-LABEL: @group_non_uniform_fmul_clustered_reduce
 func @group_non_uniform_fmul_clustered_reduce(%val: vector<2xf32>) -> vector<2xf32> {
-  %four = spv.constant 4 : i32
+  %four = spv.Constant 4 : i32
   // CHECK: %{{.+}} = spv.GroupNonUniformFMul "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xf32>
   %0 = spv.GroupNonUniformFMul "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xf32>
   return %0: vector<2xf32>
@@ -167,7 +167,7 @@ func @group_non_uniform_iadd_reduce(%val: i32) -> i32 {
 
 // CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce
 func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
-  %four = spv.constant 4 : i32
+  %four = spv.Constant 4 : i32
   // CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
   %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
   return %0: vector<2xi32>
@@ -200,7 +200,7 @@ func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>, %size: i32) -
 // -----
 
 func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
-  %five = spv.constant 5 : i32
+  %five = spv.Constant 5 : i32
   // expected-error @+1 {{cluster size operand must be a power of two}}
   %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%five) : vector<2xi32>
   return %0: vector<2xi32>
@@ -221,7 +221,7 @@ func @group_non_uniform_imul_reduce(%val: i32) -> i32 {
 
 // CHECK-LABEL: @group_non_uniform_imul_clustered_reduce
 func @group_non_uniform_imul_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> {
-  %four = spv.constant 4 : i32
+  %four = spv.Constant 4 : i32
   // CHECK: %{{.+}} = spv.GroupNonUniformIMul "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
   %0 = spv.GroupNonUniformIMul "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
   return %0: vector<2xi32>
index 9b8e66b..dbabaed 100644 (file)
@@ -7,7 +7,7 @@
 spv.module Logical GLSL450 {
   spv.globalVariable @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
   spv.func @access_chain() -> () "None" {
-    %0 = spv.constant 1: i32
+    %0 = spv.Constant 1: i32
     // CHECK: [[VAR1:%.*]] = spv.mlir.addressof @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
     // CHECK-NEXT: spv.AccessChain [[VAR1]][{{.*}}, {{.*}}] : !spv.ptr<!spv.struct<(f32, !spv.array<4 x f32>)>, Input>
     %1 = spv.mlir.addressof @var1 : !spv.ptr<!spv.struct<(f32, !spv.array<4xf32>)>, Input>
@@ -49,29 +49,29 @@ spv.module Logical GLSL450 {
 // -----
 
 //===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
 //===----------------------------------------------------------------------===//
 
 func @const() -> () {
-  // CHECK: spv.constant true
-  // CHECK: spv.constant 42 : i32
-  // CHECK: spv.constant 5.000000e-01 : f32
-  // CHECK: spv.constant dense<[2, 3]> : vector<2xi32>
-  // CHECK: spv.constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>>
-  // CHECK: spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
-  // CHECK: spv.constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
-  // CHECK: spv.constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
-  // CHECK: spv.constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
-
-  %0 = spv.constant true
-  %1 = spv.constant 42 : i32
-  %2 = spv.constant 0.5 : f32
-  %3 = spv.constant dense<[2, 3]> : vector<2xi32>
-  %4 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
-  %5 = spv.constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
-  %6 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
-  %7 = spv.constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
-  %8 = spv.constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+  // CHECK: spv.Constant true
+  // CHECK: spv.Constant 42 : i32
+  // CHECK: spv.Constant 5.000000e-01 : f32
+  // CHECK: spv.Constant dense<[2, 3]> : vector<2xi32>
+  // CHECK: spv.Constant [dense<3.000000e+00> : vector<2xf32>] : !spv.array<1 x vector<2xf32>>
+  // CHECK: spv.Constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  // CHECK: spv.Constant dense<1.000000e+00> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+  // CHECK: spv.Constant dense<{{\[}}[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  // CHECK: spv.Constant dense<{{\[}}[1.000000e+00, 2.000000e+00, 3.000000e+00], [4.000000e+00, 5.000000e+00, 6.000000e+00]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+
+  %0 = spv.Constant true
+  %1 = spv.Constant 42 : i32
+  %2 = spv.Constant 0.5 : f32
+  %3 = spv.Constant dense<[2, 3]> : vector<2xi32>
+  %4 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.array<1xvector<2xf32>>
+  %5 = spv.Constant dense<1> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  %6 = spv.Constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
+  %7 = spv.Constant dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> : !spv.array<2 x !spv.array<3 x i32>>
+  %8 = spv.Constant dense<[[1.0, 2.0, 3.0], [4.0, 5.0, 6.0]]> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x f32>>
   return
 }
 
@@ -79,7 +79,7 @@ func @const() -> () {
 
 func @unaccepted_std_attr() -> () {
   // expected-error @+1 {{cannot have value of type 'none'}}
-  %0 = spv.constant unit : none
+  %0 = spv.Constant unit : none
   return
 }
 
@@ -87,7 +87,7 @@ func @unaccepted_std_attr() -> () {
 
 func @array_constant() -> () {
   // expected-error @+1 {{has array element whose type ('vector<2xi32>') does not match the result element type ('vector<2xf32>')}}
-  %0 = spv.constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spv.array<2xvector<2xf32>>
+  %0 = spv.Constant [dense<3.0> : vector<2xf32>, dense<4> : vector<2xi32>] : !spv.array<2xvector<2xf32>>
   return
 }
 
@@ -95,7 +95,7 @@ func @array_constant() -> () {
 
 func @array_constant() -> () {
   // expected-error @+1 {{must have spv.array result type for array value}}
-  %0 = spv.constant [dense<3.0> : vector<2xf32>] : !spv.rtarray<vector<2xf32>>
+  %0 = spv.Constant [dense<3.0> : vector<2xf32>] : !spv.rtarray<vector<2xf32>>
   return
 }
 
@@ -103,7 +103,7 @@ func @array_constant() -> () {
 
 func @non_nested_array_constant() -> () {
   // expected-error @+1 {{only support nested array result type}}
-  %0 = spv.constant dense<3.0> : tensor<2x2xf32> : !spv.array<2xvector<2xf32>>
+  %0 = spv.Constant dense<3.0> : tensor<2x2xf32> : !spv.array<2xvector<2xf32>>
   return
 }
 
@@ -111,21 +111,21 @@ func @non_nested_array_constant() -> () {
 
 func @value_result_type_mismatch() -> () {
   // expected-error @+1 {{must have spv.array result type for array value}}
-  %0 = "spv.constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>)
+  %0 = "spv.Constant"() {value = dense<0> : tensor<4xi32>} : () -> (vector<4xi32>)
 }
 
 // -----
 
 func @value_result_type_mismatch() -> () {
   // expected-error @+1 {{result element type ('i32') does not match value element type ('f32')}}
-  %0 = spv.constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>>
+  %0 = spv.Constant dense<1.0> : tensor<2x3xf32> : !spv.array<2 x !spv.array<3 x i32>>
 }
 
 // -----
 
 func @value_result_num_elements_mismatch() -> () {
   // expected-error @+1 {{result number of elements (6) does not match value number of elements (4)}}
-  %0 = spv.constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>>
+  %0 = spv.Constant dense<1.0> : tensor<2x2xf32> : !spv.array<2 x !spv.array<3 x f32>>
   return
 }
 
@@ -308,7 +308,7 @@ spv.module Logical GLSL450 {
 
 // TODO: Fix test case after initialization with normal constant is addressed
 // spv.module Logical GLSL450 {
-//   %0 = spv.constant 4.0 : f32
+//   %0 = spv.Constant 4.0 : f32
 //   // CHECK1: spv.Variable init(%0) : !spv.ptr<f32, Private>
 //   spv.globalVariable @var1 init(%0) : !spv.ptr<f32, Private>
 // }
@@ -337,7 +337,7 @@ spv.module Logical GLSL450 {
 
 // TODO: Fix test case after initialization with constant is addressed
 // spv.module Logical GLSL450 {
-//   %0 = spv.constant 4.0 : f32
+//   %0 = spv.Constant 4.0 : f32
 //   // CHECK1: spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 //   spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr<f32, Private>
 // }
@@ -480,7 +480,7 @@ spv.module Logical GLSL450 {
 // expected-error@+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.module Logical GLSL450 {
   spv.func @compute() -> f32 "None" {
     // CHECK: spv.mlir.referenceof @sc3 : f32
     %0 = spv.mlir.referenceof @sc3 : f32
-    %1 = spv.constant 6.0 : f32
+    %1 = spv.Constant 6.0 : f32
     %2 = spv.FAdd %0, %1 : f32
     spv.ReturnValue %2 : f32
   }
@@ -801,10 +801,10 @@ spv.module Logical GLSL450 {
 
 spv.module Logical GLSL450 {
   spv.func @foo() -> i32 "None" {
-    // CHECK: [[LHS:%.*]] = spv.constant
-    %0 = spv.constant 1: i32
-    // CHECK: [[RHS:%.*]] = spv.constant
-    %1 = spv.constant 1: i32
+    // CHECK: [[LHS:%.*]] = spv.Constant
+    %0 = spv.Constant 1: i32
+    // CHECK: [[RHS:%.*]] = spv.Constant
+    %1 = spv.Constant 1: i32
 
     // CHECK: spv.SpecConstantOperation wraps "spv.IAdd"([[LHS]], [[RHS]]) : (i32, i32) -> i32
     %2 = spv.SpecConstantOperation wraps "spv.IAdd"(%0, %1) : (i32, i32) -> i32
@@ -831,7 +831,7 @@ spv.module Logical GLSL450 {
 
 spv.module Logical GLSL450 {
   spv.func @foo() -> i32 "None" {
-    %0 = spv.constant 1: i32
+    %0 = spv.Constant 1: i32
     // expected-error @+1 {{op expects parent op 'spv.SpecConstantOperation'}}
     spv.mlir.yield %0 : i32
   }
index d9c7406..e0c70f7 100644 (file)
@@ -18,7 +18,7 @@ spv.module Logical GLSL450 {
   attributes {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ARG1:%.*]] = spv.mlir.addressof [[VAR1]]
     // CHECK: [[ADDRESSARG0:%.*]] = spv.mlir.addressof [[VAR0]]
-    // CHECK: [[CONST0:%.*]] = spv.constant 0 : i32
+    // CHECK: [[CONST0:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG0PTR:%.*]] = spv.AccessChain [[ADDRESSARG0]]{{\[}}[[CONST0]]
     // CHECK: [[ARG0:%.*]] = spv.Load "StorageBuffer" [[ARG0PTR]]
     // CHECK: spv.Return
index 7db0f88..c2a56d2 100644 (file)
@@ -40,19 +40,19 @@ spv.module Logical GLSL450 {
     {spv.interface_var_abi = #spv.interface_var_abi<(0, 6), StorageBuffer>}) "None"
   attributes  {spv.entry_point_abi = {local_size = dense<[32, 1, 1]> : vector<3xi32>}} {
     // CHECK: [[ADDRESSARG6:%.*]] = spv.mlir.addressof [[VAR6]]
-    // CHECK: [[CONST6:%.*]] = spv.constant 0 : i32
+    // CHECK: [[CONST6:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG6PTR:%.*]] = spv.AccessChain [[ADDRESSARG6]]{{\[}}[[CONST6]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG6PTR]]
     // CHECK: [[ADDRESSARG5:%.*]] = spv.mlir.addressof [[VAR5]]
-    // CHECK: [[CONST5:%.*]] = spv.constant 0 : i32
+    // CHECK: [[CONST5:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG5PTR:%.*]] = spv.AccessChain [[ADDRESSARG5]]{{\[}}[[CONST5]]
     // CHECK: {{%.*}} = spv.Load "StorageBuffer" [[ARG5PTR]]
     // CHECK: [[ADDRESSARG4:%.*]] = spv.mlir.addressof [[VAR4]]
-    // CHECK: [[CONST4:%.*]] = spv.constant 0 : i32
+    // CHECK: [[CONST4:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG4PTR:%.*]] = spv.AccessChain [[ADDRESSARG4]]{{\[}}[[CONST4]]
     // CHECK: [[ARG4:%.*]] = spv.Load "StorageBuffer" [[ARG4PTR]]
     // CHECK: [[ADDRESSARG3:%.*]] = spv.mlir.addressof [[VAR3]]
-    // CHECK: [[CONST3:%.*]] = spv.constant 0 : i32
+    // CHECK: [[CONST3:%.*]] = spv.Constant 0 : i32
     // CHECK: [[ARG3PTR:%.*]] = spv.AccessChain [[ADDRESSARG3]]{{\[}}[[CONST3]]
     // CHECK: [[ARG3:%.*]] = spv.Load "StorageBuffer" [[ARG3PTR]]
     // CHECK: [[ADDRESSARG2:%.*]] = spv.mlir.addressof [[VAR2]]
@@ -102,7 +102,7 @@ spv.module Logical GLSL450 {
     // CHECK: spv.IAdd [[ARG4]]
     %37 = spv.IAdd %arg4, %11 : i32
     // CHECK: spv.AccessChain [[ARG0]]
-    %c0 = spv.constant 0 : i32
+    %c0 = spv.Constant 0 : i32
     %38 = spv.AccessChain %arg0[%c0, %36, %37] : !spv.ptr<!spv.struct<(!spv.array<12 x !spv.array<4 x f32>>)>, StorageBuffer>, i32, i32, i32
     %39 = spv.Load "StorageBuffer" %38 : f32
     // CHECK: spv.AccessChain [[ARG1]]
index cb955b4..66c1ddc 100644 (file)
@@ -5,11 +5,11 @@
 //===----------------------------------------------------------------------===//
 
 func @combine_full_access_chain() -> f32 {
-  // CHECK: %[[INDEX:.*]] = spv.constant 0
+  // CHECK: %[[INDEX:.*]] = spv.Constant 0
   // CHECK-NEXT: %[[VAR:.*]] = spv.Variable
   // CHECK-NEXT: %[[PTR:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]], %[[INDEX]]]
   // CHECK-NEXT: spv.Load "Function" %[[PTR]]
-  %c0 = spv.constant 0: i32
+  %c0 = spv.Constant 0: i32
   %0 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
   %1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>, i32
   %2 = spv.AccessChain %1[%c0, %c0] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, i32, i32
@@ -20,13 +20,13 @@ func @combine_full_access_chain() -> f32 {
 // -----
 
 func @combine_access_chain_multi_use() -> !spv.array<4xf32> {
-  // CHECK: %[[INDEX:.*]] = spv.constant 0
+  // CHECK: %[[INDEX:.*]] = spv.Constant 0
   // CHECK-NEXT: %[[VAR:.*]] = spv.Variable
   // CHECK-NEXT: %[[PTR_0:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]]]
   // CHECK-NEXT: %[[PTR_1:.*]] = spv.AccessChain %[[VAR]][%[[INDEX]], %[[INDEX]], %[[INDEX]]]
   // CHECK-NEXT: spv.Load "Function" %[[PTR_0]]
   // CHECK-NEXT: spv.Load "Function" %[[PTR_1]]
-  %c0 = spv.constant 0: i32
+  %c0 = spv.Constant 0: i32
   %0 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
   %1 = spv.AccessChain %0[%c0] : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>, i32
   %2 = spv.AccessChain %1[%c0] : !spv.ptr<!spv.array<4x!spv.array<4xf32>>, Function>, i32
@@ -39,14 +39,14 @@ func @combine_access_chain_multi_use() -> !spv.array<4xf32> {
 // -----
 
 func @dont_combine_access_chain_without_common_base() -> !spv.array<4xi32> {
-  // CHECK: %[[INDEX:.*]] = spv.constant 1
+  // CHECK: %[[INDEX:.*]] = spv.Constant 1
   // CHECK-NEXT: %[[VAR_0:.*]] = spv.Variable
   // CHECK-NEXT: %[[VAR_1:.*]] = spv.Variable
   // CHECK-NEXT: %[[VAR_0_PTR:.*]] = spv.AccessChain %[[VAR_0]][%[[INDEX]]]
   // CHECK-NEXT: %[[VAR_1_PTR:.*]] = spv.AccessChain %[[VAR_1]][%[[INDEX]]]
   // CHECK-NEXT: spv.Load "Function" %[[VAR_0_PTR]]
   // CHECK-NEXT: spv.Load "Function" %[[VAR_1_PTR]]
-  %c1 = spv.constant 1: i32
+  %c1 = spv.Constant 1: i32
   %0 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
   %1 = spv.Variable : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>
   %2 = spv.AccessChain %0[%c1] : !spv.ptr<!spv.struct<(!spv.array<4x!spv.array<4xf32>>, !spv.array<4xi32>)>, Function>, i32
@@ -92,10 +92,10 @@ func @convert_bitcast_multi_use(%arg0 : vector<2xf32>, %arg1 : !spv.ptr<i64, Uni
 
 // CHECK-LABEL: extract_vector
 func @extract_vector() -> (i32, i32, i32) {
-  // CHECK: spv.constant 42 : i32
-  // CHECK: spv.constant -33 : i32
-  // CHECK: spv.constant 6 : i32
-  %0 = spv.constant dense<[42, -33, 6]> : vector<3xi32>
+  // CHECK: spv.Constant 42 : i32
+  // CHECK: spv.Constant -33 : i32
+  // CHECK: spv.Constant 6 : i32
+  %0 = spv.Constant dense<[42, -33, 6]> : vector<3xi32>
   %1 = spv.CompositeExtract %0[0 : i32] : vector<3xi32>
   %2 = spv.CompositeExtract %0[1 : i32] : vector<3xi32>
   %3 = spv.CompositeExtract %0[2 : i32] : vector<3xi32>
@@ -106,9 +106,9 @@ func @extract_vector() -> (i32, i32, i32) {
 
 // CHECK-LABEL: extract_array_final
 func @extract_array_final() -> (i32, i32) {
-  // CHECK: spv.constant 4 : i32
-  // CHECK: spv.constant -5 : i32
-  %0 = spv.constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+  // CHECK: spv.Constant 4 : i32
+  // CHECK: spv.Constant -5 : i32
+  %0 = spv.Constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
   %1 = spv.CompositeExtract %0[0 : i32, 0 : i32] : !spv.array<1 x vector<2 x i32>>
   %2 = spv.CompositeExtract %0[0 : i32, 1 : i32] : !spv.array<1 x vector<2 x i32>>
   return %1, %2 : i32, i32
@@ -118,8 +118,8 @@ func @extract_array_final() -> (i32, i32) {
 
 // CHECK-LABEL: extract_array_interm
 func @extract_array_interm() -> (vector<2xi32>) {
-  // CHECK: spv.constant dense<[4, -5]> : vector<2xi32>
-  %0 = spv.constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+  // CHECK: spv.Constant dense<[4, -5]> : vector<2xi32>
+  %0 = spv.Constant [dense<[4, -5]> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
   %1 = spv.CompositeExtract %0[0 : i32] : !spv.array<1 x vector<2 x i32>>
   return %1 : vector<2xi32>
 }
@@ -138,15 +138,15 @@ func @extract_from_not_constant() -> i32 {
 // -----
 
 //===----------------------------------------------------------------------===//
-// spv.constant
+// spv.Constant
 //===----------------------------------------------------------------------===//
 
 // TODO: test constants in 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_scalar_constant() -> (i32, i32) {
 // -----
 
 func @deduplicate_vector_constant() -> (vector<3xi32>, vector<3xi32>) {
-  // CHECK: %[[CST:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %0 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[CST:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %0 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
   // CHECK-NEXT: return %[[CST]], %[[CST]]
   return %0, %1 : vector<3xi32>, vector<3xi32>
 }
@@ -164,9 +164,9 @@ func @deduplicate_vector_constant() -> (vector<3xi32>, vector<3xi32>) {
 // -----
 
 func @deduplicate_composite_constant() -> (!spv.array<1 x vector<2xi32>>, !spv.array<1 x vector<2xi32>>) {
-  // CHECK: %[[CST:.*]] = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
-  %0 = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
-  %1 = spv.constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+  // CHECK: %[[CST:.*]] = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+  %0 = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
+  %1 = spv.Constant [dense<5> : vector<2xi32>] : !spv.array<1 x vector<2xi32>>
   // CHECK-NEXT: return %[[CST]], %[[CST]]
   return %0, %1 : !spv.array<1 x vector<2xi32>>, !spv.array<1 x vector<2xi32>>
 }
@@ -180,7 +180,7 @@ func @deduplicate_composite_constant() -> (!spv.array<1 x vector<2xi32>>, !spv.a
 // CHECK-LABEL: @iadd_zero
 // CHECK-SAME: (%[[ARG:.*]]: i32)
 func @iadd_zero(%arg0: i32) -> (i32, i32) {
-  %zero = spv.constant 0 : i32
+  %zero = spv.Constant 0 : i32
   %0 = spv.IAdd %arg0, %zero : i32
   %1 = spv.IAdd %zero, %arg0 : i32
   // CHECK: return %[[ARG]], %[[ARG]]
@@ -189,12 +189,12 @@ func @iadd_zero(%arg0: i32) -> (i32, i32) {
 
 // CHECK-LABEL: @const_fold_scalar_iadd_normal
 func @const_fold_scalar_iadd_normal() -> (i32, i32, i32) {
-  %c5 = spv.constant 5 : i32
-  %cn8 = spv.constant -8 : i32
+  %c5 = spv.Constant 5 : i32
+  %cn8 = spv.Constant -8 : i32
 
-  // CHECK: spv.constant 10
-  // CHECK: spv.constant -16
-  // CHECK: spv.constant -3
+  // CHECK: spv.Constant 10
+  // CHECK: spv.Constant -16
+  // CHECK: spv.Constant -3
   %0 = spv.IAdd %c5, %c5 : i32
   %1 = spv.IAdd %cn8, %cn8 : i32
   %2 = spv.IAdd %c5, %cn8 : i32
@@ -203,34 +203,34 @@ func @const_fold_scalar_iadd_normal() -> (i32, i32, i32) {
 
 // CHECK-LABEL: @const_fold_scalar_iadd_flow
 func @const_fold_scalar_iadd_flow() -> (i32, i32, i32, i32) {
-  %c1 = spv.constant 1 : i32
-  %c2 = spv.constant 2 : i32
-  %c3 = spv.constant 4294967295 : i32  // 2^32 - 1: 0xffff ffff
-  %c4 = spv.constant -2147483648 : i32 // -2^31   : 0x8000 0000
-  %c5 = spv.constant -1 : i32          //         : 0xffff ffff
-  %c6 = spv.constant -2 : i32          //         : 0xffff fffe
+  %c1 = spv.Constant 1 : i32
+  %c2 = spv.Constant 2 : i32
+  %c3 = spv.Constant 4294967295 : i32  // 2^32 - 1: 0xffff ffff
+  %c4 = spv.Constant -2147483648 : i32 // -2^31   : 0x8000 0000
+  %c5 = spv.Constant -1 : i32          //         : 0xffff ffff
+  %c6 = spv.Constant -2 : i32          //         : 0xffff fffe
 
   // 0x0000 0001 + 0xffff ffff = 0x1 0000 0000 -> 0x0000 0000
-  // CHECK: spv.constant 0
+  // CHECK: spv.Constant 0
   %0 = spv.IAdd %c1, %c3 : i32
   // 0x0000 0002 + 0xffff ffff = 0x1 0000 0001 -> 0x0000 0001
-  // CHECK: spv.constant 1
+  // CHECK: spv.Constant 1
   %1 = spv.IAdd %c2, %c3 : i32
   // 0x8000 0000 + 0xffff ffff = 0x1 7fff ffff -> 0x7fff ffff
-  // CHECK: spv.constant 2147483647
+  // CHECK: spv.Constant 2147483647
   %2 = spv.IAdd %c4, %c5 : i32
   // 0x8000 0000 + 0xffff fffe = 0x1 7fff fffe -> 0x7fff fffe
-  // CHECK: spv.constant 2147483646
+  // CHECK: spv.Constant 2147483646
   %3 = spv.IAdd %c4, %c6 : i32
   return %0, %1, %2, %3: i32, i32, i32, i32
 }
 
 // CHECK-LABEL: @const_fold_vector_iadd
 func @const_fold_vector_iadd() -> vector<3xi32> {
-  %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32>
-  %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32>
+  %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32>
+  %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32>
 
-  // CHECK: spv.constant dense<[39, -70, 155]>
+  // CHECK: spv.Constant dense<[39, -70, 155]>
   %0 = spv.IAdd %vc1, %vc2 : vector<3xi32>
   return %0: vector<3xi32>
 }
@@ -244,9 +244,9 @@ func @const_fold_vector_iadd() -> vector<3xi32> {
 // CHECK-LABEL: @imul_zero_one
 // CHECK-SAME: (%[[ARG:.*]]: i32)
 func @imul_zero_one(%arg0: i32) -> (i32, i32) {
-  // CHECK: %[[ZERO:.*]] = spv.constant 0
-  %zero = spv.constant 0 : i32
-  %one = spv.constant 1: i32
+  // CHECK: %[[ZERO:.*]] = spv.Constant 0
+  %zero = spv.Constant 0 : i32
+  %one = spv.Constant 1: i32
   %0 = spv.IMul %arg0, %zero : i32
   %1 = spv.IMul %one, %arg0 : i32
   // CHECK: return %[[ZERO]], %[[ARG]]
@@ -255,13 +255,13 @@ func @imul_zero_one(%arg0: i32) -> (i32, i32) {
 
 // CHECK-LABEL: @const_fold_scalar_imul_normal
 func @const_fold_scalar_imul_normal() -> (i32, i32, i32) {
-  %c5 = spv.constant 5 : i32
-  %cn8 = spv.constant -8 : i32
-  %c7 = spv.constant 7 : i32
+  %c5 = spv.Constant 5 : i32
+  %cn8 = spv.Constant -8 : i32
+  %c7 = spv.Constant 7 : i32
 
-  // CHECK: spv.constant 35
-  // CHECK: spv.constant -40
-  // CHECK: spv.constant -56
+  // CHECK: spv.Constant 35
+  // CHECK: spv.Constant -40
+  // CHECK: spv.Constant -56
   %0 = spv.IMul %c7, %c5 : i32
   %1 = spv.IMul %c5, %cn8 : i32
   %2 = spv.IMul %cn8, %c7 : i32
@@ -270,18 +270,18 @@ func @const_fold_scalar_imul_normal() -> (i32, i32, i32) {
 
 // CHECK-LABEL: @const_fold_scalar_imul_flow
 func @const_fold_scalar_imul_flow() -> (i32, i32, i32) {
-  %c1 = spv.constant 2 : i32
-  %c2 = spv.constant 4 : i32
-  %c3 = spv.constant 4294967295 : i32  // 2^32 - 1 : 0xffff ffff
-  %c4 = spv.constant 2147483647 : i32  // 2^31 - 1 : 0x7fff ffff
+  %c1 = spv.Constant 2 : i32
+  %c2 = spv.Constant 4 : i32
+  %c3 = spv.Constant 4294967295 : i32  // 2^32 - 1 : 0xffff ffff
+  %c4 = spv.Constant 2147483647 : i32  // 2^31 - 1 : 0x7fff ffff
 
   // (0xffff ffff << 1) = 0x1 ffff fffe -> 0xffff fffe
-  // CHECK: %[[CST2:.*]] = spv.constant -2
+  // CHECK: %[[CST2:.*]] = spv.Constant -2
   %0 = spv.IMul %c1, %c3 : i32
   // (0x7fff ffff << 1) = 0x0 ffff fffe -> 0xffff fffe
   %1 = spv.IMul %c1, %c4 : i32
   // (0x7fff ffff << 2) = 0x1 ffff fffc -> 0xffff fffc
-  // CHECK: %[[CST4:.*]] = spv.constant -4
+  // CHECK: %[[CST4:.*]] = spv.Constant -4
   %2 = spv.IMul %c4, %c2 : i32
   // CHECK: return %[[CST2]], %[[CST2]], %[[CST4]]
   return %0, %1, %2: i32, i32, i32
@@ -290,10 +290,10 @@ func @const_fold_scalar_imul_flow() -> (i32, i32, i32) {
 
 // CHECK-LABEL: @const_fold_vector_imul
 func @const_fold_vector_imul() -> vector<3xi32> {
-  %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32>
-  %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32>
+  %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32>
+  %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32>
 
-  // CHECK: spv.constant dense<[-126, 825, 3556]>
+  // CHECK: spv.Constant dense<[-126, 825, 3556]>
   %0 = spv.IMul %vc1, %vc2 : vector<3xi32>
   return %0: vector<3xi32>
 }
@@ -306,20 +306,20 @@ func @const_fold_vector_imul() -> vector<3xi32> {
 
 // CHECK-LABEL: @isub_x_x
 func @isub_x_x(%arg0: i32) -> i32 {
-  // CHECK: spv.constant 0
+  // CHECK: spv.Constant 0
   %0 = spv.ISub %arg0, %arg0: i32
   return %0: i32
 }
 
 // CHECK-LABEL: @const_fold_scalar_isub_normal
 func @const_fold_scalar_isub_normal() -> (i32, i32, i32) {
-  %c5 = spv.constant 5 : i32
-  %cn8 = spv.constant -8 : i32
-  %c7 = spv.constant 7 : i32
+  %c5 = spv.Constant 5 : i32
+  %cn8 = spv.Constant -8 : i32
+  %c7 = spv.Constant 7 : i32
 
-  // CHECK: spv.constant 2
-  // CHECK: spv.constant 13
-  // CHECK: spv.constant -15
+  // CHECK: spv.Constant 2
+  // CHECK: spv.Constant 13
+  // CHECK: spv.Constant -15
   %0 = spv.ISub %c7, %c5 : i32
   %1 = spv.ISub %c5, %cn8 : i32
   %2 = spv.ISub %cn8, %c7 : i32
@@ -328,34 +328,34 @@ func @const_fold_scalar_isub_normal() -> (i32, i32, i32) {
 
 // CHECK-LABEL: @const_fold_scalar_isub_flow
 func @const_fold_scalar_isub_flow() -> (i32, i32, i32, i32) {
-  %c1 = spv.constant 0 : i32
-  %c2 = spv.constant 1 : i32
-  %c3 = spv.constant 4294967295 : i32  // 2^32 - 1 : 0xffff ffff
-  %c4 = spv.constant 2147483647 : i32  // 2^31     : 0x7fff ffff
-  %c5 = spv.constant -1 : i32          //          : 0xffff ffff
-  %c6 = spv.constant -2 : i32          //          : 0xffff fffe
+  %c1 = spv.Constant 0 : i32
+  %c2 = spv.Constant 1 : i32
+  %c3 = spv.Constant 4294967295 : i32  // 2^32 - 1 : 0xffff ffff
+  %c4 = spv.Constant 2147483647 : i32  // 2^31     : 0x7fff ffff
+  %c5 = spv.Constant -1 : i32          //          : 0xffff ffff
+  %c6 = spv.Constant -2 : i32          //          : 0xffff fffe
 
   // 0x0000 0000 - 0xffff ffff -> 0x0000 0000 + 0x0000 0001 = 0x0000 0001
-  // CHECK: spv.constant 1
+  // CHECK: spv.Constant 1
   %0 = spv.ISub %c1, %c3 : i32
   // 0x0000 0001 - 0xffff ffff -> 0x0000 0001 + 0x0000 0001 = 0x0000 0002
-  // CHECK: spv.constant 2
+  // CHECK: spv.Constant 2
   %1 = spv.ISub %c2, %c3 : i32
   // 0xffff ffff - 0x7fff ffff -> 0xffff ffff + 0x8000 0001 = 0x1 8000 0000
-  // CHECK: spv.constant -2147483648
+  // CHECK: spv.Constant -2147483648
   %2 = spv.ISub %c5, %c4 : i32
   // 0xffff fffe - 0x7fff ffff -> 0xffff fffe + 0x8000 0001 = 0x1 7fff ffff
-  // CHECK: spv.constant 2147483647
+  // CHECK: spv.Constant 2147483647
   %3 = spv.ISub %c6, %c4 : i32
   return %0, %1, %2, %3: i32, i32, i32, i32
 }
 
 // CHECK-LABEL: @const_fold_vector_isub
 func @const_fold_vector_isub() -> vector<3xi32> {
-  %vc1 = spv.constant dense<[42, -55, 127]> : vector<3xi32>
-  %vc2 = spv.constant dense<[-3, -15, 28]> : vector<3xi32>
+  %vc1 = spv.Constant dense<[42, -55, 127]> : vector<3xi32>
+  %vc2 = spv.Constant dense<[-3, -15, 28]> : vector<3xi32>
 
-  // CHECK: spv.constant dense<[45, -40, 99]>
+  // CHECK: spv.Constant dense<[45, -40, 99]>
   %0 = spv.ISub %vc1, %vc2 : vector<3xi32>
   return %0: vector<3xi32>
 }
@@ -369,9 +369,9 @@ func @const_fold_vector_isub() -> vector<3xi32> {
 // CHECK-LABEL: @convert_logical_and_true_false_scalar
 // CHECK-SAME: %[[ARG:.+]]: i1
 func @convert_logical_and_true_false_scalar(%arg: i1) -> (i1, i1) {
-  %true = spv.constant true
-  // CHECK: %[[FALSE:.+]] = spv.constant false
-  %false = spv.constant false
+  %true = spv.Constant true
+  // CHECK: %[[FALSE:.+]] = spv.Constant false
+  %false = spv.Constant false
   %0 = spv.LogicalAnd %true, %arg: i1
   %1 = spv.LogicalAnd %arg, %false: i1
   // CHECK: return %[[ARG]], %[[FALSE]]
@@ -381,9 +381,9 @@ func @convert_logical_and_true_false_scalar(%arg: i1) -> (i1, i1) {
 // CHECK-LABEL: @convert_logical_and_true_false_vector
 // CHECK-SAME: %[[ARG:.+]]: vector<3xi1>
 func @convert_logical_and_true_false_vector(%arg: vector<3xi1>) -> (vector<3xi1>, vector<3xi1>) {
-  %true = spv.constant dense<true> : vector<3xi1>
-  // CHECK: %[[FALSE:.+]] = spv.constant dense<false>
-  %false = spv.constant dense<false> : vector<3xi1>
+  %true = spv.Constant dense<true> : vector<3xi1>
+  // CHECK: %[[FALSE:.+]] = spv.Constant dense<false>
+  %false = spv.Constant dense<false> : vector<3xi1>
   %0 = spv.LogicalAnd %true, %arg: vector<3xi1>
   %1 = spv.LogicalAnd %arg, %false: vector<3xi1>
   // CHECK: return %[[ARG]], %[[FALSE]]
@@ -456,9 +456,9 @@ func @convert_logical_not_to_logical_equal(%arg0: vector<3xi1>, %arg1: vector<3x
 // CHECK-LABEL: @convert_logical_or_true_false_scalar
 // CHECK-SAME: %[[ARG:.+]]: i1
 func @convert_logical_or_true_false_scalar(%arg: i1) -> (i1, i1) {
-  // CHECK: %[[TRUE:.+]] = spv.constant true
-  %true = spv.constant true
-  %false = spv.constant false
+  // CHECK: %[[TRUE:.+]] = spv.Constant true
+  %true = spv.Constant true
+  %false = spv.Constant false
   %0 = spv.LogicalOr %true, %arg: i1
   %1 = spv.LogicalOr %arg, %false: i1
   // CHECK: return %[[TRUE]], %[[ARG]]
@@ -468,9 +468,9 @@ func @convert_logical_or_true_false_scalar(%arg: i1) -> (i1, i1) {
 // CHECK-LABEL: @convert_logical_or_true_false_vector
 // CHECK-SAME: %[[ARG:.+]]: vector<3xi1>
 func @convert_logical_or_true_false_vector(%arg: vector<3xi1>) -> (vector<3xi1>, vector<3xi1>) {
-  // CHECK: %[[TRUE:.+]] = spv.constant dense<true>
-  %true = spv.constant dense<true> : vector<3xi1>
-  %false = spv.constant dense<false> : vector<3xi1>
+  // CHECK: %[[TRUE:.+]] = spv.Constant dense<true>
+  %true = spv.Constant dense<true> : vector<3xi1>
+  %false = spv.Constant dense<false> : vector<3xi1>
   %0 = spv.LogicalOr %true, %arg: vector<3xi1>
   %1 = spv.LogicalOr %arg, %false: vector<3xi1>
   // CHECK: return %[[TRUE]], %[[ARG]]
@@ -484,11 +484,11 @@ func @convert_logical_or_true_false_vector(%arg: vector<3xi1>) -> (vector<3xi1>,
 //===----------------------------------------------------------------------===//
 
 func @canonicalize_selection_op_scalar_type(%cond: i1) -> () {
-  %0 = spv.constant 0: i32
-  // CHECK: %[[TRUE_VALUE:.*]] = spv.constant 1 : i32
-  %1 = spv.constant 1: i32
-  // CHECK: %[[FALSE_VALUE:.*]] = spv.constant 2 : i32
-  %2 = spv.constant 2: i32
+  %0 = spv.Constant 0: i32
+  // CHECK: %[[TRUE_VALUE:.*]] = spv.Constant 1 : i32
+  %1 = spv.Constant 1: i32
+  // CHECK: %[[FALSE_VALUE:.*]] = spv.Constant 2 : i32
+  %2 = spv.Constant 2: i32
   // CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<i32, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<i32, Function>
 
@@ -515,11 +515,11 @@ func @canonicalize_selection_op_scalar_type(%cond: i1) -> () {
 // -----
 
 func @canonicalize_selection_op_vector_type(%cond: i1) -> () {
-  %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
-  // CHECK: %[[TRUE_VALUE:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  // CHECK: %[[FALSE_VALUE:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+  %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+  // CHECK: %[[TRUE_VALUE:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[FALSE_VALUE:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
   // CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
 
@@ -547,11 +547,11 @@ func @canonicalize_selection_op_vector_type(%cond: i1) -> () {
 
 // Store to a different variables.
 func @cannot_canonicalize_selection_op_0(%cond: i1) -> () {
-  %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+  %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
   // CHECK: %[[DST_VAR_0:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
   // CHECK: %[[DST_VAR_1:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
@@ -584,11 +584,11 @@ func @cannot_canonicalize_selection_op_0(%cond: i1) -> () {
 
 // A conditional block consists of more than 2 operations.
 func @cannot_canonicalize_selection_op_1(%cond: i1) -> () {
-  %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+  %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
   // CHECK: %[[DST_VAR_0:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
   // CHECK: %[[DST_VAR_1:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
@@ -620,11 +620,11 @@ func @cannot_canonicalize_selection_op_1(%cond: i1) -> () {
 
 // A control-flow goes into `^then` block from `^else` block.
 func @cannot_canonicalize_selection_op_2(%cond: i1) -> () {
-  %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+  %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
   // CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
 
@@ -652,11 +652,11 @@ func @cannot_canonicalize_selection_op_2(%cond: i1) -> () {
 
 // `spv.Return` as a block terminator.
 func @cannot_canonicalize_selection_op_3(%cond: i1) -> () {
-  %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+  %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
   // CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
 
@@ -684,11 +684,11 @@ func @cannot_canonicalize_selection_op_3(%cond: i1) -> () {
 
 // Different memory access attributes.
 func @cannot_canonicalize_selection_op_4(%cond: i1) -> () {
-  %0 = spv.constant dense<[0, 1, 2]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_0:.*]] = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  %1 = spv.constant dense<[1, 2, 3]> : vector<3xi32>
-  // CHECK: %[[SRC_VALUE_1:.*]] = spv.constant dense<[2, 3, 4]> : vector<3xi32>
-  %2 = spv.constant dense<[2, 3, 4]> : vector<3xi32>
+  %0 = spv.Constant dense<[0, 1, 2]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_0:.*]] = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  %1 = spv.Constant dense<[1, 2, 3]> : vector<3xi32>
+  // CHECK: %[[SRC_VALUE_1:.*]] = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
+  %2 = spv.Constant dense<[2, 3, 4]> : vector<3xi32>
   // CHECK: %[[DST_VAR:.*]] = spv.Variable init({{%.*}}) : !spv.ptr<vector<3xi32>, Function>
   %3 = spv.Variable init(%0) : !spv.ptr<vector<3xi32>, Function>
 
index 90e9b85..8f31ba7 100644 (file)
@@ -2,10 +2,10 @@
 
 // CHECK: func @clamp_fordlessthan(%[[INPUT:.*]]: f32)
 func @clamp_fordlessthan(%input: f32) -> f32 {
-  // CHECK: %[[MIN:.*]] = spv.constant
-  %min = spv.constant 0.5 : f32
-  // CHECK: %[[MAX:.*]] = spv.constant
-  %max = spv.constant 1.0 : f32
+  // CHECK: %[[MIN:.*]] = spv.Constant
+  %min = spv.Constant 0.5 : f32
+  // CHECK: %[[MAX:.*]] = spv.Constant
+  %max = spv.Constant 1.0 : f32
 
   // CHECK: [[RES:%.*]] = spv.GLSL.FClamp %[[INPUT]], %[[MIN]], %[[MAX]]
   %0 = spv.FOrdLessThan %min, %input : f32
@@ -21,10 +21,10 @@ func @clamp_fordlessthan(%input: f32) -> f32 {
 
 // CHECK: func @clamp_fordlessthanequal(%[[INPUT:.*]]: f32)
 func @clamp_fordlessthanequal(%input: f32) -> f32 {
-  // CHECK: %[[MIN:.*]] = spv.constant
-  %min = spv.constant 0.5 : f32
-  // CHECK: %[[MAX:.*]] = spv.constant
-  %max = spv.constant 1.0 : f32
+  // CHECK: %[[MIN:.*]] = spv.Constant
+  %min = spv.Constant 0.5 : f32
+  // CHECK: %[[MAX:.*]] = spv.Constant
+  %max = spv.Constant 1.0 : f32
 
   // CHECK: [[RES:%.*]] = spv.GLSL.FClamp %[[INPUT]], %[[MIN]], %[[MAX]]
   %0 = spv.FOrdLessThanEqual %min, %input : f32
@@ -40,10 +40,10 @@ func @clamp_fordlessthanequal(%input: f32) -> f32 {
 
 // CHECK: func @clamp_slessthan(%[[INPUT:.*]]: si32)
 func @clamp_slessthan(%input: si32) -> si32 {
-  // CHECK: %[[MIN:.*]] = spv.constant
-  %min = spv.constant 0 : si32
-  // CHECK: %[[MAX:.*]] = spv.constant
-  %max = spv.constant 10 : si32
+  // CHECK: %[[MIN:.*]] = spv.Constant
+  %min = spv.Constant 0 : si32
+  // CHECK: %[[MAX:.*]] = spv.Constant
+  %max = spv.Constant 10 : si32
 
   // CHECK: [[RES:%.*]] = spv.GLSL.SClamp %[[INPUT]], %[[MIN]], %[[MAX]]
   %0 = spv.SLessThan %min, %input : si32
@@ -59,10 +59,10 @@ func @clamp_slessthan(%input: si32) -> si32 {
 
 // CHECK: func @clamp_slessthanequal(%[[INPUT:.*]]: si32)
 func @clamp_slessthanequal(%input: si32) -> si32 {
-  // CHECK: %[[MIN:.*]] = spv.constant
-  %min = spv.constant 0 : si32
-  // CHECK: %[[MAX:.*]] = spv.constant
-  %max = spv.constant 10 : si32
+  // CHECK: %[[MIN:.*]] = spv.Constant
+  %min = spv.Constant 0 : si32
+  // CHECK: %[[MAX:.*]] = spv.Constant
+  %max = spv.Constant 10 : si32
 
   // CHECK: [[RES:%.*]] = spv.GLSL.SClamp %[[INPUT]], %[[MIN]], %[[MAX]]
   %0 = spv.SLessThanEqual %min, %input : si32
@@ -78,10 +78,10 @@ func @clamp_slessthanequal(%input: si32) -> si32 {
 
 // CHECK: func @clamp_ulessthan(%[[INPUT:.*]]: i32)
 func @clamp_ulessthan(%input: i32) -> i32 {
-  // CHECK: %[[MIN:.*]] = spv.constant
-  %min = spv.constant 0 : i32
-  // CHECK: %[[MAX:.*]] = spv.constant
-  %max = spv.constant 10 : i32
+  // CHECK: %[[MIN:.*]] = spv.Constant
+  %min = spv.Constant 0 : i32
+  // CHECK: %[[MAX:.*]] = spv.Constant
+  %max = spv.Constant 10 : i32
 
   // CHECK: [[RES:%.*]] = spv.GLSL.UClamp %[[INPUT]], %[[MIN]], %[[MAX]]
   %0 = spv.ULessThan %min, %input : i32
@@ -97,10 +97,10 @@ func @clamp_ulessthan(%input: i32) -> i32 {
 
 // CHECK: func @clamp_ulessthanequal(%[[INPUT:.*]]: i32)
 func @clamp_ulessthanequal(%input: i32) -> i32 {
-  // CHECK: %[[MIN:.*]] = spv.constant
-  %min = spv.constant 0 : i32
-  // CHECK: %[[MAX:.*]] = spv.constant
-  %max = spv.constant 10 : i32
+  // CHECK: %[[MIN:.*]] = spv.Constant
+  %min = spv.Constant 0 : i32
+  // CHECK: %[[MAX:.*]] = spv.Constant
+  %max = spv.Constant 10 : i32
 
   // CHECK: [[RES:%.*]] = spv.GLSL.UClamp %[[INPUT]], %[[MIN]], %[[MAX]]
   %0 = spv.ULessThanEqual %min, %input : i32
index 983fc26..7b66359 100644 (file)
@@ -17,13 +17,13 @@ spv.module Logical GLSL450 {
 
 spv.module Logical GLSL450 {
   spv.func @callee() -> i32 "None" {
-    %0 = spv.constant 42 : i32
+    %0 = spv.Constant 42 : i32
     spv.ReturnValue %0 : i32
   }
 
   // CHECK-LABEL: @calling_single_block_retval_func
   spv.func @calling_single_block_retval_func() -> i32 "None" {
-    // CHECK-NEXT: %[[CST:.*]] = spv.constant 42
+    // CHECK-NEXT: %[[CST:.*]] = spv.Constant 42
     %0 = spv.FunctionCall @callee() : () -> (i32)
     // CHECK-NEXT: spv.ReturnValue %[[CST]]
     spv.ReturnValue %0 : i32
@@ -36,12 +36,12 @@ spv.module Logical GLSL450 {
   spv.globalVariable @data bind(0, 0) : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
   spv.func @callee() "None" {
     %0 = spv.mlir.addressof @data : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>
-    %1 = spv.constant 0: i32
+    %1 = spv.Constant 0: i32
     %2 = spv.AccessChain %0[%1, %1] : !spv.ptr<!spv.struct<(!spv.rtarray<i32> [0])>, StorageBuffer>, i32, i32
     spv.Branch ^next
 
   ^next:
-    %3 = spv.constant 42: i32
+    %3 = spv.Constant 42: i32
     spv.Store "StorageBuffer" %2, %3 : i32
     spv.Return
   }
@@ -49,11 +49,11 @@ spv.module Logical GLSL450 {
   // CHECK-LABEL: @calling_multi_block_ret_func
   spv.func @calling_multi_block_ret_func() "None" {
     // CHECK-NEXT:   spv.mlir.addressof
-    // CHECK-NEXT:   spv.constant 0
+    // CHECK-NEXT:   spv.Constant 0
     // CHECK-NEXT:   spv.AccessChain
     // CHECK-NEXT:   spv.Branch ^bb1
     // CHECK-NEXT: ^bb1:
-    // CHECK-NEXT:   spv.constant
+    // CHECK-NEXT:   spv.Constant
     // CHECK-NEXT:   spv.Store
     // CHECK-NEXT:   spv.Branch ^bb2
     spv.FunctionCall @callee() : () -> ()
@@ -81,7 +81,7 @@ spv.module Logical GLSL450 {
 
   // CHECK-LABEL: @calling_selection_ret_func
   spv.func @calling_selection_ret_func() "None" {
-    %0 = spv.constant true
+    %0 = spv.Constant true
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
@@ -104,8 +104,8 @@ spv.module Logical GLSL450 {
 
   // CHECK-LABEL: @calling_selection_no_ret_func
   spv.func @calling_selection_no_ret_func() "None" {
-    // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true
-    %0 = spv.constant true
+    // CHECK-NEXT: %[[TRUE:.*]] = spv.Constant true
+    %0 = spv.Constant true
     // CHECK-NEXT: spv.selection
     // CHECK-NEXT:   spv.BranchConditional %[[TRUE]], ^bb1, ^bb2
     // CHECK-NEXT: ^bb1:
@@ -137,7 +137,7 @@ spv.module Logical GLSL450 {
 
   // CHECK-LABEL: @calling_loop_ret_func
   spv.func @calling_loop_ret_func() "None" {
-    %0 = spv.constant true
+    %0 = spv.Constant true
     // CHECK: spv.FunctionCall
     spv.FunctionCall @callee(%0) : (i1) -> ()
     spv.Return
@@ -164,8 +164,8 @@ spv.module Logical GLSL450 {
 
   // CHECK-LABEL: @calling_loop_no_ret_func
   spv.func @calling_loop_no_ret_func() "None" {
-    // CHECK-NEXT: %[[TRUE:.*]] = spv.constant true
-    %0 = spv.constant true
+    // CHECK-NEXT: %[[TRUE:.*]] = spv.Constant true
+    %0 = spv.Constant true
     // CHECK-NEXT: spv.loop
     // CHECK-NEXT:   spv.Branch ^bb1
     // CHECK-NEXT: ^bb1:
@@ -189,7 +189,7 @@ spv.module Logical GLSL450 {
 
   // CHECK: @inline_into_selection_region
   spv.func @inline_into_selection_region() "None" {
-    %1 = spv.constant 0 : i32
+    %1 = spv.Constant 0 : i32
     // CHECK-DAG: [[ADDRESS_ARG0:%.*]] = spv.mlir.addressof @arg_0
     // CHECK-DAG: [[ADDRESS_ARG1:%.*]] = spv.mlir.addressof @arg_1
     // CHECK-DAG: [[LOADPTR:%.*]] = spv.AccessChain [[ADDRESS_ARG0]]
index 2a7eb65..f21f97c 100644 (file)
@@ -20,7 +20,7 @@ spv.module Logical GLSL450 {
   spv.globalVariable @var5 bind(1,3) : !spv.ptr<!spv.struct<(!spv.array<256xf32>)>, StorageBuffer>
 
   spv.func @kernel() -> () "None" {
-    %c0 = spv.constant 0 : i32
+    %c0 = spv.Constant 0 : i32
     // CHECK: {{%.*}} = spv.mlir.addressof @var0 : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
     %0 = spv.mlir.addressof @var0 : !spv.ptr<!spv.struct<(i32, !spv.struct<(f32, i32)>, f32)>, Uniform>
     // CHECK:  {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.struct<(i32 [0], !spv.struct<(f32 [0], i32 [4])> [4], f32 [12])>, Uniform>
index 5206692..e9955ce 100644 (file)
@@ -3,10 +3,10 @@
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @bool_const
   spv.func @bool_const() -> () "None" {
-    // CHECK: spv.constant true
-    %0 = spv.constant true
-    // CHECK: spv.constant false
-    %1 = spv.constant false
+    // CHECK: spv.Constant true
+    %0 = spv.Constant true
+    // CHECK: spv.Constant false
+    %1 = spv.Constant false
 
     %2 = spv.Variable init(%0): !spv.ptr<i1, Function>
     %3 = spv.Variable init(%1): !spv.ptr<i1, Function>
@@ -15,12 +15,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @i32_const
   spv.func @i32_const() -> () "None" {
-    // CHECK: spv.constant 0 : i32
-    %0 = spv.constant  0 : i32
-    // CHECK: spv.constant 10 : i32
-    %1 = spv.constant 10 : i32
-    // CHECK: spv.constant -5 : i32
-    %2 = spv.constant -5 : i32
+    // CHECK: spv.Constant 0 : i32
+    %0 = spv.Constant  0 : i32
+    // CHECK: spv.Constant 10 : i32
+    %1 = spv.Constant 10 : i32
+    // CHECK: spv.Constant -5 : i32
+    %2 = spv.Constant -5 : i32
 
     %3 = spv.IAdd %0, %1 : i32
     %4 = spv.IAdd %2, %3 : i32
@@ -29,12 +29,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @si32_const
   spv.func @si32_const() -> () "None" {
-    // CHECK: spv.constant 0 : si32
-    %0 = spv.constant  0 : si32
-    // CHECK: spv.constant 10 : si32
-    %1 = spv.constant 10 : si32
-    // CHECK: spv.constant -5 : si32
-    %2 = spv.constant -5 : si32
+    // CHECK: spv.Constant 0 : si32
+    %0 = spv.Constant  0 : si32
+    // CHECK: spv.Constant 10 : si32
+    %1 = spv.Constant 10 : si32
+    // CHECK: spv.Constant -5 : si32
+    %2 = spv.Constant -5 : si32
 
     %3 = spv.IAdd %0, %1 : si32
     %4 = spv.IAdd %2, %3 : si32
@@ -46,12 +46,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // because they all use 1 as the signedness bit. So we always treat them
   // as signless integers.
   spv.func @ui32_const() -> () "None" {
-    // CHECK: spv.constant 0 : i32
-    %0 = spv.constant  0 : ui32
-    // CHECK: spv.constant 10 : i32
-    %1 = spv.constant 10 : ui32
-    // CHECK: spv.constant -5 : i32
-    %2 = spv.constant 4294967291 : ui32
+    // CHECK: spv.Constant 0 : i32
+    %0 = spv.Constant  0 : ui32
+    // CHECK: spv.Constant 10 : i32
+    %1 = spv.Constant 10 : ui32
+    // CHECK: spv.Constant -5 : i32
+    %2 = spv.Constant 4294967291 : ui32
 
     %3 = spv.IAdd %0, %1 : ui32
     %4 = spv.IAdd %2, %3 : ui32
@@ -60,14 +60,14 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @i64_const
   spv.func @i64_const() -> () "None" {
-    // CHECK: spv.constant 4294967296 : i64
-    %0 = spv.constant           4294967296 : i64 //  2^32
-    // CHECK: spv.constant -4294967296 : i64
-    %1 = spv.constant          -4294967296 : i64 // -2^32
-    // CHECK: spv.constant 9223372036854775807 : i64
-    %2 = spv.constant  9223372036854775807 : i64 //  2^63 - 1
-    // CHECK: spv.constant -9223372036854775808 : i64
-    %3 = spv.constant -9223372036854775808 : i64 // -2^63
+    // CHECK: spv.Constant 4294967296 : i64
+    %0 = spv.Constant           4294967296 : i64 //  2^32
+    // CHECK: spv.Constant -4294967296 : i64
+    %1 = spv.Constant          -4294967296 : i64 // -2^32
+    // CHECK: spv.Constant 9223372036854775807 : i64
+    %2 = spv.Constant  9223372036854775807 : i64 //  2^63 - 1
+    // CHECK: spv.Constant -9223372036854775808 : i64
+    %3 = spv.Constant -9223372036854775808 : i64 // -2^63
 
     %4 = spv.IAdd %0, %1 : i64
     %5 = spv.IAdd %2, %3 : i64
@@ -76,10 +76,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @i16_const
   spv.func @i16_const() -> () "None" {
-    // CHECK: spv.constant -32768 : i16
-    %0 = spv.constant -32768 : i16 // -2^15
-    // CHECK: spv.constant 32767 : i16
-    %1 = spv.constant 32767 : i16 //  2^15 - 1
+    // CHECK: spv.Constant -32768 : i16
+    %0 = spv.Constant -32768 : i16 // -2^15
+    // CHECK: spv.Constant 32767 : i16
+    %1 = spv.Constant 32767 : i16 //  2^15 - 1
 
     %2 = spv.IAdd %0, %1 : i16
     spv.Return
@@ -87,18 +87,18 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @float_const
   spv.func @float_const() -> () "None" {
-    // CHECK: spv.constant 0.000000e+00 : f32
-    %0 = spv.constant 0. : f32
-    // CHECK: spv.constant 1.000000e+00 : f32
-    %1 = spv.constant 1. : f32
-    // CHECK: spv.constant -0.000000e+00 : f32
-    %2 = spv.constant -0. : f32
-    // CHECK: spv.constant -1.000000e+00 : f32
-    %3 = spv.constant -1. : f32
-    // CHECK: spv.constant 7.500000e-01 : f32
-    %4 = spv.constant 0.75 : f32
-    // CHECK: spv.constant -2.500000e-01 : f32
-    %5 = spv.constant -0.25 : f32
+    // CHECK: spv.Constant 0.000000e+00 : f32
+    %0 = spv.Constant 0. : f32
+    // CHECK: spv.Constant 1.000000e+00 : f32
+    %1 = spv.Constant 1. : f32
+    // CHECK: spv.Constant -0.000000e+00 : f32
+    %2 = spv.Constant -0. : f32
+    // CHECK: spv.Constant -1.000000e+00 : f32
+    %3 = spv.Constant -1. : f32
+    // CHECK: spv.Constant 7.500000e-01 : f32
+    %4 = spv.Constant 0.75 : f32
+    // CHECK: spv.Constant -2.500000e-01 : f32
+    %5 = spv.Constant -0.25 : f32
 
     %6 = spv.FAdd %0, %1 : f32
     %7 = spv.FAdd %2, %3 : f32
@@ -109,10 +109,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // CHECK-LABEL: @double_const
   spv.func @double_const() -> () "None" {
     // TODO: test range boundary values
-    // CHECK: spv.constant 1.024000e+03 : f64
-    %0 = spv.constant 1024. : f64
-    // CHECK: spv.constant -1.024000e+03 : f64
-    %1 = spv.constant -1024. : f64
+    // CHECK: spv.Constant 1.024000e+03 : f64
+    %0 = spv.Constant 1024. : f64
+    // CHECK: spv.Constant -1.024000e+03 : f64
+    %1 = spv.Constant -1024. : f64
 
     %2 = spv.FAdd %0, %1 : f64
     spv.Return
@@ -120,10 +120,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @half_const
   spv.func @half_const() -> () "None" {
-    // CHECK: spv.constant 5.120000e+02 : f16
-    %0 = spv.constant 512. : f16
-    // CHECK: spv.constant -5.120000e+02 : f16
-    %1 = spv.constant -512. : f16
+    // CHECK: spv.Constant 5.120000e+02 : f16
+    %0 = spv.Constant 512. : f16
+    // CHECK: spv.Constant -5.120000e+02 : f16
+    %1 = spv.Constant -512. : f16
 
     %2 = spv.FAdd %0, %1 : f16
     spv.Return
@@ -131,12 +131,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @bool_vector_const
   spv.func @bool_vector_const() -> () "None" {
-    // CHECK: spv.constant dense<false> : vector<2xi1>
-    %0 = spv.constant dense<false> : vector<2xi1>
-    // CHECK: spv.constant dense<[true, true, true]> : vector<3xi1>
-    %1 = spv.constant dense<true> : vector<3xi1>
-    // CHECK: spv.constant dense<[false, true]> : vector<2xi1>
-    %2 = spv.constant dense<[false, true]> : vector<2xi1>
+    // CHECK: spv.Constant dense<false> : vector<2xi1>
+    %0 = spv.Constant dense<false> : vector<2xi1>
+    // CHECK: spv.Constant dense<[true, true, true]> : vector<3xi1>
+    %1 = spv.Constant dense<true> : vector<3xi1>
+    // CHECK: spv.Constant dense<[false, true]> : vector<2xi1>
+    %2 = spv.Constant dense<[false, true]> : vector<2xi1>
 
     %3 = spv.Variable init(%0): !spv.ptr<vector<2xi1>, Function>
     %4 = spv.Variable init(%1): !spv.ptr<vector<3xi1>, Function>
@@ -146,12 +146,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @int_vector_const
   spv.func @int_vector_const() -> () "None" {
-    // CHECK: spv.constant dense<0> : vector<3xi32>
-    %0 = spv.constant dense<0> : vector<3xi32>
-    // CHECK: spv.constant dense<1> : vector<3xi32>
-    %1 = spv.constant dense<1> : vector<3xi32>
-    // CHECK: spv.constant dense<[2, -3, 4]> : vector<3xi32>
-    %2 = spv.constant dense<[2, -3, 4]> : vector<3xi32>
+    // CHECK: spv.Constant dense<0> : vector<3xi32>
+    %0 = spv.Constant dense<0> : vector<3xi32>
+    // CHECK: spv.Constant dense<1> : vector<3xi32>
+    %1 = spv.Constant dense<1> : vector<3xi32>
+    // CHECK: spv.Constant dense<[2, -3, 4]> : vector<3xi32>
+    %2 = spv.Constant dense<[2, -3, 4]> : vector<3xi32>
 
     %3 = spv.IAdd %0, %1 : vector<3xi32>
     %4 = spv.IAdd %2, %3 : vector<3xi32>
@@ -160,12 +160,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @fp_vector_const
   spv.func @fp_vector_const() -> () "None" {
-    // CHECK: spv.constant dense<0.000000e+00> : vector<4xf32>
-    %0 = spv.constant dense<0.> : vector<4xf32>
-    // CHECK: spv.constant dense<-1.500000e+01> : vector<4xf32>
-    %1 = spv.constant dense<-15.> : vector<4xf32>
-    // CHECK: spv.constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32>
-    %2 = spv.constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32>
+    // CHECK: spv.Constant dense<0.000000e+00> : vector<4xf32>
+    %0 = spv.Constant dense<0.> : vector<4xf32>
+    // CHECK: spv.Constant dense<-1.500000e+01> : vector<4xf32>
+    %1 = spv.Constant dense<-15.> : vector<4xf32>
+    // CHECK: spv.Constant dense<[7.500000e-01, -2.500000e-01, 1.000000e+01, 4.200000e+01]> : vector<4xf32>
+    %2 = spv.Constant dense<[0.75, -0.25, 10., 42.]> : vector<4xf32>
 
     %3 = spv.FAdd %0, %1 : vector<4xf32>
     %4 = spv.FAdd %2, %3 : vector<4xf32>
@@ -174,51 +174,51 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @ui64_array_const
   spv.func @ui64_array_const() -> (!spv.array<3xui64>) "None" {
-    // CHECK: spv.constant [5, 6, 7] : !spv.array<3 x i64>
-    %0 = spv.constant [5 : ui64, 6 : ui64, 7 : ui64] : !spv.array<3 x ui64>
+    // CHECK: spv.Constant [5, 6, 7] : !spv.array<3 x i64>
+    %0 = spv.Constant [5 : ui64, 6 : ui64, 7 : ui64] : !spv.array<3 x ui64>
 
     spv.ReturnValue %0: !spv.array<3xui64>
   }
 
   // CHECK-LABEL: @si32_array_const
   spv.func @si32_array_const() -> (!spv.array<3xsi32>) "None" {
-    // CHECK: spv.constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
-    %0 = spv.constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
+    // CHECK: spv.Constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
+    %0 = spv.Constant [5 : si32, 6 : si32, 7 : si32] : !spv.array<3 x si32>
 
     spv.ReturnValue %0 : !spv.array<3xsi32>
   }
   // CHECK-LABEL: @float_array_const
   spv.func @float_array_const() -> (!spv.array<2 x vector<2xf32>>) "None" {
-    // CHECK: spv.constant [dense<3.000000e+00> : vector<2xf32>, dense<[4.000000e+00, 5.000000e+00]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
-    %0 = spv.constant [dense<3.0> : vector<2xf32>, dense<[4., 5.]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
+    // CHECK: spv.Constant [dense<3.000000e+00> : vector<2xf32>, dense<[4.000000e+00, 5.000000e+00]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
+    %0 = spv.Constant [dense<3.0> : vector<2xf32>, dense<[4., 5.]> : vector<2xf32>] : !spv.array<2 x vector<2xf32>>
 
     spv.ReturnValue %0 : !spv.array<2 x vector<2xf32>>
   }
 
   // CHECK-LABEL: @ignore_not_used_const
   spv.func @ignore_not_used_const() -> () "None" {
-    %0 = spv.constant false
+    %0 = spv.Constant false
     // CHECK-NEXT: spv.Return
     spv.Return
   }
 
   // CHECK-LABEL: @materialize_const_at_each_use
   spv.func @materialize_const_at_each_use() -> (i32) "None" {
-    // CHECK: %[[USE1:.*]] = spv.constant 42 : i32
-    // CHECK: %[[USE2:.*]] = spv.constant 42 : i32
+    // CHECK: %[[USE1:.*]] = spv.Constant 42 : i32
+    // CHECK: %[[USE2:.*]] = spv.Constant 42 : i32
     // CHECK: spv.IAdd %[[USE1]], %[[USE2]]
-    %0 = spv.constant 42 : i32
+    %0 = spv.Constant 42 : i32
     %1 = spv.IAdd %0, %0 : i32
     spv.ReturnValue %1 : i32
   }
 
   // CHECK-LABEL: @const_variable
   spv.func @const_variable(%arg0 : i32, %arg1 : i32) -> () "None" {
-    // CHECK: %[[CONST:.*]] = spv.constant 5 : i32
+    // CHECK: %[[CONST:.*]] = spv.Constant 5 : i32
     // CHECK: spv.Variable init(%[[CONST]]) : !spv.ptr<i32, Function>
     // CHECK: spv.IAdd %arg0, %arg1
     %0 = spv.IAdd %arg0, %arg1 : i32
-    %1 = spv.constant 5 : i32
+    %1 = spv.Constant 5 : i32
     %2 = spv.Variable init(%1) : !spv.ptr<i32, Function>
     %3 = spv.Load "Function" %2 : i32
     %4 = spv.IAdd %0, %3 : i32
@@ -227,15 +227,15 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @multi_dimensions_const
   spv.func @multi_dimensions_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" {
-    // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
-    %0 = spv.constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    // CHECK: spv.Constant {{\[}}{{\[}}[1 : i32, 2 : i32, 3 : i32], [4 : i32, 5 : i32, 6 : i32]], {{\[}}[7 : i32, 8 : i32, 9 : i32], [10 : i32, 11 : i32, 12 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    %0 = spv.Constant dense<[[[1, 2, 3], [4, 5, 6]], [[7, 8, 9], [10, 11, 12]]]> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
     spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
   }
 
   // CHECK-LABEL: @multi_dimensions_splat_const
   spv.func @multi_dimensions_splat_const() -> (!spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>) "None" {
-    // CHECK: spv.constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
-    %0 = spv.constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    // CHECK: spv.Constant {{\[}}{{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]], {{\[}}[1 : i32, 1 : i32, 1 : i32], [1 : i32, 1 : i32, 1 : i32]]] : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
+    %0 = spv.Constant dense<1> : tensor<2x2x3xi32> : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
     spv.ReturnValue %0 : !spv.array<2 x !spv.array<2 x !spv.array<3 x i32, stride=4>, stride=12>, stride=24>
   }
 }
index 3f8c4bf..bf7deeb 100644 (file)
@@ -94,7 +94,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [CooperativeMatrixNV], [SPV_N
 
   // CHECK-LABEL: @cooperative_matrix_access_chain
   spv.func @cooperative_matrix_access_chain(%a : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>) -> !spv.ptr<f32, Function> "None" {
-    %0 = spv.constant 0: i32
+    %0 = spv.Constant 0: i32
     // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
     %1 = spv.AccessChain %a[%0] : !spv.ptr<!spv.coopmatrix<8x16xf32, Subgroup>, Function>, i32
     spv.ReturnValue %1 : !spv.ptr<f32, Function>
index 1b4abef..88fb919 100644 (file)
@@ -44,7 +44,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   }
 
   spv.func @local_var() "None" {
-    %zero = spv.constant 0: i32
+    %zero = spv.Constant 0: i32
     // CHECK: loc({{".*debug.mlir"}}:49:12)
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     spv.Return
@@ -68,8 +68,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   }
 
   spv.func @loop(%count : i32) -> () "None" {
-    %zero = spv.constant 0: i32
-    %one = spv.constant 1: i32
+    %zero = spv.Constant 0: i32
+    %one = spv.Constant 1: i32
     %ivar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     %jvar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     spv.loop {
@@ -121,9 +121,9 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   }
 
   spv.func @selection(%cond: i1) -> () "None" {
-    %zero = spv.constant 0: i32
-    %one = spv.constant 1: i32
-    %two = spv.constant 2: i32
+    %zero = spv.Constant 0: i32
+    %one = spv.Constant 1: i32
+    %two = spv.Constant 2: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     spv.selection {
       // CHECK: loc({{".*debug.mlir"}}:128:5)
index 04d375a..92be451 100644 (file)
@@ -3,7 +3,7 @@
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @var1 : !spv.ptr<!spv.array<4xf32>, Input>
   spv.func @fmain() -> i32 "None" {
-    %0 = spv.constant 16 : i32
+    %0 = spv.Constant 16 : i32
     %1 = spv.mlir.addressof @var1 : !spv.ptr<!spv.array<4xf32>, Input>
     // CHECK: {{%.*}} = spv.FunctionCall @f_0({{%.*}}) : (i32) -> i32
     %3 = spv.FunctionCall @f_0(%0) : (i32) -> i32
@@ -24,7 +24,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   }
 
   spv.func @f_loop_with_function_call(%count : i32) -> () "None" {
-    %zero = spv.constant 0: i32
+    %zero = spv.Constant 0: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     spv.loop {
       spv.Branch ^header
@@ -44,7 +44,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
     spv.Return
   }
   spv.func @f_inc(%arg0 : !spv.ptr<i32, Function>) -> () "None" {
-      %one = spv.constant 1 : i32
+      %one = spv.Constant 1 : i32
       %0 = spv.Load "Function" %arg0 : i32
       %1 = spv.IAdd %0, %one : i32
       spv.Store "Function" %arg0, %1 : i32
index 1cdba58..16c1061 100644 (file)
@@ -28,7 +28,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() "None" {
     // CHECK: %[[ADDR:.*]] = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
     %0 = spv.mlir.addressof @globalInvocationID : !spv.ptr<vector<3xi32>, Input>
-    %1 = spv.constant 0: i32
+    %1 = spv.Constant 0: i32
     // CHECK: spv.AccessChain %[[ADDR]]
     %2 = spv.AccessChain %0[%1] : !spv.ptr<vector<3xi32>, Input>, i32
     spv.Return
index 7be965d..6a5d532 100644 (file)
@@ -93,16 +93,16 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.SpecConstant @condition_scalar = true
   spv.func @select() -> () "None" {
-    %0 = spv.constant 4.0 : f32
-    %1 = spv.constant 5.0 : f32
+    %0 = spv.Constant 4.0 : f32
+    %1 = spv.Constant 5.0 : f32
     %2 = spv.mlir.referenceof @condition_scalar : i1
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, f32
     %3 = spv.Select %2, %0, %1 : i1, f32
-    %4 = spv.constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32>
-    %5 = spv.constant dense<[6.0, 7.0, 8.0, 9.0]> : vector<4xf32>
+    %4 = spv.Constant dense<[2.0, 3.0, 4.0, 5.0]> : vector<4xf32>
+    %5 = spv.Constant dense<[6.0, 7.0, 8.0, 9.0]> : vector<4xf32>
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : i1, vector<4xf32>
     %6 = spv.Select %2, %4, %5 : i1, vector<4xf32>
-    %7 = spv.constant dense<[true, true, true, true]> : vector<4xi1>
+    %7 = spv.Constant dense<[true, true, true, true]> : vector<4xi1>
     // CHECK: spv.Select {{.*}}, {{.*}}, {{.*}} : vector<4xi1>, vector<4xf32>
     %8 = spv.Select %7, %4, %5 : vector<4xi1>, vector<4xf32>
     spv.Return
index d54754e..a18a420 100644 (file)
@@ -5,8 +5,8 @@
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   // for (int i = 0; i < count; ++i) {}
   spv.func @loop(%count : i32) -> () "None" {
-    %zero = spv.constant 0: i32
-    %one = spv.constant 1: i32
+    %zero = spv.Constant 0: i32
+    %one = spv.Constant 1: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
 // CHECK:        spv.Branch ^bb1
@@ -35,7 +35,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
     ^continue:
 // CHECK-NEXT:     spv.Load
       %val1 = spv.Load "Function" %var : i32
-// CHECK-NEXT:     spv.constant 1
+// CHECK-NEXT:     spv.Constant 1
 // CHECK-NEXT:     spv.IAdd
       %add = spv.IAdd %val1, %one : i32
 // CHECK-NEXT:     spv.Store
@@ -64,13 +64,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @GV2 bind(0, 1) : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
   spv.func @loop_kernel() "None" {
     %0 = spv.mlir.addressof @GV1 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
-    %1 = spv.constant 0 : i32
+    %1 = spv.Constant 0 : i32
     %2 = spv.AccessChain %0[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
     %3 = spv.mlir.addressof @GV2 : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>
     %5 = spv.AccessChain %3[%1] : !spv.ptr<!spv.struct<(!spv.array<10 x f32, stride=4> [0])>, StorageBuffer>, i32
-    %6 = spv.constant 4 : i32
-    %7 = spv.constant 42 : i32
-    %8 = spv.constant 2 : i32
+    %6 = spv.Constant 4 : i32
+    %7 = spv.Constant 42 : i32
+    %8 = spv.Constant 2 : i32
 // CHECK:        spv.Branch ^bb1(%{{.*}} : i32)
 // CHECK-NEXT: ^bb1(%[[OUTARG:.*]]: i32):
 // CHECK-NEXT:   spv.loop {
@@ -112,8 +112,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   //   for (int j = 0; j < count; ++j) { }
   // }
   spv.func @loop(%count : i32) -> () "None" {
-    %zero = spv.constant 0: i32
-    %one = spv.constant 1: i32
+    %zero = spv.Constant 0: i32
+    %one = spv.Constant 1: i32
     %ivar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
     %jvar = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
@@ -135,7 +135,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 // CHECK-NEXT:   ^bb2:
     ^body:
-// CHECK-NEXT:     spv.constant 0
+// CHECK-NEXT:     spv.Constant 0
 // CHECK-NEXT:                  spv.Store
       spv.Store "Function" %jvar, %zero : i32
 // CHECK-NEXT:     spv.Branch ^bb3
@@ -164,7 +164,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
       ^continue:
 // CHECK-NEXT:       spv.Load
         %jval1 = spv.Load "Function" %jvar : i32
-// CHECK-NEXT:       spv.constant 1
+// CHECK-NEXT:       spv.Constant 1
 // CHECK-NEXT:       spv.IAdd
         %add = spv.IAdd %jval1, %one : i32
 // CHECK-NEXT:       spv.Store
@@ -185,7 +185,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
     ^continue:
 // CHECK-NEXT:     spv.Load
       %ival1 = spv.Load "Function" %ivar : i32
-// CHECK-NEXT:     spv.constant 1
+// CHECK-NEXT:     spv.Constant 1
 // CHECK-NEXT:     spv.IAdd
       %add = spv.IAdd %ival1, %one : i32
 // CHECK-NEXT:     spv.Store
index 8bf0217..8730044 100644 (file)
@@ -30,13 +30,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @load_store_zero_rank_float(%arg0: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>) "None" {
     // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>
     // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : f32
-    %0 = spv.constant 0 : i32
+    %0 = spv.Constant 0 : i32
     %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
     %2 = spv.Load "StorageBuffer" %1 : f32
 
     // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>
     // CHECK-NEXT: spv.Store "StorageBuffer" [[STORE_PTR]], [[VAL]] : f32
-    %3 = spv.constant 0 : i32
+    %3 = spv.Constant 0 : i32
     %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<(!spv.array<1 x f32, stride=4> [0])>, StorageBuffer>, i32, i32
     spv.Store "StorageBuffer" %4, %2 : f32
     spv.Return
@@ -45,13 +45,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @load_store_zero_rank_int(%arg0: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, %arg1: !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>) "None" {
     // CHECK: [[LOAD_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>
     // CHECK-NEXT: [[VAL:%.*]] = spv.Load "StorageBuffer" [[LOAD_PTR]] : i32
-    %0 = spv.constant 0 : i32
+    %0 = spv.Constant 0 : i32
     %1 = spv.AccessChain %arg0[%0, %0] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
     %2 = spv.Load "StorageBuffer" %1 : i32
 
     // CHECK: [[STORE_PTR:%.*]] = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>
     // CHECK-NEXT: spv.Store "StorageBuffer" [[STORE_PTR]], [[VAL]] : i32
-    %3 = spv.constant 0 : i32
+    %3 = spv.Constant 0 : i32
     %4 = spv.AccessChain %arg1[%3, %3] : !spv.ptr<!spv.struct<(!spv.array<1 x i32, stride=4> [0])>, StorageBuffer>, i32, i32
     spv.Store "StorageBuffer" %4, %2 : i32
     spv.Return
index 5cb035f..e0d576f 100644 (file)
@@ -10,7 +10,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @group_non_uniform_broadcast
   spv.func @group_non_uniform_broadcast(%value: f32) -> f32 "None" {
-    %one = spv.constant 1 : i32
+    %one = spv.Constant 1 : i32
     // CHECK: spv.GroupNonUniformBroadcast Subgroup %{{.*}}, %{{.*}} : f32, i32
     %0 = spv.GroupNonUniformBroadcast Subgroup %value, %one : f32, i32
     spv.ReturnValue %0: f32
@@ -60,7 +60,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // CHECK-LABEL: @group_non_uniform_iadd_clustered_reduce
   spv.func @group_non_uniform_iadd_clustered_reduce(%val: vector<2xi32>) -> vector<2xi32> "None" {
-    %four = spv.constant 4 : i32
+    %four = spv.Constant 4 : i32
     // CHECK: %{{.+}} = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %{{.+}} cluster_size(%{{.+}}) : vector<2xi32>
     %0 = spv.GroupNonUniformIAdd "Workgroup" "ClusteredReduce" %val cluster_size(%four) : vector<2xi32>
     spv.ReturnValue %0: vector<2xi32>
index bb94ecc..ba320b3 100644 (file)
@@ -4,8 +4,8 @@
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
-// CHECK:        %[[CST:.*]] = spv.constant 0
-    %zero = spv.constant 0 : i32
+// CHECK:        %[[CST:.*]] = spv.Constant 0
+    %zero = spv.Constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[CST]] : i32)
     spv.Branch ^bb1(%zero : i32)
 // CHECK-NEXT: ^bb1(%{{.*}}: i32):
@@ -25,10 +25,10 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
-// CHECK:        %[[ZERO:.*]] = spv.constant 0
-    %zero = spv.constant 0 : i32
-// CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
-    %one = spv.constant 1.0 : f32
+// CHECK:        %[[ZERO:.*]] = spv.Constant 0
+    %zero = spv.Constant 0 : i32
+// CHECK-NEXT:   %[[ONE:.*]] = spv.Constant 1
+    %one = spv.Constant 1.0 : f32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[ZERO]], %[[ONE]] : i32, f32)
     spv.Branch ^bb1(%zero, %one : i32, f32)
 
@@ -49,8 +49,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @foo() -> () "None" {
-// CHECK:        %[[CST0:.*]] = spv.constant 0
-    %zero = spv.constant 0 : i32
+// CHECK:        %[[CST0:.*]] = spv.Constant 0
+    %zero = spv.Constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb1(%[[CST0]] : i32)
     spv.Branch ^bb1(%zero : i32)
 
@@ -58,7 +58,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   ^bb1(%arg0: i32):
 // CHECK-NEXT:   %[[ADD:.*]] = spv.IAdd %[[ARG]], %[[ARG]] : i32
     %0 = spv.IAdd %arg0, %arg0 : i32
-// CHECK-NEXT:   %[[CST1:.*]] = spv.constant 0
+// CHECK-NEXT:   %[[CST1:.*]] = spv.Constant 0
 // CHECK-NEXT:   spv.Branch ^bb2(%[[CST1]], %[[ADD]] : i32, i32)
     spv.Branch ^bb2(%zero, %0 : i32, i32)
 
@@ -83,8 +83,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
     spv.Branch ^bb1
 
 // CHECK-NEXT: ^bb1:
-// CHECK-NEXT:   %[[ZERO:.*]] = spv.constant 0
-// CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
+// CHECK-NEXT:   %[[ZERO:.*]] = spv.Constant 0
+// CHECK-NEXT:   %[[ONE:.*]] = spv.Constant 1
 // CHECK-NEXT:   spv.Branch ^bb2(%[[ZERO]], %[[ONE]] : i32, f32)
 
 // CHECK-NEXT: ^bb2(%{{.*}}: i32, %{{.*}}: f32):
@@ -94,8 +94,8 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   // This block is reordered to follow domination order.
   ^bb1:
-    %zero = spv.constant 0 : i32
-    %one = spv.constant 1.0 : f32
+    %zero = spv.Constant 0 : i32
+    %one = spv.Constant 1.0 : f32
     spv.Branch ^bb2(%zero, %one : i32, f32)
   }
 
@@ -115,21 +115,21 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 // CHECK:      spv.selection
     spv.selection {
-      %true = spv.constant true
+      %true = spv.Constant true
 // CHECK:        spv.BranchConditional %{{.*}}, ^bb1, ^bb2
       spv.BranchConditional %true, ^true, ^false
 
 // CHECK-NEXT: ^bb1:
     ^true:
-// CHECK-NEXT:   %[[ZERO:.*]] = spv.constant 0
-      %zero = spv.constant 0 : i32
+// CHECK-NEXT:   %[[ZERO:.*]] = spv.Constant 0
+      %zero = spv.Constant 0 : i32
 // CHECK-NEXT:   spv.Branch ^bb3(%[[ZERO]] : i32)
       spv.Branch ^phi(%zero: i32)
 
 // CHECK-NEXT: ^bb2:
     ^false:
-// CHECK-NEXT:   %[[ONE:.*]] = spv.constant 1
-      %one = spv.constant 1 : i32
+// CHECK-NEXT:   %[[ONE:.*]] = spv.Constant 1
+      %one = spv.Constant 1 : i32
 // CHECK-NEXT:   spv.Branch ^bb3(%[[ONE]] : i32)
       spv.Branch ^phi(%one: i32)
 
@@ -162,9 +162,9 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.globalVariable @__builtin_var_NumWorkgroups__ built_in("NumWorkgroups") : !spv.ptr<vector<3xi32>, Input>
   spv.globalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi32>, Input>
   spv.func @fmul_kernel() "None" {
-    %3 = spv.constant 12 : i32
-    %4 = spv.constant 32 : i32
-    %5 = spv.constant 4 : i32
+    %3 = spv.Constant 12 : i32
+    %4 = spv.Constant 32 : i32
+    %5 = spv.Constant 4 : i32
     %6 = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi32>, Input>
     %7 = spv.Load "Input" %6 : vector<3xi32>
     %8 = spv.CompositeExtract %7[0 : i32] : vector<3xi32>
@@ -243,12 +243,12 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @fmul_kernel() "None" {
-    %cst4 = spv.constant 4 : i32
+    %cst4 = spv.Constant 4 : i32
 
-    %val1 = spv.constant 43 : i32
-    %val2 = spv.constant 44 : i32
+    %val1 = spv.Constant 43 : i32
+    %val2 = spv.Constant 44 : i32
 
-// CHECK:        spv.constant 43
+// CHECK:        spv.Constant 43
 // CHECK-NEXT:   spv.Branch ^[[BB1:.+]](%{{.+}} : i32)
 // CHECK-NEXT: ^[[BB1]](%{{.+}}: i32):
 // CHECK-NEXT:   spv.loop
@@ -264,7 +264,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
       spv.mlir.merge
     }
 
-// CHECK:        spv.constant 44
+// CHECK:        spv.Constant 44
 // CHECK-NEXT:   spv.Branch ^[[BB2:.+]](%{{.+}} : i32)
 // CHECK-NEXT: ^[[BB2]](%{{.+}}: i32):
 // CHECK-NEXT:   spv.loop
index e05f4a1..31dfcbc 100644 (file)
@@ -6,13 +6,13 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
   spv.func @selection(%cond: i1) -> () "None" {
 // CHECK:        spv.Branch ^bb1
 // CHECK-NEXT: ^bb1:
-    %zero = spv.constant 0: i32
-    %one = spv.constant 1: i32
-    %two = spv.constant 2: i32
+    %zero = spv.Constant 0: i32
+    %one = spv.Constant 1: i32
+    %two = spv.Constant 2: i32
     %var = spv.Variable init(%zero) : !spv.ptr<i32, Function>
 
 // CHECK-NEXT:   spv.selection control(Flatten)
-// CHECK-NEXT:     spv.constant 0
+// CHECK-NEXT:     spv.Constant 0
 // CHECK-NEXT:     spv.Variable
     spv.selection control(Flatten) {
 // CHECK-NEXT: spv.BranchConditional %{{.*}} [5, 10], ^bb1, ^bb2
@@ -20,7 +20,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 // CHECK-NEXT:   ^bb1:
     ^then:
-// CHECK-NEXT:     spv.constant 1
+// CHECK-NEXT:     spv.Constant 1
 // CHECK-NEXT:     spv.Store
       spv.Store "Function" %var, %one : i32
 // CHECK-NEXT:     spv.Branch ^bb3
@@ -28,7 +28,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 // CHECK-NEXT:   ^bb2:
     ^else:
-// CHECK-NEXT:     spv.constant 2
+// CHECK-NEXT:     spv.Constant 2
 // CHECK-NEXT:     spv.Store
       spv.Store "Function" %var, %two : i32
 // CHECK-NEXT:     spv.Branch ^bb3
@@ -67,7 +67,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
 // CHECK:        ^bb1:
     ^then:
-      %zero = spv.constant 0 : i32
+      %zero = spv.Constant 0 : i32
       spv.ReturnValue  %zero : i32
 
 // CHECK:        ^bb2:
@@ -76,7 +76,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
       spv.mlir.merge
     }
 
-    %one = spv.constant 1 : i32
+    %one = spv.Constant 1 : i32
     spv.ReturnValue  %one : i32
   }
 
index 70b6e32..a69834b 100644 (file)
@@ -94,17 +94,17 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
 
   spv.func @use_composite() -> (i32) "None" {
     // CHECK: [[USE1:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32
-    // CHECK: [[USE2:%.*]] = spv.constant 0 : i32
+    // CHECK: [[USE2:%.*]] = spv.Constant 0 : i32
 
     // CHECK: [[RES1:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE1]], [[USE2]]) : (i32, i32) -> i32
 
     // CHECK: [[USE3:%.*]] = spv.mlir.referenceof @sc_i32_1 : i32
-    // CHECK: [[USE4:%.*]] = spv.constant 0 : i32
+    // CHECK: [[USE4:%.*]] = spv.Constant 0 : i32
 
     // CHECK: [[RES2:%.*]] = spv.SpecConstantOperation wraps "spv.ISub"([[USE3]], [[USE4]]) : (i32, i32) -> i32
 
     %0 = spv.mlir.referenceof @sc_i32_1 : i32
-    %1 = spv.constant 0 : i32
+    %1 = spv.Constant 0 : i32
     %2 = spv.SpecConstantOperation wraps "spv.ISub"(%0, %1) : (i32, i32) -> i32
 
     // CHECK: [[RES3:%.*]] = spv.SpecConstantOperation wraps "spv.IMul"([[RES1]], [[RES2]]) : (i32, i32) -> i32
index 6a287b2..efea010 100644 (file)
@@ -15,7 +15,7 @@ spv.module Logical GLSL450 requires #spv.vce<v1.0, [Shader], []> {
     %6 = spv.CompositeExtract %5[1 : i32, 2 : i32] : !spv.array<4x!spv.array<4xi32>>
     // CHECK: {{%.*}} = spv.undef : !spv.ptr<!spv.struct<(f32)>, StorageBuffer>
     %7 = spv.undef : !spv.ptr<!spv.struct<(f32)>, StorageBuffer>
-    %8 = spv.constant 0 : i32
+    %8 = spv.Constant 0 : i32
     %9 = spv.AccessChain %7[%8] : !spv.ptr<!spv.struct<(f32)>, StorageBuffer>, i32
     spv.Return
   }