| 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) |
| } |
| |
| In0 = struct @align(4) { |
| data_in0:array<i32, 512> @offset(0) |
| } |
| |
| In1 = struct @align(4) { |
| data_in1:array<i32, 512> @offset(0) |
| } |
| |
| $B1: { # root |
| %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, In0, read> = var @binding_point(0, 0) |
| %x_19:ptr<storage, In1, read> = var @binding_point(0, 1) |
| } |
| |
| %main_1 = func(%tint_wgid:vec3<u32>):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 = access %tint_wgid, 0u |
| %18:u32 = mul 128u, %17 |
| store %base_index_in, %18 |
| %19:u32 = access %tint_wgid, 0u |
| %20:u32 = mul 256u, %19 |
| store %base_index_out, %20 |
| store %index_in0, 0i |
| store %index_in1, -128i |
| store %index_out0, 0i |
| store %index_out1, -128i |
| store %condition_index, 0i |
| store %i, 0i |
| loop [b: $B3, c: $B4] { # loop_1 |
| $B3: { # body |
| %21:i32 = load %i |
| %22:bool = lt %21, 256i |
| if %22 [t: $B5, f: $B6] { # if_1 |
| $B5: { # true |
| exit_if # if_1 |
| } |
| $B6: { # false |
| exit_loop # loop_1 |
| } |
| } |
| %23:i32 = load %condition_index |
| %24:ptr<storage, i32, read> = access %x_13, 0u, %23 |
| %25:i32 = load %24 |
| %26:bool = eq %25, 0i |
| if %26 [t: $B7, f: $B8] { # if_2 |
| $B7: { # true |
| %27:u32 = load %base_index_out |
| %x_75:u32 = let %27 |
| %29:i32 = load %index_out0 |
| %x_76:i32 = let %29 |
| %31:u32 = bitcast %x_76 |
| %32:u32 = add %x_75, %31 |
| %33:ptr<storage, i32, read_write> = access %x_15, 0u, %32 |
| %34:ptr<storage, i32, read_write> = let %33 |
| %35:u32 = load %base_index_in |
| %36:u32 = let %35 |
| %37:i32 = load %index_in0 |
| %38:u32 = bitcast %37 |
| %39:u32 = add %36, %38 |
| %40:ptr<storage, i32, read> = access %x_17, 0u, %39 |
| %41:i32 = load %40 |
| store %34, %41 |
| %42:i32 = load %index_out0 |
| %43:i32 = add %42, 1i |
| store %index_out0, %43 |
| %44:i32 = load %index_in1 |
| %45:i32 = add %44, 1i |
| store %index_in1, %45 |
| exit_if # if_2 |
| } |
| $B8: { # false |
| %46:u32 = load %base_index_out |
| %x_90:u32 = let %46 |
| %48:i32 = load %index_out1 |
| %x_91:i32 = let %48 |
| %50:u32 = bitcast %x_91 |
| %51:u32 = add %x_90, %50 |
| %52:ptr<storage, i32, read_write> = access %x_15, 0u, %51 |
| %53:ptr<storage, i32, read_write> = let %52 |
| %54:u32 = load %base_index_in |
| %55:u32 = let %54 |
| %56:i32 = load %index_in1 |
| %57:u32 = bitcast %56 |
| %58:u32 = add %55, %57 |
| %59:ptr<storage, i32, read> = access %x_19, 0u, %58 |
| %60:i32 = load %59 |
| store %53, %60 |
| %61:i32 = load %index_out1 |
| %62:i32 = add %61, 1i |
| store %index_out1, %62 |
| %63:i32 = load %index_in1 |
| %64:i32 = add %63, 1i |
| store %index_in1, %64 |
| exit_if # if_2 |
| } |
| } |
| %65:i32 = load %condition_index |
| %66:i32 = load %condition_index |
| %67:i32 = add %66, 1i |
| %68:ptr<storage, i32, read> = access %x_13, 0u, %67 |
| %69:i32 = load %68 |
| %70:i32 = add %65, %69 |
| store %condition_index, %70 |
| %71:i32 = load %index_in0 |
| store %temp0, %71 |
| %72:i32 = load %index_in1 |
| store %index_in0, %72 |
| %73:i32 = load %temp0 |
| store %index_in1, %73 |
| %74:i32 = load %index_out0 |
| store %temp1, %74 |
| %75:i32 = load %index_out1 |
| store %index_out0, %75 |
| %76:i32 = load %temp1 |
| store %index_out1, %76 |
| continue # -> $B4 |
| } |
| $B4: { # continuing |
| %77:i32 = load %i |
| %78:i32 = add %77, 1i |
| store %i, %78 |
| next_iteration # -> $B3 |
| } |
| } |
| ret |
| } |
| } |
| %tint_symbol = @compute @workgroup_size(4, 1, 1) func(%gl_WorkGroupID_param:vec3<u32> [@workgroup_id]):void { |
| $B9: { |
| %81:void = call %main_1, %gl_WorkGroupID_param |
| 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. * |
| ******************************************************************** |