spirv-reader: spec op iadd, isub, imul
Fixed: 402726681 402726207 402726051
Change-Id: If0c9ba13bcfda7b15688b7c72f5a70c7aec8d3a5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/249595
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index e963325..65af473 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -378,6 +378,15 @@
case spv::Op::OpSNegate:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSNegate, 3);
break;
+ case spv::Op::OpIAdd:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kAdd, 3);
+ break;
+ case spv::Op::OpISub:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSub, 3);
+ break;
+ case spv::Op::OpIMul:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kMul, 3);
+ break;
default:
TINT_ICE() << "Unknown spec constant operation: " << op;
}
@@ -3295,6 +3304,7 @@
Emit(binary, inst.result_id());
}
+ /// Emits the logical negation of the result of the given SPIR-V instruction.
/// @param inst the SPIR-V instruction
/// @param op the binary operator to use
void EmitInvertedBinary(const spvtools::opt::Instruction& inst, core::BinaryOp op) {
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index c3b3cda..e4eeb04 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -1273,7 +1273,7 @@
)");
}
-TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Not) {
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Not_ResultI32_i32) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -1309,6 +1309,117 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Not_ResultI32_u32) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %false = OpSpecConstantFalse %bool
+ %true = OpSpecConstantTrue %bool
+ %one = OpConstant %u32 1
+ %myconst = OpSpecConstantOp %i32 Not %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.not<i32> 1u
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Not_ResultU32_i32) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %false = OpSpecConstantFalse %bool
+ %true = OpSpecConstantTrue %bool
+ %one = OpConstant %i32 1
+ %myconst = OpSpecConstantOp %u32 Not %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.not<u32> 1i
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_Not_ResultU32_u32) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %false = OpSpecConstantFalse %bool
+ %true = OpSpecConstantTrue %bool
+ %one = OpConstant %u32 1
+ %myconst = OpSpecConstantOp %u32 Not %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.not<u32> 1u
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
TEST_F(SpirvParserTest, Var_OpSpecConstantOp_FConvert) {
EXPECT_IR(R"(
OpCapability Shader
@@ -1344,10 +1455,9 @@
)");
}
-TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SNegate) {
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SNegate_ResultI32_i32) {
EXPECT_IR(R"(
OpCapability Shader
- OpCapability Float16
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
@@ -1378,6 +1488,537 @@
)");
}
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SNegate_ResultI32_u32) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %u32 1
+ %myconst = OpSpecConstantOp %i32 SNegate %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.s_negate<i32> 1u
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SNegate_ResultU32_i32) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %myconst = OpSpecConstantOp %u32 SNegate %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.s_negate<u32> 1i
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_SNegate_ResultU32_u32) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %u32 1
+ %myconst = OpSpecConstantOp %u32 SNegate %one
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.s_negate<u32> 1u
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IAdd_ResultI32_i32_i32) {
+ // All signed
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %i32 IAdd %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.add<i32> 1i, 2i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IAdd_ResultI32_i32_u32) {
+ // Last arg different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %i32 IAdd %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.add<i32> 1i, 2u
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IAdd_ResultI32_u32_i32) {
+ // Second arg different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %u32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %i32 IAdd %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.add<i32> 1u, 2i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IAdd_ResultU32_i32_i32) {
+ // Result different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %u32 IAdd %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.add<u32> 1i, 2i
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ISub_ResultI32_i32_i32) {
+ // All signed
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %i32 ISub %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.sub<i32> 1i, 2i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ISub_ResultI32_i32_u32) {
+ // Last arg different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %i32 ISub %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.sub<i32> 1i, 2u
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ISub_ResultI32_u32_i32) {
+ // Second arg different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %u32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %i32 ISub %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.sub<i32> 1u, 2i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_ISub_ResultU32_i32_i32) {
+ // Result different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %u32 ISub %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.sub<u32> 1i, 2i
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IMul_ResultI32_i32_i32) {
+ // All signed
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %i32 IMul %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.mul<i32> 1i, 2i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IMul_ResultI32_i32_u32) {
+ // Last arg different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %two = OpConstant %u32 2
+ %myconst = OpSpecConstantOp %i32 IMul %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.mul<i32> 1i, 2u
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IMul_ResultI32_u32_i32) {
+ // Second arg different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %u32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %i32 IMul %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %i32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:i32 = spirv.mul<i32> 1u, 2i
+ %myconst:i32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:i32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Var_OpSpecConstantOp_IMul_ResultU32_i32_i32) {
+ // Result different
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %myconst "myconst"
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %u32 = OpTypeInt 32 0
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %myconst = OpSpecConstantOp %u32 IMul %one %two
+ %voidfn = OpTypeFunction %void
+ %main = OpFunction %void None %voidfn
+ %main_entry = OpLabel
+ %1 = OpCopyObject %u32 %myconst
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:u32 = spirv.mul<u32> 1i, 2i
+ %myconst:u32 = override %1
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:u32 = let %myconst
+ ret
+ }
+}
+)");
+}
+
// In the case of all literals, SPIR-V opt treats the `OpSpecConstantComposite` as an
// `OpConstantComposite` so it appears in the constant manager already. This then needs no handling
// on our side.
