wgpu icon indicating copy to clipboard operation
wgpu copied to clipboard

Relax vector, matrix and array OOB validation

Open teoxoy opened this issue 3 years ago • 4 comments

Relax vector, matrix and array OOB validation by handling OOB accesses at runtime.

The following example should work; however, validation currently errors because the index is out of bounds.

@compute @workgroup_size(1)
fn main() {
    let idx = 9;
    let v = vec2<f32>();
    let m = mat2x2<f32>();
    let a = array<f32,1>();
    let vi = v[idx];
    let mi = m[idx];
    let ai = a[idx];
}

[edit] Note that WGSL says that let-bound identifiers cannot be constant expressions, so the shader translation time bounds checking rules do not apply to this example.

Related

  • gfx-rs/wgpu#4381
  • gfx-rs/wgpu#4389
  • #6396

teoxoy avatar May 14 '22 17:05 teoxoy

Note that this issue only covers the excessive validation. There may be missing bounds checks, and those should be filed as separate issues.

jimblandy avatar Dec 12 '23 15:12 jimblandy

See also: https://github.com/gfx-rs/wgpu/issues/4337

nical avatar Jan 03 '24 10:01 nical

Curious, why should this work? Do you just mean it should pass validation? And then the bounds-check policy is used to determine what should happen with vi, mi, etc.?

I think this would be a regression. If you can detect OOB accesses, as these clearly are, at shader translation time, why not report them?

bavalpey avatar Oct 03 '24 15:10 bavalpey

@bavalpey: It's for compliance with the WebGPU spec., which forces safe (but implementation-defined) behavior: https://www.w3.org/TR/WGSL/#out-of-bounds-access-sec

ErichDonGubler avatar Oct 03 '24 17:10 ErichDonGubler