| 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 |
| %11:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %12:f32 = load_vector_element %11, 1u |
| %13:bool = gt %10, %12 |
| if %13 [t: $B3] { # if_1 |
| $B3: { # true |
| %14:void = msl.threadgroup_barrier 4u |
| exit_if # if_1 |
| } |
| } |
| %15:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %16:f32 = load_vector_element %15, 0u |
| %17:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %18:f32 = load_vector_element %17, 1u |
| %19:bool = gt %16, %18 |
| if %19 [t: $B4] { # if_2 |
| $B4: { # true |
| %20:void = msl.threadgroup_barrier 4u |
| exit_if # if_2 |
| } |
| } |
| %21:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %22:f32 = load_vector_element %21, 1u |
| %23:i32 = call %tint_f32_to_i32, %22 |
| store %i, %23 |
| loop [b: $B5, c: $B6] { # loop_1 |
| $B5: { # body |
| %25:i32 = load %i |
| %26:bool = gt %25, 0i |
| if %26 [t: $B7, f: $B8] { # if_3 |
| $B7: { # true |
| exit_if # if_3 |
| } |
| $B8: { # false |
| exit_loop # loop_1 |
| } |
| } |
| %27:void = msl.threadgroup_barrier 4u |
| continue # -> $B6 |
| } |
| $B6: { # continuing |
| %28:i32 = load %i |
| %29:i32 = sub %28, 1i |
| store %i, %29 |
| 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 |
| %30:i32 = load %i_1 |
| %31:bool = lt %30, 2i |
| if %31 [t: $B13, f: $B14] { # if_4 |
| $B13: { # true |
| exit_if # if_4 |
| } |
| $B14: { # false |
| exit_loop # loop_3 |
| } |
| } |
| %32:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %33:f32 = load_vector_element %32, 0u |
| %34:bool = gt %33, 1.0f |
| if %34 [t: $B15] { # if_5 |
| $B15: { # true |
| %35:void = msl.threadgroup_barrier 4u |
| exit_if # if_5 |
| } |
| } |
| %36:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %37:f32 = load_vector_element %36, 0u |
| %38:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %39:f32 = load_vector_element %38, 1u |
| %40:bool = gt %37, %39 |
| if %40 [t: $B16] { # if_6 |
| $B16: { # true |
| %41:void = msl.threadgroup_barrier 4u |
| exit_if # if_6 |
| } |
| } |
| %42:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %43:f32 = load_vector_element %42, 1u |
| %44:i32 = call %tint_f32_to_i32, %43 |
| store %z, %44 |
| loop [b: $B17, c: $B18] { # loop_4 |
| $B17: { # body |
| %45:i32 = load %z |
| %46:bool = gt %45, 0i |
| if %46 [t: $B19, f: $B20] { # if_7 |
| $B19: { # true |
| exit_if # if_7 |
| } |
| $B20: { # false |
| exit_loop # loop_4 |
| } |
| } |
| %47:f32 = load %GLF_live3s |
| %48:f32 = add %47, 1.0f |
| store %GLF_live3s, %48 |
| continue # -> $B18 |
| } |
| $B18: { # continuing |
| %49:i32 = load %z |
| %50:i32 = sub %49, 1i |
| store %z, %50 |
| next_iteration # -> $B17 |
| } |
| } |
| %51:i32 = load %i_1 |
| %52:bool = gte %51, 1i |
| if %52 [t: $B21] { # if_8 |
| $B21: { # true |
| %53:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %54:f32 = load_vector_element %53, 0u |
| %55:bool = gt %54, 1.0f |
| if %55 [t: $B22] { # if_9 |
| $B22: { # true |
| %56:void = msl.threadgroup_barrier 4u |
| exit_if # if_9 |
| } |
| } |
| exit_if # if_8 |
| } |
| } |
| continue # -> $B12 |
| } |
| $B12: { # continuing |
| %57:i32 = load %i_1 |
| %58:i32 = add %57, 1i |
| store %i_1, %58 |
| next_iteration # -> $B11 |
| } |
| } |
| continue # -> $B10 |
| } |
| $B10: { # continuing |
| %59:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %60:f32 = load_vector_element %59, 0u |
| %x_111:f32 = let %60 |
| %62:ptr<uniform, vec2<f32>, read> = access %x_6, 0u |
| %63:f32 = load_vector_element %62, 1u |
| %x_113:f32 = let %63 |
| %65:bool = gt %x_111, %x_113 |
| %66:bool = eq %65, false |
| break_if %66 # -> [t: exit_loop loop_2, f: $B9] |
| } |
| } |
| ret |
| } |
| } |
| %tint_symbol = @compute @workgroup_size(1, 1, 1) func():void { |
| $B23: { |
| %68:void = call %main_1 |
| ret |
| } |
| } |
| %tint_f32_to_i32 = func(%value:f32):i32 { |
| $B24: { |
| %70:i32 = convert %value |
| %71:bool = gte %value, -2147483648.0f |
| %72:i32 = select -2147483648i, %70, %71 |
| %73:bool = lte %value, 2147483520.0f |
| %74:i32 = select 2147483647i, %72, %73 |
| ret %74 |
| } |
| } |
| |
| 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. * |
| ******************************************************************** |
