Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Div

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
ec7dd38
·
verified ·
1 Parent(s): eebf847

sync 91d990483a17

Browse files
Files changed (8) hide show
  1. README.md +12 -8
  2. build/webgpu/bench.json +0 -1
  3. build/webgpu/binary-broadcast-vec4.wgsl.jinja +33 -37
  4. build/webgpu/binary-broadcast.wgsl.jinja +20 -16
  5. build/webgpu/binary-vec4.wgsl.jinja +37 -8
  6. build/webgpu/manifest.json +148 -186
  7. build/webgpu/metadata.json +19 -9
  8. build/webgpu/test.json +47 -5
README.md CHANGED
@@ -18,16 +18,16 @@ See the [ONNX `Div` spec](https://onnx.ai/onnx/operators/onnx__Div.html) for the
18
 
19
  ## Inputs
20
 
21
- | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
  | --- | --- | --- | --- | --- | --- | --- |
23
- | `A` | `a` | `T` | — | — | First operand (dividend). | required |
24
- | `B` | `b` | `T` | — | — | Second operand (divisor). | required |
25
 
26
  ## Outputs
27
 
28
- | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
29
  | --- | --- | --- | --- | --- | --- | --- |
30
- | `C` | `c` | `T` | derived | broadcast result of `A` and `B` | Result of elementwise division; same element type as the inputs. | required |
31
 
32
  ## Type constraints
33
 
@@ -37,7 +37,7 @@ See the [ONNX `Div` spec](https://onnx.ai/onnx/operators/onnx__Div.html) for the
37
 
38
  ## Files
39
 
40
- - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
41
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
42
  - [`test.json`](build/webgpu/test.json) — correctness cases
43
  - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
@@ -47,10 +47,14 @@ See the [ONNX `Div` spec](https://onnx.ai/onnx/operators/onnx__Div.html) for the
47
 
48
  ## Use with `@huggingface/kernels`
49
 
50
- The loader derives every required output's shape and logical dtype from the manifest contract and this call.
51
- It then allocates the result tensors automatically.
 
 
 
52
 
53
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
 
54
 
55
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
56
 
 
18
 
19
  ## Inputs
20
 
21
+ | Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
22
  | --- | --- | --- | --- | --- | --- | --- |
23
+ | `a` | `A` | `T` | — | — | First operand (dividend). | required |
24
+ | `b` | `B` | `T` | — | — | Second operand (divisor). | required |
25
 
26
  ## Outputs
27
 
28
+ | Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
29
  | --- | --- | --- | --- | --- | --- | --- |
30
+ | `c` | `C` | `T` | derived | broadcast result of `a` and `b` | Result of elementwise division; same element type as the inputs. | required |
31
 
32
  ## Type constraints
33
 
 
37
 
38
  ## Files
39
 
40
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
41
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
42
  - [`test.json`](build/webgpu/test.json) — correctness cases
43
  - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 
47
 
48
  ## Use with `@huggingface/kernels`
49
 
50
+ ```sh
51
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.2
52
+ ```
53
+
54
+ Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
55
 
56
  The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
57
+ It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
58
 
59
  Replace each `*Data` placeholder with a typed array containing the corresponding input data.
60
 
build/webgpu/bench.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "ai.onnx.Div",
3
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
  "cases": [
5
  {
 
1
  {
 
2
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
3
  "cases": [
4
  {
build/webgpu/binary-broadcast-vec4.wgsl.jinja CHANGED
@@ -42,14 +42,10 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
42
  {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
43
  {%- endmacro %}
44
 
45
- {% if usesF16 %}
46
- enable f16;
47
- {% endif %}
48
  {{ env.wgsl.resourceDeclarations }}
49
 
50
- // Vec4 broadcast binary op. The same compiled-stride folding as
51
- // binary-broadcast (offsets are compile-time-strength-reducible), but each
52
- // thread writes a vec4 of the output. If the output's innermost axis is a
53
  // multiple of four, each operand uses one base offset. For an
54
  // odd innermost extent, a vec4 may cross a row boundary: same-shape operands
55
  // remain contiguous, while broadcast operands compute the four lane offsets
@@ -63,36 +59,36 @@ enable f16;
63
  // operand is vector-bound and loaded directly; broadcast operands remain
64
  // scalar-bound so their independent lane offsets stay addressable.
65
  {% set a_numel = namespace(value=1) %}
66
- {% for d in source.aShape %}
67
  {% set a_numel.value = a_numel.value * d %}
68
  {% endfor %}
69
  {% set b_numel = namespace(value=1) %}
70
- {% for d in source.bShape %}
71
  {% set b_numel.value = b_numel.value * d %}
72
  {% endfor %}
73
  {% set c_numel = namespace(value=1) %}
74
- {% for d in source.cShape %}
75
  {% set c_numel.value = c_numel.value * d %}
76
  {% endfor %}
77
- {% set a_same = namespace(value=(source.aRank == source.cRank)) %}
78
  {% if a_same.value %}
79
- {% for axis in range(source.cRank) %}
80
- {% if source.aShape[axis] != source.cShape[axis] %}
81
  {% set a_same.value = false %}
82
  {% endif %}
83
  {% endfor %}
84
  {% endif %}
85
- {% set b_same = namespace(value=(source.bRank == source.cRank)) %}
86
  {% if b_same.value %}
87
- {% for axis in range(source.cRank) %}
88
- {% if source.bShape[axis] != source.cShape[axis] %}
89
  {% set b_same.value = false %}
90
  {% endif %}
91
  {% endfor %}
92
  {% endif %}
93
- {% set a_inner = source.aShape[source.aRank - 1] if source.aRank >= 1 else 1 %}
94
- {% set b_inner = source.bShape[source.bRank - 1] if source.bRank >= 1 else 1 %}
95
- {% set c_inner = source.cShape[source.cRank - 1] %}
96
  {% set crosses_inner_rows = c_inner % 4 != 0 %}
97
  {% if a_numel.value == 1 %}{% set a_mode = "scalar" %}
98
  {% elif a_inner == 1 %}{% set a_mode = "splat" %}
@@ -102,31 +98,31 @@ enable f16;
102
  {% else %}{% set b_mode = "contig" %}{% endif %}
103
 
104
  {% if not a_same.value and a_mode != "scalar" %}
105
- {{ offset_fn("a_offset", source.aShape, source.aRank, false, a_numel.value, source.cShape, source.cRank, c_numel.value) }}
106
  {% endif %}
107
 
108
  {% if not b_same.value and b_mode != "scalar" %}
109
- {{ offset_fn("b_offset", source.bShape, source.bRank, false, b_numel.value, source.cShape, source.cRank, c_numel.value) }}
110
  {% endif %}
111
 
112
  {% set is_int = scalar == "i32" or scalar == "u32" %}
113
  {% set acc = scalar if is_int else "f32" %}
114
  // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
115
  // use i32/u32 storage.
116
- {% if source.cDtype == "int8" %}
117
  fn wrap_dtype(v: vec4<i32>) -> vec4<i32> { return (v << vec4<u32>(24u)) >> vec4<u32>(24u); }
118
  {% set wrap = "wrap_dtype" %}
119
- {% elif source.cDtype == "uint8" %}
120
  fn wrap_dtype(v: vec4<u32>) -> vec4<u32> { return v & vec4<u32>(0xFFu); }
121
  {% set wrap = "wrap_dtype" %}
122
  {% else %}
123
  {% set wrap = "" %}
124
  {% endif %}
125
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
126
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
127
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
128
- // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
129
- let i4 = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
130
  if (i4 >= params.count) {
131
  return;
132
  }
@@ -140,15 +136,15 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
140
  let av = vec4<{{ acc }}>({{ acc }}(a[0]));
141
  {% elif crosses_inner_rows %}
142
  let av = vec4<{{ acc }}>(
143
- {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }}]),
144
- {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base + 1u") }}]),
145
- {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base + 2u") }}]),
146
- {{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base + 3u") }}])
147
  );
148
  {% elif a_mode == "splat" %}
149
- let av = vec4<{{ acc }}>({{ acc }}(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }}]));
150
  {% else %}
151
- let ao = {{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "base") }};
152
  let av = vec4<{{ acc }}>({{ acc }}(a[ao]), {{ acc }}(a[ao + 1u]), {{ acc }}(a[ao + 2u]), {{ acc }}(a[ao + 3u]));
153
  {% endif %}
154
  {% if b_same.value %}
@@ -157,15 +153,15 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
157
  let bv = vec4<{{ acc }}>({{ acc }}(b[0]));
158
  {% elif crosses_inner_rows %}
159
  let bv = vec4<{{ acc }}>(
160
- {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }}]),
161
- {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base + 1u") }}]),
162
- {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base + 2u") }}]),
163
- {{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base + 3u") }}])
164
  );
165
  {% elif b_mode == "splat" %}
166
- let bv = vec4<{{ acc }}>({{ acc }}(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }}]));
167
  {% else %}
168
- let bo = {{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "base") }};
169
  let bv = vec4<{{ acc }}>({{ acc }}(b[bo]), {{ acc }}(b[bo + 1u]), {{ acc }}(b[bo + 2u]), {{ acc }}(b[bo + 3u]));
170
  {% endif %}
171
  {% if is_int %}
 
42
  {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
43
  {%- endmacro %}
44
 
 
 
 
45
  {{ env.wgsl.resourceDeclarations }}
46
 
47
+ // Vec4 broadcast binary op. Offsets use compile-time strides, and each
48
+ // invocation writes four output values. If the output's innermost axis is a
 
49
  // multiple of four, each operand uses one base offset. For an
50
  // odd innermost extent, a vec4 may cross a row boundary: same-shape operands
51
  // remain contiguous, while broadcast operands compute the four lane offsets
 
59
  // operand is vector-bound and loaded directly; broadcast operands remain
60
  // scalar-bound so their independent lane offsets stay addressable.
61
  {% set a_numel = namespace(value=1) %}
62
+ {% for d in aShape %}
63
  {% set a_numel.value = a_numel.value * d %}
64
  {% endfor %}
65
  {% set b_numel = namespace(value=1) %}
66
+ {% for d in bShape %}
67
  {% set b_numel.value = b_numel.value * d %}
68
  {% endfor %}
69
  {% set c_numel = namespace(value=1) %}
70
+ {% for d in cShape %}
71
  {% set c_numel.value = c_numel.value * d %}
72
  {% endfor %}
73
+ {% set a_same = namespace(value=(aRank == cRank)) %}
74
  {% if a_same.value %}
75
+ {% for axis in range(cRank) %}
76
+ {% if aShape[axis] != cShape[axis] %}
77
  {% set a_same.value = false %}
78
  {% endif %}
79
  {% endfor %}
80
  {% endif %}
81
+ {% set b_same = namespace(value=(bRank == cRank)) %}
82
  {% if b_same.value %}
83
+ {% for axis in range(cRank) %}
84
+ {% if bShape[axis] != cShape[axis] %}
85
  {% set b_same.value = false %}
86
  {% endif %}
87
  {% endfor %}
88
  {% endif %}
89
+ {% set a_inner = aShape[aRank - 1] if aRank >= 1 else 1 %}
90
+ {% set b_inner = bShape[bRank - 1] if bRank >= 1 else 1 %}
91
+ {% set c_inner = cShape[cRank - 1] %}
92
  {% set crosses_inner_rows = c_inner % 4 != 0 %}
93
  {% if a_numel.value == 1 %}{% set a_mode = "scalar" %}
94
  {% elif a_inner == 1 %}{% set a_mode = "splat" %}
 
98
  {% else %}{% set b_mode = "contig" %}{% endif %}
99
 
100
  {% if not a_same.value and a_mode != "scalar" %}
101
+ {{ offset_fn("a_offset", aShape, aRank, false, a_numel.value, cShape, cRank, c_numel.value) }}
102
  {% endif %}
103
 
104
  {% if not b_same.value and b_mode != "scalar" %}
105
+ {{ offset_fn("b_offset", bShape, bRank, false, b_numel.value, cShape, cRank, c_numel.value) }}
106
  {% endif %}
107
 
108
  {% set is_int = scalar == "i32" or scalar == "u32" %}
109
  {% set acc = scalar if is_int else "f32" %}
110
  // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
111
  // use i32/u32 storage.
112
+ {% if cDtype == "int8" %}
113
  fn wrap_dtype(v: vec4<i32>) -> vec4<i32> { return (v << vec4<u32>(24u)) >> vec4<u32>(24u); }
114
  {% set wrap = "wrap_dtype" %}
115
+ {% elif cDtype == "uint8" %}
116
  fn wrap_dtype(v: vec4<u32>) -> vec4<u32> { return v & vec4<u32>(0xFFu); }
117
  {% set wrap = "wrap_dtype" %}
118
  {% else %}
119
  {% set wrap = "" %}
120
  {% endif %}
121
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
122
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
123
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
124
+ // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
125
+ let i4 = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
126
  if (i4 >= params.count) {
127
  return;
128
  }
 
136
  let av = vec4<{{ acc }}>({{ acc }}(a[0]));
137
  {% elif crosses_inner_rows %}
138
  let av = vec4<{{ acc }}>(
139
+ {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base") }}]),
140
+ {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base + 1u") }}]),
141
+ {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base + 2u") }}]),
142
+ {{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base + 3u") }}])
143
  );
144
  {% elif a_mode == "splat" %}
145
+ let av = vec4<{{ acc }}>({{ acc }}(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "base") }}]));
146
  {% else %}
147
+ let ao = {{ broadcast_offset_call("a_offset", aShape, cShape, "base") }};
148
  let av = vec4<{{ acc }}>({{ acc }}(a[ao]), {{ acc }}(a[ao + 1u]), {{ acc }}(a[ao + 2u]), {{ acc }}(a[ao + 3u]));
149
  {% endif %}
150
  {% if b_same.value %}
 
153
  let bv = vec4<{{ acc }}>({{ acc }}(b[0]));
154
  {% elif crosses_inner_rows %}
155
  let bv = vec4<{{ acc }}>(
156
+ {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base") }}]),
157
+ {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base + 1u") }}]),
158
+ {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base + 2u") }}]),
159
+ {{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base + 3u") }}])
160
  );
161
  {% elif b_mode == "splat" %}
162
+ let bv = vec4<{{ acc }}>({{ acc }}(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "base") }}]));
163
  {% else %}
164
+ let bo = {{ broadcast_offset_call("b_offset", bShape, cShape, "base") }};
165
  let bv = vec4<{{ acc }}>({{ acc }}(b[bo]), {{ acc }}(b[bo + 1u]), {{ acc }}(b[bo + 2u]), {{ acc }}(b[bo + 3u]));
166
  {% endif %}
167
  {% if is_int %}
build/webgpu/binary-broadcast.wgsl.jinja CHANGED
@@ -1,13 +1,14 @@
1
  {% macro flat_tail_open() %}
2
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
4
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
- // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
6
- let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
7
  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
8
  // the logical tensor length (or its storage binding) to be vec4 aligned.
9
- let begin = invocation * {{ source.itemsPerInvocation }}u;
10
- let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
 
11
  for (var i = begin; i < end; i = i + 1u) {
12
  {%- endmacro %}
13
  {% macro flat_tail_close() %}
@@ -75,14 +76,17 @@ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif
75
  {% endif %}
76
  {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
77
  {%- endmacro %}{% macro binary_broadcast_offsets() %}
78
- {{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.cShape, source.cRank) }}
 
 
 
 
 
 
79
 
80
- {{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.cShape, source.cRank) }}
81
  {%- endmacro %}
82
 
83
- {% if usesF16 %}
84
- enable f16;
85
- {% endif %}
86
  {{ env.wgsl.resourceDeclarations }}
87
 
88
 
@@ -90,21 +94,21 @@ enable f16;
90
 
91
  {{ flat_tail_open() }}
92
  {% if scalar == "i32" or scalar == "u32" %}
93
- let av = a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}];
94
- let bv = b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}];
95
  let r = av / bv;
96
  // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
97
  // use i32/u32 storage.
98
- {% if source.cDtype == "int8" %}
99
  c[i] = (r << 24u) >> 24u;
100
- {% elif source.cDtype == "uint8" %}
101
  c[i] = r & 0xFFu;
102
  {% else %}
103
  c[i] = r;
104
  {% endif %}
105
  {% else %}
106
- let av = f32(a[{{ broadcast_offset_call("a_offset", source.aShape, source.cShape, "i") }}]);
107
- let bv = f32(b[{{ broadcast_offset_call("b_offset", source.bShape, source.cShape, "i") }}]);
108
  let quotient = av / bv;
109
  let rounded = round(quotient);
110
  // GPU f32 division can land slightly off an exact integer quotient. Snap only
 
1
  {% macro flat_tail_open() %}
2
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
3
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
4
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
5
+ // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
6
+ let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
7
  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
8
  // the logical tensor length (or its storage binding) to be vec4 aligned.
9
+ {% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
10
+ let begin = invocation * {{ itemsPerInvocation }}u;
11
+ let end = min(begin + {{ itemsPerInvocation }}u, params.count);
12
  for (var i = begin; i < end; i = i + 1u) {
13
  {%- endmacro %}
14
  {% macro flat_tail_close() %}
 
76
  {% endif %}
77
  {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
78
  {%- endmacro %}{% macro binary_broadcast_offsets() %}
79
+ {% set aShape = aShape | default([]) %}
80
+ {% set aRank = aRank | default(0) %}
81
+ {% set bShape = bShape | default([]) %}
82
+ {% set bRank = bRank | default(0) %}
83
+ {% set cShape = cShape | default([]) %}
84
+ {% set cRank = cRank | default(0) %}
85
+ {{ broadcast_offset_fn("a_offset", aShape, aRank, cShape, cRank) }}
86
 
87
+ {{ broadcast_offset_fn("b_offset", bShape, bRank, cShape, cRank) }}
88
  {%- endmacro %}
89
 
 
 
 
90
  {{ env.wgsl.resourceDeclarations }}
91
 
92
 
 
94
 
95
  {{ flat_tail_open() }}
96
  {% if scalar == "i32" or scalar == "u32" %}
97
+ let av = a[{{ broadcast_offset_call("a_offset", aShape, cShape, "i") }}];
98
+ let bv = b[{{ broadcast_offset_call("b_offset", bShape, cShape, "i") }}];
99
  let r = av / bv;
100
  // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
101
  // use i32/u32 storage.
102
+ {% if cDtype == "int8" %}
103
  c[i] = (r << 24u) >> 24u;
104
+ {% elif cDtype == "uint8" %}
105
  c[i] = r & 0xFFu;
106
  {% else %}
107
  c[i] = r;
108
  {% endif %}
109
  {% else %}
110
+ let av = f32(a[{{ broadcast_offset_call("a_offset", aShape, cShape, "i") }}]);
111
+ let bv = f32(b[{{ broadcast_offset_call("b_offset", bShape, cShape, "i") }}]);
112
  let quotient = av / bv;
113
  let rounded = round(quotient);
114
  // GPU f32 division can land slightly off an exact integer quotient. Snap only
build/webgpu/binary-vec4.wgsl.jinja CHANGED
@@ -1,25 +1,51 @@
1
- {% if usesF16 %}
2
- enable f16;
3
- {% endif %}
4
  {{ env.wgsl.resourceDeclarations }}
5
 
 
 
 
 
 
6
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
7
- fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
8
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
9
- // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
10
- let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
 
 
 
 
 
 
 
 
 
 
 
 
11
  if (i >= params.count) {
12
  return;
13
  }
 
 
 
 
 
 
 
14
  let av = a[i];
 
 
 
 
15
  let bv = b[i];
 
16
  {% if scalar == "i32" or scalar == "u32" %}
17
  let r = av / bv;
18
  // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
19
  // use i32/u32 storage.
20
- {% if source.cDtype == "int8" %}
21
  c[i] = (r << vec4<u32>(24u)) >> vec4<u32>(24u);
22
- {% elif source.cDtype == "uint8" %}
23
  c[i] = r & vec4<u32>(0xFFu);
24
  {% else %}
25
  c[i] = r;
@@ -49,4 +75,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
49
  select(quotient.w, rounded.w, snapProduct.w && snapClose.w)
50
  );
51
  {% endif %}
 
 
 
52
  }
 
 
 
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
+ {% set vec4PerThread = vec4PerThread %}
4
+ {% if vec4PerThread > 1 %}
5
+ const ITEMS: u32 = {{ vec4PerThread }}u;
6
+ {% endif %}
7
+
8
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
9
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
10
  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
11
+ // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
12
+ {% if vec4PerThread > 1 %}
13
+ // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
14
+ // access consecutive words on every step, while each lane can keep several
15
+ // independent loads in flight.
16
+ let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
17
+ let span = (params.count + ITEMS - 1u) / ITEMS;
18
+ for (var j = 0u; j < ITEMS; j = j + 1u) {
19
+ let i = tid + j * span;
20
+ if (i >= params.count) {
21
+ break;
22
+ }
23
+ {% else %}
24
+ let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
25
  if (i >= params.count) {
26
  return;
27
  }
28
+ {% endif %}
29
+
30
+ {% set scalarOperand = scalarOperand if scalarOperand is defined else "" %}
31
+ {% if scalarOperand == "a" %}
32
+ // One-element operand: read once and splat across the vector.
33
+ let av = {{ vectorScalar }}(a[0]);
34
+ {% else %}
35
  let av = a[i];
36
+ {% endif %}
37
+ {% if scalarOperand == "b" %}
38
+ let bv = {{ vectorScalar }}(b[0]);
39
+ {% else %}
40
  let bv = b[i];
41
+ {% endif %}
42
  {% if scalar == "i32" or scalar == "u32" %}
43
  let r = av / bv;
44
  // Narrow integer operations wrap modulo the logical dtype width; int8/uint8
45
  // use i32/u32 storage.
46
+ {% if cDtype == "int8" %}
47
  c[i] = (r << vec4<u32>(24u)) >> vec4<u32>(24u);
48
+ {% elif cDtype == "uint8" %}
49
  c[i] = r & vec4<u32>(0xFFu);
50
  {% else %}
51
  c[i] = r;
 
75
  select(quotient.w, rounded.w, snapProduct.w && snapClose.w)
76
  );
77
  {% endif %}
78
+ {% if vec4PerThread > 1 %}
79
+ }
80
+ {% endif %}
81
  }
build/webgpu/manifest.json CHANGED
@@ -2,234 +2,196 @@
2
  "domain": "ai.onnx",
3
  "name": "Div",
4
  "sinceVersion": 14,
5
- "description": "Performs elementwise binary division of two tensors with NumPy-style multidirectional broadcasting. For integer types, division truncates toward zero.",
6
- "inputs": [
7
- { "role": "A", "dtype": "T", "description": "First operand (dividend)." },
8
- { "role": "B", "dtype": "T", "description": "Second operand (divisor)." }
9
- ],
10
- "outputs": [
11
- {
12
- "role": "C",
13
- "dtype": "T",
14
- "rank": "max(ranks.A, ranks.B)",
15
- "description": "Result of elementwise division; same element type as the inputs.",
16
- "shape": "broadcastShape(shapes.A, shapes.B)"
17
- }
18
- ],
19
- "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int8", "uint8"] },
20
- "args": {
21
- "a": { "kind": "tensor", "semantic": "A", "role": "input" },
22
- "b": { "kind": "tensor", "semantic": "B", "role": "input" },
23
- "c": { "kind": "tensor", "semantic": "C", "role": "output" }
24
  },
25
- "tunables": { "WORKGROUP_SIZE": 256 },
 
26
  "variants": [
27
  {
28
  "id": "same_shape_vec4",
29
  "priority": 20,
30
- "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 == 0", "f16Ok(dtypes.T)"],
31
- "constants": {
32
- "scalar": "dtypes.T",
33
- "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
34
- "usesF16": "dtypes.T == \"f16\""
35
- },
36
  "passes": [
37
  {
38
  "id": "main",
39
  "name": "Div.vec4",
40
- "source": { "shader": "binary-vec4.wgsl.jinja", "inputs": { "op": "\"div\"", "cDtype": "tensorDtypes.C" } },
41
- "bindings": [
42
- {
43
- "name": "a",
44
- "arg": "a",
45
- "semantic": "A",
46
- "buffer": { "type": "read-only-storage" },
47
- "elementType": "$vectorScalar"
48
- },
49
- {
50
- "name": "b",
51
- "arg": "b",
52
- "semantic": "B",
53
- "buffer": { "type": "read-only-storage" },
54
- "elementType": "$vectorScalar"
55
- },
56
- {
57
- "name": "c",
58
- "arg": "c",
59
- "semantic": "C",
60
- "buffer": { "type": "storage" },
61
- "elementType": "$vectorScalar"
62
- },
63
- {
64
- "name": "params",
65
- "semantic": "kernel.params",
66
- "buffer": { "type": "uniform" },
67
- "struct": {
68
- "name": "Params",
69
- "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
70
- }
71
- }
72
- ],
73
- "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
74
  }
75
  ]
76
  },
77
  {
78
  "id": "broadcast_vec4",
79
- "when": ["ranks.A <= ranks.C", "ranks.B <= ranks.C", "ranks.C >= 1", "numel(shapes.C) % 4 == 0", "numel(shapes.C) >= 4", "f16Ok(dtypes.T)"],
80
- "constants": {
 
81
  "scalar": "dtypes.T",
82
- "aElement": "\"vec4<\" ~ dtypes.T ~ \">\" if sameShape(shapes.A, shapes.C) else dtypes.T",
83
- "bElement": "\"vec4<\" ~ dtypes.T ~ \">\" if sameShape(shapes.B, shapes.C) else dtypes.T",
84
- "vec4Scalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
85
- "usesF16": "dtypes.T == \"f16\""
86
  },
87
  "passes": [
88
  {
89
  "id": "main",
90
  "name": "Div",
91
- "source": {
92
- "shader": "binary-broadcast-vec4.wgsl.jinja",
93
- "inputs": {
94
- "aShape": "shapes.A",
95
- "bShape": "shapes.B",
96
- "cShape": "shapes.C",
97
- "aRank": "ranks.A",
98
- "bRank": "ranks.B",
99
- "cRank": "ranks.C",
100
- "op": "\"div\"",
101
- "cDtype": "tensorDtypes.C"
102
- }
103
  },
104
- "bindings": [
105
- {
106
- "name": "a",
107
- "arg": "a",
108
- "semantic": "A",
109
- "buffer": { "type": "read-only-storage" },
110
- "elementType": "$aElement"
111
- },
112
- {
113
- "name": "b",
114
- "arg": "b",
115
- "semantic": "B",
116
- "buffer": { "type": "read-only-storage" },
117
- "elementType": "$bElement"
118
- },
119
- { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$vec4Scalar" },
120
- {
121
- "name": "params",
122
- "semantic": "kernel.params",
123
- "buffer": { "type": "uniform" },
124
- "struct": {
125
- "name": "Params",
126
- "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C) / 4" }]
127
- }
128
- }
129
- ],
130
- "dispatch": { "threads": "numel(shapes.C) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
131
  }
132
- ],
133
- "priority": 10
134
  },
135
  {
136
  "id": "same_shape_scalar_x4",
137
  "priority": 15,
138
- "when": ["sameShape(shapes.A, shapes.C)", "sameShape(shapes.B, shapes.C)", "numel(shapes.C) > 0", "numel(shapes.C) % 4 != 0", "f16Ok(dtypes.T)"],
139
- "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
140
  "passes": [
141
  {
142
  "id": "main",
143
  "name": "Div",
144
- "source": {
145
- "shader": "binary-broadcast.wgsl.jinja",
146
- "inputs": {
147
- "aShape": "shapes.A",
148
- "bShape": "shapes.B",
149
- "cShape": "shapes.C",
150
- "aRank": "ranks.A",
151
- "bRank": "ranks.B",
152
- "cRank": "ranks.C",
153
- "op": "\"div\"",
154
- "cDtype": "tensorDtypes.C",
155
- "itemsPerInvocation": 4
156
- }
157
  },
158
- "bindings": [
159
- {
160
- "name": "a",
161
- "arg": "a",
162
- "semantic": "A",
163
- "buffer": { "type": "read-only-storage" },
164
- "elementType": "$scalar"
165
- },
166
- {
167
- "name": "b",
168
- "arg": "b",
169
- "semantic": "B",
170
- "buffer": { "type": "read-only-storage" },
171
- "elementType": "$scalar"
172
- },
173
- { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$scalar" },
174
- {
175
- "name": "params",
176
- "semantic": "kernel.params",
177
- "buffer": { "type": "uniform" },
178
- "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
179
- }
180
- ],
181
- "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
182
  }
183
  ]
184
  },
185
  {
186
  "id": "broadcast",
187
- "when": ["ranks.A <= ranks.C", "ranks.B <= ranks.C", "f16Ok(dtypes.T)"],
188
- "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
189
  "passes": [
190
  {
191
  "id": "main",
192
  "name": "Div",
193
- "source": {
194
- "shader": "binary-broadcast.wgsl.jinja",
195
- "inputs": {
196
- "aShape": "shapes.A",
197
- "bShape": "shapes.B",
198
- "cShape": "shapes.C",
199
- "aRank": "ranks.A",
200
- "bRank": "ranks.B",
201
- "cRank": "ranks.C",
202
- "op": "\"div\"",
203
- "cDtype": "tensorDtypes.C",
204
- "itemsPerInvocation": 4
205
- }
206
  },
207
- "bindings": [
208
- {
209
- "name": "a",
210
- "arg": "a",
211
- "semantic": "A",
212
- "buffer": { "type": "read-only-storage" },
213
- "elementType": "$scalar"
214
- },
215
- {
216
- "name": "b",
217
- "arg": "b",
218
- "semantic": "B",
219
- "buffer": { "type": "read-only-storage" },
220
- "elementType": "$scalar"
221
- },
222
- { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "storage" }, "elementType": "$scalar" },
223
- {
224
- "name": "params",
225
- "semantic": "kernel.params",
226
- "buffer": { "type": "uniform" },
227
- "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.C)" }] }
228
- }
229
- ],
230
- "dispatch": { "threads": "ceilDiv(numel(shapes.C), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
231
  }
232
  ]
233
  }
234
- ]
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
235
  }
 
2
  "domain": "ai.onnx",
3
  "name": "Div",
4
  "sinceVersion": 14,
5
+ "inputs": { "a": { "onnx": "A", "dtype": "T" }, "b": { "onnx": "B", "dtype": "T" } },
6
+ "outputs": {
7
+ "c": { "onnx": "C", "dtype": "T", "rank": "max(ranks.a, ranks.b)", "shape": "broadcastShape(shapes.a, shapes.b)" }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
8
  },
9
+ "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int8", "uint8"] },
10
+ "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
11
  "variants": [
12
  {
13
  "id": "same_shape_vec4",
14
  "priority": 20,
15
+ "when": ["sameShape(shapes.a, shapes.c)", "sameShape(shapes.b, shapes.c)", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
16
+ "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
 
 
 
 
17
  "passes": [
18
  {
19
  "id": "main",
20
  "name": "Div.vec4",
21
+ "shader": "binary-vec4.wgsl.jinja",
22
+ "derive": {
23
+ "op": "\"div\"",
24
+ "cDtype": "tensorDtypes.c",
25
+ "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
26
+ },
27
+ "bindings": ["a", "b", "c_binary", "params"],
28
+ "dispatch": {
29
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
30
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
31
+ "z": 1
32
+ }
33
+ }
34
+ ]
35
+ },
36
+ {
37
+ "id": "scalar_b_vec4",
38
+ "priority": 18,
39
+ "when": ["sameShape(shapes.a, shapes.c)", "numel(shapes.b) == 1", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
40
+ "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
41
+ "passes": [
42
+ {
43
+ "id": "main",
44
+ "name": "Div.scalarBVec4",
45
+ "shader": "binary-vec4.wgsl.jinja",
46
+ "derive": {
47
+ "op": "\"div\"",
48
+ "cDtype": "tensorDtypes.c",
49
+ "scalarOperand": "\"b\"",
50
+ "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
51
+ },
52
+ "bindings": ["a", "b_2", "c_binary", "params"],
53
+ "dispatch": {
54
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
55
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
56
+ "z": 1
57
+ }
58
+ }
59
+ ]
60
+ },
61
+ {
62
+ "id": "scalar_a_vec4",
63
+ "priority": 18,
64
+ "when": ["sameShape(shapes.b, shapes.c)", "numel(shapes.a) == 1", "numel(shapes.c) > 0", "numel(shapes.c) % 4 == 0", "f16Ok(dtypes.T)"],
65
+ "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
66
+ "passes": [
67
+ {
68
+ "id": "main",
69
+ "name": "Div.scalarAVec4",
70
+ "shader": "binary-vec4.wgsl.jinja",
71
+ "derive": {
72
+ "op": "\"div\"",
73
+ "cDtype": "tensorDtypes.c",
74
+ "scalarOperand": "\"a\"",
75
+ "vec4PerThread": "4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1"
76
+ },
77
+ "bindings": ["a_2", "b", "c_binary", "params"],
78
+ "dispatch": {
79
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
80
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c) / 4, 4 if numel(shapes.c) * dtypeBytes(tensorDtypes.c) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
81
+ "z": 1
82
+ }
83
  }
84
  ]
85
  },
86
  {
87
  "id": "broadcast_vec4",
88
+ "priority": 10,
89
+ "when": ["ranks.a <= ranks.c", "ranks.b <= ranks.c", "ranks.c >= 1", "numel(shapes.c) % 4 == 0", "numel(shapes.c) >= 4", "f16Ok(dtypes.T)"],
90
+ "derive": {
91
  "scalar": "dtypes.T",
92
+ "aElement": "\"vec4<\" ~ dtypes.T ~ \">\" if sameShape(shapes.a, shapes.c) else dtypes.T",
93
+ "bElement": "\"vec4<\" ~ dtypes.T ~ \">\" if sameShape(shapes.b, shapes.c) else dtypes.T",
94
+ "vec4Scalar": "\"vec4<\" ~ dtypes.T ~ \">\""
 
95
  },
96
  "passes": [
97
  {
98
  "id": "main",
99
  "name": "Div",
100
+ "shader": "binary-broadcast-vec4.wgsl.jinja",
101
+ "derive": {
102
+ "aShape": "shapes.a",
103
+ "bShape": "shapes.b",
104
+ "cShape": "shapes.c",
105
+ "aRank": "ranks.a",
106
+ "bRank": "ranks.b",
107
+ "cRank": "ranks.c",
108
+ "op": "\"div\"",
109
+ "cDtype": "tensorDtypes.c"
 
 
110
  },
111
+ "bindings": ["a_3", "b_3", "c_2_binary", "params"],
112
+ "dispatch": {
113
+ "x": "min(ceilDiv((numel(shapes.c) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
114
+ "y": "ceilDiv(ceilDiv((numel(shapes.c) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
115
+ "z": 1
116
+ }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
117
  }
118
+ ]
 
119
  },
120
  {
121
  "id": "same_shape_scalar_x4",
122
  "priority": 15,
123
+ "when": ["sameShape(shapes.a, shapes.c)", "sameShape(shapes.b, shapes.c)", "numel(shapes.c) > 0", "numel(shapes.c) % 4 != 0", "f16Ok(dtypes.T)"],
124
+ "derive": { "scalar": "dtypes.T" },
125
  "passes": [
126
  {
127
  "id": "main",
128
  "name": "Div",
129
+ "shader": "binary-broadcast.wgsl.jinja",
130
+ "derive": {
131
+ "aShape": "shapes.a",
132
+ "bShape": "shapes.b",
133
+ "cShape": "shapes.c",
134
+ "aRank": "ranks.a",
135
+ "bRank": "ranks.b",
136
+ "cRank": "ranks.c",
137
+ "op": "\"div\"",
138
+ "cDtype": "tensorDtypes.c",
139
+ "itemsPerInvocation": 4
 
 
140
  },
141
+ "bindings": ["a_2", "b_2", "c_3", "params_2"],
142
+ "dispatch": {
143
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
144
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
145
+ "z": 1
146
+ }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
147
  }
148
  ]
149
  },
150
  {
151
  "id": "broadcast",
152
+ "when": ["ranks.a <= ranks.c", "ranks.b <= ranks.c", "f16Ok(dtypes.T)"],
153
+ "derive": { "scalar": "dtypes.T" },
154
  "passes": [
155
  {
156
  "id": "main",
157
  "name": "Div",
158
+ "shader": "binary-broadcast.wgsl.jinja",
159
+ "derive": {
160
+ "aShape": "shapes.a",
161
+ "bShape": "shapes.b",
162
+ "cShape": "shapes.c",
163
+ "aRank": "ranks.a",
164
+ "bRank": "ranks.b",
165
+ "cRank": "ranks.c",
166
+ "op": "\"div\"",
167
+ "cDtype": "tensorDtypes.c",
168
+ "itemsPerInvocation": 4
 
 
169
  },
170
+ "bindings": ["a_2", "b_2", "c_3", "params_2"],
171
+ "dispatch": {
172
+ "x": "min(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
173
+ "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.c), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
174
+ "z": 1
175
+ }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
176
  }
177
  ]
178
  }
179
+ ],
180
+ "bindings": {
181
+ "a": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
182
+ "b": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
183
+ "c_binary": { "buffer": "storage", "elementType": "$vectorScalar", "name": "c" },
184
+ "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c) / 4" }] },
185
+ "b_2": { "buffer": "read-only-storage", "name": "b", "elementType": "$scalar" },
186
+ "a_2": { "buffer": "read-only-storage", "name": "a", "elementType": "$scalar" },
187
+ "a_3": { "buffer": "read-only-storage", "name": "a", "elementType": "$aElement" },
188
+ "b_3": { "buffer": "read-only-storage", "name": "b", "elementType": "$bElement" },
189
+ "c_2_binary": { "buffer": "storage", "name": "c", "elementType": "$vec4Scalar" },
190
+ "c_3": { "buffer": "storage", "name": "c", "elementType": "$scalar" },
191
+ "params_2": {
192
+ "buffer": "uniform",
193
+ "name": "params",
194
+ "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.c)" }]
195
+ }
196
+ }
197
  }
build/webgpu/metadata.json CHANGED
@@ -1,20 +1,30 @@
1
  {
2
  "name": "ai.onnx.Div",
3
- "id": "_ai_onnx_div_webgpu_af5d5b7",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "T8ahwVvFuvykKXC3r8ulXGaNXUQDbPI1PbKp5rqstSQ=",
11
- "binary-broadcast-vec4.wgsl.jinja": "U3eGI7MMFuwzIsyzIo+9YCReBwWqdMojY4wfH4yOWOU=",
12
- "binary-broadcast.wgsl.jinja": "ZA6kUovcMlxJK9HzmH0+wH0TlN2L0iGq57j5ayF+Sl8=",
13
- "binary-vec4.wgsl.jinja": "d/SCw5hd7ZTCW1+V/zDeZmzTq2Cvy/VGcup4xlfnW+g=",
14
- "manifest.json": "fbw8UcqY7Gfrzy2KqagX1azKQlVfXg5RUivgVT5pxL8=",
15
- "test.json": "nI5UD3dutlCko0fqrCnHpndsor6lw9SEhDbEPGMZst4="
16
  }
17
  },
18
- "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
19
- "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Div" }
 
 
 
 
 
 
 
 
 
 
20
  }
 
1
  {
2
  "name": "ai.onnx.Div",
3
+ "id": "_ai_onnx_div_webgpu_efa12f3",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "P9M5ABblrxZXFebRyjROuEvlMZ9Dj/gfBV536o28y0c=",
11
+ "binary-broadcast-vec4.wgsl.jinja": "D7MgNRBwqYQWAr1FO01kz0tlYlPVD61Eq3ySgWx0wlk=",
12
+ "binary-broadcast.wgsl.jinja": "PzhDFTkzRTci0/u3PfHb1K1UelW/2yZ6eEg5R3WMgQ8=",
13
+ "binary-vec4.wgsl.jinja": "FcjDYUYpxhws/abgE+IVpMJdVpWzaK3y2C7e+npskHA=",
14
+ "manifest.json": "RDnwe6I7AOEJULY/+EcZDTaiTzWh4gmmXPZUezNzOPU=",
15
+ "test.json": "kCokVoivZqeEvXu4trVF4PN6Ev6adrofbxfY3l5Q+NU="
16
  }
17
  },
18
+ "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
19
+ "webgpu": {
20
+ "manifestSpec": "2.0",
21
+ "variants": {
22
+ "same_shape_vec4": ["binary-vec4.wgsl.jinja"],
23
+ "scalar_b_vec4": ["binary-vec4.wgsl.jinja"],
24
+ "scalar_a_vec4": ["binary-vec4.wgsl.jinja"],
25
+ "broadcast_vec4": ["binary-broadcast-vec4.wgsl.jinja"],
26
+ "same_shape_scalar_x4": ["binary-broadcast.wgsl.jinja"],
27
+ "broadcast": ["binary-broadcast.wgsl.jinja"]
28
+ }
29
+ }
30
  }
build/webgpu/test.json CHANGED
@@ -1,5 +1,4 @@
1
  {
2
- "op": "ai.onnx.Div",
3
  "fixtureArrays": {
4
  "onnx_backend_div_input_a": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
5
  },
@@ -62,7 +61,7 @@
62
  "provenance": {
63
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
64
  "test": "MathOpTest.Div",
65
- "notes": "Scalar-path companion: equal subnormal magnitudes divide to ordinary finite values."
66
  },
67
  "inputs": {
68
  "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-38, -1e-38, 4e-39] } },
@@ -458,11 +457,15 @@
458
  {
459
  "name": "f32_div_by_infinity",
460
  "inputs": {
461
- "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, -1.0, 0.0, 100.0] } },
 
 
 
 
462
  "b": {
463
  "dtype": "float32",
464
  "shape": [4],
465
- "data": { "kind": "values", "values": ["Infinity", "Infinity", "Infinity", "-Infinity"] }
466
  }
467
  },
468
  "outputs": {
@@ -470,7 +473,8 @@
470
  "dtype": "float32",
471
  "shape": [4],
472
  "tolerance": 0,
473
- "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] }
 
474
  }
475
  }
476
  },
@@ -504,6 +508,44 @@
504
  }
505
  },
506
  "outputs": { "c": { "dtype": "float16", "shape": [2, 8], "tolerance": 0 } }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
507
  }
508
  ]
509
  }
 
1
  {
 
2
  "fixtureArrays": {
3
  "onnx_backend_div_input_a": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
4
  },
 
61
  "provenance": {
62
  "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
63
  "test": "MathOpTest.Div",
64
+ "notes": "On an unaligned scalar path, equal subnormal magnitudes must divide to ordinary finite values."
65
  },
66
  "inputs": {
67
  "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-38, -1e-38, 4e-39] } },
 
457
  {
458
  "name": "f32_div_by_infinity",
459
  "inputs": {
460
+ "a": {
461
+ "dtype": "float32",
462
+ "shape": [4],
463
+ "data": { "kind": "values", "values": [1.0, "Infinity", -1.0, "-Infinity"] }
464
+ },
465
  "b": {
466
  "dtype": "float32",
467
  "shape": [4],
468
+ "data": { "kind": "values", "values": ["Infinity", "Infinity", "-Infinity", "-Infinity"] }
469
  }
470
  },
471
  "outputs": {
 
473
  "dtype": "float32",
474
  "shape": [4],
475
  "tolerance": 0,
476
+ "allowNaN": true,
477
+ "data": { "kind": "values", "values": [0.0, "NaN", 0.0, "NaN"] }
478
  }
479
  }
480
  },
 
508
  }
509
  },
510
  "outputs": { "c": { "dtype": "float16", "shape": [2, 8], "tolerance": 0 } }
511
+ },
512
+ {
513
+ "name": "scalar_b_vec4_route",
514
+ "provenance": {
515
+ "notes": "Route lock for the one-element-operand vec4 kernel: the other operand matches the output and the vector count is a multiple of four, so the scalar reads once and splats."
516
+ },
517
+ "inputs": {
518
+ "a": {
519
+ "dtype": "float32",
520
+ "shape": [2, 8],
521
+ "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 1.5 }
522
+ },
523
+ "b": {
524
+ "dtype": "float32",
525
+ "shape": [1],
526
+ "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.11, "scale": 1.5 }
527
+ }
528
+ },
529
+ "outputs": { "c": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000001 } }
530
+ },
531
+ {
532
+ "name": "scalar_a_vec4_route",
533
+ "provenance": {
534
+ "notes": "Route lock for the one-element-operand vec4 kernel: the other operand matches the output and the vector count is a multiple of four, so the scalar reads once and splats."
535
+ },
536
+ "inputs": {
537
+ "a": {
538
+ "dtype": "float32",
539
+ "shape": [1],
540
+ "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23, "scale": 1.5 }
541
+ },
542
+ "b": {
543
+ "dtype": "float32",
544
+ "shape": [2, 8],
545
+ "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.11, "scale": 1.5 }
546
+ }
547
+ },
548
+ "outputs": { "c": { "dtype": "float32", "shape": [2, 8], "tolerance": 0.000001 } }
549
  }
550
  ]
551
  }