| SKIP: FAILED |
| |
| ../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: In2 = struct @align(4) { |
| data_in2:array<i32, 8> @offset(0) |
| } |
| |
| Out0 = struct @align(4) { |
| data_out0:array<i32, 1024> @offset(0) |
| } |
| |
| In1 = struct @align(4) { |
| data_in1:array<i32, 512> @offset(0) |
| } |
| |
| In0 = struct @align(4) { |
| data_in0:array<i32, 512> @offset(0) |
| } |
| |
| $B1: { # root |
| %gl_WorkGroupID:ptr<private, vec3<u32>, read_write> = var |
| %x_13:ptr<storage, In2, read> = var @binding_point(0, 2) |
| %x_15:ptr<storage, Out0, read_write> = var @binding_point(0, 3) |
| %x_17:ptr<storage, In1, read> = var @binding_point(0, 1) |
| %x_19:ptr<storage, In0, read> = var @binding_point(0, 0) |
| } |
| |
| %main_1 = func():void { |
| $B2: { |
| %base_index_in:ptr<function, u32, read_write> = var |
| %base_index_out:ptr<function, u32, read_write> = var |
| %index_in0:ptr<function, i32, read_write> = var |
| %index_in1:ptr<function, i32, read_write> = var |
| %index_out0:ptr<function, i32, read_write> = var |
| %index_out1:ptr<function, i32, read_write> = var |
| %condition_index:ptr<function, i32, read_write> = var |
| %i:ptr<function, i32, read_write> = var |
| %temp0:ptr<function, i32, read_write> = var |
| %temp1:ptr<function, i32, read_write> = var |
| %17:u32 = load_vector_element %gl_WorkGroupID, 0u |
| %x_58:u32 = let %17 |
| %19:u32 = mul 128u, %x_58 |
| store %base_index_in, %19 |
| %20:u32 = load_vector_element %gl_WorkGroupID, 0u |
| %x_61:u32 = let %20 |
| %22:u32 = mul 256u, %x_61 |
| store %base_index_out, %22 |
| store %index_in0, 127i |
| store %index_in1, 383i |
| store %index_out0, 255i |
| store %index_out1, 383i |
| store %condition_index, 0i |
| store %i, 0i |
| loop [b: $B3, c: $B4] { # loop_1 |
| $B3: { # body |
| %23:i32 = load %i |
| %x_67:i32 = let %23 |
| %25:bool = lt %x_67, 256i |
| if %25 [t: $B5, f: $B6] { # if_1 |
| $B5: { # true |
| exit_if # if_1 |
| } |
| $B6: { # false |
| exit_loop # loop_1 |
| } |
| } |
| %26:i32 = load %condition_index |
| %x_70:i32 = let %26 |
| %28:ptr<storage, i32, read> = access %x_13, 0u, %x_70 |
| %29:i32 = load %28 |
| %x_72:i32 = let %29 |
| %31:bool = eq %x_72, 0i |
| if %31 [t: $B7, f: $B8] { # if_2 |
| $B7: { # true |
| %32:u32 = load %base_index_out |
| %x_77:u32 = let %32 |
| %34:i32 = load %index_out0 |
| %x_78:i32 = let %34 |
| %36:u32 = load %base_index_in |
| %x_81:u32 = let %36 |
| %38:i32 = load %index_in0 |
| %x_82:i32 = let %38 |
| %40:u32 = bitcast %x_82 |
| %41:u32 = add %x_81, %40 |
| %42:ptr<storage, i32, read> = access %x_17, 0u, %41 |
| %43:i32 = load %42 |
| %x_86:i32 = let %43 |
| %45:u32 = bitcast %x_78 |
| %46:u32 = add %x_77, %45 |
| %47:ptr<storage, i32, read_write> = access %x_15, 0u, %46 |
| store %47, %x_86 |
| %48:i32 = load %index_out0 |
| %x_88:i32 = let %48 |
| %50:i32 = sub %x_88, 1i |
| store %index_out0, %50 |
| %51:i32 = load %index_in1 |
| %x_90:i32 = let %51 |
| %53:i32 = sub %x_90, 1i |
| store %index_in1, %53 |
| exit_if # if_2 |
| } |
| $B8: { # false |
| %54:u32 = load %base_index_out |
| %x_92:u32 = let %54 |
| %56:i32 = load %index_out1 |
| %x_93:i32 = let %56 |
| %58:u32 = load %base_index_in |
| %x_96:u32 = let %58 |
| %60:i32 = load %index_in1 |
| %x_97:i32 = let %60 |
| %62:u32 = bitcast %x_97 |
| %63:u32 = add %x_96, %62 |
| %64:ptr<storage, i32, read> = access %x_19, 0u, %63 |
| %65:i32 = load %64 |
| %x_101:i32 = let %65 |
| %67:u32 = bitcast %x_93 |
| %68:u32 = add %x_92, %67 |
| %69:ptr<storage, i32, read_write> = access %x_15, 0u, %68 |
| store %69, %x_101 |
| %70:i32 = load %index_out1 |
| %x_103:i32 = let %70 |
| %72:i32 = sub %x_103, 1i |
| store %index_out1, %72 |
| %73:i32 = load %index_in1 |
| %x_105:i32 = let %73 |
| %75:i32 = sub %x_105, 1i |
| store %index_in1, %75 |
| exit_if # if_2 |
| } |
| } |
| %76:i32 = load %condition_index |
| %x_107:i32 = let %76 |
| %78:i32 = add %x_107, 1i |
| %79:ptr<storage, i32, read> = access %x_13, 0u, %78 |
| %80:i32 = load %79 |
| %x_110:i32 = let %80 |
| %82:i32 = load %condition_index |
| %x_111:i32 = let %82 |
| %84:i32 = add %x_111, %x_110 |
| store %condition_index, %84 |
| %85:i32 = load %index_in0 |
| %x_113:i32 = let %85 |
| store %temp0, %x_113 |
| %87:i32 = load %index_in1 |
| %x_114:i32 = let %87 |
| store %index_in0, %x_114 |
| %89:i32 = load %temp0 |
| %x_115:i32 = let %89 |
| store %index_in1, %x_115 |
| %91:i32 = load %index_out0 |
| %x_116:i32 = let %91 |
| store %temp1, %x_116 |
| %93:i32 = load %index_out1 |
| %x_117:i32 = let %93 |
| store %index_out0, %x_117 |
| %95:i32 = load %temp1 |
| %x_118:i32 = let %95 |
| store %index_out1, %x_118 |
| continue # -> $B4 |
| } |
| $B4: { # continuing |
| %97:i32 = load %i |
| %x_119:i32 = let %97 |
| %99:i32 = add %x_119, 1i |
| store %i, %99 |
| next_iteration # -> $B3 |
| } |
| } |
| ret |
| } |
| } |
| %tint_symbol = @compute @workgroup_size(4, 1, 1) func(%gl_WorkGroupID_param:vec3<u32> [@workgroup_id]):void { |
| $B9: { |
| store %gl_WorkGroupID, %gl_WorkGroupID_param |
| %102:void = call %main_1 |
| ret |
| } |
| } |
| |
| 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. * |
| ******************************************************************** |
