[spirv-reader][ir] Enable disabled tests. The remainder of the SPIR-V IR reader disabled tests can be enabled as all the needed support is available. Bug: 42250952 Change-Id: I6a0aa273a16ae14847d1e22a902f828cbd0db048 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/251199 Commit-Queue: dan sinclair <dsinclair@chromium.org> Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc index 8c55c24..1cee8c8 100644 --- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc +++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -658,7 +658,7 @@ )"); } -TEST_F(SpirvReaderTest, DISABLED_RectifyOperandsAndResult_FindUMsb) { +TEST_F(SpirvReaderTest, RectifyOperandsAndResult_FindUMsb) { // Check signedness conversion of arguments and results. // SPIR-V signed arg -> cast arg to unsigned // signed result -> cast result to signed @@ -693,14 +693,23 @@ R"( %main = @compute @workgroup_size(1u, 1u, 1u) func():void { $B1: { - let x_1 = bitcast<i32>(firstLeadingBit(bitcast<u32>(i1))); - let x_2 = bitcast<vec2i>(firstLeadingBit(bitcast<vec2u>(v2i1))); - let x_3 = firstLeadingBit(bitcast<u32>(i1)); - let x_4 = firstLeadingBit(bitcast<vec2u>(v2i1)); - let x_5 = bitcast<i32>(firstLeadingBit(u1)); - let x_6 = bitcast<vec2i>(firstLeadingBit(v2u1)); - let x_7 = firstLeadingBit(u1); - let x_8 = firstLeadingBit(v2u1); + %2:u32 = bitcast 30i + %3:u32 = firstLeadingBit %2 + %4:i32 = bitcast %3 + %5:vec2<u32> = bitcast vec2<i32>(30i, 40i) + %6:vec2<u32> = firstLeadingBit %5 + %7:vec2<i32> = bitcast %6 + %8:u32 = bitcast 30i + %9:u32 = firstLeadingBit %8 + %10:vec2<u32> = bitcast vec2<i32>(30i, 40i) + %11:vec2<u32> = firstLeadingBit %10 + %12:u32 = firstLeadingBit 10u + %13:i32 = bitcast %12 + %14:vec2<u32> = firstLeadingBit vec2<u32>(10u, 20u) + %15:vec2<i32> = bitcast %14 + %16:u32 = firstLeadingBit 10u + %17:vec2<u32> = firstLeadingBit vec2<u32>(10u, 20u) + ret } } )");
diff --git a/src/tint/lang/spirv/reader/lower/texture_test.cc b/src/tint/lang/spirv/reader/lower/texture_test.cc index 56941a0..481039c 100644 --- a/src/tint/lang/spirv/reader/lower/texture_test.cc +++ b/src/tint/lang/spirv/reader/lower/texture_test.cc
@@ -222,7 +222,7 @@ ASSERT_EQ(expect, str()); } -TEST_F(SpirvReader_TextureTest, DISABLED_Type_Image_DepthUnknown) { +TEST_F(SpirvReader_TextureTest, Type_Image_DepthUnknown) { b.Append(mod.root_block, [&] { auto* v = b.Var("wg", ty.ptr(handle, ty.Get<spirv::type::Image>(
diff --git a/src/tint/lang/spirv/reader/parser/atomics_test.cc b/src/tint/lang/spirv/reader/parser/atomics_test.cc index 28f51de..20d181f 100644 --- a/src/tint/lang/spirv/reader/parser/atomics_test.cc +++ b/src/tint/lang/spirv/reader/parser/atomics_test.cc
@@ -791,67 +791,6 @@ )"); } -// TODO(dsinclair): Requires support for variable pointers -TEST_F(SpirvParser_AtomicsTest, DISABLED_FunctionParam_subpointer) { - EXPECT_IR(R"( - OpCapability Shader - OpCapability VariablePointers - OpExtension "SPV_KHR_variable_pointers" - OpMemoryModel Logical GLSL450 - OpEntryPoint GLCompute %main "main" - OpExecutionMode %main LocalSize 1 1 1 - OpName %wg "wg" - OpName %main "main" - %int = OpTypeInt 32 1 - %uint = OpTypeInt 32 0 - %uint_0 = OpConstant %uint 0 - %uint_1 = OpConstant %uint 1 - %uint_2 = OpConstant %uint 2 - %uint_4 = OpConstant %uint 4 - %int_1 = OpConstant %int 1 - %arr = OpTypeArray %uint %uint_4 - %ptr_arr = OpTypePointer Workgroup %arr - %ptr_uint = OpTypePointer Workgroup %uint - %void = OpTypeVoid - %10 = OpTypeFunction %void - %11 = OpTypeFunction %void %ptr_uint - %wg = OpVariable %ptr_arr Workgroup - - %foo = OpFunction %void None %11 - %param = OpFunctionParameter %ptr_uint - %foo_start = OpLabel - OpAtomicStore %param %uint_2 %uint_0 %uint_1 - OpReturn - OpFunctionEnd - - %main = OpFunction %void None %10 - %45 = OpLabel - %42 = OpAccessChain %ptr_uint %wg %int_1 - %44 = OpFunctionCall %void %foo %42 - OpReturn - OpFunctionEnd -)", - R"( -$B1: { # root - %wg:ptr<workgroup, array<u32, 4>, read_write> = var undef -} - -%2 = func(%3:ptr<workgroup, array<u32, 4>, read_write>):void { - $B2: { - %4:ptr<workgroup, u32, read_write> = access %3, 1i - %5:void = spirv.atomic_store %4, 2u, 0u, 1u - ret - } -} -%main = @compute @workgroup_size(1u, 1u, 1u) func():void { - $B3: { - %7:void = call %2, %wg - ret - } -} -)"); -} - TEST_F(SpirvParser_AtomicsTest, AtomicAdd) { EXPECT_IR(R"( OpCapability Shader
diff --git a/src/tint/lang/spirv/reader/parser/import_test.cc b/src/tint/lang/spirv/reader/parser/import_test.cc index 0f7ac6a..df38ffb 100644 --- a/src/tint/lang/spirv/reader/parser/import_test.cc +++ b/src/tint/lang/spirv/reader/parser/import_test.cc
@@ -127,8 +127,7 @@ SPV_ENV_UNIVERSAL_1_3); } -// TODO(dsinclair): internal compiler error: TINT_UNIMPLEMENTED unhandled SPIR-V type: [uint32] -TEST_F(SpirvParserTest, DISABLED_Import_NonSemantic_IgnoredExtInsts) { +TEST_F(SpirvParserTest, Import_NonSemantic_IgnoredExtInsts) { // This is the clspv-compiled output of this OpenCL C: // kernel void foo(global int*A) { A=A; } // It emits NonSemantic.ClspvReflection.1 extended instructions. @@ -184,8 +183,20 @@ %28 = OpExtInst %void %20 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 )", R"( -%main = @compute @workgroup_size(1u, 1u, 1u) func():void { - $B1: { +tint_symbol_1 = struct @align(4) { + tint_symbol:array<u32> @offset(0) +} + +$B1: { # root + %1:u32 = override 1u @id(0) + %2:u32 = override 1u @id(1) + %3:u32 = override 1u @id(2) + %4:ptr<storage, tint_symbol_1, read_write> = var undef @binding_point(0, 0) +} + +%foo = @compute @workgroup_size(1u, 1u, 1u) func():void { + $B2: { + %6:ptr<storage, u32, read_write> = access %4, 0u, 0u ret } }
diff --git a/src/tint/lang/spirv/reader/parser/misc_test.cc b/src/tint/lang/spirv/reader/parser/misc_test.cc index b866b88..310e765 100644 --- a/src/tint/lang/spirv/reader/parser/misc_test.cc +++ b/src/tint/lang/spirv/reader/parser/misc_test.cc
@@ -275,8 +275,7 @@ )"); } -// TODO(dsinclair): Requires OpSelectionMerge and OpBranchConditional -TEST_F(SpirvParserTest, DISABLED_OpUnreachable_InsideIf) { +TEST_F(SpirvParserTest, OpUnreachable_InsideIf) { EXPECT_IR(R"( OpCapability Shader OpMemoryModel Logical GLSL450 @@ -299,16 +298,21 @@ R"( %main = @compute @workgroup_size(1u, 1u, 1u) func():void { $B1: { - %2:void = if true [t: $B2] { - unreachable + if true [t: $B2, f: $B3] { # if_1 + $B2: { # true + unreachable + } + $B3: { # false + exit_if # if_1 + } } + ret } } )"); } -// TODO(dsinclair): Requires OpBranch -TEST_F(SpirvParserTest, DISABLED_OpUnreachable_InsideLoop) { +TEST_F(SpirvParserTest, OpUnreachable_InsideLoop) { EXPECT_IR(R"( OpCapability Shader OpMemoryModel Logical GLSL450 @@ -335,9 +339,24 @@ R"( %main = @compute @workgroup_size(1u, 1u, 1u) func():void { $B1: { - %2:void = loop %true [] - unreachable + loop [b: $B2, c: $B3] { # loop_1 + $B2: { # body + %2:bool = or true, true + if %2 [t: $B4, f: $B5] { # if_1 + $B4: { # true + unreachable + } + $B5: { # false + unreachable + } + } + unreachable + } + $B3: { # continuing + next_iteration # -> $B2 + } } + ret } } )"); @@ -423,8 +442,7 @@ )"); } -// TODO(dsinclair): Requires OpSelectionMerge and OpBranchConditional -TEST_F(SpirvParserTest, DISABLED_OpKill_InsideIf) { +TEST_F(SpirvParserTest, OpKill_InsideIf) { EXPECT_IR(R"( OpCapability Shader OpMemoryModel Logical GLSL450 @@ -447,9 +465,14 @@ R"( %main = @fragment func():void { $B1: { - %2:void = if true [t: $B2] { - discard - ret + if true [t: $B2, f: $B3] { # if_1 + $B2: { # true + discard + ret + } + $B3: { # false + exit_if # if_1 + } } discard ret @@ -458,8 +481,7 @@ )"); } -// TODO(dsinclair): Requires OpBranch -TEST_F(SpirvParserTest, DISABLED_OpKill_InsideLoop) { +TEST_F(SpirvParserTest, OpKill_InsideLoop) { EXPECT_IR(R"( OpCapability Shader OpMemoryModel Logical GLSL450 @@ -486,9 +508,23 @@ R"( %main = @fragment func():void { $B1: { - %2:void = loop %true [] - discard - ret + loop [b: $B2, c: $B3] { # loop_1 + $B2: { # body + %2:bool = or true, true + if %2 [t: $B4, f: $B5] { # if_1 + $B4: { # true + discard + ret + } + $B5: { # false + unreachable + } + } + unreachable + } + $B3: { # continuing + next_iteration # -> $B2 + } } discard ret