Bugzilla

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:

    ```wgsl

        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:

    ```
    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.
    ```

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:

    ```wgsl

        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.
    ```