sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +66 -0
build/webgpu/bench.json +79 -0
build/webgpu/manifest.json +225 -0
build/webgpu/metadata.json +19 -0
build/webgpu/mod-vec4.wgsl.jinja +110 -0
build/webgpu/mod.wgsl.jinja +161 -0
build/webgpu/test.json +773 -0

README.md CHANGED Viewed

@@ -1,3 +1,69 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Mod
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Performs elementwise binary modulo on tensors `A` and `B` with multidirectional broadcasting. When `fmod` is `0` (default), applies Python-style `%` with the sign of the divisor; when `fmod` is `1`, applies C-style `fmod` with the sign of the dividend.
+See the [ONNX `Mod` spec](https://onnx.ai/onnx/operators/onnx__Mod.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `A` | `a` | `T` | — | — | Dividend tensor. | required |
+| `B` | `b` | `T` | — | — | Divisor tensor. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `C` | `c` | `T` | derived | broadcast result of `A` and `B` | Remainder tensor; same shape as the broadcast result of A and B. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `fmod` | `0` | Controls the modulo mode: `0` (default) uses Python-style integer mod (sign of divisor); `1` uses C-style `fmod` (sign of dividend, floating-point types only). |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `uint32`, `int16`, `int8`, `uint8` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`mod-vec4.wgsl.jinja`](build/webgpu/mod-vec4.wgsl.jinja)
+- [`mod.wgsl.jinja`](build/webgpu/mod.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Mod", { version: 1 });
+const { c } = await kernel({ a: { data: aData, shape: [3] }, b: { data: bData, shape: [3] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,79 @@

+{
+  "op": "ai.onnx.Mod",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "8m_u32_scalar_div",
+      "preset": "smoke",
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [8388608] },
+        "b": { "dtype": "uint32", "shape": [1], "data": { "kind": "constant", "value": 17 } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [8388608] } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "8388608 * 4 * 2" }] }
+    },
+    {
+      "name": "1m_u32",
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [1048576] },
+        "b": { "dtype": "uint32", "shape": [1], "data": { "kind": "constant", "value": 17 } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [1048576] } }
+    },
+    {
+      "name": "1m_f32_fmod_same_shape",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [1048576] },
+        "b": { "dtype": "float32", "shape": [1048576], "data": { "kind": "constant", "value": 3.7 } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [1048576] } }
+    },
+    {
+      "name": "f32_broadcast_per_channel_fallback",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [2048, 2049], "dist": "normal", "seed": 711, "scale": 2 },
+        "b": { "dtype": "float32", "shape": [2049], "data": { "kind": "constant", "value": 3.0 } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [2048, 2049] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2048 * 2049 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "1m_f16_fmod_same_shape",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 1048576 },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float16", "shape": [1048576], "dist": "normal", "seed": 712, "scale": 8 },
+        "b": { "dtype": "float16", "shape": [1048576], "data": { "kind": "constant", "value": 3.5 } }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [1048576] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }] }
+    },
+    {
+      "name": "1m_i32_floor_mod_same_shape",
+      "preset": "smoke",
+      "vars": { "dtype": "int32", "count": 1048576 },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [1048576], "dist": "normal", "seed": 713, "scale": 1000 },
+        "b": { "dtype": "int32", "shape": [1048576], "data": { "kind": "constant", "value": 7 } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [1048576] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }] }
+    },
+    {
+      "name": "i32_floor_mod_broadcast_non_vec4_shape",
+      "preset": "stress",
+      "attrs": { "fmod": 0 },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [2049, 2051], "dist": "normal", "seed": 801, "scale": 1000 },
+        "b": { "dtype": "int32", "shape": [2051], "data": { "kind": "constant", "value": 7 } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [2049, 2051] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2049 * 2051 * 4 * 2" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,225 @@

+{
+  "domain": "ai.onnx",
+  "name": "Mod",
+  "sinceVersion": 13,
+  "description": "Performs elementwise binary modulo on tensors `A` and `B` with multidirectional broadcasting. When `fmod` is `0` (default), applies Python-style `%` with the sign of the divisor; when `fmod` is `1`, applies C-style `fmod` with the sign of the dividend.",
+  "inputs": [
+    { "role": "A", "dtype": "T", "description": "Dividend tensor." },
+    { "role": "B", "dtype": "T", "description": "Divisor tensor." }
+  ],
+  "outputs": [
+    {
+      "role": "C",
+      "dtype": "T",
+      "rank": "max(ranks.A, ranks.B)",
+      "description": "Remainder tensor; same shape as the broadcast result of A and B.",
+      "shape": "broadcastShape(shapes.A, shapes.B)"
+    }
+  ],
+  "attributes": { "fmod": 0 },
+  "attributeDescriptions": {
+    "fmod": "Controls the modulo mode: `0` (default) uses Python-style integer mod (sign of divisor); `1` uses C-style `fmod` (sign of dividend, floating-point types only)."
+  },
+  "attributeConstraints": { "fmod": { "values": [0, 1] } },
+  "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int16", "int8", "uint8"] },
+  "args": {
+    "a": { "kind": "tensor", "semantic": "A", "role": "input" },
+    "b": { "kind": "tensor", "semantic": "B", "role": "input" },
+    "c": { "kind": "tensor", "semantic": "C", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": { "floatFmodContract": "attrs.fmod != 0 or not (dtypes.T == \"f32\" or dtypes.T == \"f16\")" },
+  "constants": {
+    "scalar": "dtypes.T",
+    "usesF16": "dtypes.T == \"f16\"",
+    "isU32": "dtypes.T == \"u32\"",
+    "isI32": "dtypes.T == \"i32\"",
+    "fmod": "attrs.fmod != 0"
+  },
+  "variants": [
+    {
+      "id": "last_axis_broadcast_vec4",
+      "priority": 25,
+      "when": ["ranks.C >= 1", "sameShape(shapes.A, shapes.C)", "ranks.B == 1", "dim(shapes.B, 0) == dim(shapes.C, ranks.C - 1)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0", "f16Ok(dtypes.T)", "floatFmodContract"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Mod.lastAxisBroadcastVec4",
+          "source": { "shader": "mod-vec4.wgsl.jinja", "inputs": { "perAxisB": true, "axisDim": "dim(shapes.B, 0)" } },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "c",
+              "arg": "c",
+              "semantic": "C",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "scalar_b_vec4",
+      "priority": 30,
+      "when": ["sameShape(shapes.A, shapes.C)", "numel(shapes.B) == 1", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0", "f16Ok(dtypes.T)", "floatFmodContract"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Mod.scalarBVec4",
+          "source": { "shader": "mod-vec4.wgsl.jinja", "inputs": { "scalarB": true } },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 1
+            },
+            {
+              "name": "c",
+              "arg": "c",
+              "semantic": "C",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "same_shape_vec4",
+      "priority": 26,
+      "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0", "f16Ok(dtypes.T)", "floatFmodContract"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Mod.vec4",
+          "source": { "shader": "mod-vec4.wgsl.jinja", "inputs": {} },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "c",
+              "arg": "c",
+              "semantic": "C",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast",
+      "when": ["ranks.A >= 0", "ranks.B >= 0", "ranks.C >= 0", "ranks.A <= ranks.C", "ranks.B <= ranks.C", "f16Ok(dtypes.T)", "floatFmodContract"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Mod",
+          "source": {
+            "shader": "mod.wgsl.jinja",
+            "inputs": {
+              "aShape": "shapes.A",
+              "bShape": "shapes.B",
+              "cShape": "shapes.C",
+              "aRank": "ranks.A",
+              "bRank": "ranks.B",
+              "cRank": "ranks.C"
+            }
+          },
+          "bindings": [
+            {
+              "name": "a",
+              "arg": "a",
+              "semantic": "A",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "b",
+              "arg": "b",
+              "semantic": "B",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.C)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "name": "ai.onnx.Mod",
+  "id": "_ai_onnx_mod_webgpu_4a7f8df",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "QdfVC09wyipoUAu6hPHgQsOE7mtlr2T9OZJffvCh8YQ=",
+      "manifest.json": "EWrTl0E4Q3pRgojKEehDCETQSaGUYd6LvDy+6DtWFCI=",
+      "mod-vec4.wgsl.jinja": "OoOWti3TFUtnneaDtonReshC71gzRLiUQOOK4IZ9P7w=",
+      "mod.wgsl.jinja": "R3tnz/Y9NdD2Ut/jujKAUl2H2Y2CYSu9LpKXdUYJsz4=",
+      "test.json": "3XigfrMBzVhWpcZHr5WWptcT2N637Wb7iEjzEValb4A="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Mod" }
+}

build/webgpu/mod-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,110 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if not isU32 and not isI32 %}
+// Truncated floating remainder (C fmod / sign of dividend). The naive
+// `a - b*trunc(a/b)` loses all precision once |a| is large enough that the f32
+// product `b*trunc(a/b)` rounds back to `a` (e.g. fmod(1e20, 3) cancels to 0).
+// This uses an exact Sterbenz scaled-subtraction reduction: each
+// `x - scaled` is exact because `scaled <= x < 2*scaled`, and scaling by powers
+// of two is exact, so the final remainder carries no rounding error.
+fn fmod_trunc(a: f32, b: f32) -> f32 {
+  let aa = abs(a);
+  let ab = abs(b);
+  // b == 0, or a non-finite: defer to the hardware formula for IEEE NaN/Inf.
+  if (!(ab > 0.0) || bitcast<u32>(aa) >= 0x7f800000u) {
+    return a - b * trunc(a / b);
+  }
+  if (aa < ab) {
+    return a;
+  }
+  var scaled = ab;
+  loop {
+    let nxt = scaled + scaled;
+    if (nxt > aa) { break; }
+    scaled = nxt;
+  }
+  var x = aa;
+  loop {
+    if (x >= scaled) { x = x - scaled; }
+    if (scaled == ab) { break; }
+    scaled = scaled * 0.5;
+  }
+  return select(x, -x, a < 0.0);
+}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d(note="vec4-limit") }}
+{% if source.perAxisB is defined and source.perAxisB %}
+  // B broadcasts over the innermost axis. A/C stay 128-bit packed even when
+  // an odd row width makes vec4 groups cross a row boundary; derive each
+  // lane's wrapped B coordinate from the flat scalar position.
+  let b0 = (i * 4u) % {{ source.axisDim }}u;
+  let b1 = (b0 + 1u) % {{ source.axisDim }}u;
+  let b2 = (b0 + 2u) % {{ source.axisDim }}u;
+  let b3 = (b0 + 3u) % {{ source.axisDim }}u;
+{% endif %}
+{% if isU32 %}
+  c[i] = a[i] % {% if source.scalarB is defined and source.scalarB %}vec4<u32>(b[0]){% elif source.perAxisB is defined and source.perAxisB %}vec4<u32>(b[b0], b[b1], b[b2], b[b3]){% else %}b[i]{% endif %};
+{% elif isI32 %}
+  let av = a[i];
+  let bv = {% if source.scalarB is defined and source.scalarB %}vec4<i32>(b[0]){% elif source.perAxisB is defined and source.perAxisB %}vec4<i32>(b[b0], b[b1], b[b2], b[b3]){% else %}b[i]{% endif %};
+  let rt = av % bv;
+{% if fmod %}
+  c[i] = rt;
+{% else %}
+  c[i] = select(rt, rt + bv, (rt != vec4<i32>(0)) & ((rt < vec4<i32>(0)) != (bv < vec4<i32>(0))));
+{% endif %}
+{% else %}
+  let av = vec4<f32>(a[i]);
+{% if source.perAxisB is defined and source.perAxisB %}
+  let bv = vec4<f32>(f32(b[b0]), f32(b[b1]), f32(b[b2]), f32(b[b3]));
+{% else %}
+  let bv = {% if source.scalarB is defined and source.scalarB %}vec4<f32>(f32(b[0])){% else %}vec4<f32>(b[i]){% endif %};
+{% endif %}
+  let rt = vec4<f32>(
+    fmod_trunc(av.x, bv.x),
+    fmod_trunc(av.y, bv.y),
+    fmod_trunc(av.z, bv.z),
+    fmod_trunc(av.w, bv.w)
+  );
+  c[i] = vec4<{{ scalar }}>(rt);
+{% endif %}
+}

build/webgpu/mod.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,161 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
+fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
+{% if out_numel == 0 %}
+  return 0u;
+{% elif op_numel == 1 %}
+  return 0u;
+{% elif op_same %}
+  return out_index;
+{% else %}
+  var offset = 0u;
+{% for axis in range(outRank) %}
+{% set op_axis = axis - (outRank - opRank) %}
+{% if op_axis >= 0 and opShape[op_axis] != 1 %}
+{% set c_stride = namespace(value=1) %}
+{% for j in range(axis + 1, outRank) %}
+{% set c_stride.value = c_stride.value * outShape[j] %}
+{% endfor %}
+{% set op_stride = namespace(value=1) %}
+{% for j in range(op_axis + 1, opRank) %}
+{% set op_stride.value = op_stride.value * opShape[j] %}
+{% endfor %}
+{% if c_stride.value == 1 %}
+  let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
+{% else %}
+  let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
+{% endif %}
+{% if op_stride.value == 1 %}
+  offset = offset + coord{{ axis }};
+{% else %}
+  offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
+{% endif %}
+{% endif %}
+{% endfor %}
+  return offset;
+{% endif %}
+}
+{%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
+{% set op_numel = namespace(value=1) %}
+{% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
+{% set out_numel = namespace(value=1) %}
+{% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
+{{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
+{%- endmacro %}{% macro broadcast_offset_fn(fn_name, opShape, opRank, outShape, outRank) %}
+{% set op_numel = namespace(value=1) %}
+{% for d in opShape %}
+{% set op_numel.value = op_numel.value * d %}
+{% endfor %}
+{% set out_numel = namespace(value=1) %}
+{% for d in outShape %}
+{% set out_numel.value = out_numel.value * d %}
+{% endfor %}
+{% set op_same = namespace(value=(opRank == outRank)) %}
+{% if op_same.value %}
+{% for axis in range(outRank) %}
+{% if opShape[axis] != outShape[axis] %}
+{% set op_same.value = false %}
+{% endif %}
+{% endfor %}
+{% endif %}
+{{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
+{%- endmacro %}{% macro binary_broadcast_offsets() %}
+{{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.cShape, source.cRank) }}
+{{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.cShape, source.cRank) }}
+{%- endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{{ binary_broadcast_offsets() }}
+{% if not isU32 and not isI32 %}
+// Truncated floating remainder (C fmod / sign of dividend). The naive
+// `a - b*trunc(a/b)` loses all precision once |a| is large enough that the f32
+// product `b*trunc(a/b)` rounds back to `a` (e.g. fmod(1e20, 3) cancels to 0).
+// This uses an exact Sterbenz scaled-subtraction reduction: each
+// `x - scaled` is exact because `scaled <= x < 2*scaled`, and scaling by powers
+// of two is exact, so the final remainder carries no rounding error.
+fn fmod_trunc(a: f32, b: f32) -> f32 {
+  let aa = abs(a);
+  let ab = abs(b);
+  // b == 0, or a non-finite: defer to the hardware formula for IEEE NaN/Inf.
+  if (!(ab > 0.0) || bitcast<u32>(aa) >= 0x7f800000u) {
+    return a - b * trunc(a / b);
+  }
+  if (aa < ab) {
+    return a;
+  }
+  var scaled = ab;
+  loop {
+    let nxt = scaled + scaled;
+    if (nxt > aa) { break; }
+    scaled = nxt;
+  }
+  var x = aa;
+  loop {
+    if (x >= scaled) { x = x - scaled; }
+    if (scaled == ab) { break; }
+    scaled = scaled * 0.5;
+  }
+  return select(x, -x, a < 0.0);
+}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d() }}
+{% if isU32 %}
+  c[i] = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}] % b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
+{% elif isI32 %}
+  let av = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}];
+  let bv = b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
+  let rt = av % bv;
+{% if fmod %}
+  c[i] = rt;
+{% else %}
+  c[i] = select(rt, rt + bv, rt != 0 && (rt < 0) != (bv < 0));
+{% endif %}
+{% else %}
+  let av = f32(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}]);
+  let bv = f32(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}]);
+  let rt = fmod_trunc(av, bv);
+  c[i] = {{ scalar }}(rt);
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,773 @@

+{
+  "op": "ai.onnx.Mod",
+  "fixtureArrays": {
+    "ort_int32_rank3_scalar_divisor_broadcast_input_a": [0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29]
+  },
+  "cases": [
+    {
+      "name": "int16_floor_mod_boundaries",
+      "inputs": {
+        "a": { "dtype": "int16", "shape": [4], "data": { "kind": "values", "values": [-32768, -32768, 32767, -32767] } },
+        "b": { "dtype": "int16", "shape": [4], "data": { "kind": "values", "values": [-1, 32767, -32768, -32768] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int16",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 32766, -1, -32767] }
+        }
+      }
+    },
+    {
+      "name": "dispatch_cliff_broadcast",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [16776961],
+          "data": { "kind": "cycle", "values": [5.5, 6.5, 7.5, 8.5, 9.5] }
+        },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [3.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [16776961], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "u32_broadcast",
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [5, 6, 7, 8, 9, 10] } },
+        "b": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 4, 5] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [2, 3] } }
+    },
+    {
+      "name": "f32_broadcast",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [5.5, 6.5, 7.5, 8.5] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [4] } }
+    },
+    {
+      "name": "f32_fmod_subnormal_dividend_scalar_b_vec4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Fmod_float_mixed_sign",
+        "notes": "For |x| < |y|, fmod(x, y) returns x; finite subnormal dividends must not be flushed to zero."
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 4e-39, -4e-39] }
+        },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 4e-39, -4e-39] }
+        }
+      }
+    },
+    {
+      "name": "f32_fmod_subnormal_dividend_scalar_path",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Fmod_float_mixed_sign",
+        "notes": "Scalar broadcast companion: subnormal dividends smaller than the divisor are valid nonzero remainders."
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 4e-39] } },
+        "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 4e-39] }
+        }
+      }
+    },
+    {
+      "name": "f32_fmod_negative_operands",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-5.5, 5.5, -5.5, 5.5] } },
+        "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [2.0, 2.0, -2.0, -2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_scalar_fmod_negative_dividend",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.5] } },
+        "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_fmod_large_dividend_cancellation",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": {
+            "kind": "values",
+            "values": [100000000000000000000.0, 10000000000000.0, -100000000000000000000.0, 100000000000000000000.0, 2500000000.0, 123456789.0]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [3.0, 7.0, 3.0, 7.0, 100.0, 1000.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [6], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_f32_fmod_mixed_sign",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Fmod_float_mixed_sign"
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-4.3, 7.2, 5.0, 4.3, -7.2, 8.0] }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [2.1, -3.4, 8.0, -2.1, 3.4, 5.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_f16_fmod_mixed_sign",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Fmod_float16_mixed_sign"
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-4.3, 7.2, 5.0, 4.3, -7.2, 8.0] }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [2.1, -3.4, 8.0, -2.1, 3.4, 5.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [6], "tolerance": 0.002 } }
+    },
+    {
+      "name": "ort_uint32_mod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.UInt32_mod"
+      },
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [6], "data": { "kind": "values", "values": [4, 7, 5, 4, 7, 8] } },
+        "b": { "dtype": "uint32", "shape": [6], "data": { "kind": "values", "values": [2, 3, 8, 2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_mixed_sign_floor_mod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int32_mixed_sign"
+      },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_mixed_sign_fmod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int32_mixed_sign_fmod"
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_rank3_scalar_divisor_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int32_mod_bcast"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "int32",
+          "shape": [3, 2, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_int32_rank3_scalar_divisor_broadcast_input_a" }
+          }
+        },
+        "b": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [7] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int32",
+          "shape": [3, 2, 5],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [0, 1, 2, 3, 4, 5, 6, 0, 1, 2, 3, 4, 5, 6, 0, 1, 2, 3, 4, 5, 6, 0, 1, 2, 3, 4, 5, 6, 0, 1]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_int8_mixed_sign_floor_mod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int8_mixed_sign"
+      },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int8",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0, -2, 5, 0, 2, 3] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_int8_mixed_sign_fmod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int8_mixed_sign_fmod"
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int8",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0, 1, 5, 0, -1, 3] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_uint8_mod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.UInt8_mod"
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [6], "data": { "kind": "values", "values": [4, 7, 5, 4, 7, 8] } },
+        "b": { "dtype": "uint8", "shape": [6], "data": { "kind": "values", "values": [2, 3, 8, 2, 3, 5] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "uint8",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0, 1, 5, 0, 1, 3] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_int8_mixed_sign_floor_mod_vec4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int8_mixed_sign",
+        "notes": "Extended to 8 elements to exercise the vectorized WebGPU path while preserving ORT's mixed-sign floor-mod semantics."
+      },
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [8], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8, -9, 9] } },
+        "b": { "dtype": "int8", "shape": [8], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5, 4, -4] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, -2, 5, 0, 2, 3, 3, -3] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_uint8_mod_vec4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.UInt8_mod",
+        "notes": "Extended to 8 elements to exercise the vectorized WebGPU path while preserving ORT's unsigned mod semantics."
+      },
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [8], "data": { "kind": "values", "values": [4, 7, 5, 4, 7, 8, 255, 128] } },
+        "b": { "dtype": "uint8", "shape": [8], "data": { "kind": "values", "values": [2, 3, 8, 2, 3, 5, 10, 127] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "uint8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, 1, 5, 0, 1, 3, 5, 1] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_int16_mixed_sign_floor_mod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int16_mixed_sign"
+      },
+      "inputs": {
+        "a": { "dtype": "int16", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int16", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int16", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int16_mixed_sign_fmod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int16_mixed_sign_fmod"
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "int16", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int16", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int16", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_scalar_divisor_broadcast",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int32_mod_bcast"
+      },
+      "inputs": {
+        "a": {
+          "dtype": "int32",
+          "shape": [3, 2, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_int32_rank3_scalar_divisor_broadcast_input_a" }
+          }
+        },
+        "b": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [7] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [3, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_large_values_no_float_cast",
+      "inputs": {
+        "a": {
+          "dtype": "int32",
+          "shape": [6],
+          "data": {
+            "kind": "values",
+            "values": [2147483647, 2000000001, 16777217, -2147483647, -2000000001, -16777217]
+          }
+        },
+        "b": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [3, 97, 3, 3, 97, 3] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "u32_large_values_no_float_cast",
+      "inputs": {
+        "a": {
+          "dtype": "uint32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [4294967295, 4294967294, 4000000001, 16777217, 16777216] }
+        },
+        "b": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [97, 65535, 4000000000, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "zero_sized_broadcast_noop",
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [0, 3], "data": { "kind": "values", "values": [] } },
+        "b": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [0, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_zero_broadcast_loop_span",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.DimWithZeroHandling",
+        "notes": "Int32 projection of ORT's zero-dimension Mod BroadcastLoopSpan case."
+      },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [2, 2, 0], "data": { "kind": "values", "values": [] } },
+        "b": { "dtype": "int32", "shape": [2, 1], "data": { "kind": "values", "values": [1, 2] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [2, 2, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_mod_broadcast",
+      "inputs": {
+        "a": {
+          "dtype": "int32",
+          "shape": [3, 2, 5],
+          "data": {
+            "kind": "values",
+            "values": { "$ref": "#/fixtureArrays/ort_int32_rank3_scalar_divisor_broadcast_input_a" }
+          }
+        },
+        "b": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [7] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [3, 2, 5] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_broadcast" }
+    },
+    {
+      "name": "onnx_backend_mod_int64_fmod",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [6] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_int64_fmod",
+        "notes": "ONNX int64 tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_mod_mixed_sign_float16",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-4.30078125, 7.19921875, 5.0, 4.30078125, -7.19921875, 8.0] }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [6],
+          "data": { "kind": "values", "values": [2.099609375, -3.400390625, 8.0, -2.099609375, 3.400390625, 5.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [6] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_mixed_sign_float16" }
+    },
+    {
+      "name": "onnx_backend_mod_mixed_sign_float32",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": {
+            "kind": "values",
+            "values": [-4.300000190734863, 7.199999809265137, 5.0, 4.300000190734863, -7.199999809265137, 8.0]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": {
+            "kind": "values",
+            "values": [2.0999999046325684, -3.4000000953674316, 8.0, -2.0999999046325684, 3.4000000953674316, 5.0]
+          }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [6] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_mixed_sign_float32" }
+    },
+    {
+      "name": "onnx_backend_mod_mixed_sign_int32",
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int32", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [6] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_mixed_sign_int32" }
+    },
+    {
+      "name": "onnx_backend_mod_mixed_sign_int8",
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [-4, 7, 5, 4, -7, 8] } },
+        "b": { "dtype": "int8", "shape": [6], "data": { "kind": "values", "values": [2, -3, 8, -2, 3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "int8", "shape": [6] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_mixed_sign_int8" }
+    },
+    {
+      "name": "onnx_backend_mod_uint32",
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 7, 5] } },
+        "b": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 3, 8] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [3] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_uint32" }
+    },
+    {
+      "name": "onnx_backend_mod_uint8",
+      "inputs": {
+        "a": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [4, 7, 5] } },
+        "b": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [2, 3, 8] } }
+      },
+      "outputs": { "c": { "dtype": "uint8", "shape": [3] } },
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mod_uint8" }
+    },
+    {
+      "name": "vec4_u32_same_shape",
+      "inputs": {
+        "a": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [4, 7, 5, 4, 7, 8, 9, 10] } },
+        "b": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [2, 3, 8, 3, 4, 5, 7, 6] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "uint32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [0, 1, 5, 1, 3, 3, 2, 4] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "i32_scalar_divisor_vec4_floor_mod",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int32_mod_bcast",
+        "notes": "A vec4-aligned int32 tensor with a scalar negative divisor exercises scalar-divisor broadcasting and drives floor-mod correction on dividends on both sides of zero."
+      },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [-9, 7, -7, 15, 0, -1, 22, -22] } },
+        "b": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-7] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int32",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-2, 0, 0, -6, 0, -1, -6, -1] }
+        }
+      }
+    },
+    {
+      "name": "i32_scalar_divisor_vec4_fmod",
+      "attrs": { "fmod": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "ModOpTest.Int32_mod_bcast",
+        "notes": "The fmod=1 half of the same route: truncated remainder, so the sign follows the dividend and the floor-mod adjust must NOT fire. Left unpinned so the reference is exercised on the arm its pinned sibling covers only in the kernel."
+      },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [-9, 7, -7, 15, 0, -1, 22, -22] } },
+        "b": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [-7] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_i32_floor_mod_mixed_sign",
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [-4, 7, 5, -7, 8, -9, 10, -3] } },
+        "b": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [3, -3, 3, 3, -5, 5, 7, -2] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [2, -2, 2, 2, -2, 1, 3, -1] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "vec4_i32_fmod_mixed_sign",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [-4, 7, 5, -7, 8, -9, 10, -3] } },
+        "b": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [3, -3, 3, 3, -5, 5, 7, -2] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "int32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-1, 1, 2, -1, 3, -4, 3, -1] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "vec4_f16_fmod_mixed_sign",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [5.5, -5.5, 7.25, -7.25, 9.5, -9.5, 3.75, -3.75] }
+        },
+        "b": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [2.0, 2.0, 3.0, 3.0, 4.0, -4.0, 1.5, 1.5] }
+        }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.5, -1.5, 1.25, -1.25, 1.5, -1.5, 0.75, -0.75] },
+          "tolerance": 0.001
+        }
+      }
+    },
+    {
+      "name": "vec4_f32_fmod_large_dividend_cancellation",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": {
+            "kind": "values",
+            "values": [100000000000000000000.0, 10000000000000.0, -100000000000000000000.0, 100000000000000000000.0, 2500000000.0, 123456789.0, -10000000000000.0, 999999999.0]
+          }
+        },
+        "b": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [3.0, 7.0, 3.0, 7.0, 100.0, 1000.0, 7.0, 10.0] }
+        }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "i32_floor_mod_non_vec4_numel_negative_divisor",
+      "inputs": {
+        "a": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [5, -5, 7] } },
+        "b": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [-3, -3, -3] } }
+      },
+      "outputs": { "c": { "dtype": "int32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_fmod_values_near_f16_max",
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [65504.0, 65504.0, -65504.0, 1024.0] }
+        },
+        "b": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [3.0, 0.5, 3.0, 0.125] } }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [4], "tolerance": 0.5 } }
+    },
+    {
+      "name": "u32_scalar_b_non_vec4_large_values",
+      "inputs": {
+        "a": {
+          "dtype": "uint32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [4294967295, 4294967294, 100, 7, 3] }
+        },
+        "b": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4294967293] } }
+      },
+      "outputs": {
+        "c": {
+          "dtype": "uint32",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [2, 1, 100, 7, 3] }
+        }
+      }
+    },
+    {
+      "name": "i8_floor_mod_divisor_equals_dividend",
+      "inputs": {
+        "a": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, -128, -1, 127] } },
+        "b": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, 127, -1, -127] } }
+      },
+      "outputs": {
+        "c": { "dtype": "int8", "shape": [4], "tolerance": 0, "data": { "kind": "values", "values": [0, 126, 0, 0] } }
+      }
+    },
+    {
+      "name": "rank7_broadcast_scalar_tail",
+      "inputs": {
+        "a": {
+          "dtype": "uint32",
+          "shape": [2, 1, 2, 1, 2, 1, 3],
+          "data": { "kind": "cycle", "values": [5, 7, 11, 13] }
+        },
+        "b": { "dtype": "uint32", "shape": [1, 2, 1, 2, 1, 2, 1], "data": { "kind": "cycle", "values": [2, 3, 4] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [2, 2, 2, 2, 2, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "last_axis_broadcast_vec4_odd_row_crossing",
+      "provenance": {
+        "notes": "Four vec4 groups cross odd-width row boundaries, exercising wrapped per-lane B indices. Expected values come from the TypeScript Mod reference."
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float32",
+          "shape": [4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "scale": 17.0, "offset": -2.0 }
+        },
+        "b": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [2.0, 3.0, 4.0, 5.0, 6.0] } }
+      },
+      "outputs": { "c": { "dtype": "float32", "shape": [4, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "last_axis_broadcast_vec4_f16_odd_row_crossing",
+      "provenance": {
+        "notes": "f16 companion for the per-lane broadcast path; locks explicit widening of scalar B lanes before fmod arithmetic."
+      },
+      "attrs": { "fmod": 1 },
+      "inputs": {
+        "a": {
+          "dtype": "float16",
+          "shape": [4, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.41, "scale": 17.0, "offset": -2.0 }
+        },
+        "b": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [2.0, 3.0, 4.0, 5.0, 6.0] } }
+      },
+      "outputs": { "c": { "dtype": "float16", "shape": [4, 5], "tolerance": 0.01 } }
+    },
+    {
+      "name": "rank8_broadcast_alternating",
+      "attrs": {},
+      "inputs": {
+        "a": {
+          "dtype": "uint32",
+          "shape": [2, 1, 2, 1, 2, 1, 2, 3],
+          "data": { "kind": "cycle", "values": [7, 9, 13, 20] }
+        },
+        "b": { "dtype": "uint32", "shape": [1, 2, 1, 2, 1, 2, 1, 1], "data": { "kind": "cycle", "values": [3, 5] } }
+      },
+      "outputs": { "c": { "dtype": "uint32", "shape": [2, 2, 2, 2, 2, 2, 2, 3], "tolerance": 0 } }
+    }
+  ]
+}