[JS/Web] Added Uniforms support to binary ops.

[satyajandhyala]

Description

Added Uniform support to binary ops

Motivation and Context

To improve performance

[satyajandhyala]

For the test_add with both inputs same shape I got the following shader

      struct Uniforms { vec_size:u32 };
      @group(0) @binding(3) var<uniform> uniforms: Uniforms;
const aData_shape = vec3<u32>(3,4,5);
const aData_strides = vec3<u32>(20,5,1);
const bData_shape = vec3<u32>(3,4,5);
const bData_strides = vec3<u32>(20,5,1);
const outputData_shape = vec3<u32>(3,4,5);
const outputData_strides = vec3<u32>(20,5,1);

        @group(0) @binding(0) var<storage, read> aData: array<vec4<f32>>;
@group(0) @binding(1) var<storage, read> bData: array<vec4<f32>>;
@group(0) @binding(2) var<storage, read_write> outputData: array<vec4<f32>>;




        @compute @workgroup_size(64, 1, 1)
  fn main(@builtin(global_invocation_id) global_id : vec3<u32>,
    @builtin(local_invocation_id) local_id : vec3<u32>) {
    let global_idx = global_id.x;

        if (global_idx >= uniforms.vec_size) { return; }
        outputData[global_idx]=aData[global_idx]+bData[global_idx];
      }
      

[satyajandhyala]

For the test add with different input shapes I got the following shader

      struct Uniforms { vec_size:u32, aData_shape:vec3<u32>, aData_strides:vec3<u32>, bData_shape:u32, bData_strides:u32, outputData_shape:vec3<u32>, outputData_strides:vec3<u32> };
      @group(0) @binding(3) var<uniform> uniforms: Uniforms;


  fn o2i_outputData(offset: u32) -> vec3<u32> {
    var indices: vec3<u32>;
    var current = offset;

    let dim0 = current / uniforms.outputData_strides[0];
    let rest0 = current % uniforms.outputData_strides[0];
    indices[0] = dim0;
    current = rest0;

    let dim1 = current / uniforms.outputData_strides[1];
    let rest1 = current % uniforms.outputData_strides[1];
    indices[1] = dim1;
    current = rest1;
    indices[2] = current;
    return indices;
  }

        @group(0) @binding(0) var<storage, read> aData: array<vec4<f32>>;
@group(0) @binding(1) var<storage, read> bData: array<vec4<f32>>;
@group(0) @binding(2) var<storage, read_write> outputData: array<vec4<f32>>;




          fn calcOffsetaData(outputIndices: vec3<u32>) -> u32{
            var retval : u32 = 0;
            for (var i = 3 - 1; i >= 0; i = i - 1) {
              retval = retval + uniforms.aData_strides[i] * (outputIndices[i] % uniforms.aData_shape[i]);
            }
            return retval;
          }

          fn calcOffsetbData(outputIndices: vec3<u32>) -> u32{
            var retval : u32 = 0;
            retval = uniforms.bData_strides * (outputIndices[3 - 1] % uniforms.bData_shape);
            return retval;
          }


        @compute @workgroup_size(64, 1, 1)
  fn main(@builtin(global_invocation_id) global_id : vec3<u32>,
    @builtin(local_invocation_id) local_id : vec3<u32>) {
    let global_idx = global_id.x;

        if (global_idx >= uniforms.vec_size) { return; }


            let outputIndices0 = o2i_outputData(global_idx * 4u + 0u);
            let offsetA0 = calcOffsetaData(outputIndices0);
            let offsetB0 = calcOffsetbData(outputIndices0);
            let indexA0 = offsetA0 / 4u;
            let indexB0 = offsetB0 / 4u;
            let componentA0 = offsetA0 % 4u;
            let componentB0 = offsetB0 % 4u;
            outputData[global_idx][0] = (aData[indexA0][componentA0]+bData[indexB0][componentB0]);


            let outputIndices1 = o2i_outputData(global_idx * 4u + 1u);
            let offsetA1 = calcOffsetaData(outputIndices1);
            let offsetB1 = calcOffsetbData(outputIndices1);
            let indexA1 = offsetA1 / 4u;
            let indexB1 = offsetB1 / 4u;
            let componentA1 = offsetA1 % 4u;
            let componentB1 = offsetB1 % 4u;
            outputData[global_idx][1] = (aData[indexA1][componentA1]+bData[indexB1][componentB1]);


            let outputIndices2 = o2i_outputData(global_idx * 4u + 2u);
            let offsetA2 = calcOffsetaData(outputIndices2);
            let offsetB2 = calcOffsetbData(outputIndices2);
            let indexA2 = offsetA2 / 4u;
            let indexB2 = offsetB2 / 4u;
            let componentA2 = offsetA2 % 4u;
            let componentB2 = offsetB2 % 4u;
            outputData[global_idx][2] = (aData[indexA2][componentA2]+bData[indexB2][componentB2]);


            let outputIndices3 = o2i_outputData(global_idx * 4u + 3u);
            let offsetA3 = calcOffsetaData(outputIndices3);
            let offsetB3 = calcOffsetbData(outputIndices3);
            let indexA3 = offsetA3 / 4u;
            let indexB3 = offsetB3 / 4u;
            let componentA3 = offsetA3 % 4u;
            let componentB3 = offsetB3 % 4u;
            outputData[global_idx][3] = (aData[indexA3][componentA3]+bData[indexB3][componentB3]);


      }

[qjia7]

LGTM, thanks!