Naga Fails to Validate Out Texture Argument

[lisyarus]

Description
In my C++ project, that uses wgpu-native-v22.1.0.5, shader compilation fails on MacOS on M1 hardware with the following error:

thread '<unnamed>' panicked at /Users/runner/.cargo/git/checkouts/wgpu-53e70f8674b08dd4/5c5c8b1/naga/src/back/msl/writer.rs:249:29:
internal error: entered unreachable code: module is not valid
stack backtrace:
   0:        0x100dcf2e8 - std::backtrace_rs::backtrace::libunwind::trace::hbebc8679d47bdc2c
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/../../backtrace/src/backtrace/libunwind.rs:116:5
   1:        0x100dcf2e8 - std::backtrace_rs::backtrace::trace_unsynchronized::h3a2e9637943241aa
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/../../backtrace/src/backtrace/mod.rs:66:5
   2:        0x100dcf2e8 - std::sys::backtrace::_print_fmt::he430849680584674
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/sys/backtrace.rs:65:5
   3:        0x100dcf2e8 - <std::sys::backtrace::BacktraceLock::print::DisplayBacktrace as core::fmt::Display>::fmt::h243268f17d714c7f
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/sys/backtrace.rs:40:26
   4:        0x100e0ef68 - core::fmt::rt::Argument::fmt::h0d339881c25f3c31
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/core/src/fmt/rt.rs:173:76
   5:        0x100e0ef68 - core::fmt::write::hb3cfb8a30e72d7ff
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/core/src/fmt/mod.rs:1182:21
   6:        0x100dc52cc - std::io::Write::write_fmt::hfb2314975de9ecf1
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/io/mod.rs:1827:15
   7:        0x100dd17f8 - std::sys::backtrace::BacktraceLock::print::he14461129ccbfef5
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/sys/backtrace.rs:43:9
   8:        0x100dd17f8 - std::panicking::default_hook::{{closure}}::h14c7718ccf39d316
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/panicking.rs:269:22
   9:        0x100dd141c - std::panicking::default_hook::hc62e60da3be2f352
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/panicking.rs:296:9
  10:        0x100dd23d4 - std::panicking::rust_panic_with_hook::h09e8a656f11e82b2
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/panicking.rs:800:13
  11:        0x100dd1cfc - std::panicking::begin_panic_handler::{{closure}}::h1230eb3cc91b241c
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/panicking.rs:667:13
  12:        0x100dcf774 - std::sys::backtrace::__rust_end_short_backtrace::hc3491307aceda2c2
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/sys/backtrace.rs:168:18
  13:        0x100dd19ec - rust_begin_unwind
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/std/src/panicking.rs:665:5
  14:        0x100e47a18 - core::panicking::panic_fmt::ha4b80a05b9fff47a
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/core/src/panicking.rs:74:14
  15:        0x100bd743c - <naga::back::msl::writer::TypeContext as core::fmt::Display>::fmt::ha54a986cde67cb26
  16:        0x100e0ef68 - core::fmt::rt::Argument::fmt::h0d339881c25f3c31
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/core/src/fmt/rt.rs:173:76
  17:        0x100e0ef68 - core::fmt::write::hb3cfb8a30e72d7ff
                               at /rustc/eeb90cda1969383f56a2637cbd3037bdf598841c/library/core/src/fmt/mod.rs:1182:21
  18:        0x100beae7c - naga::back::msl::writer::Writer<W>::write_functions::hd0ee49a746763fbb
  19:        0x100be6890 - naga::back::msl::writer::Writer<W>::write::h2d46705c3337cc50
  20:        0x100cfaec4 - naga::back::msl::write_string::h8a026bbc33459419
  21:        0x100b8ded0 - wgpu_hal::metal::device::<impl wgpu_hal::metal::Device>::load_shader::h3a4964a03ce918ff
  22:        0x100b98e3c - objc::rc::autorelease::autoreleasepool::hebbb0043551780d7
  23:        0x100ab7d40 - wgpu_core::device::resource::Device<A>::create_render_pipeline::h76ddb7fcd2e3bf60
  24:        0x100a5c914 - wgpu_core::device::global::<impl wgpu_core::global::Global>::device_create_render_pipeline::h595b34c9c52c9974
  25:        0x100a3fc08 - _wgpuDeviceCreateRenderPipeline
  26:        0x1009f9ff0 - __ZN15PreviewPipelineC2EP14WGPUDeviceImplR14ShaderRegistry17WGPUTextureFormatP23WGPUBindGroupLayoutImplS6_
  27:        0x1009fa7e8 - __ZN8Renderer4ImplC2EP14WGPUDeviceImplP13WGPUQueueImpl17WGPUTextureFormatR14ShaderRegistry
  28:        0x1009fad84 - __ZN8RendererC1EP14WGPUDeviceImplP13WGPUQueueImpl17WGPUTextureFormatR14ShaderRegistry
  29:        0x1009f7b24 - _main

Looking at the naga/src/back/msl/writer.rs:249 from the wgpu-22.1.0 tag, it seems that the compiler somehow failed to parse texture access flags or something like that.

Here is the full shader source that triggers the error:

const PI = 3.14159265358979323846;

// Heaviside step function
fn chiPlus(x : f32) -> f32 {
	return step(0.0, x);
}

fn completeBasis(Z : vec3f) -> mat3x3f {
	var X = select(vec3f(1.0, 0.0, 0.0), vec3f(0.0, 1.0, 0.0), abs(Z.x) > 0.5);
	X -= Z * dot(X, Z);
	X = normalize(X);
	let Y = cross(Z, X);
	return mat3x3f(X, Y, Z);
}

fn perspectiveDivide(v : vec4f) -> vec3f {
	return v.xyz / v.w;
}

fn det(v0 : vec3f, v1 : vec3f, v2 : vec3f) -> f32 {
	return dot(v0, cross(v1, v2));
}

fn solve3x3Equation(matrix : mat3x3f, rhs : vec3f) -> vec3f {
	// Use Cramer's rule to solve the system
	let det = det(matrix[0], matrix[1], matrix[2]);
	let d0 = det(rhs, matrix[1], matrix[2]);
	let d1 = det(matrix[0], rhs, matrix[2]);
	let d2 = det(matrix[0], matrix[1], rhs);

	return vec3f(d0, d1, d2) / det;
}
struct Camera
{
	viewProjectionMatrix : mat4x4f,
	viewProjectionInverseMatrix : mat4x4f,
	position: vec3f,
	screenSize : vec2u,
	frameID : u32,
	globalFrameID : u32,
}
struct Material
{
	baseColorFactorAndTransmission : vec4f,
	metallicRoughnessFactorAndIor : vec4f,
	emissiveFactor : vec4f,
	textureLayers : vec4u,
}


const MAX_ENV_MAP_INTENSITY = 100.0;

fn sampleEnvMap(environmentMap: texture_storage_2d<rgba32float, read>, direction : vec3f) -> vec3f {
	let dimensions = vec2f(textureDimensions(environmentMap));
	let x = atan2(direction.z, direction.x) / PI * 0.5 + 0.5;
	let y = -atan2(direction.y, length(direction.xz)) / PI + 0.5;

	return min(vec3f(MAX_ENV_MAP_INTENSITY), textureLoad(environmentMap, vec2u(dimensions * vec2f(x, y))).xyz);
}

fn acesToneMap(color : vec3f) -> vec3f {
	let A = 2.51;
	let B = 0.03;
	let C = 2.43;
	let D = 0.59;
	let E = 0.14;

	return saturate((color * (A * color + B)) / (color * (C * color + D) + E));
}

fn gammaCorrect(color : vec3f) -> vec3f {
	return pow(color, vec3f(1.0 / 2.2));
}

@group(0) @binding(0) var<uniform> camera : Camera;

@group(1) @binding(0) var<storage, read> materials : array<Material>;
@group(1) @binding(1) var environmentMap : texture_storage_2d<rgba32float, read>;
@group(1) @binding(2) var textureSampler : sampler;
@group(1) @binding(3) var albedoTexture : texture_2d_array<f32>;

struct VertexInput {
	@builtin(vertex_index) index : u32,
	@location(0) position : vec3f,
	@location(1) normal : vec3f,
	@location(2) materialID : u32,
	@location(3) texcoord : vec2f,
}

struct VertexOutput {
	@builtin(position) position : vec4f,
	@location(0) worldPosition : vec3f,
	@location(1) normal : vec3f,
	@location(2) texcoord : vec2f,
	@location(3) @interpolate(flat) materialID : u32,
}

@vertex
fn vertexMain(in : VertexInput) -> VertexOutput {
	return VertexOutput(
		camera.viewProjectionMatrix * vec4f(in.position, 1.0),
		in.position,
		in.normal,
		in.texcoord,
		in.materialID,
	);
}

@fragment
fn fragmentMain(in : VertexOutput, @builtin(front_facing) front_facing : bool) -> @location(0) vec4f {
	let material = materials[in.materialID];

	let normal = normalize(select(-in.normal, in.normal, front_facing));

	let lightDirection = normalize(vec3f(1.0, 3.0, 2.0));

	let albedoSample = textureSampleLevel(albedoTexture, textureSampler, in.texcoord, material.textureLayers.x, 0.0);

	if (albedoSample.a < 0.5) {
		discard;
	}

	let albedo = material.baseColorFactorAndTransmission.rgb * albedoSample.rgb;

	let litColor = albedo * (0.5 + 0.5 * dot(normal, lightDirection)) + material.emissiveFactor.rgb;

	let cameraDirection = normalize(camera.position - in.worldPosition);
	let reflectedDirection = 2.0 * normal * dot(normal, cameraDirection) - cameraDirection;

	//let reflectedColor = sampleEnvMap(environmentMap, reflectedDirection) * albedo;
	let reflectedColor = vec3f(0.0);

	let metallic = material.metallicRoughnessFactorAndIor.b;

	let color = mix(litColor, reflectedColor, metallic);

	return vec4f(gammaCorrect(acesToneMap(color)), 1.0);
}

struct BackgroundVertexOutput
{
	@builtin(position) position : vec4f,
	@location(0) worldSpacePosition : vec3f,
}

@vertex
fn backgroundVertexMain(@builtin(vertex_index) index : u32) -> BackgroundVertexOutput
{
	var ndc = vec4f(-1.0, -1.0, 0.0, 1.0);
	if (index == 1u) {
		ndc = vec4f(3.0, -1.0, 0.0, 1.0);
	}
	else if (index == 2u) {
		ndc = vec4f(-1.0, 3.0, 0.0, 1.0);
	}

	return BackgroundVertexOutput(
		ndc,
		perspectiveDivide(camera.viewProjectionInverseMatrix * ndc)
	);
}

@fragment
fn backgroundFragmentMain(in : BackgroundVertexOutput) -> @location(0) vec4f
{
	let direction = in.worldSpacePosition - camera.position;
	//let color = sampleEnvMap(environmentMap, direction);
	let color = vec3f(0.0);
	return vec4f(gammaCorrect(acesToneMap(color)), 1.0);
}

Specifically the fn sampleEnvMap(environmentMap: texture_storage_2d<rgba32float, read>, direction : vec3f) function looks like it might be the problem. Commenting it out (and removing its usages) fixes the problem (though another shader compilation crashes after that, but it's a different bug).

This code works on Linux with wgpu-native v0.19.4.1. I've initially tried it on MacOS with the same wgpu-native version, and the same error emerged.

Repro steps
Unfortunately, my only reproducible example is rather large: clone this repo, build it, and run, supplying any scene from the test_scenes directory.

Expected vs observed behavior
Expected behavior: shader compiles properly
Observed behavior: shader compilation crashes

Extra materials
I've tried enabling METAL_DEVICE_WRAPPER_TYPE=1 but I didn't see any extra logs.

Platform
Device: Apple M1 Pro
OS: macOS Ventura 13.3
wgpu-native version: v0.19.4.1 and v22.1.0.5
OS window created using SDL2 and a small wrapper to create a wgpu surface.

[cwfitzgerald]

I think this may be a validator issue as I don't think a texture argument should have made it all the way to the backend, though some other maintainer should validate.

Can you try running your shader against trunk naga cli cargo run --bin naga your_shader.wgsl

[lisyarus]

@cwfitzgerald Sure, here's the result:

lisyarus@i111354133 ~/wgpu $ cargo run --bin naga ../preview.wgsl
warning: /Users/lisyarus/wgpu/deno_webgpu/Cargo.toml: unused manifest key: target.cfg(not(target_arch = "wasm32")).dependencies.wgt.package
    Finished `dev` profile [unoptimized + debuginfo] target(s) in 0.69s
     Running `target/debug/naga ../preview.wgsl`
Validation successful

[dj2]

I tried it with tint and it passed validation and generated an MSL shader once I fixed:

out/s.wgsl:27:11 error: cannot use let det as call target
  let d0 = det(rhs, matrix[1], matrix[2]);
           ^^^

out/s.wgsl:26:6 note: let det declared here
  let det = det(matrix[0], matrix[1], matrix[2]);
      ^^^

[cwfitzgerald]

Hmm, will need a maintainer who understands more about wgsl to chime in.

[lisyarus]

I tried it with tint and it passed validation and generated an MSL shader once I fixed:

out/s.wgsl:27:11 error: cannot use let det as call target
  let d0 = det(rhs, matrix[1], matrix[2]);
           ^^^

out/s.wgsl:26:6 note: let det declared here
  let det = det(matrix[0], matrix[1], matrix[2]);
      ^^^

Oh, I didn't even notice I had this name clash. I guess WGSL allows it, since everything works on Vulkan?
Anyway, fixing it doesn't remove the original problem, sadly.

[dj2]

I'm guessing that it's getting dropped because it's dead code. When I tried it in Tint, I just compiled everything so that function gets compiled. The browser will, probably, ignore it because it isn't reachable from any of the entry points.

The first call let det = det(matrix[0], matrix[1], matrix[2]); is allowed because the let comes into existance at the ;. So it understands the det() function at that point. After that, the det local is in existance so the calls to det() fail, it's no longer a method at that scope.