blob: 66a945bf426260a35b075bbde6f898b2e90c1861 [file] [log] [blame]
SKIP: crbug.com/tint/833 loop emission is broken
#include <metal_stdlib>
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;
}