| SKIP: FAILED |
| |
| ../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: S = struct @align(4) { |
| a:i32 @offset(0) |
| b:i32 @offset(4) |
| } |
| |
| $B1: { # root |
| %s:ptr<storage, i32, read_write> = var @binding_point(0, 0) |
| %g1:ptr<workgroup, atomic<i32>, read_write> = var |
| } |
| |
| %accept_value = func(%val:i32):i32 { |
| $B2: { |
| ret %val |
| } |
| } |
| %accept_ptr_deref_call_func = func(%val_1:ptr<function, i32, read_write>):i32 { # %val_1: 'val' |
| $B3: { |
| %7:i32 = load %val_1 |
| %8:i32 = let %7 |
| %9:i32 = load %val_1 |
| %10:i32 = call %accept_value, %9 |
| %11:i32 = add %8, %10 |
| ret %11 |
| } |
| } |
| %accept_ptr_deref_pass_through = func(%val_2:ptr<function, i32, read_write>):i32 { # %val_2: 'val' |
| $B4: { |
| %14:i32 = load %val_2 |
| %15:i32 = let %14 |
| %16:i32 = call %accept_ptr_deref_call_func, %val_2 |
| %17:i32 = add %15, %16 |
| ret %17 |
| } |
| } |
| %accept_ptr_to_struct_and_access = func(%val_3:ptr<function, S, read_write>):i32 { # %val_3: 'val' |
| $B5: { |
| %20:ptr<function, i32, read_write> = access %val_3, 0u |
| %21:i32 = load %20 |
| %22:ptr<function, i32, read_write> = access %val_3, 1u |
| %23:i32 = load %22 |
| %24:i32 = add %21, %23 |
| ret %24 |
| } |
| } |
| %accept_ptr_to_struct_access_pass_ptr = func(%val_4:ptr<function, S, read_write>):i32 { # %val_4: 'val' |
| $B6: { |
| %27:ptr<function, i32, read_write> = access %val_4, 0u |
| %b:ptr<function, i32, read_write> = let %27 |
| store %b, 2i |
| %29:i32 = load %b |
| ret %29 |
| } |
| } |
| %accept_ptr_vec_access_elements = func(%v1:ptr<function, vec3<f32>, read_write>):i32 { |
| $B7: { |
| %32:vec3<f32> = load %v1 |
| %33:vec3<f32> = load %v1 |
| %34:vec3<f32> = cross %32, %33 |
| %35:f32 = access %34, 0u |
| store_vector_element %v1, 0u, %35 |
| %36:f32 = load_vector_element %v1, 0u |
| %37:i32 = call %tint_f32_to_i32, %36 |
| ret %37 |
| } |
| } |
| %call_builtin_with_mod_scope_ptr = func():i32 { |
| $B8: { |
| %40:i32 = atomicLoad %g1 |
| ret %40 |
| } |
| } |
| %tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void { |
| $B9: { |
| %43:bool = eq %tint_local_index, 0u |
| if %43 [t: $B10] { # if_1 |
| $B10: { # true |
| %44:void = atomicStore %g1, 0i |
| exit_if # if_1 |
| } |
| } |
| %45:void = msl.threadgroup_barrier 4u |
| %v1_1:ptr<function, i32, read_write> = var, 0i # %v1_1: 'v1' |
| %v2:ptr<function, S, read_write> = var, S(0i) |
| %v3:ptr<function, S, read_write> = let %v2 |
| %v4:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(0.0f) |
| %50:i32 = atomicLoad %g1 |
| %t1:i32 = let %50 |
| %52:i32 = call %accept_ptr_deref_pass_through, %v1_1 |
| %53:i32 = let %52 |
| %54:i32 = call %accept_ptr_to_struct_and_access, %v2 |
| %55:i32 = add %53, %54 |
| %56:i32 = let %55 |
| %57:i32 = call %accept_ptr_to_struct_and_access, %v3 |
| %58:i32 = add %56, %57 |
| %59:i32 = let %58 |
| %60:i32 = call %accept_ptr_vec_access_elements, %v4 |
| %61:i32 = add %59, %60 |
| %62:i32 = let %61 |
| %63:i32 = call %accept_ptr_to_struct_access_pass_ptr, %v2 |
| %64:i32 = add %62, %63 |
| %65:i32 = let %64 |
| %66:i32 = call %call_builtin_with_mod_scope_ptr |
| %67:i32 = add %65, %66 |
| %68:i32 = add %67, %t1 |
| store %s, %68 |
| ret |
| } |
| } |
| %tint_f32_to_i32 = func(%value:f32):i32 { |
| $B11: { |
| %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. * |
| ******************************************************************** |