blob: bd202cb5ef9c591b00d26b66fa575e87d8e15228 [file] [log] [blame]
SKIP: FAILED
../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: Inner = struct @align(16) {
scalar_f32:f32 @offset(0)
scalar_i32:i32 @offset(4)
scalar_u32:u32 @offset(8)
scalar_f16:f16 @offset(12)
vec2_f32:vec2<f32> @offset(16)
vec2_i32:vec2<i32> @offset(24)
vec2_u32:vec2<u32> @offset(32)
vec2_f16:vec2<f16> @offset(40)
vec3_f32:vec3<f32> @offset(48)
vec3_i32:vec3<i32> @offset(64)
vec3_u32:vec3<u32> @offset(80)
vec3_f16:vec3<f16> @offset(96)
vec4_f32:vec4<f32> @offset(112)
vec4_i32:vec4<i32> @offset(128)
vec4_u32:vec4<u32> @offset(144)
vec4_f16:vec4<f16> @offset(160)
mat2x2_f32:mat2x2<f32> @offset(168)
mat2x3_f32:mat2x3<f32> @offset(192)
mat2x4_f32:mat2x4<f32> @offset(224)
mat3x2_f32:mat3x2<f32> @offset(256)
mat3x3_f32:mat3x3<f32> @offset(288)
mat3x4_f32:mat3x4<f32> @offset(336)
mat4x2_f32:mat4x2<f32> @offset(384)
mat4x3_f32:mat4x3<f32> @offset(416)
mat4x4_f32:mat4x4<f32> @offset(480)
mat2x2_f16:mat2x2<f16> @offset(544)
mat2x3_f16:mat2x3<f16> @offset(552)
mat2x4_f16:mat2x4<f16> @offset(568)
mat3x2_f16:mat3x2<f16> @offset(584)
mat3x3_f16:mat3x3<f16> @offset(600)
mat3x4_f16:mat3x4<f16> @offset(624)
mat4x2_f16:mat4x2<f16> @offset(648)
mat4x3_f16:mat4x3<f16> @offset(664)
mat4x4_f16:mat4x4<f16> @offset(696)
arr2_vec3_f32:array<vec3<f32>, 2> @offset(736)
arr2_mat4x2_f16:array<mat4x2<f16>, 2> @offset(768)
}
S = struct @align(16) {
arr:array<Inner> @offset(0)
}
$B1: { # root
%sb:ptr<storage, S, read_write> = var @binding_point(0, 0)
}
%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%idx:u32 [@local_invocation_index]):void {
$B2: {
%4:ptr<storage, f32, read_write> = access %sb, 0u, %idx, 0u
store %4, 0.0f
%5:ptr<storage, i32, read_write> = access %sb, 0u, %idx, 1u
store %5, 0i
%6:ptr<storage, u32, read_write> = access %sb, 0u, %idx, 2u
store %6, 0u
%7:ptr<storage, f16, read_write> = access %sb, 0u, %idx, 3u
store %7, 0.0h
%8:ptr<storage, vec2<f32>, read_write> = access %sb, 0u, %idx, 4u
store %8, vec2<f32>(0.0f)
%9:ptr<storage, vec2<i32>, read_write> = access %sb, 0u, %idx, 5u
store %9, vec2<i32>(0i)
%10:ptr<storage, vec2<u32>, read_write> = access %sb, 0u, %idx, 6u
store %10, vec2<u32>(0u)
%11:ptr<storage, vec2<f16>, read_write> = access %sb, 0u, %idx, 7u
store %11, vec2<f16>(0.0h)
%12:ptr<storage, vec3<f32>, read_write> = access %sb, 0u, %idx, 8u
store %12, vec3<f32>(0.0f)
%13:ptr<storage, vec3<i32>, read_write> = access %sb, 0u, %idx, 9u
store %13, vec3<i32>(0i)
%14:ptr<storage, vec3<u32>, read_write> = access %sb, 0u, %idx, 10u
store %14, vec3<u32>(0u)
%15:ptr<storage, vec3<f16>, read_write> = access %sb, 0u, %idx, 11u
store %15, vec3<f16>(0.0h)
%16:ptr<storage, vec4<f32>, read_write> = access %sb, 0u, %idx, 12u
store %16, vec4<f32>(0.0f)
%17:ptr<storage, vec4<i32>, read_write> = access %sb, 0u, %idx, 13u
store %17, vec4<i32>(0i)
%18:ptr<storage, vec4<u32>, read_write> = access %sb, 0u, %idx, 14u
store %18, vec4<u32>(0u)
%19:ptr<storage, vec4<f16>, read_write> = access %sb, 0u, %idx, 15u
store %19, vec4<f16>(0.0h)
%20:ptr<storage, mat2x2<f32>, read_write> = access %sb, 0u, %idx, 16u
store %20, mat2x2<f32>(vec2<f32>(0.0f))
%21:ptr<storage, mat2x3<f32>, read_write> = access %sb, 0u, %idx, 17u
%22:void = call %tint_store_and_preserve_padding, %21, mat2x3<f32>(vec3<f32>(0.0f))
%24:ptr<storage, mat2x4<f32>, read_write> = access %sb, 0u, %idx, 18u
store %24, mat2x4<f32>(vec4<f32>(0.0f))
%25:ptr<storage, mat3x2<f32>, read_write> = access %sb, 0u, %idx, 19u
store %25, mat3x2<f32>(vec2<f32>(0.0f))
%26:ptr<storage, mat3x3<f32>, read_write> = access %sb, 0u, %idx, 20u
%27:void = call %tint_store_and_preserve_padding_1, %26, mat3x3<f32>(vec3<f32>(0.0f))
%29:ptr<storage, mat3x4<f32>, read_write> = access %sb, 0u, %idx, 21u
store %29, mat3x4<f32>(vec4<f32>(0.0f))
%30:ptr<storage, mat4x2<f32>, read_write> = access %sb, 0u, %idx, 22u
store %30, mat4x2<f32>(vec2<f32>(0.0f))
%31:ptr<storage, mat4x3<f32>, read_write> = access %sb, 0u, %idx, 23u
%32:void = call %tint_store_and_preserve_padding_2, %31, mat4x3<f32>(vec3<f32>(0.0f))
%34:ptr<storage, mat4x4<f32>, read_write> = access %sb, 0u, %idx, 24u
store %34, mat4x4<f32>(vec4<f32>(0.0f))
%35:ptr<storage, mat2x2<f16>, read_write> = access %sb, 0u, %idx, 25u
store %35, mat2x2<f16>(vec2<f16>(0.0h))
%36:ptr<storage, mat2x3<f16>, read_write> = access %sb, 0u, %idx, 26u
%37:void = call %tint_store_and_preserve_padding_3, %36, mat2x3<f16>(vec3<f16>(0.0h))
%39:ptr<storage, mat2x4<f16>, read_write> = access %sb, 0u, %idx, 27u
store %39, mat2x4<f16>(vec4<f16>(0.0h))
%40:ptr<storage, mat3x2<f16>, read_write> = access %sb, 0u, %idx, 28u
store %40, mat3x2<f16>(vec2<f16>(0.0h))
%41:ptr<storage, mat3x3<f16>, read_write> = access %sb, 0u, %idx, 29u
%42:void = call %tint_store_and_preserve_padding_4, %41, mat3x3<f16>(vec3<f16>(0.0h))
%44:ptr<storage, mat3x4<f16>, read_write> = access %sb, 0u, %idx, 30u
store %44, mat3x4<f16>(vec4<f16>(0.0h))
%45:ptr<storage, mat4x2<f16>, read_write> = access %sb, 0u, %idx, 31u
store %45, mat4x2<f16>(vec2<f16>(0.0h))
%46:ptr<storage, mat4x3<f16>, read_write> = access %sb, 0u, %idx, 32u
%47:void = call %tint_store_and_preserve_padding_5, %46, mat4x3<f16>(vec3<f16>(0.0h))
%49:ptr<storage, mat4x4<f16>, read_write> = access %sb, 0u, %idx, 33u
store %49, mat4x4<f16>(vec4<f16>(0.0h))
%50:ptr<storage, array<vec3<f32>, 2>, read_write> = access %sb, 0u, %idx, 34u
%51:void = call %tint_store_and_preserve_padding_6, %50, array<vec3<f32>, 2>(vec3<f32>(0.0f))
%53:ptr<storage, array<mat4x2<f16>, 2>, read_write> = access %sb, 0u, %idx, 35u
store %53, array<mat4x2<f16>, 2>(mat4x2<f16>(vec2<f16>(0.0h)))
ret
}
}
%tint_store_and_preserve_padding = func(%target:ptr<storage, mat2x3<f32>, read_write>, %value_param:mat2x3<f32>):void {
$B3: {
%56:ptr<storage, vec3<f32>, read_write> = access %target, 0u
%57:vec3<f32> = access %value_param, 0u
store %56, %57
%58:ptr<storage, vec3<f32>, read_write> = access %target, 1u
%59:vec3<f32> = access %value_param, 1u
store %58, %59
ret
}
}
%tint_store_and_preserve_padding_1 = func(%target_1:ptr<storage, mat3x3<f32>, read_write>, %value_param_1:mat3x3<f32>):void { # %tint_store_and_preserve_padding_1: 'tint_store_and_preserve_padding', %target_1: 'target', %value_param_1: 'value_param'
$B4: {
%62:ptr<storage, vec3<f32>, read_write> = access %target_1, 0u
%63:vec3<f32> = access %value_param_1, 0u
store %62, %63
%64:ptr<storage, vec3<f32>, read_write> = access %target_1, 1u
%65:vec3<f32> = access %value_param_1, 1u
store %64, %65
%66:ptr<storage, vec3<f32>, read_write> = access %target_1, 2u
%67:vec3<f32> = access %value_param_1, 2u
store %66, %67
ret
}
}
%tint_store_and_preserve_padding_2 = func(%target_2:ptr<storage, mat4x3<f32>, read_write>, %value_param_2:mat4x3<f32>):void { # %tint_store_and_preserve_padding_2: 'tint_store_and_preserve_padding', %target_2: 'target', %value_param_2: 'value_param'
$B5: {
%70:ptr<storage, vec3<f32>, read_write> = access %target_2, 0u
%71:vec3<f32> = access %value_param_2, 0u
store %70, %71
%72:ptr<storage, vec3<f32>, read_write> = access %target_2, 1u
%73:vec3<f32> = access %value_param_2, 1u
store %72, %73
%74:ptr<storage, vec3<f32>, read_write> = access %target_2, 2u
%75:vec3<f32> = access %value_param_2, 2u
store %74, %75
%76:ptr<storage, vec3<f32>, read_write> = access %target_2, 3u
%77:vec3<f32> = access %value_param_2, 3u
store %76, %77
ret
}
}
%tint_store_and_preserve_padding_3 = func(%target_3:ptr<storage, mat2x3<f16>, read_write>, %value_param_3:mat2x3<f16>):void { # %tint_store_and_preserve_padding_3: 'tint_store_and_preserve_padding', %target_3: 'target', %value_param_3: 'value_param'
$B6: {
%80:ptr<storage, vec3<f16>, read_write> = access %target_3, 0u
%81:vec3<f16> = access %value_param_3, 0u
store %80, %81
%82:ptr<storage, vec3<f16>, read_write> = access %target_3, 1u
%83:vec3<f16> = access %value_param_3, 1u
store %82, %83
ret
}
}
%tint_store_and_preserve_padding_4 = func(%target_4:ptr<storage, mat3x3<f16>, read_write>, %value_param_4:mat3x3<f16>):void { # %tint_store_and_preserve_padding_4: 'tint_store_and_preserve_padding', %target_4: 'target', %value_param_4: 'value_param'
$B7: {
%86:ptr<storage, vec3<f16>, read_write> = access %target_4, 0u
%87:vec3<f16> = access %value_param_4, 0u
store %86, %87
%88:ptr<storage, vec3<f16>, read_write> = access %target_4, 1u
%89:vec3<f16> = access %value_param_4, 1u
store %88, %89
%90:ptr<storage, vec3<f16>, read_write> = access %target_4, 2u
%91:vec3<f16> = access %value_param_4, 2u
store %90, %91
ret
}
}
%tint_store_and_preserve_padding_5 = func(%target_5:ptr<storage, mat4x3<f16>, read_write>, %value_param_5:mat4x3<f16>):void { # %tint_store_and_preserve_padding_5: 'tint_store_and_preserve_padding', %target_5: 'target', %value_param_5: 'value_param'
$B8: {
%94:ptr<storage, vec3<f16>, read_write> = access %target_5, 0u
%95:vec3<f16> = access %value_param_5, 0u
store %94, %95
%96:ptr<storage, vec3<f16>, read_write> = access %target_5, 1u
%97:vec3<f16> = access %value_param_5, 1u
store %96, %97
%98:ptr<storage, vec3<f16>, read_write> = access %target_5, 2u
%99:vec3<f16> = access %value_param_5, 2u
store %98, %99
%100:ptr<storage, vec3<f16>, read_write> = access %target_5, 3u
%101:vec3<f16> = access %value_param_5, 3u
store %100, %101
ret
}
}
%tint_store_and_preserve_padding_6 = func(%target_6:ptr<storage, array<vec3<f32>, 2>, read_write>, %value_param_6:array<vec3<f32>, 2>):void { # %tint_store_and_preserve_padding_6: 'tint_store_and_preserve_padding', %target_6: 'target', %value_param_6: 'value_param'
$B9: {
loop [i: $B10, b: $B11, c: $B12] { # loop_1
$B10: { # initializer
next_iteration 0u # -> $B11
}
$B11 (%idx_1:u32): { # body
%105:bool = gte %idx_1, 2u
if %105 [t: $B13] { # if_1
$B13: { # true
exit_loop # loop_1
}
}
%106:ptr<storage, vec3<f32>, read_write> = access %target_6, %idx_1
%107:vec3<f32> = access %value_param_6, %idx_1
store %106, %107
continue # -> $B12
}
$B12: { # continuing
%108:u32 = add %idx_1, 1u
next_iteration %108 # -> $B11
}
}
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. *
********************************************************************