| SKIP: FAILED |
| |
| <dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: Params = struct @align(4) { |
| filterDim:u32 @offset(0) |
| blockDim:u32 @offset(4) |
| } |
| |
| Flip = struct @align(4) { |
| value:u32 @offset(0) |
| } |
| |
| %b1 = block { # root |
| %samp:ptr<handle, sampler, read> = var @binding_point(0, 0) |
| %params:ptr<uniform, Params, read> = var @binding_point(0, 1) |
| %inputTex:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 1) |
| %outputTex:ptr<handle, texture_storage_2d<rgba8unorm, write>, read> = var @binding_point(1, 2) |
| %flip:ptr<uniform, Flip, read> = var @binding_point(1, 3) |
| %tile:ptr<workgroup, array<array<vec3<f32>, 256>, 4>, read_write> = var |
| } |
| |
| %tint_symbol = @compute @workgroup_size(64, 1, 1) func(%WorkGroupID:vec3<u32> [@workgroup_id], %LocalInvocationID:vec3<u32> [@local_invocation_id]):void -> %b2 { |
| %b2 = block { |
| %10:ptr<uniform, u32, read> = access %params, 0u |
| %11:u32 = load %10 |
| %12:u32 = sub %11, 1u |
| %filterOffset:u32 = div %12, 2u |
| %14:texture_2d<f32> = load %inputTex |
| %dims:vec2<u32> = textureDimensions %14, 0i |
| %16:vec2<u32> = swizzle %WorkGroupID, xy |
| %17:ptr<uniform, u32, read> = access %params, 1u |
| %18:u32 = load %17 |
| %19:vec2<u32> = construct %18, 4u |
| %20:vec2<u32> = mul %16, %19 |
| %21:vec2<u32> = swizzle %LocalInvocationID, xy |
| %22:vec2<u32> = mul %21, vec2<u32>(4u, 1u) |
| %23:vec2<u32> = add %20, %22 |
| %24:vec2<u32> = construct %filterOffset, 0u |
| %baseIndex:vec2<u32> = sub %23, %24 |
| loop [i: %b3, b: %b4, c: %b5] { # loop_1 |
| %b3 = block { # initializer |
| %r:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b4 |
| } |
| %b4 = block { # body |
| %27:u32 = load %r |
| %28:bool = lt %27, 4u |
| if %28 [t: %b6, f: %b7] { # if_1 |
| %b6 = block { # true |
| exit_if # if_1 |
| } |
| %b7 = block { # false |
| exit_loop # loop_1 |
| } |
| } |
| loop [i: %b8, b: %b9, c: %b10] { # loop_2 |
| %b8 = block { # initializer |
| %c:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b9 |
| } |
| %b9 = block { # body |
| %30:u32 = load %c |
| %31:bool = lt %30, 4u |
| if %31 [t: %b11, f: %b12] { # if_2 |
| %b11 = block { # true |
| exit_if # if_2 |
| } |
| %b12 = block { # false |
| exit_loop # loop_2 |
| } |
| } |
| %32:u32 = load %c |
| %33:u32 = load %r |
| %34:vec2<u32> = construct %32, %33 |
| %35:vec2<u32> = add %baseIndex, %34 |
| %loadIndex:ptr<function, vec2<u32>, read_write> = var, %35 |
| %37:ptr<uniform, u32, read> = access %flip, 0u |
| %38:u32 = load %37 |
| %39:bool = neq %38, 0u |
| if %39 [t: %b13] { # if_3 |
| %b13 = block { # true |
| %40:vec2<u32> = load %loadIndex |
| %41:vec2<u32> = swizzle %40, yx |
| store %loadIndex, %41 |
| exit_if # if_3 |
| } |
| } |
| %42:u32 = load %r |
| %43:u32 = access %LocalInvocationID, 0u |
| %44:u32 = mul 4u, %43 |
| %45:u32 = load %c |
| %46:u32 = add %44, %45 |
| %47:ptr<workgroup, vec3<f32>, read_write> = access %tile, %42, %46 |
| %48:texture_2d<f32> = load %inputTex |
| %49:sampler = load %samp |
| %50:vec2<u32> = load %loadIndex |
| %51:vec2<f32> = convert %50 |
| %52:vec2<f32> = add %51, vec2<f32>(0.25f) |
| %53:vec2<f32> = convert %dims |
| %54:vec2<f32> = div %52, %53 |
| %55:vec4<f32> = textureSampleLevel %48, %49, %54, 0.0f |
| %56:vec3<f32> = swizzle %55, xyz |
| store %47, %56 |
| continue %b10 |
| } |
| %b10 = block { # continuing |
| %57:u32 = load %c |
| %58:u32 = add %57, 1u |
| store %c, %58 |
| next_iteration %b9 |
| } |
| } |
| continue %b5 |
| } |
| %b5 = block { # continuing |
| %59:u32 = load %r |
| %60:u32 = add %59, 1u |
| store %r, %60 |
| next_iteration %b4 |
| } |
| } |
| %61:void = workgroupBarrier |
| loop [i: %b14, b: %b15, c: %b16] { # loop_3 |
| %b14 = block { # initializer |
| %r_1:ptr<function, u32, read_write> = var, 0u # %r_1: 'r' |
| next_iteration %b15 |
| } |
| %b15 = block { # body |
| %63:u32 = load %r_1 |
| %64:bool = lt %63, 4u |
| if %64 [t: %b17, f: %b18] { # if_4 |
| %b17 = block { # true |
| exit_if # if_4 |
| } |
| %b18 = block { # false |
| exit_loop # loop_3 |
| } |
| } |
| loop [i: %b19, b: %b20, c: %b21] { # loop_4 |
| %b19 = block { # initializer |
| %c_1:ptr<function, u32, read_write> = var, 0u # %c_1: 'c' |
| next_iteration %b20 |
| } |
| %b20 = block { # body |
| %66:u32 = load %c_1 |
| %67:bool = lt %66, 4u |
| if %67 [t: %b22, f: %b23] { # if_5 |
| %b22 = block { # true |
| exit_if # if_5 |
| } |
| %b23 = block { # false |
| exit_loop # loop_4 |
| } |
| } |
| %68:u32 = load %c_1 |
| %69:u32 = load %r_1 |
| %70:vec2<u32> = construct %68, %69 |
| %71:vec2<u32> = add %baseIndex, %70 |
| %writeIndex:ptr<function, vec2<u32>, read_write> = var, %71 |
| %73:ptr<uniform, u32, read> = access %flip, 0u |
| %74:u32 = load %73 |
| %75:bool = neq %74, 0u |
| if %75 [t: %b24] { # if_6 |
| %b24 = block { # true |
| %76:vec2<u32> = load %writeIndex |
| %77:vec2<u32> = swizzle %76, yx |
| store %writeIndex, %77 |
| exit_if # if_6 |
| } |
| } |
| %78:u32 = access %LocalInvocationID, 0u |
| %79:u32 = mul 4u, %78 |
| %80:u32 = load %c_1 |
| %center:u32 = add %79, %80 |
| %82:bool = gte %center, %filterOffset |
| %83:bool = if %82 [t: %b25, f: %b26] { # if_7 |
| %b25 = block { # true |
| %84:u32 = sub 256u, %filterOffset |
| %85:bool = lt %center, %84 |
| exit_if %85 # if_7 |
| } |
| %b26 = block { # false |
| exit_if false # if_7 |
| } |
| } |
| %86:bool = if %83 [t: %b27, f: %b28] { # if_8 |
| %b27 = block { # true |
| %87:vec2<u32> = load %writeIndex |
| %88:vec2<bool> = lt %87, %dims |
| %89:bool = all %88 |
| exit_if %89 # if_8 |
| } |
| %b28 = block { # false |
| exit_if false # if_8 |
| } |
| } |
| if %86 [t: %b29] { # if_9 |
| %b29 = block { # true |
| %acc:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(0.0f) |
| loop [i: %b30, b: %b31, c: %b32] { # loop_5 |
| %b30 = block { # initializer |
| %f:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b31 |
| } |
| %b31 = block { # body |
| %92:u32 = load %f |
| %93:ptr<uniform, u32, read> = access %params, 0u |
| %94:u32 = load %93 |
| %95:bool = lt %92, %94 |
| if %95 [t: %b33, f: %b34] { # if_10 |
| %b33 = block { # true |
| exit_if # if_10 |
| } |
| %b34 = block { # false |
| exit_loop # loop_5 |
| } |
| } |
| %96:u32 = load %f |
| %97:u32 = add %center, %96 |
| %98:u32 = sub %97, %filterOffset |
| %i:ptr<function, u32, read_write> = var, %98 |
| %100:vec3<f32> = load %acc |
| %101:ptr<uniform, u32, read> = access %params, 0u |
| %102:u32 = load %101 |
| %103:f32 = convert %102 |
| %104:f32 = div 1.0f, %103 |
| %105:u32 = load %r_1 |
| %106:u32 = load %i |
| %107:ptr<workgroup, vec3<f32>, read_write> = access %tile, %105, %106 |
| %108:vec3<f32> = load %107 |
| %109:vec3<f32> = mul %104, %108 |
| %110:vec3<f32> = add %100, %109 |
| store %acc, %110 |
| continue %b32 |
| } |
| %b32 = block { # continuing |
| %111:u32 = load %f |
| %112:u32 = add %111, 1u |
| store %f, %112 |
| next_iteration %b31 |
| } |
| } |
| %113:texture_storage_2d<rgba8unorm, write> = load %outputTex |
| %114:vec2<u32> = load %writeIndex |
| %115:vec3<f32> = load %acc |
| %116:vec4<f32> = construct %115, 1.0f |
| %117:void = textureStore %113, %114, %116 |
| exit_if # if_9 |
| } |
| } |
| continue %b21 |
| } |
| %b21 = block { # continuing |
| %118:u32 = load %c_1 |
| %119:u32 = add %118, 1u |
| store %c_1, %119 |
| next_iteration %b20 |
| } |
| } |
| continue %b16 |
| } |
| %b16 = block { # continuing |
| %120:u32 = load %r_1 |
| %121:u32 = add %120, 1u |
| store %r_1, %121 |
| next_iteration %b15 |
| } |
| } |
| 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. * |
| ******************************************************************** |
