| SKIP: FAILED |
| |
| <dawn>/src/tint/lang/msl/writer/printer/printer.cc:355 internal compiler error: Matrix = struct @align(4) { |
| numbers:array<f32> @offset(0) |
| } |
| |
| Uniforms = struct @align(4) { |
| dimAOuter:u32 @offset(0) |
| dimInner:u32 @offset(4) |
| dimBOuter:u32 @offset(8) |
| } |
| |
| %b1 = block { # root |
| %firstMatrix:ptr<storage, Matrix, read> = var @binding_point(0, 0) |
| %secondMatrix:ptr<storage, Matrix, read> = var @binding_point(0, 1) |
| %resultMatrix:ptr<storage, Matrix, read_write> = var @binding_point(0, 2) |
| %uniforms:ptr<uniform, Uniforms, read> = var @binding_point(0, 3) |
| %mm_Asub:ptr<workgroup, array<array<f32, 64>, 64>, read_write> = var |
| %mm_Bsub:ptr<workgroup, array<array<f32, 64>, 64>, read_write> = var |
| } |
| |
| %mm_readA = func(%row:u32, %col:u32):f32 -> %b2 { |
| %b2 = block { |
| %10:ptr<uniform, u32, read> = access %uniforms, 0u |
| %11:u32 = load %10 |
| %12:bool = lt %row, %11 |
| %13:bool = if %12 [t: %b3, f: %b4] { # if_1 |
| %b3 = block { # true |
| %14:ptr<uniform, u32, read> = access %uniforms, 1u |
| %15:u32 = load %14 |
| %16:bool = lt %col, %15 |
| exit_if %16 # if_1 |
| } |
| %b4 = block { # false |
| exit_if false # if_1 |
| } |
| } |
| if %13 [t: %b5] { # if_2 |
| %b5 = block { # true |
| %17:ptr<uniform, u32, read> = access %uniforms, 1u |
| %18:u32 = load %17 |
| %19:u32 = mul %row, %18 |
| %20:u32 = add %19, %col |
| %21:ptr<storage, f32, read> = access %firstMatrix, 0u, %20 |
| %result:f32 = load %21 |
| ret %result |
| } |
| } |
| ret 0.0f |
| } |
| } |
| %mm_readB = func(%row_1:u32, %col_1:u32):f32 -> %b6 { # %row_1: 'row', %col_1: 'col' |
| %b6 = block { |
| %26:ptr<uniform, u32, read> = access %uniforms, 1u |
| %27:u32 = load %26 |
| %28:bool = lt %row_1, %27 |
| %29:bool = if %28 [t: %b7, f: %b8] { # if_3 |
| %b7 = block { # true |
| %30:ptr<uniform, u32, read> = access %uniforms, 2u |
| %31:u32 = load %30 |
| %32:bool = lt %col_1, %31 |
| exit_if %32 # if_3 |
| } |
| %b8 = block { # false |
| exit_if false # if_3 |
| } |
| } |
| if %29 [t: %b9] { # if_4 |
| %b9 = block { # true |
| %33:ptr<uniform, u32, read> = access %uniforms, 2u |
| %34:u32 = load %33 |
| %35:u32 = mul %row_1, %34 |
| %36:u32 = add %35, %col_1 |
| %37:ptr<storage, f32, read> = access %secondMatrix, 0u, %36 |
| %result_1:f32 = load %37 # %result_1: 'result' |
| ret %result_1 |
| } |
| } |
| ret 0.0f |
| } |
| } |
| %mm_write = func(%row_2:u32, %col_2:u32, %value:f32):void -> %b10 { # %row_2: 'row', %col_2: 'col' |
| %b10 = block { |
| %43:ptr<uniform, u32, read> = access %uniforms, 0u |
| %44:u32 = load %43 |
| %45:bool = lt %row_2, %44 |
| %46:bool = if %45 [t: %b11, f: %b12] { # if_5 |
| %b11 = block { # true |
| %47:ptr<uniform, u32, read> = access %uniforms, 2u |
| %48:u32 = load %47 |
| %49:bool = lt %col_2, %48 |
| exit_if %49 # if_5 |
| } |
| %b12 = block { # false |
| exit_if false # if_5 |
| } |
| } |
| if %46 [t: %b13] { # if_6 |
| %b13 = block { # true |
| %50:ptr<uniform, u32, read> = access %uniforms, 2u |
| %51:u32 = load %50 |
| %52:u32 = mul %row_2, %51 |
| %index:u32 = add %col_2, %52 |
| %54:ptr<storage, f32, read_write> = access %resultMatrix, 0u, %index |
| store %54, %value |
| exit_if # if_6 |
| } |
| } |
| ret |
| } |
| } |
| %tint_symbol = @compute @workgroup_size(16, 16, 1) func(%local_id:vec3<u32> [@local_invocation_id], %global_id:vec3<u32> [@global_invocation_id]):void -> %b14 { |
| %b14 = block { |
| %58:u32 = access %local_id, 1u |
| %tileRow:u32 = mul %58, 4u |
| %60:u32 = access %local_id, 0u |
| %tileCol:u32 = mul %60, 4u |
| %62:u32 = access %global_id, 1u |
| %globalRow:u32 = mul %62, 4u |
| %64:u32 = access %global_id, 0u |
| %globalCol:u32 = mul %64, 4u |
| %66:ptr<uniform, u32, read> = access %uniforms, 1u |
| %67:u32 = load %66 |
| %68:u32 = sub %67, 1u |
| %69:u32 = div %68, 64u |
| %numTiles:u32 = add %69, 1u |
| %acc:ptr<function, array<f32, 16>, read_write> = var |
| %ACached:ptr<function, f32, read_write> = var |
| %BCached:ptr<function, array<f32, 4>, read_write> = var |
| loop [i: %b15, b: %b16, c: %b17] { # loop_1 |
| %b15 = block { # initializer |
| %index_1:ptr<function, u32, read_write> = var, 0u # %index_1: 'index' |
| next_iteration %b16 |
| } |
| %b16 = block { # body |
| %75:u32 = load %index_1 |
| %76:bool = lt %75, 16u |
| if %76 [t: %b18, f: %b19] { # if_7 |
| %b18 = block { # true |
| exit_if # if_7 |
| } |
| %b19 = block { # false |
| exit_loop # loop_1 |
| } |
| } |
| %77:u32 = load %index_1 |
| %78:ptr<function, f32, read_write> = access %acc, %77 |
| store %78, 0.0f |
| continue %b17 |
| } |
| %b17 = block { # continuing |
| %79:u32 = load %index_1 |
| %80:u32 = add %79, 1u |
| store %index_1, %80 |
| next_iteration %b16 |
| } |
| } |
| %ColPerThreadA:u32 = let 4u |
| %82:u32 = access %local_id, 0u |
| %tileColA:u32 = mul %82, %ColPerThreadA |
| %RowPerThreadB:u32 = let 4u |
| %85:u32 = access %local_id, 1u |
| %tileRowB:u32 = mul %85, %RowPerThreadB |
| loop [i: %b20, b: %b21, c: %b22] { # loop_2 |
| %b20 = block { # initializer |
| %t:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b21 |
| } |
| %b21 = block { # body |
| %88:u32 = load %t |
| %89:bool = lt %88, %numTiles |
| if %89 [t: %b23, f: %b24] { # if_8 |
| %b23 = block { # true |
| exit_if # if_8 |
| } |
| %b24 = block { # false |
| exit_loop # loop_2 |
| } |
| } |
| loop [i: %b25, b: %b26, c: %b27] { # loop_3 |
| %b25 = block { # initializer |
| %innerRow:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b26 |
| } |
| %b26 = block { # body |
| %91:u32 = load %innerRow |
| %92:bool = lt %91, 4u |
| if %92 [t: %b28, f: %b29] { # if_9 |
| %b28 = block { # true |
| exit_if # if_9 |
| } |
| %b29 = block { # false |
| exit_loop # loop_3 |
| } |
| } |
| loop [i: %b30, b: %b31, c: %b32] { # loop_4 |
| %b30 = block { # initializer |
| %innerCol:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b31 |
| } |
| %b31 = block { # body |
| %94:u32 = load %innerCol |
| %95:bool = lt %94, %ColPerThreadA |
| if %95 [t: %b33, f: %b34] { # if_10 |
| %b33 = block { # true |
| exit_if # if_10 |
| } |
| %b34 = block { # false |
| exit_loop # loop_4 |
| } |
| } |
| %96:u32 = load %innerRow |
| %inputRow:u32 = add %tileRow, %96 |
| %98:u32 = load %innerCol |
| %inputCol:u32 = add %tileColA, %98 |
| %100:ptr<workgroup, f32, read_write> = access %mm_Asub, %inputRow, %inputCol |
| %101:u32 = load %innerRow |
| %102:u32 = add %globalRow, %101 |
| %103:u32 = load %t |
| %104:u32 = mul %103, 64u |
| %105:u32 = add %104, %inputCol |
| %106:f32 = call %mm_readA, %102, %105 |
| store %100, %106 |
| continue %b32 |
| } |
| %b32 = block { # continuing |
| %107:u32 = load %innerCol |
| %108:u32 = add %107, 1u |
| store %innerCol, %108 |
| next_iteration %b31 |
| } |
| } |
| continue %b27 |
| } |
| %b27 = block { # continuing |
| %109:u32 = load %innerRow |
| %110:u32 = add %109, 1u |
| store %innerRow, %110 |
| next_iteration %b26 |
| } |
| } |
| loop [i: %b35, b: %b36, c: %b37] { # loop_5 |
| %b35 = block { # initializer |
| %innerRow_1:ptr<function, u32, read_write> = var, 0u # %innerRow_1: 'innerRow' |
| next_iteration %b36 |
| } |
| %b36 = block { # body |
| %112:u32 = load %innerRow_1 |
| %113:bool = lt %112, %RowPerThreadB |
| if %113 [t: %b38, f: %b39] { # if_11 |
| %b38 = block { # true |
| exit_if # if_11 |
| } |
| %b39 = block { # false |
| exit_loop # loop_5 |
| } |
| } |
| loop [i: %b40, b: %b41, c: %b42] { # loop_6 |
| %b40 = block { # initializer |
| %innerCol_1:ptr<function, u32, read_write> = var, 0u # %innerCol_1: 'innerCol' |
| next_iteration %b41 |
| } |
| %b41 = block { # body |
| %115:u32 = load %innerCol_1 |
| %116:bool = lt %115, 4u |
| if %116 [t: %b43, f: %b44] { # if_12 |
| %b43 = block { # true |
| exit_if # if_12 |
| } |
| %b44 = block { # false |
| exit_loop # loop_6 |
| } |
| } |
| %117:u32 = load %innerRow_1 |
| %inputRow_1:u32 = add %tileRowB, %117 # %inputRow_1: 'inputRow' |
| %119:u32 = load %innerCol_1 |
| %inputCol_1:u32 = add %tileCol, %119 # %inputCol_1: 'inputCol' |
| %121:u32 = load %innerCol_1 |
| %122:ptr<workgroup, f32, read_write> = access %mm_Bsub, %121, %inputCol_1 |
| %123:u32 = load %t |
| %124:u32 = mul %123, 64u |
| %125:u32 = add %124, %inputRow_1 |
| %126:u32 = load %innerCol_1 |
| %127:u32 = add %globalCol, %126 |
| %128:f32 = call %mm_readB, %125, %127 |
| store %122, %128 |
| continue %b42 |
| } |
| %b42 = block { # continuing |
| %129:u32 = load %innerCol_1 |
| %130:u32 = add %129, 1u |
| store %innerCol_1, %130 |
| next_iteration %b41 |
| } |
| } |
| continue %b37 |
| } |
| %b37 = block { # continuing |
| %131:u32 = load %innerRow_1 |
| %132:u32 = add %131, 1u |
| store %innerRow_1, %132 |
| next_iteration %b36 |
| } |
| } |
| %133:void = workgroupBarrier |
| loop [i: %b45, b: %b46, c: %b47] { # loop_7 |
| %b45 = block { # initializer |
| %k:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b46 |
| } |
| %b46 = block { # body |
| %135:u32 = load %k |
| %136:bool = lt %135, 64u |
| if %136 [t: %b48, f: %b49] { # if_13 |
| %b48 = block { # true |
| exit_if # if_13 |
| } |
| %b49 = block { # false |
| exit_loop # loop_7 |
| } |
| } |
| loop [i: %b50, b: %b51, c: %b52] { # loop_8 |
| %b50 = block { # initializer |
| %inner:ptr<function, u32, read_write> = var, 0u |
| next_iteration %b51 |
| } |
| %b51 = block { # body |
| %138:u32 = load %inner |
| %139:bool = lt %138, 4u |
| if %139 [t: %b53, f: %b54] { # if_14 |
| %b53 = block { # true |
| exit_if # if_14 |
| } |
| %b54 = block { # false |
| exit_loop # loop_8 |
| } |
| } |
| %140:u32 = load %inner |
| %141:ptr<function, f32, read_write> = access %BCached, %140 |
| %142:u32 = load %k |
| %143:u32 = load %inner |
| %144:u32 = add %tileCol, %143 |
| %145:ptr<workgroup, f32, read_write> = access %mm_Bsub, %142, %144 |
| %146:f32 = load %145 |
| store %141, %146 |
| continue %b52 |
| } |
| %b52 = block { # continuing |
| %147:u32 = load %inner |
| %148:u32 = add %147, 1u |
| store %inner, %148 |
| next_iteration %b51 |
| } |
| } |
| loop [i: %b55, b: %b56, c: %b57] { # loop_9 |
| %b55 = block { # initializer |
| %innerRow_2:ptr<function, u32, read_write> = var, 0u # %innerRow_2: 'innerRow' |
| next_iteration %b56 |
| } |
| %b56 = block { # body |
| %150:u32 = load %innerRow_2 |
| %151:bool = lt %150, 4u |
| if %151 [t: %b58, f: %b59] { # if_15 |
| %b58 = block { # true |
| exit_if # if_15 |
| } |
| %b59 = block { # false |
| exit_loop # loop_9 |
| } |
| } |
| %152:u32 = load %innerRow_2 |
| %153:u32 = add %tileRow, %152 |
| %154:u32 = load %k |
| %155:ptr<workgroup, f32, read_write> = access %mm_Asub, %153, %154 |
| %156:f32 = load %155 |
| store %ACached, %156 |
| loop [i: %b60, b: %b61, c: %b62] { # loop_10 |
| %b60 = block { # initializer |
| %innerCol_2:ptr<function, u32, read_write> = var, 0u # %innerCol_2: 'innerCol' |
| next_iteration %b61 |
| } |
| %b61 = block { # body |
| %158:u32 = load %innerCol_2 |
| %159:bool = lt %158, 4u |
| if %159 [t: %b63, f: %b64] { # if_16 |
| %b63 = block { # true |
| exit_if # if_16 |
| } |
| %b64 = block { # false |
| exit_loop # loop_10 |
| } |
| } |
| %160:u32 = load %innerRow_2 |
| %161:u32 = mul %160, 4u |
| %162:u32 = load %innerCol_2 |
| %index_2:u32 = add %161, %162 # %index_2: 'index' |
| %164:ptr<function, f32, read_write> = access %acc, %index_2 |
| %165:ptr<function, f32, read_write> = access %acc, %index_2 |
| %166:f32 = load %165 |
| %167:f32 = load %ACached |
| %168:u32 = load %innerCol_2 |
| %169:ptr<function, f32, read_write> = access %BCached, %168 |
| %170:f32 = load %169 |
| %171:f32 = mul %167, %170 |
| %172:f32 = add %166, %171 |
| store %164, %172 |
| continue %b62 |
| } |
| %b62 = block { # continuing |
| %173:u32 = load %innerCol_2 |
| %174:u32 = add %173, 1u |
| store %innerCol_2, %174 |
| next_iteration %b61 |
| } |
| } |
| continue %b57 |
| } |
| %b57 = block { # continuing |
| %175:u32 = load %innerRow_2 |
| %176:u32 = add %175, 1u |
| store %innerRow_2, %176 |
| next_iteration %b56 |
| } |
| } |
| continue %b47 |
| } |
| %b47 = block { # continuing |
| %177:u32 = load %k |
| %178:u32 = add %177, 1u |
| store %k, %178 |
| next_iteration %b46 |
| } |
| } |
| %179:void = workgroupBarrier |
| continue %b22 |
| } |
| %b22 = block { # continuing |
| %180:u32 = load %t |
| %181:u32 = add %180, 1u |
| store %t, %181 |
| next_iteration %b21 |
| } |
| } |
| loop [i: %b65, b: %b66, c: %b67] { # loop_11 |
| %b65 = block { # initializer |
| %innerRow_3:ptr<function, u32, read_write> = var, 0u # %innerRow_3: 'innerRow' |
| next_iteration %b66 |
| } |
| %b66 = block { # body |
| %183:u32 = load %innerRow_3 |
| %184:bool = lt %183, 4u |
| if %184 [t: %b68, f: %b69] { # if_17 |
| %b68 = block { # true |
| exit_if # if_17 |
| } |
| %b69 = block { # false |
| exit_loop # loop_11 |
| } |
| } |
| loop [i: %b70, b: %b71, c: %b72] { # loop_12 |
| %b70 = block { # initializer |
| %innerCol_3:ptr<function, u32, read_write> = var, 0u # %innerCol_3: 'innerCol' |
| next_iteration %b71 |
| } |
| %b71 = block { # body |
| %186:u32 = load %innerCol_3 |
| %187:bool = lt %186, 4u |
| if %187 [t: %b73, f: %b74] { # if_18 |
| %b73 = block { # true |
| exit_if # if_18 |
| } |
| %b74 = block { # false |
| exit_loop # loop_12 |
| } |
| } |
| %188:u32 = load %innerRow_3 |
| %189:u32 = mul %188, 4u |
| %190:u32 = load %innerCol_3 |
| %index_3:u32 = add %189, %190 # %index_3: 'index' |
| %192:u32 = load %innerRow_3 |
| %193:u32 = add %globalRow, %192 |
| %194:u32 = load %innerCol_3 |
| %195:u32 = add %globalCol, %194 |
| %196:ptr<function, f32, read_write> = access %acc, %index_3 |
| %197:f32 = load %196 |
| %198:void = call %mm_write, %193, %195, %197 |
| continue %b72 |
| } |
| %b72 = block { # continuing |
| %199:u32 = load %innerCol_3 |
| %200:u32 = add %199, 1u |
| store %innerCol_3, %200 |
| next_iteration %b71 |
| } |
| } |
| continue %b67 |
| } |
| %b67 = block { # continuing |
| %201:u32 = load %innerRow_3 |
| %202:u32 = add %201, 1u |
| store %innerRow_3, %202 |
| next_iteration %b66 |
| } |
| } |
| 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. * |
| ******************************************************************** |
