#include <metal_stdlib>
#include <simd/simd.h>
#include <metal_simdgroup_matrix>

using namespace metal;

struct SSBO
{
    bfloat data[1];
};

kernel void main0(device SSBO& ssbo [[buffer(0)]])
{
    simdgroup_bfloat8x8 _21;
    simdgroup_load(_21, &ssbo.data[0u], 8u);
    simdgroup_bfloat8x8 _22;
    simdgroup_load(_22, &ssbo.data[0u], 8u);
    simdgroup_bfloat8x8 _23;
    simdgroup_load(_23, &ssbo.data[0u], 8u);
    simdgroup_bfloat8x8 _24;
    simdgroup_multiply_accumulate(_24, _21, _22, _23);
    simdgroup_store(_24, &ssbo.data[0u], 8u);
}

