xref: /aosp_15_r20/external/skia/tests/sksl/compute/MatrixMultiply.metal (revision c8dee2aa9b3f27cf6c858bd81872bdeb2c07ed17)

#include <metal_stdlib>
#include <simd/simd.h>
#ifdef __clang__
#pragma clang diagnostic ignored "-Wall"
#endif
using namespace metal;
struct Inputs {
    uint3 sk_GlobalInvocationID;
};
struct sizeBuffer {
    int2 sizes[1];
};
struct inputs1 {
    float data1[1];
};
struct inputs2 {
    float data2[1];
};
struct result {
    float resultData[1];
};
struct Globals {
    device sizeBuffer* _anonInterface0;
    const device inputs1* _anonInterface1;
    const device inputs2* _anonInterface2;
    device result* _anonInterface3;
};
kernel 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)]]) {
    Globals _globals{&_anonInterface0, &_anonInterface1, &_anonInterface2, &_anonInterface3};
    (void)_globals;
    Inputs _in = { sk_GlobalInvocationID };
    _globals._anonInterface0->sizes[2] = int2(_globals._anonInterface0->sizes[0].x, _globals._anonInterface0->sizes[1].y);
    int2 resultCell = int2(int(_in.sk_GlobalInvocationID.x), int(_in.sk_GlobalInvocationID.y));
    float result = 0.0;
    for (int i = 0;i < _globals._anonInterface0->sizes[0].y; ++i) {
        int a = i + resultCell.x * _globals._anonInterface0->sizes[0].y;
        int b = resultCell.y + i * _globals._anonInterface0->sizes[1].y;
        result += _globals._anonInterface1->data1[a] * _globals._anonInterface2->data2[b];
    }
    int index = resultCell.y + resultCell.x * _globals._anonInterface0->sizes[1].y;
    _globals._anonInterface3->resultData[index] = result;
    return;
}