alias number = f32;

struct Scalar {
  size: vec4<u32>,
  data: array<number>,
}

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

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

@group(0) @binding(0)
var<storage, read> m: MatrixBatch;
@group(0) @binding(1)
var<storage, read> v: UVectorBatch;
@group(0) @binding(2)
var<storage, read> lr: Scalar;
@group(0) @binding(3)
var<storage, read_write> grs: MatrixBatch;

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
  grs.size = m.size;

  var vi = global_id.z * v.size.x + global_id.y;
  var gi = vi * grs.size.x + global_id.x;
  grs.data[gi] = lr.data[0] * (m.data[gi] - select(0f, 1f, v.data[vi] == global_id.x));
}