SKIP: crbug.com/tint/833 loop emission is broken #include using namespace metal; struct Uniforms { /* 0x0000 */ packed_uint2 aShape; /* 0x0008 */ packed_uint2 bShape; /* 0x0010 */ packed_uint2 outShape; }; struct Matrix { /* 0x0000 */ uint numbers[1]; }; struct tint_symbol_2 { uint3 global_id [[thread_position_in_grid]]; }; kernel void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) { uint3 const global_id = tint_symbol_1.global_id; uint2 const resultCell = uint2(global_id.y, global_id.x); uint const dimInner = uniforms.aShape.y; uint const dimOutter = uniforms.outShape.y; uint result = 0u; { uint i = 0u; { bool tint_msl_is_first_1 = true; uint const a; uint const b; for(;;) { if (!tint_msl_is_first_1) { i = (i + 1u); } tint_msl_is_first_1 = false; if (!((i < dimInner))) { break; } a = (i + (resultCell.x * dimInner)); b = (resultCell.y + (i * dimOutter)); result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b])); } } } uint const index = (resultCell.y + (resultCell.x * dimOutter)); resultMatrix.numbers[index] = result; return; }