struct PackedMatrixBatch {
  size: vec4<u32>,
  data: array<vec4<f32>>,
}

struct MatrixBatch {
  size: vec4<u32>,
  data: array<f32>,
}

@group(0) @binding(0)
var<storage, read> e: PackedMatrixBatch;
@group(0) @binding(1)
var<storage, read> tw: PackedMatrixBatch;
@group(0) @binding(2)
var<storage, read_write> aw: MatrixBatch;

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
  aw.size.x = tw.size.y;
  aw.size.y = e.size.y;
  aw.size.z = e.size.z;

  var s = 0f;
  let vc = e.size.x / 4;
  let lo = (global_id.z * e.size.y + global_id.y) * vc;
  let ro = (global_id.z * tw.size.y + global_id.x) * vc;
  for (var i = 0u; i < vc; i++) {
    s += dot(e.data[lo + i], tw.data[ro + i]);
  }

  aw.data[(global_id.z * aw.size.y + global_id.y) * aw.size.x + global_id.x] = s;
}