wgpu icon indicating copy to clipboard operation
wgpu copied to clipboard

SPIRV validation error in test with double continue switches

Open Imberflur opened this issue 1 year ago • 0 comments

I ran into a SPIRV validation issue in a test case I added in my PR https://github.com/gfx-rs/wgpu/pull/5654#issuecomment-2092368818. I've been able to reduce it down to:

This input wgsl

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    loop {
        switch 0 {
            default: {
                continue;
            }
        }
        switch 0 {
            default: {
                continue;
            }
        }
    }
}

This output spirv

; SPIR-V
; Version: 1.6
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 22
; Schema: 0
               OpCapability Shader
          %1 = OpExtInstImport "GLSL.std.450"
               OpMemoryModel Logical GLSL450
               OpEntryPoint GLCompute %2 "main" %gl_GlobalInvocationID
               OpExecutionMode %2 LocalSize 1 1 1
               OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
       %void = OpTypeVoid
       %uint = OpTypeInt 32 0
     %v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
          %8 = OpTypeFunction %void
        %int = OpTypeInt 32 1
      %int_0 = OpConstant %int 0
          %2 = OpFunction %void None %8
         %11 = OpLabel
         %12 = OpLoad %v3uint %gl_GlobalInvocationID
               OpBranch %13
         %13 = OpLabel
               OpBranch %14
         %14 = OpLabel
               OpLoopMerge %15 %16 None
               OpBranch %17
         %17 = OpLabel
               OpSelectionMerge %18 None
               OpSwitch %int_0 %19
         %19 = OpLabel
               OpBranch %16
         %18 = OpLabel
               OpSelectionMerge %20 None
               OpSwitch %int_0 %21
         %21 = OpLabel
               OpBranch %16
         %20 = OpLabel
               OpBranch %16
         %16 = OpLabel
               OpBranch %14
         %15 = OpLabel
               OpReturn
               OpFunctionEnd

This error:

error: line 32: Case construct that targets '21[%21]' has invalid branch to block '16[%16]' (not another case construct, corresponding merge, outer loop merge or outer loop continue)
  %21 = OpLabel

line 32 is the OpBranch %16 in %21 = OpLabel

Afaict %16 is actually an outer loop continue, is this a spirv validator bug?

Imberflur avatar May 04 '24 06:05 Imberflur