Vectorize [un]pack4x{I, U}8[Clamp] on spv and msl

[robamler]

Connections
Related to, but independent of:

• Use specialized intrinsics for dot4{I, U}8Packed on SPIR-V and HLSL #7574
• Potentially optimize dot4{I,U}8Packed on Metal #7653

Description
Emits simpler code (using vectorized ops) for unpack4xI8, unpack4xU8, pack4xI8, pack4xU8, pack4xI8Clamp, and pack4xU8Clamp on SPIR-V (if capability "Int8" is available) and Metal (with MSL 2.1+).

• on SPIR-V (if capability "Int8" is available): exploit that SClamp, UClamp, and OpUConvert accept vector arguments (in which case results are computed per component), and that OpBitcast can cast between vectors and scalars, with a well-defined bit order that matches that required by the WGSL spec (see below for details).
• on Metal (MSL 2.1+ only): translate (see shader playground)
  • unpack4x{I, U}8(x) → [u]int4(as_type<packed_[u]char4>(x)),
  • pack4x{I, U}8(x) → as_type<uint>(packed_[u]char4(x)),
  • pack4xI8Clamp(x) → as_type<uint>(packed_char4(metal::clamp(x, -128, 127))), and
  • pack4xU8Clamp(x) → as_type<uint>(packed_uchar4(metal::clamp(x, 0, 255))).

Regarding byte order (this was not relevant for the dot product in #7653 but is relevant now):

• WGSL states:

  Component e[i] of the input is mapped to bits 8 x i through 8 x i + 7 of the result.

• This corresponds to little-endian byte order, which is what Apple uses for Metal.
• SPIR-V states the same byte order in more complicated words:

  Within this mapping, any single component of S [remark: a 32-bit int in our case] (mapping to multiple components of L [remark: a 4-vector of 8-bit ints in our case]) maps its lower-ordered bits to the lower-numbered components of L

Testing

• SPIR-V part was already covered by snapshot tests.
• Added a snapshot test for MSL 2.1+
• It would be nice to add tests for correctness (rather than just validity) of the MSL and SPIR-V output here. I.e., run something like the test in wgpu_gpu::shader::data_builtins::create_unpack4xU8_test on both the optimized and the unoptimized generated code. I'm not sure how to do this but I'd welcome any hints.

Squash or Rebase?
Each commit should pass CI.

Checklist

• Run cargo fmt.
• Run taplo format.
• Run cargo clippy --tests. If applicable, add:
  • --target wasm32-unknown-unknown
• Run cargo xtask test to run tests.
• If this contains user-facing changes, add a CHANGELOG.md entry.

[robamler]

FWIW, the optimization of pack4xU8 speeds up our motivating use case (same as in #7595 (comment)) by between 1.1% and 1.7%, with a standard error of at most 0.06% across all benchmarks. So the improvement here is much smaller than what we found for dot4I8Packed, but it's definitely statistically significant. Curiously, the impact of the optimizations of unpack4xU8 is within the error bars, even though this function is also used in our shader, and I verified manually that the vectorized code for it gets emitted. These are very ad-hoc benchmarks, and I ran them only on SPIR-V for now.

[robamler]

I squashed some fixes to make this PR easier to review. Each commit should now separately pass CI.

[teoxoy]

Looks good overall, left a few small comments.

[robamler]

Thank you so much for the review, this is great! I addressed your comments.

[teoxoy]

It would be nice to add tests for correctness (rather than just validity) of the MSL and SPIR-V output here. I.e., run something like the test in wgpu_gpu::shader::data_builtins::create_unpack4xU8_test on both the optimized and the unoptimized generated code. I'm not sure how to do this but I'd welcome any hints.

This is blocked on #4582.

[teoxoy]

Thanks for addressing the comments, let's land it!

[teoxoy]

Ah it would probably be nice to land this without squashing, could you move the code changes in the last commit to their respective commits?

[robamler]

No problem, done.

[teoxoy]

Thank you!