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

@group(0) @binding(0)
var<storage, read> a: MatrixBatch;
@group(0) @binding(1)
var<storage, read> dldo: MatrixBatch;
@group(0) @binding(2)
var<storage, read_write> gtout: MatrixBatch;

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
  gtout.size = a.size;
  gtout.size.y = dldo.size.x;

  var s = 0f;
  var ai = global_id.z * a.size.y * a.size.x + global_id.x;
  var di = global_id.z * dldo.size.x * dldo.size.y + global_id.y;
  for (var i = 0u; i < a.size.y; i++) {
    s += a.data[ai] * dldo.data[di];
    ai += a.size.x;
    di += dldo.size.x;
  }

  gtout.data[global_id.z * gtout.size.y * gtout.size.x + global_id.y * gtout.size.x + global_id.x] = s;
}