| SKIP: FAILED |
| |
| ../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: RenderParams = struct @align(16) { |
| modelViewProjectionMatrix:mat4x4<f32> @offset(0) |
| right:vec3<f32> @offset(64) |
| up:vec3<f32> @offset(80) |
| } |
| |
| VertexOutput = struct @align(16) { |
| position:vec4<f32> @offset(0), @builtin(position) |
| color:vec4<f32> @offset(16), @location(0) |
| quad_pos:vec2<f32> @offset(32), @location(1) |
| } |
| |
| VertexInput = struct @align(16) { |
| position:vec3<f32> @offset(0), @location(0) |
| color:vec4<f32> @offset(16), @location(1) |
| quad_pos:vec2<f32> @offset(32), @location(2) |
| } |
| |
| SimulationParams = struct @align(16) { |
| deltaTime:f32 @offset(0) |
| seed:vec4<f32> @offset(16) |
| } |
| |
| Particle = struct @align(16) { |
| position:vec3<f32> @offset(0) |
| lifetime:f32 @offset(12) |
| color:vec4<f32> @offset(16) |
| velocity:vec2<f32> @offset(32) |
| } |
| |
| Particles = struct @align(16) { |
| particles:array<Particle> @offset(0) |
| } |
| |
| UBO = struct @align(4) { |
| width:u32 @offset(0) |
| } |
| |
| Buffer = struct @align(4) { |
| weights:array<f32> @offset(0) |
| } |
| |
| $B1: { # root |
| %rand_seed:ptr<private, vec2<f32>, read_write> = var |
| %render_params:ptr<uniform, RenderParams, read> = var @binding_point(0, 5) |
| %sim_params:ptr<uniform, SimulationParams, read> = var @binding_point(0, 0) |
| %data:ptr<storage, Particles, read_write> = var @binding_point(0, 1) |
| %tint_symbol:ptr<handle, texture_1d<f32>, read> = var @binding_point(0, 2) |
| %ubo:ptr<uniform, UBO, read> = var @binding_point(0, 3) |
| %buf_in:ptr<storage, Buffer, read> = var @binding_point(0, 4) |
| %buf_out:ptr<storage, Buffer, read_write> = var @binding_point(0, 5) |
| %tex_in:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 6) |
| %tex_out:ptr<handle, texture_storage_2d<rgba8unorm, write>, read> = var @binding_point(0, 7) |
| } |
| |
| %asinh_468a48 = func():void { |
| $B2: { |
| %arg_0:ptr<function, f16, read_write> = var, 0.0h |
| %13:f16 = load %arg_0 |
| %14:f16 = asinh %13 |
| %res:ptr<function, f16, read_write> = var, %14 |
| ret |
| } |
| } |
| %vertex_main = @vertex func():vec4<f32> [@position] { |
| $B3: { |
| %17:void = call %asinh_468a48 |
| ret vec4<f32>(0.0f) |
| } |
| } |
| %fragment_main = @fragment func():void { |
| $B4: { |
| %19:void = call %asinh_468a48 |
| ret |
| } |
| } |
| %rgba32uintin = @compute @workgroup_size(1, 1, 1) func():void { |
| $B5: { |
| %21:void = call %asinh_468a48 |
| ret |
| } |
| } |
| %vs_main = @vertex func(%in:VertexInput):VertexOutput { |
| $B6: { |
| %24:ptr<uniform, vec3<f32>, read> = access %render_params, 1u |
| %25:vec3<f32> = load %24 |
| %26:ptr<uniform, vec3<f32>, read> = access %render_params, 2u |
| %27:vec3<f32> = load %26 |
| %28:mat2x3<f32> = construct %25, %27 |
| %29:vec2<f32> = access %in, 2u |
| %30:vec3<f32> = mul %28, %29 |
| %quad_pos:ptr<function, vec3<f32>, read_write> = var, %30 |
| %32:vec3<f32> = access %in, 0u |
| %33:vec3<f32> = load %quad_pos |
| %34:vec3<f32> = add %33, 0.00999999977648258209f |
| %35:vec3<f32> = sub %32, %34 |
| %position:ptr<function, vec3<f32>, read_write> = var, %35 |
| %out:ptr<function, VertexOutput, read_write> = var |
| %38:ptr<function, vec4<f32>, read_write> = access %out, 0u |
| %39:ptr<uniform, mat4x4<f32>, read> = access %render_params, 0u |
| %40:mat4x4<f32> = load %39 |
| %41:mat4x4<f32> = let %40 |
| %42:vec3<f32> = load %position |
| %43:vec4<f32> = construct %42, 1.0f |
| %44:vec4<f32> = mul %41, %43 |
| store %38, %44 |
| %45:ptr<function, vec4<f32>, read_write> = access %out, 1u |
| %46:vec4<f32> = access %in, 1u |
| store %45, %46 |
| %47:ptr<function, vec2<f32>, read_write> = access %out, 2u |
| %48:vec2<f32> = access %in, 2u |
| store %47, %48 |
| %49:VertexOutput = load %out |
| ret %49 |
| } |
| } |
| %simulate = @compute @workgroup_size(64, 1, 1) func(%GlobalInvocationID:vec3<u32> [@global_invocation_id]):void { |
| $B7: { |
| %52:ptr<uniform, vec4<f32>, read> = access %sim_params, 1u |
| %53:vec4<f32> = load %52 |
| %54:vec2<f32> = swizzle %53, xy |
| %55:vec2<f32> = let %54 |
| %56:vec2<u32> = swizzle %GlobalInvocationID, xy |
| %57:vec2<f32> = convert %56 |
| %58:vec2<f32> = mul %55, %57 |
| %59:ptr<uniform, vec4<f32>, read> = access %sim_params, 1u |
| %60:vec4<f32> = load %59 |
| %61:vec2<f32> = swizzle %60, zw |
| %62:vec2<f32> = mul %58, %61 |
| store %rand_seed, %62 |
| %63:u32 = access %GlobalInvocationID, 0u |
| %idx:u32 = let %63 |
| %65:ptr<storage, Particle, read_write> = access %data, 0u, %idx |
| %66:Particle = load %65 |
| %particle:ptr<function, Particle, read_write> = var, %66 |
| %68:ptr<storage, Particle, read_write> = access %data, 0u, %idx |
| %69:Particle = load %particle |
| %70:void = call %tint_store_and_preserve_padding, %68, %69 |
| ret |
| } |
| } |
| %export_level = @compute @workgroup_size(64, 1, 1) func(%coord:vec3<u32> [@global_invocation_id]):void { |
| $B8: { |
| %74:vec2<u32> = swizzle %coord, xy |
| %75:texture_storage_2d<rgba8unorm, write> = load %tex_out |
| %76:vec2<u32> = textureDimensions %75 |
| %77:vec2<u32> = construct %76 |
| %78:vec2<bool> = lt %74, %77 |
| %79:bool = all %78 |
| if %79 [t: $B9] { # if_1 |
| $B9: { # true |
| %80:u32 = access %coord, 0u |
| %81:u32 = access %coord, 1u |
| %82:ptr<uniform, u32, read> = access %ubo, 0u |
| %83:u32 = load %82 |
| %84:u32 = mul %81, %83 |
| %85:u32 = and %84, 31u |
| %86:u32 = shl %80, %85 |
| %dst_offset:u32 = let %86 |
| %88:u32 = access %coord, 0u |
| %89:u32 = sub %88, 2u |
| %90:u32 = access %coord, 1u |
| %91:u32 = and 2u, 31u |
| %92:u32 = shr %90, %91 |
| %93:ptr<uniform, u32, read> = access %ubo, 0u |
| %94:u32 = load %93 |
| %95:u32 = mul %92, %94 |
| %96:u32 = add %89, %95 |
| %src_offset:u32 = let %96 |
| %98:u32 = and 0u, 31u |
| %99:u32 = shl %src_offset, %98 |
| %100:ptr<storage, f32, read> = access %buf_in, 0u, %99 |
| %101:f32 = load %100 |
| %a:f32 = let %101 |
| %103:u32 = add %src_offset, 1u |
| %104:ptr<storage, f32, read> = access %buf_in, 0u, %103 |
| %105:f32 = load %104 |
| %b:f32 = let %105 |
| %107:u32 = add %src_offset, 1u |
| %108:ptr<uniform, u32, read> = access %ubo, 0u |
| %109:u32 = load %108 |
| %110:u32 = add %107, %109 |
| %111:ptr<storage, f32, read> = access %buf_in, 0u, %110 |
| %112:f32 = load %111 |
| %c:f32 = let %112 |
| %114:u32 = add %src_offset, 1u |
| %115:ptr<uniform, u32, read> = access %ubo, 0u |
| %116:u32 = load %115 |
| %117:u32 = add %114, %116 |
| %118:ptr<storage, f32, read> = access %buf_in, 0u, %117 |
| %119:f32 = load %118 |
| %d:f32 = let %119 |
| %121:vec4<f32> = construct %a, %b, %c, %d |
| %122:f32 = dot %121, vec4<f32>(1.0f) |
| %sum:f32 = let %122 |
| %124:ptr<storage, f32, read_write> = access %buf_out, 0u, %dst_offset |
| %125:f32 = mod %sum, 4.0f |
| store %124, %125 |
| %126:f32 = mul %a, %b |
| %127:f32 = div %a, %b |
| %128:f32 = add %127, %c |
| %129:vec4<f32> = construct %a, %126, %128, %sum |
| %130:vec4<f32> = let %129 |
| %131:f32 = max %sum, 0.0f |
| %132:vec4<f32> = add %130, %131 |
| %probabilities:vec4<f32> = let %132 |
| %134:texture_storage_2d<rgba8unorm, write> = load %tex_out |
| %135:texture_storage_2d<rgba8unorm, write> = let %134 |
| %136:vec2<u32> = swizzle %coord, xy |
| %137:vec2<i32> = convert %136 |
| %138:void = textureStore %135, %137, %probabilities |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| %tint_store_and_preserve_padding = func(%target:ptr<storage, Particle, read_write>, %value_param:Particle):void { |
| $B10: { |
| %141:ptr<storage, vec3<f32>, read_write> = access %target, 0u |
| %142:vec3<f32> = access %value_param, 0u |
| store %141, %142 |
| %143:ptr<storage, f32, read_write> = access %target, 1u |
| %144:f32 = access %value_param, 1u |
| store %143, %144 |
| %145:ptr<storage, vec4<f32>, read_write> = access %target, 2u |
| %146:vec4<f32> = access %value_param, 2u |
| store %145, %146 |
| %147:ptr<storage, vec2<f32>, read_write> = access %target, 3u |
| %148:vec2<f32> = access %value_param, 3u |
| store %147, %148 |
| 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. * |
| ******************************************************************** |
