1915297 - WebGPU: `Remove-bg` renders blank images

Status: NEW

(Core :: Graphics: WebGPU, defect, P3)

Description

[Erich Gubler [:ErichDonGubler]]

I noticed this page used WebGPU while browsing HN, and that it didn't work yet due to shader compilation failures: https://news.ycombinator.com/item?id=41358490

Comment 1

[Erich Gubler [:ErichDonGubler]]

There seem to be two shader compilation issues in my testing on my M1 Macbook Pro:

• A shader named Concat is blocked on dynamic indexing into constant array values (bug 1913424); source:

  
      struct Uniforms { outputSize:u32, sizeInConcatAxis0:u32, sizeInConcatAxis1:u32, input0_shape:vec4<u32>, input0_strides:vec4<u32>, input1_shape:vec4<u32>, input1_strides:vec4<u32>, output_shape:vec4<u32>, output_strides:vec4<u32> };
      @group(0) @binding(3) var<uniform> uniforms: Uniforms;
  fn i2o_input0(indices: vec4<u32>) -> u32 {
    return uniforms.input0_strides[3] * (indices[3])+uniforms.input0_strides[2] * (indices[2])+uniforms.input0_strides[1] * (indices[1])+uniforms.input0_strides[0] * (indices[0]);
  }
  
  fn get_input0ByIndices(indices: vec4<u32>) -> f32 {
    return input0[i2o_input0(indices)];
  }
  
  fn i2o_input1(indices: vec4<u32>) -> u32 {
    return uniforms.input1_strides[3] * (indices[3])+uniforms.input1_strides[2] * (indices[2])+uniforms.input1_strides[1] * (indices[1])+uniforms.input1_strides[0] * (indices[0]);
  }
  
  fn get_input1ByIndices(indices: vec4<u32>) -> f32 {
    return input1[i2o_input1(indices)];
  }
  
  fn o2i_output(offset: u32) -> vec4<u32> {
    var indices: vec4<u32>;
    var current = offset;
  
    let dim0 = current / uniforms.output_strides[0];
    let rest0 = current % uniforms.output_strides[0];
    indices[0] = dim0;
    current = rest0;
  
    let dim1 = current / uniforms.output_strides[1];
    let rest1 = current % uniforms.output_strides[1];
    indices[1] = dim1;
    current = rest1;
  
    let dim2 = current / uniforms.output_strides[2];
    let rest2 = current % uniforms.output_strides[2];
    indices[2] = dim2;
    current = rest2;
    indices[3] = current;
    return indices;
  }
  
  
  @group(0) @binding(0) var<storage, read> input0: array<f32>;
  @group(0) @binding(1) var<storage, read> input1: array<f32>;
  @group(0) @binding(2) var<storage, read_write> output: array<f32>;
  
  
  fn calculateInputIndex(index: u32) -> u32 {
    let sizeInConcatAxis = array<u32, 2u>(uniforms.sizeInConcatAxis0,uniforms.sizeInConcatAxis1);
    for (var i: u32 = 0u; i < 2; i += 1u ) {
      if (index < sizeInConcatAxis[i]) {
        return i;
      }
    }
    return 2u;
  }
  
  @compute @workgroup_size(64, 1, 1)
  fn main(@builtin(global_invocation_id) global_id : vec3<u32>,
    @builtin(workgroup_id) workgroup_id : vec3<u32>,
    @builtin(local_invocation_id) local_id : vec3<u32>) {
    let global_idx = global_id.x; let local_idx = local_id.x;
  
    if (global_idx >= uniforms.outputSize) { return; }
  
    var indices = o2i_output(global_idx);
  
    let inputIndex = calculateInputIndex(indices[3]);
    if (inputIndex != 0u) {
      let sizeInConcatAxis = array<u32, 2u>(uniforms.sizeInConcatAxis0,uniforms.sizeInConcatAxis1);
      indices[3] -= sizeInConcatAxis[inputIndex - 1u];
    }
  
    if (inputIndex == 0u) { output[global_idx]=get_input0ByIndices(indices); }
  else { output[global_idx]=get_input1ByIndices(indices); }
  }
  

  ...where the error returned by Naga states:

  error:
     ┌─ Concat.wgsl:49:3
     │
  49 │ ╭   fn calculateInputIndex(index: u32) -> u32 {
  50 │ │     let sizeInConcatAxis = array<u32, 2u>(uniforms.sizeInConcatAxis0,uniforms.sizeInConcatAxis1);
  51 │ │     for (var i: u32 = 0u; i < 2; i += 1u ) {
  52 │ │       if (index < sizeInConcatAxis[i]) {
     │ │                   ^^^^^^^^^^^^^^^^^^^ naga::Expression [14]
     · │
  55 │ │     }
  56 │ │     return 2u;
     │ ╰──────────────^ naga::Function [5]
  
  Function [5] 'calculateInputIndex' is invalid:
          Expression [14] is invalid
          The expression [7] may only be indexed by a constant
  

• Another shader titled Resize has this source:

  
  
      struct Uniforms { output_size:u32, scales:vec4<f32>, roi:array<vec4<f32>, 2>, input_shape:vec4<u32>, input_strides:vec4<u32>, output_shape:vec4<u32>, output_strides:vec4<u32> };
      @group(0) @binding(2) var<uniform> uniforms: Uniforms;
  fn i2o_input(indices: vec4<u32>) -> u32 {
    return uniforms.input_strides[3] * (indices[3])+uniforms.input_strides[2] * (indices[2])+uniforms.input_strides[1] * (indices[1])+uniforms.input_strides[0] * (indices[0]);
  }
  
  fn get_inputByIndices(indices: vec4<u32>) -> f32 {
    return input[i2o_input(indices)];
  }
  
  fn o2i_output(offset: u32) -> vec4<u32> {
    var indices: vec4<u32>;
    var current = offset;
  
    let dim0 = current / uniforms.output_strides[0];
    let rest0 = current % uniforms.output_strides[0];
    indices[0] = dim0;
    current = rest0;
  
    let dim1 = current / uniforms.output_strides[1];
    let rest1 = current % uniforms.output_strides[1];
    indices[1] = dim1;
    current = rest1;
  
    let dim2 = current / uniforms.output_strides[2];
    let rest2 = current % uniforms.output_strides[2];
    indices[2] = dim2;
    current = rest2;
    indices[3] = current;
    return indices;
  }
  
  
      fn getOriginalCoordinateFromResizedCoordinate(xResized: u32, xScale: f32, lengthResized: u32,
     lengthOriginal: u32, roiStart: f32, roiEnd: f32) -> f32 { return ((f32(xResized) + 0.5) / f32(xScale)) - 0.5;};
  
  
    fn calculateOriginalIndicesFromOutputIndices(output_indices: vec4<u32>) -> array<f32, 4> {
      var original_indices: array<f32, 4>;
      for (var i:u32 = 0; i < 4; i++) {
        var output_index = output_indices[i];
        var scale = uniforms.scales[i];
        var roi_low = uniforms.roi[(i) / 4][(i) % 4];
        var roi_hi = uniforms.roi[(i + 4) / 4][(i + 4) % 4];
        if (scale == 1.0) {
          original_indices[i] = f32(output_index);
        } else {
          var input_shape_i = uniforms.input_shape[i];
          var output_shape_i = uniforms.output_shape[i];
          original_indices[i] = getOriginalCoordinateFromResizedCoordinate(output_index, scale, output_shape_i,
                                                                           input_shape_i, roi_low, roi_hi);
        }
      }
      return original_indices;
    };
  
    fn getInputValue(batch: u32, channel: u32, row: u32, col: u32) -> f32 {
      var input_indices: vec4<u32>;
      input_indices[2]=max(0, min(row, 16 - 1));;
      input_indices[3]=max(0, min(col, 16 - 1));;
  
    input_indices[1]=channel;;
    input_indices[0]=batch;;
  
      return get_inputByIndices(input_indices);
    }
  
    fn bilinearInterpolation(output_indices: vec4<u32>) -> f32 {
      var originalIndices = calculateOriginalIndicesFromOutputIndices(output_indices);
      var row:f32 = originalIndices[2];
      var col:f32 = originalIndices[3];
      ;
      row = max(0, min(row, 16 - 1));
      col = max(0, min(col, 16 - 1));
      var row1: u32 = u32(row);
      var col1: u32 = u32(col);
      var row2: u32 = u32(row + 1);
      var col2: u32 = u32(col + 1);
      var channel: u32 = u32(originalIndices[1]);
      var batch: u32 =  u32(originalIndices[0]);
      var x11: f32 = getInputValue(batch, channel, row1, col1);
      var x12: f32 = getInputValue(batch, channel, row1, col2);
      var x21: f32 = getInputValue(batch, channel, row2, col1);
      var x22: f32 = getInputValue(batch, channel, row2, col2);
      var dx1: f32 = abs(row - f32(row1));
      var dx2: f32 = abs(f32(row2) - row);
      var dy1: f32 = abs(col - f32(col1));
      var dy2: f32 = abs(f32(col2) - col);
      if (row1 == row2) {
        dx1 = 0.5;
        dx2 = 0.5;
      }
      if (col1 == col2) {
        dy1 = 0.5;
        dy2 = 0.5;
      }
      return (x11 * dx2 * dy2 + x12 * dx2 * dy1 + x21 * dx1 * dy2 + x22 * dx1 * dy1);
    };
            ;
  
      @group(0) @binding(0) var<storage, read> input: array<f32>;
  @group(0) @binding(1) var<storage, read_write> output: array<f32>;
      @compute @workgroup_size(64, 1, 1)
  fn main(@builtin(global_invocation_id) global_id : vec3<u32>,
    @builtin(workgroup_id) workgroup_id : vec3<u32>,
    @builtin(local_invocation_id) local_id : vec3<u32>) {
    let global_idx = global_id.x; let local_idx = local_id.x;
  
        if (global_idx >= uniforms.output_size) { return; }
  
        let output_indices = o2i_output(global_idx);
        var input_indices: vec4<u32>;
        output[global_idx] = bilinearInterpolation(output_indices);;
  
      }
  

  ...with this error in Naga (almost certainly bug 1829879):

  error:
     ┌─ Resize.wgsl:59:5
     │
  59 │ ╭     fn getInputValue(batch: u32, channel: u32, row: u32, col: u32) -> f32 {
  60 │ │       var input_indices: vec4<u32>;
  61 │ │       input_indices[2]=max(0, min(row, 16 - 1));;
     │ │                               ^^^ naga::Expression [8]
  62 │ │       input_indices[3]=max(0, min(col, 16 - 1));;
     · │
  66 │ │
  67 │ │       return get_inputByIndices(input_indices);
     │ ╰───────────────────────────────────────────────^ naga::Function [5]
  
  Function [5] 'getInputValue' is invalid:
          Expression [8] is invalid
          Argument [1] to Min as expression [7] has an invalid type.