Compute shader and buffer support

[tychedelia]

Adds initial support for Compute and Buffer API objects.

Some important notes:

• Like Image, our Buffer type wraps a ShaderBuffer asset and uses the CPU-side asset data as a kind of cache. We try to optimize reads so that readback happens only when the CPU-side data is stale.
• Buffers bound as read_write storage to any pipeline (compute or material) are automatically marked stale after each frame tick via a system in Last.
• Our Python interface is very high level and looks just like an array-like thing. Because buffers in WebGPU are typed (i.e. no heterogeneous or reinterpret_cast style buffers), we do validation to ensure that you are only storing the same type of element in the buffer.
• Compute wraps a cached compute pipeline built from a Shader with a @compute entry point. We use the same dynamic shader stuff here as materials.
• dispatch(x, y, z) runs app.update() to flush pending buffer writes to the GPU, then submits the compute pass directly on the render sub-app.
• Uniform, storage buffer, and texture bindings are set via shader reflection, same as materials.
• Custom materials can now accept ShaderValue::Buffer for storage bindings
• MaterialValue -> ShaderValue to represent it being shared by any shader op.

You can test this in IPython via something like:

  from mewnala import Graphics, Shader, Compute, Buffer

  g = Graphics.new_offscreen(1, 1, "", None)
  g.begin_draw()


  shader = Shader("""
  @group(0) @binding(0)
  var<storage, read_write> data: array<f32>;

  @compute @workgroup_size(4)
  fn main(@builtin(global_invocation_id) id: vec3<u32>) {
      data[id.x] = data[id.x] * 2.0;
  }
  """)

  buf = Buffer(data=[1.0, 2.0, 3.0, 4.0])
  compute = Compute(shader)
  compute.set(data=buf)
  compute.dispatch(1, 1, 1)

  print(buf[0], buf[1], buf[2], buf[3])

  compute.dispatch(1, 1, 1)
  print(buf.read())

  buf[0] = 100.0
  compute.dispatch(1, 1, 1)
  print(buf[0])

  g.end_draw()