Description

naga's MSL backend emits an assignment sq_1 = intersection_query<...>{}; when a ray_query variable is declared inside a nested scope (e.g. a loop body) and another ray_query with the same name exists in an enclosing scope. Since Apple's Metal SDK (32023+) deleted operator= on metal::raytracing::intersection_query, this fails to compile.

Generated Metal (from a shader with two ray_query declarations sharing a function scope):

sq_1 = metal::raytracing::intersection_query<metal::raytracing::instancing, metal::raytracing::triangle_data> {};

Error from Metal compiler:

error: overload resolution selected deleted operator '='
note: candidate function has been explicitly deleted
  METAL_FUNC thread intersection_query &operator=(const thread intersection_query &) thread = delete;

Repro steps

Create a WGSL compute shader with two ray_query variables in one function scope:

enable wgpu_ray_query;

@group(0) @binding(0) var acc: acceleration_structure;

@compute @workgroup_size(1)
fn main() {
    // First ray_query declaration
    var sq: ray_query;
    rayQueryInitialize(&sq, acc, RayDesc(0x01u, 0xFFu, 0.001, 1000.0, vec3(0.0), vec3(0.0, -1.0, 0.0)));
    rayQueryProceed(&sq);

    // Second ray_query in a loop — naga renames to sq_1 and emits assignment
    for (var i: u32 = 0u; i < 4u; i++) {
        var sq: ray_query;
        rayQueryInitialize(&sq, acc, RayDesc(0x01u, 0xFFu, 0.001, 1000.0, vec3(f32(i)), vec3(0.0, -1.0, 0.0)));
        rayQueryProceed(&sq);
    }
}

A standalone reproduction project is at https://github.com/tristanpoland/naga-ray-query-repro.

Compile with wgpu 30.0.0 on macOS Sequoia (Metal SDK 32023+). Device::create_compute_pipeline will fail with the error above.

Note: a single var rq: ray_query; (no duplicate name in scope) compiles fine — the = {} copy-list-initialization for the declaration alone is accepted. The bug specifically requires two ray_query variables sharing a function scope.

Expected vs observed behavior

Expected: Pipeline creation succeeds. The second ray_query variable inside the loop body should either be declared fresh (using Metal block scoping) or left uninitialized until rayQueryInitialize (which calls .reset()).

Observed: naga hoists all local variables to function scope with unique names. For the inner-scope ray_query, it emits a declaration at the top:

intersection_query<...> sq_1;

Then at the original declaration point (inside the loop body), it emits a reassignment:

sq_1 = intersection_query<...> {};

This triggers the deleted operator= on newer Metal SDKs.

Root cause

naga/src/back/msl/writer.rs — when the MSL backend writes a local variable block, it hoists all declarations to function scope. For variables originally declared in nested scopes, the declaration is emitted at the top, and a plain assignment name = Type{}; is emitted at the original declaration point. For ray_query this assignment is both unnecessary (.reset() is called by the subsequent rayQueryInitialize) and harmful (the Metal SDK deleted operator=).

There are two locations in naga/src/back/msl/writer.rs that need fixing:

1. Variable hoisting (around line 1170): When a ray_query variable has no initializer and is hoisted to function scope, the = {} should be omitted (emit Type name; instead of Type name = {};). This part is benign — it compiles — but is still worth fixing for correctness.

2. Reassignment at nested scope (the real bug): When a variable originally declared in an inner scope is hoisted, naga emits a reassignment name = Type{}; at the original declaration point. For ray_query types, this assignment should be elided entirely, since the subsequent rayQueryInitialize(&name, ...) call already uses .reset().

The fix: add a check for RayQuery types before emitting the reassignment, and skip it.

Platform

• wgpu 30.0.0 (naga 30.0.0) — also confirmed present on trunk in the wgpu monorepo
• macOS 15.x (Sequoia) — Metal SDK version 32023
• /System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/32023/Libraries/lib/clang/32023.622/include/metal/metal_raytracing:3134
• Apple Silicon (M4)