xref: /aosp_15_r20/external/skia/tests/sksl/compute/MatrixMultiply.metal (revision c8dee2aa9b3f27cf6c858bd81872bdeb2c07ed17)
1#include <metal_stdlib>
2#include <simd/simd.h>
3#ifdef __clang__
4#pragma clang diagnostic ignored "-Wall"
5#endif
6using namespace metal;
7struct Inputs {
8    uint3 sk_GlobalInvocationID;
9};
10struct sizeBuffer {
11    int2 sizes[1];
12};
13struct inputs1 {
14    float data1[1];
15};
16struct inputs2 {
17    float data2[1];
18};
19struct result {
20    float resultData[1];
21};
22struct Globals {
23    device sizeBuffer* _anonInterface0;
24    const device inputs1* _anonInterface1;
25    const device inputs2* _anonInterface2;
26    device result* _anonInterface3;
27};
28kernel void computeMain(uint3 sk_GlobalInvocationID [[thread_position_in_grid]], device sizeBuffer& _anonInterface0 [[buffer(0)]], const device inputs1& _anonInterface1 [[buffer(1)]], const device inputs2& _anonInterface2 [[buffer(2)]], device result& _anonInterface3 [[buffer(3)]]) {
29    Globals _globals{&_anonInterface0, &_anonInterface1, &_anonInterface2, &_anonInterface3};
30    (void)_globals;
31    Inputs _in = { sk_GlobalInvocationID };
32    _globals._anonInterface0->sizes[2] = int2(_globals._anonInterface0->sizes[0].x, _globals._anonInterface0->sizes[1].y);
33    int2 resultCell = int2(int(_in.sk_GlobalInvocationID.x), int(_in.sk_GlobalInvocationID.y));
34    float result = 0.0;
35    for (int i = 0;i < _globals._anonInterface0->sizes[0].y; ++i) {
36        int a = i + resultCell.x * _globals._anonInterface0->sizes[0].y;
37        int b = resultCell.y + i * _globals._anonInterface0->sizes[1].y;
38        result += _globals._anonInterface1->data1[a] * _globals._anonInterface2->data2[b];
39    }
40    int index = resultCell.y + resultCell.x * _globals._anonInterface0->sizes[1].y;
41    _globals._anonInterface3->resultData[index] = result;
42    return;
43}
44