| #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]; |
| }; |
| |
| kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) { |
| 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; |
| while (true) { |
| if (!((i < dimInner))) { |
| break; |
| } |
| uint const a = (i + (resultCell.x * dimInner)); |
| uint const b = (resultCell.y + (i * dimOutter)); |
| result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b])); |
| { |
| i = (i + 1u); |
| } |
| } |
| } |
| uint const index = (resultCell.y + (resultCell.x * dimOutter)); |
| resultMatrix.numbers[index] = result; |
| return; |
| } |
| |