Unchecked shader creation is not honored on metal

[raphlinus]

Description
Even when shaders are created with create_shader_module_unchecked, bounds checks are still in place in the Metal backend.

Repro steps

1. Clone https://github.com/googlefonts/compute-shader-101, go into compute-shader-hello
2. Change the body of shader.wgsl to:

@group(0) @binding(0)
var<storage, read_write> data: array<atomic<u32>>;

@compute @workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    atomicExchange(&data[global_id.x], 1u);
}

3. Change the shader creation in main.rs to:

    let cs_module = unsafe { device.create_shader_module_unchecked(wgpu::ShaderModuleDescriptor {
        label: None,
        source: wgpu::ShaderSource::Wgsl(include_str!("shader.wgsl").into()),
    })};

4. Do cargo run

Note: this is a repro of gfx-rs/naga#1848. It was suggested to use the unchecked variant of shader creation, but that workaround is blocked by this bug.

Expected vs observed behavior

thread 'main' panicked at 'wgpu error: Validation Error

Caused by:
    In Device::create_compute_pipeline
    Internal error: Metal: program_source:27:48: error: cannot take the address of an rvalue of type 'metal::uint' (aka 'unsigned int')
    uint _e5 = metal::atomic_exchange_explicit(&uint(_e2) < 1 + (_buffer_sizes.size0 - 0 - 4) / 4 ? data[_e2] : DefaultConstructible(), 1u, metal::memory_order_relaxed);
                                               ^~~~~~~~~~


', /Users/raph/.cargo/registry/src/github.com-1ecc6299db9ec823/wgpu-0.14.0/src/backend/direct.rs:2403:5

Extra materials
I did a little digging, and it seems like there is logic in the vulkan back-end to plumb runtime_checks == false to the appropriate naga BoundsCheckPolicy values, but that appears to be missing in the metal back-end; the corresponding code seems to pass naga_options to the writer directly. Perhaps a better fix would be setting naga_options upstream rather than having to do it separately each back-end.

Platform
macOS 12.6, wgpu 0.14 and HEAD.

[raphlinus]

Also note, the naga cli behaves as expected: when the bounds check policy is not explicitly set, the shader is generated with no bounds check.

[jinleili]

This issue has disappeared on master.

I think create_shader_module_unchecked is just telling naga to uncheck, and the error above comes from Metal's internal error message when calling newLibraryWithSource:

wgpu/wgpu-hal/src/metal/device.rs

Lines 101 to 109 in 6d025a9

	let library = self
	.shared
	.device
	.lock()
	.new_library_with_source(source.as_ref(), &options)
	.map_err(|err| {
	log::warn!("Naga generated shader:\n{}", source);
	crate::PipelineError::Linkage(stage_bit, format!("Metal: {}", err))
	})?;

This validatiion error is unavoidable when the naga translated code does not conform to the MSL spec.

[raphlinus]

The translated code would not have violated the MSL spec had the unchecked request been honored.

[FL33TW00D]

Fixed and released: #3603