[wgsl-in] Error on non-constant matrix indexing

[hanawatson]

I believe the specification doesn't prevent expressions (with the correct return types) being used to access matrices. naga seems to reject this, though not for vectors/arrays.

code example:

@compute @workgroup_size(1)
fn compute_main() {
	var var0 = mat4x2<f32>()[1 + 1];
}

naga compilation error:

error: 
  ┌─ test.wgsl:3:12
  │
3 │     var var0 = mat4x2<f32>()[1 + 1];
  │               ^^^^^^^^^^^^^^^^^^^^^ naga::Expression [5]

Entry point compute_main at Compute is invalid: 
	Expression [5] is invalid
	The expression [1] may only be indexed by a constant

[hasali19]

Are you sure this works for arrays? This program seems to be rejected:

@compute
@workgroup_size(1)
fn main() {
    let x = array<i32, 4>()[1 + 1];
}

error: 
  ┌─ test.wgsl:4:12
  │
4 │     let x = array<i32, 4>()[1 + 1];
  │            ^^^^^^^^^^^^^^^^^^^^^^^ naga::Expression [5]

Entry point main at Compute is invalid: 
        Expression [5] is invalid
        The expression [1] may only be indexed by a constant

Edit: Looks like dynamic indexing works for global arrays but not local.

[hanawatson]

Yes, I agree - thanks for catching that! I suspect I might have tested them very slightly differently, and so got a false OK for the array indexing. Having run it again now, though, for both matrices and arrays, it looks like the problem only occurs when the indexing/declaration happen on the same line.

Accepted code (working very similarly for arrays):

@compute @workgroup_size(1)
fn compute_main() {
	var var0 = mat4x2<f32>();
	var var1 = var0[1 + 1];
}

Rejected code:

@compute @workgroup_size(1)
fn compute_main() {
	var var0 = mat4x2<f32>()[1 + 1];
}

[jimblandy]

The WGSL committee went back and forth on whether arrays and matrices should be allowed to be indexed by dynamically computed values, before ultimately (or at least, most recently) deciding to support it (gpuweb/gpuweb#2427).

The underlying problem, as explained in gfx-rs/wgpu#4337, is that SPIR-V doesn't support this: such values have to be spilled to memory before they can be indexed dynamically.

Since gfx-rs/wgpu#4337 has more explanation, I'm going to re-open that issue, and close this one as a duplicate.