| 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:vec3<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, 0) |
| %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_2d<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) |
| } |
| |
| %rand = func():f32 { |
| $B2: { |
| %12:vec2<f32> = load %rand_seed |
| %13:f32 = dot %12, vec2<f32>(23.1407794952392578125f, 232.6168975830078125f) |
| %14:f32 = cos %13 |
| %15:f32 = mul %14, 136.816802978515625f |
| %16:f32 = fract %15 |
| store_vector_element %rand_seed, 0u, %16 |
| %17:vec2<f32> = load %rand_seed |
| %18:f32 = dot %17, vec2<f32>(54.478565216064453125f, 345.841522216796875f) |
| %19:f32 = cos %18 |
| %20:f32 = mul %19, 534.7645263671875f |
| %21:f32 = fract %20 |
| store_vector_element %rand_seed, 1u, %21 |
| %22:f32 = load_vector_element %rand_seed, 1u |
| ret %22 |
| } |
| } |
| %vs_main = @vertex func(%in:VertexInput):VertexOutput { |
| $B3: { |
| %25:ptr<uniform, vec3<f32>, read> = access %render_params, 1u |
| %26:vec3<f32> = load %25 |
| %27:ptr<uniform, vec3<f32>, read> = access %render_params, 2u |
| %28:vec3<f32> = load %27 |
| %29:mat2x3<f32> = construct %26, %28 |
| %30:vec2<f32> = access %in, 2u |
| %31:vec3<f32> = mul %29, %30 |
| %quad_pos:ptr<function, vec3<f32>, read_write> = var, %31 |
| %33:vec3<f32> = access %in, 0u |
| %34:vec3<f32> = load %quad_pos |
| %35:vec3<f32> = mul %34, 0.00999999977648258209f |
| %36:vec3<f32> = add %33, %35 |
| %position:ptr<function, vec3<f32>, read_write> = var, %36 |
| %out:ptr<function, VertexOutput, read_write> = var |
| %39:ptr<function, vec4<f32>, read_write> = access %out, 0u |
| %40:ptr<uniform, mat4x4<f32>, read> = access %render_params, 0u |
| %41:mat4x4<f32> = load %40 |
| %42:mat4x4<f32> = let %41 |
| %43:vec3<f32> = load %position |
| %44:vec4<f32> = construct %43, 1.0f |
| %45:vec4<f32> = mul %42, %44 |
| store %39, %45 |
| %46:ptr<function, vec4<f32>, read_write> = access %out, 1u |
| %47:vec4<f32> = access %in, 1u |
| store %46, %47 |
| %48:ptr<function, vec2<f32>, read_write> = access %out, 2u |
| %49:vec2<f32> = access %in, 2u |
| store %48, %49 |
| %50:VertexOutput = load %out |
| ret %50 |
| } |
| } |
| %fs_main = @fragment func(%in_1:VertexOutput):vec4<f32> [@location(0)] { # %in_1: 'in' |
| $B4: { |
| %53:vec4<f32> = access %in_1, 1u |
| %color:ptr<function, vec4<f32>, read_write> = var, %53 |
| %55:f32 = load_vector_element %color, 3u |
| %56:f32 = let %55 |
| %57:vec2<f32> = access %in_1, 2u |
| %58:f32 = length %57 |
| %59:f32 = sub 1.0f, %58 |
| %60:f32 = max %59, 0.0f |
| %61:f32 = mul %56, %60 |
| store_vector_element %color, 3u, %61 |
| %62:vec4<f32> = load %color |
| ret %62 |
| } |
| } |
| %simulate = @compute @workgroup_size(64, 1, 1) func(%GlobalInvocationID:vec3<u32> [@global_invocation_id]):void { |
| $B5: { |
| %65:ptr<uniform, vec4<f32>, read> = access %sim_params, 1u |
| %66:vec4<f32> = load %65 |
| %67:vec2<f32> = swizzle %66, xy |
| %68:vec2<f32> = let %67 |
| %69:vec2<u32> = swizzle %GlobalInvocationID, xy |
| %70:vec2<f32> = convert %69 |
| %71:vec2<f32> = add %68, %70 |
| %72:ptr<uniform, vec4<f32>, read> = access %sim_params, 1u |
| %73:vec4<f32> = load %72 |
| %74:vec2<f32> = swizzle %73, zw |
| %75:vec2<f32> = mul %71, %74 |
| store %rand_seed, %75 |
| %76:u32 = access %GlobalInvocationID, 0u |
| %idx:u32 = let %76 |
| %78:ptr<storage, Particle, read_write> = access %data, 0u, %idx |
| %79:Particle = load %78 |
| %particle:ptr<function, Particle, read_write> = var, %79 |
| %81:ptr<function, vec3<f32>, read_write> = access %particle, 3u |
| %82:ptr<function, vec3<f32>, read_write> = access %particle, 3u |
| %83:f32 = load_vector_element %82, 2u |
| %84:ptr<uniform, f32, read> = access %sim_params, 0u |
| %85:f32 = load %84 |
| %86:f32 = mul %85, 0.5f |
| %87:f32 = sub %83, %86 |
| store_vector_element %81, 2u, %87 |
| %88:ptr<function, vec3<f32>, read_write> = access %particle, 0u |
| %89:ptr<function, vec3<f32>, read_write> = access %particle, 0u |
| %90:vec3<f32> = load %89 |
| %91:ptr<uniform, f32, read> = access %sim_params, 0u |
| %92:f32 = load %91 |
| %93:ptr<function, vec3<f32>, read_write> = access %particle, 3u |
| %94:vec3<f32> = load %93 |
| %95:vec3<f32> = mul %92, %94 |
| %96:vec3<f32> = add %90, %95 |
| store %88, %96 |
| %97:ptr<function, f32, read_write> = access %particle, 1u |
| %98:ptr<function, f32, read_write> = access %particle, 1u |
| %99:f32 = load %98 |
| %100:ptr<uniform, f32, read> = access %sim_params, 0u |
| %101:f32 = load %100 |
| %102:f32 = sub %99, %101 |
| store %97, %102 |
| %103:ptr<function, vec4<f32>, read_write> = access %particle, 2u |
| %104:ptr<function, f32, read_write> = access %particle, 1u |
| %105:f32 = load %104 |
| %106:f32 = smoothstep 0.0f, 0.5f, %105 |
| store_vector_element %103, 3u, %106 |
| %107:ptr<function, f32, read_write> = access %particle, 1u |
| %108:f32 = load %107 |
| %109:bool = lt %108, 0.0f |
| if %109 [t: $B6] { # if_1 |
| $B6: { # true |
| %coord:ptr<function, vec2<u32>, read_write> = var, vec2<u32>(0u) |
| loop [i: $B7, b: $B8, c: $B9] { # loop_1 |
| $B7: { # initializer |
| %111:texture_2d<f32> = load %tint_symbol |
| %112:u32 = textureNumLevels %111 |
| %113:u32 = sub %112, 1u |
| %level:ptr<function, u32, read_write> = var, %113 |
| next_iteration # -> $B8 |
| } |
| $B8: { # body |
| %115:u32 = load %level |
| %116:bool = gt %115, 0u |
| if %116 [t: $B10, f: $B11] { # if_2 |
| $B10: { # true |
| exit_if # if_2 |
| } |
| $B11: { # false |
| exit_loop # loop_1 |
| } |
| } |
| %117:texture_2d<f32> = load %tint_symbol |
| %118:vec2<u32> = load %coord |
| %119:u32 = load %level |
| %120:vec4<f32> = textureLoad %117, %118, %119 |
| %probabilites:vec4<f32> = let %120 |
| %122:f32 = call %rand |
| %123:vec4<f32> = construct %122 |
| %value:vec4<f32> = let %123 |
| %125:vec3<f32> = swizzle %probabilites, xyz |
| %126:vec4<f32> = construct 0.0f, %125 |
| %127:vec4<bool> = gte %value, %126 |
| %128:vec4<bool> = lt %value, %probabilites |
| %129:vec4<bool> = and %127, %128 |
| %mask:vec4<bool> = let %129 |
| %131:vec2<u32> = load %coord |
| %132:vec2<u32> = mul %131, 2u |
| store %coord, %132 |
| %133:vec2<bool> = swizzle %mask, yw |
| %134:bool = any %133 |
| %135:u32 = select 0u, 1u, %134 |
| %136:u32 = load_vector_element %coord, 0u |
| %137:u32 = add %136, %135 |
| store_vector_element %coord, 0u, %137 |
| %138:vec2<bool> = swizzle %mask, zw |
| %139:bool = any %138 |
| %140:u32 = select 0u, 1u, %139 |
| %141:u32 = load_vector_element %coord, 1u |
| %142:u32 = add %141, %140 |
| store_vector_element %coord, 1u, %142 |
| continue # -> $B9 |
| } |
| $B9: { # continuing |
| %143:u32 = load %level |
| %144:u32 = sub %143, 1u |
| store %level, %144 |
| next_iteration # -> $B8 |
| } |
| } |
| %145:vec2<u32> = load %coord |
| %146:vec2<f32> = convert %145 |
| %147:vec2<f32> = let %146 |
| %148:texture_2d<f32> = load %tint_symbol |
| %149:vec2<u32> = textureDimensions %148 |
| %150:vec2<f32> = convert %149 |
| %151:vec2<f32> = div %147, %150 |
| %uv:vec2<f32> = let %151 |
| %153:ptr<function, vec3<f32>, read_write> = access %particle, 0u |
| %154:vec2<f32> = sub %uv, 0.5f |
| %155:vec2<f32> = mul %154, 3.0f |
| %156:vec2<f32> = mul %155, vec2<f32>(1.0f, -1.0f) |
| %157:vec3<f32> = construct %156, 0.0f |
| store %153, %157 |
| %158:ptr<function, vec4<f32>, read_write> = access %particle, 2u |
| %159:texture_2d<f32> = load %tint_symbol |
| %160:vec2<u32> = load %coord |
| %161:vec4<f32> = textureLoad %159, %160, 0i |
| store %158, %161 |
| %162:ptr<function, vec3<f32>, read_write> = access %particle, 3u |
| %163:f32 = call %rand |
| %164:f32 = sub %163, 0.5f |
| %165:f32 = mul %164, 0.10000000149011611938f |
| store_vector_element %162, 0u, %165 |
| %166:ptr<function, vec3<f32>, read_write> = access %particle, 3u |
| %167:f32 = call %rand |
| %168:f32 = sub %167, 0.5f |
| %169:f32 = mul %168, 0.10000000149011611938f |
| store_vector_element %166, 1u, %169 |
| %170:ptr<function, vec3<f32>, read_write> = access %particle, 3u |
| %171:f32 = call %rand |
| %172:f32 = mul %171, 0.30000001192092895508f |
| store_vector_element %170, 2u, %172 |
| %173:ptr<function, f32, read_write> = access %particle, 1u |
| %174:f32 = call %rand |
| %175:f32 = mul %174, 2.0f |
| %176:f32 = add 0.5f, %175 |
| store %173, %176 |
| exit_if # if_1 |
| } |
| } |
| %177:ptr<storage, Particle, read_write> = access %data, 0u, %idx |
| %178:Particle = load %particle |
| %179:void = call %tint_store_and_preserve_padding, %177, %178 |
| ret |
| } |
| } |
| %import_level = @compute @workgroup_size(64, 1, 1) func(%coord_1:vec3<u32> [@global_invocation_id]):void { # %coord_1: 'coord' |
| $B12: { |
| %183:u32 = access %coord_1, 0u |
| %184:u32 = access %coord_1, 1u |
| %185:ptr<uniform, u32, read> = access %ubo, 0u |
| %186:u32 = load %185 |
| %187:u32 = mul %184, %186 |
| %188:u32 = add %183, %187 |
| %offset:u32 = let %188 |
| %190:ptr<storage, f32, read_write> = access %buf_out, 0u, %offset |
| %191:texture_2d<f32> = load %tex_in |
| %192:texture_2d<f32> = let %191 |
| %193:vec2<u32> = swizzle %coord_1, xy |
| %194:vec2<i32> = convert %193 |
| %195:vec4<f32> = textureLoad %192, %194, 0i |
| %196:f32 = access %195, 3u |
| store %190, %196 |
| ret |
| } |
| } |
| %export_level = @compute @workgroup_size(64, 1, 1) func(%coord_2:vec3<u32> [@global_invocation_id]):void { # %coord_2: 'coord' |
| $B13: { |
| %199:vec2<u32> = swizzle %coord_2, xy |
| %200:texture_storage_2d<rgba8unorm, write> = load %tex_out |
| %201:vec2<u32> = textureDimensions %200 |
| %202:vec2<u32> = construct %201 |
| %203:vec2<bool> = lt %199, %202 |
| %204:bool = all %203 |
| if %204 [t: $B14] { # if_3 |
| $B14: { # true |
| %205:u32 = access %coord_2, 0u |
| %206:u32 = access %coord_2, 1u |
| %207:ptr<uniform, u32, read> = access %ubo, 0u |
| %208:u32 = load %207 |
| %209:u32 = mul %206, %208 |
| %210:u32 = add %205, %209 |
| %dst_offset:u32 = let %210 |
| %212:u32 = access %coord_2, 0u |
| %213:u32 = mul %212, 2u |
| %214:u32 = access %coord_2, 1u |
| %215:u32 = mul %214, 2u |
| %216:ptr<uniform, u32, read> = access %ubo, 0u |
| %217:u32 = load %216 |
| %218:u32 = mul %215, %217 |
| %219:u32 = add %213, %218 |
| %src_offset:u32 = let %219 |
| %221:u32 = add %src_offset, 0u |
| %222:ptr<storage, f32, read> = access %buf_in, 0u, %221 |
| %223:f32 = load %222 |
| %a:f32 = let %223 |
| %225:u32 = add %src_offset, 1u |
| %226:ptr<storage, f32, read> = access %buf_in, 0u, %225 |
| %227:f32 = load %226 |
| %b:f32 = let %227 |
| %229:u32 = add %src_offset, 0u |
| %230:ptr<uniform, u32, read> = access %ubo, 0u |
| %231:u32 = load %230 |
| %232:u32 = add %229, %231 |
| %233:ptr<storage, f32, read> = access %buf_in, 0u, %232 |
| %234:f32 = load %233 |
| %c:f32 = let %234 |
| %236:u32 = add %src_offset, 1u |
| %237:ptr<uniform, u32, read> = access %ubo, 0u |
| %238:u32 = load %237 |
| %239:u32 = add %236, %238 |
| %240:ptr<storage, f32, read> = access %buf_in, 0u, %239 |
| %241:f32 = load %240 |
| %d:f32 = let %241 |
| %243:vec4<f32> = construct %a, %b, %c, %d |
| %244:f32 = dot %243, vec4<f32>(1.0f) |
| %sum:f32 = let %244 |
| %246:ptr<storage, f32, read_write> = access %buf_out, 0u, %dst_offset |
| %247:f32 = div %sum, 4.0f |
| store %246, %247 |
| %248:f32 = add %a, %b |
| %249:f32 = add %a, %b |
| %250:f32 = add %249, %c |
| %251:vec4<f32> = construct %a, %248, %250, %sum |
| %252:vec4<f32> = let %251 |
| %253:f32 = max %sum, 0.00009999999747378752f |
| %254:vec4<f32> = div %252, %253 |
| %probabilities:vec4<f32> = let %254 |
| %256:texture_storage_2d<rgba8unorm, write> = load %tex_out |
| %257:texture_storage_2d<rgba8unorm, write> = let %256 |
| %258:vec2<u32> = swizzle %coord_2, xy |
| %259:vec2<i32> = convert %258 |
| %260:void = textureStore %257, %259, %probabilities |
| exit_if # if_3 |
| } |
| } |
| ret |
| } |
| } |
| %tint_store_and_preserve_padding = func(%target:ptr<storage, Particle, read_write>, %value_param:Particle):void { |
| $B15: { |
| %263:ptr<storage, vec3<f32>, read_write> = access %target, 0u |
| %264:vec3<f32> = access %value_param, 0u |
| store %263, %264 |
| %265:ptr<storage, f32, read_write> = access %target, 1u |
| %266:f32 = access %value_param, 1u |
| store %265, %266 |
| %267:ptr<storage, vec4<f32>, read_write> = access %target, 2u |
| %268:vec4<f32> = access %value_param, 2u |
| store %267, %268 |
| %269:ptr<storage, vec3<f32>, read_write> = access %target, 3u |
| %270:vec3<f32> = access %value_param, 3u |
| store %269, %270 |
| 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. * |
| ******************************************************************** |
