Code that works in 23.0.1 crashes in 24.0.0 (macOS)

[peterhenryd]

Description
Calling Device::create_render_pipeline with a pipeline layout whose bind group includes a texture view array entry results in a crash when using 24.0.0, but not 23.0.1. I kept the title vague because my characterization of the issue may be incorrect.

Repro steps
I suspect this bug is only reproducible on certain macOS systems. This repository creates a simple effective environment for reproducing it on my system. It runs the same code from a single macro in two crates running versions 23.0.1 and 24.0.0 of WGPU respectively. The macro creates an instance, adapter, device, and then a render pipeline with a bind group layout with two entries (a texture view array and a sampler). No computations are performed. If you plan on trying to reproduce, run the following:

cargo run --package impl-functional
cargo run --package impl-dysfunctional

Expected vs observed behavior
On my system, the 23.0.1 example (impl-functional) works, while the 24.0.0 example (impl-dysfunctional) fails with the following error:

thread 'main' panicked at /[omitted]/wgpu-24.0.1/src/backend/wgpu_core.rs:1303:26:
wgpu error: Validation Error

Caused by:
  In Device::create_render_pipeline
    Internal error in ShaderStages(FRAGMENT) shader: mapping of ResourceBinding { group: 0, binding: 0 } is missing

Extra materials
Full stacktrace

Platform
MacBook Pro (13-inch, 2020, Two Thunderbolt 3 ports)
1.4 GHz Quad-Core Intel Core i5
Intel Iris Plus Graphics 645 1536 MB
macOS 14.7

[peterhenryd]

I believe I found the issue. It seems when there is a count present for the binding, the Metal device is assuming that the binding type is a always buffer when it could be a texture array.

[cwfitzgerald]

This may be in the right area, but that's how binding arrays are lowered to metal. They are a "argument buffer" which is basically a buffer which contains a bunch of TextureIDs/Buffer pointers in it. So when we see a binding array of any kind, we need to bind one single buffer for the entire array.

[peterhenryd]

Ah, it seems then that this would be the actual problem. I'm curious if the code should look like this instead:

// We intentionally don't dereference binding_arrays here,
// so that binding arrays fall to the buffer location.

match module.types[var.ty].inner {
  crate::TypeInner::Image { .. } => target.texture.is_some() /* ADDED: || target.buffer.is_some() */,
  crate::TypeInner::Sampler { .. } => target.sampler.is_some(),
  _ => target.buffer.is_some(),
}

I intentionally included the comment because I don't understand it but it seems relevant to what we're doing here. I have a feeling there's some implications I'm unaware of so I'll just leave this thought here for now, if anyone would like to remark.

Edit: this causes a Metal shader compilation error as the shader expects a texture attribute qualifier over the given buffer for the metal::texture2d_array. It seems this feature is just unimplemented for now?

[cwfitzgerald]

Oh! The issue is actually #3648! Your shader doesn't match your BGL - texture_2d_array doesn't use the count in the BGL. Set count to zero and everything will work.

You use count for binding_array<texture_2d<f32>> which is a different feature entirely.