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

@group(0) @binding(0)
var<storage, read> wa: MatrixBatch;
@group(0) @binding(1)
var<storage, read> otdldo: MatrixBatch;
@group(0) @binding(2)
var<storage, read_write> totdldowa: MatrixBatch;

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

  var s = 0f;
  var li = global_id.z * otdldo.size.x * otdldo.size.y + global_id.x * otdldo.size.x;
  var ri = global_id.z * wa.size.x * wa.size.y + global_id.y;
  for (var i = 0u; i < otdldo.size.x; i++) {
    s += otdldo.data[li] * wa.data[ri];
    li++;
    ri += wa.size.x;
  }

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