SHADER_I16 full support

[JMS55]

Disclaimer

Entirely LLM generated. My process was to point the LLM at the SPIRV, MSL, and HLSL specs, as well as existing issues/PRs/code related to u64, f16, etc. I also had it write a crap ton of tests covering all the places I could think of that could be affected.

Connections

Closes #4424 (16-bit integer support in naga)
Fixes #7468 (Vulkan: SHADER_I16 does not enable storage_buffer16_bit_access)
Partially addresses #7109 (i16 division/modulo overflow is now guarded in HLSL/MSL/SPIR-V backends)

Description

Implements i16/u16 16-bit integer type support in WGSL shaders, gated behind the existing Features::SHADER_I16 feature flag and a new enable wgpu_int16; WGSL enable extension.

What's included

Naga (shader compiler):

• IR types (Scalar::I16/U16, Literal::I16/U16), validation (SHADER_INT16 capability), and WGSL parsing (i16, u16, vec2<i16>, etc.) with enable wgpu_int16; gating
• All backend code generation: SPIR-V (OpTypeInt 16 + SPV_KHR_16bit_storage), MSL (short/ushort), HLSL (int16_t/uint16_t), GLSL (int16_t/uint16_t), WGSL round-trip
• Builtin math overloads (abs, min, max, clamp, sign, etc.) via ScalarSet additions
• Constant evaluation, pipeline constants, override support, float conversion bounds
• Signed division/modulo overflow guards (i16::MIN / -1) in SPIR-V, MSL, and HLSL backends
• Const-expression index evaluation (TryFrom<Literal> for u32)
• MSL: fixed put_bitcasted_expression to add static_cast truncation for sub-32-bit types, preventing invalid as_type casts due to C++ integer promotion (ushort + ushort promotes to int, making as_type<short>(int) a size mismatch)
• SPIR-V: sign-extend i16 literals into the 32-bit constant word so spirv-as can round-trip the disassembly
• MSL: use typed literals in metal::select calls within div/mod/abs helpers (e.g. short(1) instead of bare 1) to avoid overload ambiguity with sub-32-bit types
• HLSL: use type constructors instead of asuint()/asint() for 16-bit bitcasts (those intrinsics only support 32-bit types)
• WGSL writer: emit enable wgpu_int16; when the module uses i16/u16 types

Atomic validation fix (pre-existing):

• Tightened Ti::Atomic validation from Sint | Uint, width: _ to explicit width: 4 and width: 8 arms, so atomic<u16> is correctly rejected instead of silently accepted. This was a pre-existing bug that also affected atomic<u8> if anyone ever tried it via SPIR-V input.

HAL backends:

• Vulkan: SHADER_I16 reporting now requires storageBuffer16BitAccess + uniformAndStorageBuffer16BitAccess (matching SHADER_F16 behavior). VK_KHR_16bit_storage extension now requested for SHADER_I16 (not just SHADER_F16). Int16 added to available SPIR-V capabilities.
• Metal: SHADER_I16 unconditionally reported (MSL natively supports short/ushort).
• DX12: SHADER_I16 gated on SM 6.2 + Native16BitShaderOpsSupported (same condition as SHADER_F16). -enable-16bit-types now passed to DXC at runtime when SHADER_I16 is enabled (not just SHADER_F16).

wgpu-naga-bridge:

• Maps Features::SHADER_I16 to Capabilities::SHADER_INT16.

wgpu-types:

• Updated SHADER_I16 feature docs to reflect naga support and platform availability.

What's NOT included (intentional)

• Vertex formats: VertexFormat::Uint16 still maps to Scalar::U32 in shader validation, consistent with Float16 mapping to Scalar::F32. The hardware expands 16-bit vertex data to 32-bit before the shader sees it, per WebGPU spec.
• Texture formats: R16Uint/R16Sint storage textures still produce 32-bit component types in shaders, consistent with the WebGPU spec.
• HLSL struct padding: The HLSL backend's struct padding code (naga/src/back/hlsl/writer.rs:1420) uses int (4-byte) padding and divides gap sizes by 4, which can't represent sub-4-byte gaps. This is a pre-existing limitation shared with f16 — HLSL's implicit padding handles it correctly in practice, and f16 has been shipping with this behavior and passing CTS.
• Literal suffixes: No WGSL literal suffix for i16/u16 (e.g. no 42s syntax). Values are constructed via u16(42) / i16(42).
• Subgroup bitwise ops: subgroupAnd/subgroupOr/subgroupXor are rejected for 16-bit types at the naga validation level because HLSL's WaveActiveBitAnd/Or/Xor don't support them.
• Switch statements: i16/u16 cannot be used as switch selectors — the WGSL frontend's abstract int conversion consensus logic doesn't handle i16/u16 targets yet.
• Push constants/immediates: var<immediate> with i16/u16 types is rejected, matching f16 behavior. Vulkan requires StoragePushConstant16 for 16-bit push constants, which wgpu does not currently enable.
• Atomics: atomic<u16> / atomic<i16> are rejected — 16-bit atomics aren't supported by hardware.
• GL/GLES HAL: SHADER_I16 not reported as a device feature. The naga GLSL backend can generate the code (int16_t/uint16_t), but the wgpu-hal GL/GLES adapter doesn't expose the feature since GLSL 16-bit integer support requires rare extensions.

Testing

• Snapshot test (naga/tests/in/wgsl/int16.wgsl): Comprehensive coverage of types (scalars, vectors, structs, arrays), storage (uniform/storage/workgroup buffers), operators (arithmetic, bitwise, shifts, negation, comparisons), builtins (abs/min/max/clamp/sign), select, type conversions, bitcast, array indexing with u16, vector arithmetic, and subgroup operations. Output verified for SPIR-V, HLSL, MSL, and WGSL backends.
• SPIR-V validation: Generated SPIR-V passes both spirv-val and spirv-as round-trip.
• Error tests (naga/tests/naga/wgsl_errors.rs):
  • int16_capability_and_enable: Tests that u16/i16 require enable wgpu_int16; (parse error) and SHADER_INT16 capability (validation error). Covers zero-value expressions, constructor calls, and var declarations.
  • int16_in_atomic: Tests that atomic<u16> and atomic<i16> are rejected with InvalidAtomicWidth.
• DXC xtask validation: -enable-16bit-types added for SM 6.2+ profiles only.

Squash or Rebase?

Squash.

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.

[inner-daemons]

Does std140 even support i16? I was under the impression that you couldn't use these in uniforms