| SKIP: FAILED |
| |
| ../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: theSSBO = struct @align(4) { |
| out_data:i32 @offset(0) |
| } |
| |
| buf1 = struct @align(8) { |
| injectionSwitch:vec2<f32> @offset(0) |
| } |
| |
| $B1: { # root |
| %x_4:ptr<storage, theSSBO, read_write> = var @binding_point(0, 0) |
| %x_6:ptr<uniform, buf1, read> = var @binding_point(0, 1) |
| } |
| |
| %main_1 = func():void { |
| $B2: { |
| %i:ptr<function, i32, read_write> = var |
| %GLF_live3s:ptr<function, f32, read_write> = var |
| %i_1:ptr<function, i32, read_write> = var |
| %z:ptr<function, i32, read_write> = var |
| %8:ptr<storage, i32, read_write> = access %x_4, 0u |
| store %8, 42i |
| %9:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %10:f32 = load_vector_element %9, 0u |
| %x_37:f32 = let %10 |
| %12:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %13:f32 = load_vector_element %12, 1u |
| %x_39:f32 = let %13 |
| %15:bool = gt %x_37, %x_39 |
| if %15 [t: $B3] { # if_1 |
| $B3: { # true |
| %16:void = msl.threadgroup_barrier 4u |
| exit_if # if_1 |
| } |
| } |
| %17:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %18:f32 = load_vector_element %17, 0u |
| %x_44:f32 = let %18 |
| %20:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %21:f32 = load_vector_element %20, 1u |
| %x_46:f32 = let %21 |
| %23:bool = gt %x_44, %x_46 |
| if %23 [t: $B4] { # if_2 |
| $B4: { # true |
| %24:void = msl.threadgroup_barrier 4u |
| exit_if # if_2 |
| } |
| } |
| %25:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %26:f32 = load_vector_element %25, 1u |
| %x_51:f32 = let %26 |
| %28:i32 = call %tint_f32_to_i32, %x_51 |
| store %i, %28 |
| loop [b: $B5, c: $B6] { # loop_1 |
| $B5: { # body |
| %30:i32 = load %i |
| %x_57:i32 = let %30 |
| %32:bool = gt %x_57, 0i |
| if %32 [t: $B7, f: $B8] { # if_3 |
| $B7: { # true |
| exit_if # if_3 |
| } |
| $B8: { # false |
| exit_loop # loop_1 |
| } |
| } |
| %33:void = msl.threadgroup_barrier 4u |
| continue # -> $B6 |
| } |
| $B6: { # continuing |
| %34:i32 = load %i |
| %x_60:i32 = let %34 |
| %36:i32 = sub %x_60, 1i |
| store %i, %36 |
| next_iteration # -> $B5 |
| } |
| } |
| store %GLF_live3s, 0.0f |
| loop [b: $B9, c: $B10] { # loop_2 |
| $B9: { # body |
| store %i_1, 1i |
| loop [b: $B11, c: $B12] { # loop_3 |
| $B11: { # body |
| %37:i32 = load %i_1 |
| %x_70:i32 = let %37 |
| %39:bool = lt %x_70, 2i |
| if %39 [t: $B13, f: $B14] { # if_4 |
| $B13: { # true |
| exit_if # if_4 |
| } |
| $B14: { # false |
| exit_loop # loop_3 |
| } |
| } |
| %40:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %41:f32 = load_vector_element %40, 0u |
| %x_74:f32 = let %41 |
| %43:bool = gt %x_74, 1.0f |
| if %43 [t: $B15] { # if_5 |
| $B15: { # true |
| %44:void = msl.threadgroup_barrier 4u |
| exit_if # if_5 |
| } |
| } |
| %45:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %46:f32 = load_vector_element %45, 0u |
| %x_79:f32 = let %46 |
| %48:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %49:f32 = load_vector_element %48, 1u |
| %x_81:f32 = let %49 |
| %51:bool = gt %x_79, %x_81 |
| if %51 [t: $B16] { # if_6 |
| $B16: { # true |
| %52:void = msl.threadgroup_barrier 4u |
| exit_if # if_6 |
| } |
| } |
| %53:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %54:f32 = load_vector_element %53, 1u |
| %x_86:f32 = let %54 |
| %56:i32 = call %tint_f32_to_i32, %x_86 |
| store %z, %56 |
| loop [b: $B17, c: $B18] { # loop_4 |
| $B17: { # body |
| %57:i32 = load %z |
| %x_92:i32 = let %57 |
| %59:bool = gt %x_92, 0i |
| if %59 [t: $B19, f: $B20] { # if_7 |
| $B19: { # true |
| exit_if # if_7 |
| } |
| $B20: { # false |
| exit_loop # loop_4 |
| } |
| } |
| %60:f32 = load %GLF_live3s |
| %x_95:f32 = let %60 |
| %62:f32 = add %x_95, 1.0f |
| store %GLF_live3s, %62 |
| continue # -> $B18 |
| } |
| $B18: { # continuing |
| %63:i32 = load %z |
| %x_97:i32 = let %63 |
| %65:i32 = sub %x_97, 1i |
| store %z, %65 |
| next_iteration # -> $B17 |
| } |
| } |
| %66:i32 = load %i_1 |
| %x_99:i32 = let %66 |
| %68:bool = gte %x_99, 1i |
| if %68 [t: $B21] { # if_8 |
| $B21: { # true |
| %69:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %70:f32 = load_vector_element %69, 0u |
| %x_104:f32 = let %70 |
| %72:bool = gt %x_104, 1.0f |
| if %72 [t: $B22] { # if_9 |
| $B22: { # true |
| %73:void = msl.threadgroup_barrier 4u |
| exit_if # if_9 |
| } |
| } |
| exit_if # if_8 |
| } |
| } |
| continue # -> $B12 |
| } |
| $B12: { # continuing |
| %74:i32 = load %i_1 |
| %x_108:i32 = let %74 |
| %76:i32 = add %x_108, 1i |
| store %i_1, %76 |
| next_iteration # -> $B11 |
| } |
| } |
| continue # -> $B10 |
| } |
| $B10: { # continuing |
| %77:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %78:f32 = load_vector_element %77, 0u |
| %x_111:f32 = let %78 |
| %80:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %81:f32 = load_vector_element %80, 1u |
| %x_113:f32 = let %81 |
| %83:bool = gt %x_111, %x_113 |
| %84:bool = eq %83, false |
| break_if %84 # -> [t: exit_loop loop_2, f: $B9] |
| } |
| } |
| ret |
| } |
| } |
| %tint_symbol = @compute @workgroup_size(1, 1, 1) func():void { |
| $B23: { |
| %86:void = call %main_1 |
| ret |
| } |
| } |
| %tint_f32_to_i32 = func(%value:f32):i32 { |
| $B24: { |
| %88:i32 = convert %value |
| %89:bool = gte %value, -2147483648.0f |
| %90:i32 = select -2147483648i, %88, %89 |
| %91:bool = lte %value, 2147483520.0f |
| %92:i32 = select 2147483647i, %90, %91 |
| ret %92 |
| } |
| } |
| |
| unhandled variable address space |
| ******************************************************************** |
| * The tint shader compiler has encountered an unexpected error. * |
| * * |
| * Please help us fix this issue by submitting a bug report at * |
| * crbug.com/tint with the source program that triggered the bug. * |
| ******************************************************************** |
