Reference to a field causes panic

[khyperia]

struct Struct {
    thing: i32,
};

fn doSomething(x: ptr<function, i32>) {
}

fn objFunc(obj: ptr<function, Struct>) {
    let thing = &((*obj).thing);
    doSomething(thing);
    // Workaround:
    // var tmp = (*obj).thing;
    // doSomething(&tmp);
    // (*obj).thing = tmp;
}

@compute @workgroup_size(64, 1, 1) 
fn main() {
    var obj = Struct(0);
    objFunc(&obj);
}

> cargo install naga-cli --git https://github.com/gfx-rs/naga.git
> naga test.wgsl test.spv
thread 'main' panicked at 'internal error: entered unreachable code: Expression [2] is not cached!', src\back\spv\block.rs:1931:45
note: run with `RUST_BACKTRACE=1` environment variable to display a backtrace

(tested with 1650581f, which is latest at time of writing, as well as 0.10.0)

The same panic happens when running my app (where I encountered this issue) through wgpu.

[khyperia]

oh, right, I remember my rust-gpu days now, I forgot that OpAccessChain cannot be used as an argument to OpFunctionCall, that may be why this is failing? (IIRC rust-gpu inlines any function that is called with a value resulting from a OpAccessChain)

[jimblandy]

I forgot that OpAccessChain cannot be used as an argument to OpFunctionCall

Yes, and WGSL also forbids this, for the time being:

Each argument of pointer type to a user-defined function must have the same memory view as its root identifier.

Naga's validator may well be failing to enforce this rule.

[jimblandy]

for the time being

In the long term, we have plans to support such pointers through selective specialization of functions:

https://github.com/gpuweb/gpuweb/issues/3546

[jimblandy]

Confirmed that the code in the description still causes a panic in SPIR-V output:

   0: rust_begin_unwind
             at /rustc/eb26296b556cef10fb713a38f3d16b9886080f26/library/std/src/panicking.rs:593:5
   1: core::panicking::panic_fmt
             at /rustc/eb26296b556cef10fb713a38f3d16b9886080f26/library/core/src/panicking.rs:67:14
   2: <naga::back::spv::CachedExpressions as core::ops::index::Index<naga::arena::Handle<naga::Expression>>>::index
             at /home/jimb/rust/wgpu/naga/src/back/spv/mod.rs:433:13
   3: naga::back::spv::block::<impl naga::back::spv::BlockContext>::write_block
             at /home/jimb/rust/wgpu/naga/src/back/spv/block.rs:2095:45
   4: naga::back::spv::writer::<impl naga::back::spv::Writer>::write_function
             at /home/jimb/rust/wgpu/naga/src/back/spv/writer.rs:704:9
   5: naga::back::spv::writer::<impl naga::back::spv::Writer>::write_logical_layout
             at /home/jimb/rust/wgpu/naga/src/back/spv/writer.rs:1965:22
   6: naga::back::spv::writer::<impl naga::back::spv::Writer>::write
             at /home/jimb/rust/wgpu/naga/src/back/spv/writer.rs:2036:9
   7: naga::back::spv::write_vec
             at /home/jimb/rust/wgpu/naga/src/back/spv/mod.rs:740:5
   8: naga::write_output
             at /home/jimb/rust/wgpu/naga-cli/src/bin/naga.rs:537:23
   9: naga::run
             at /home/jimb/rust/wgpu/naga-cli/src/bin/naga.rs:466:9
  10: naga::main
             at /home/jimb/rust/wgpu/naga-cli/src/bin/naga.rs:220:21

[jimblandy]

I took the time to reread the spec here, and wanted to write down my understanding:

A memory view is what pointer and reference values refer to: a specific set of bytes in some address space.

A memory view of a variable is different from a memory view of one of its components: a reference to a vector, say, represents a memory view enclosing the entire matrix, whereas a reference to one of its components represents a memory view enclosing only that component.

So the restriction that the argument be the same memory view as its originating variable means that, without the unrestricted_pointer_parameters extension, you may only pass pointers to entire variables.

This is the restriction that Naga is failing to enforce.

[kvark]

Oh that's such a non-obvious limitation. Good to know it's by design, I can write my code without assuming this will come.