Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Max

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
743d5e7
·
verified ·
1 Parent(s): c80076e

sync 2e7068faf55e

Browse files
Files changed (8) hide show
  1. README.md +62 -0
  2. build/webgpu/bench.json +31 -0
  3. build/webgpu/datamove-elementwise-copy.wgsl.jinja +20 -0
  4. build/webgpu/manifest.json +477 -0
  5. build/webgpu/metadata.json +20 -0
  6. build/webgpu/minmax-broadcast.wgsl.jinja +122 -0
  7. build/webgpu/minmax-vec4.wgsl.jinja +68 -0
  8. build/webgpu/test.json +998 -0
README.md CHANGED
@@ -1,3 +1,65 @@
1
  ---
 
2
  license: apache-2.0
 
 
 
 
3
  ---
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
  ---
2
+ library_name: kernels
3
  license: apache-2.0
4
+ tags:
5
+ - kernel
6
+ - webgpu
7
+ - wgsl
8
  ---
9
+ # ai.onnx.Max
10
+
11
+ `ai.onnx` · standard ONNX operator · ONNX opset ≥ 13
12
+
13
+ ## Description
14
+
15
+ Computes the elementwise maximum across one or more input tensors with NumPy-style multidirectional broadcasting. All inputs must share the same data type, and the output has the broadcasted shape.
16
+
17
+ See the [ONNX `Max` spec](https://onnx.ai/onnx/operators/onnx__Max.html) for the reference semantics.
18
+
19
+ ## Inputs
20
+
21
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
22
+ | --- | --- | --- | --- | --- | --- | --- |
23
+ | `A` | `a` | `T` | — | — | First input tensor. | required |
24
+ | `B` | `b` | `T` | — | — | Second input tensor, broadcast-compatible with A. | optional |
25
+ | `C` | `c` | `T` | — | — | Third input tensor, broadcast-compatible with A and B. | optional |
26
+ | `D` | `d` | `T` | — | — | Fourth input tensor, broadcast-compatible with all other inputs. | optional |
27
+ | `E` | `e` | `T` | — | — | Fifth input tensor, broadcast-compatible with all other inputs. | optional |
28
+
29
+ ## Outputs
30
+
31
+ | Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
32
+ | --- | --- | --- | --- | --- | --- | --- |
33
+ | `max` | `y` | `T` | derived | derived; see description | Elementwise maximum of all input tensors. | required |
34
+
35
+ ## Type constraints
36
+
37
+ | Variable | Allowed dtypes |
38
+ | --- | --- |
39
+ | `T` | `float32`, `float16`, `int32`, `uint32`, `int16`, `int8`, `uint8` |
40
+
41
+ ## Files
42
+
43
+ - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
44
+ - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
45
+ - [`test.json`](build/webgpu/test.json) — correctness cases
46
+ - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
47
+ - [`datamove-elementwise-copy.wgsl.jinja`](build/webgpu/datamove-elementwise-copy.wgsl.jinja)
48
+ - [`minmax-broadcast.wgsl.jinja`](build/webgpu/minmax-broadcast.wgsl.jinja)
49
+ - [`minmax-vec4.wgsl.jinja`](build/webgpu/minmax-vec4.wgsl.jinja)
50
+
51
+ ## Use with `@huggingface/kernels`
52
+
53
+ The loader derives every required output's shape and logical dtype from the manifest contract and this call.
54
+ It then allocates the result tensors 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
+
58
+ Replace each `*Data` placeholder with a typed array containing the corresponding input data.
59
+
60
+ ```js
61
+ import { getKernel } from "@huggingface/kernels";
62
+
63
+ const kernel = await getKernel("webgpu-kernels/ai.onnx.Max", { version: 1 });
64
+ const { y } = await kernel({ a: { data: aData, shape: [] } });
65
+ ```
build/webgpu/bench.json ADDED
@@ -0,0 +1,31 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Max",
3
+ "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
4
+ "cases": [
5
+ {
6
+ "name": "max-f32-4m",
7
+ "preset": "smoke",
8
+ "vars": { "dtype": "float32", "count": 4194304 },
9
+ "inputs": {
10
+ "a": { "dtype": "float32", "shape": [4194304], "dist": "normal", "seed": 602, "scale": 2 },
11
+ "b": { "dtype": "float32", "shape": [4194304], "dist": "normal", "seed": 603, "scale": 2 }
12
+ },
13
+ "outputs": { "y": { "dtype": "float32", "shape": [4194304] } },
14
+ "bench": {
15
+ "primary": true,
16
+ "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }]
17
+ }
18
+ },
19
+ {
20
+ "name": "max-broadcast-scalar-non-vec4-f32",
21
+ "preset": "edge",
22
+ "vars": { "dtype": "float32", "count": 4194305 },
23
+ "inputs": {
24
+ "a": { "dtype": "float32", "shape": [4194305], "dist": "normal", "seed": 610, "scale": 2 },
25
+ "b": { "dtype": "float32", "shape": [4194305], "dist": "normal", "seed": 611, "scale": 2 }
26
+ },
27
+ "outputs": { "y": { "dtype": "float32", "shape": [4194305] } },
28
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 3" }] }
29
+ }
30
+ ]
31
+ }
build/webgpu/datamove-elementwise-copy.wgsl.jinja ADDED
@@ -0,0 +1,20 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% if usesF16 %}
2
+ enable f16;
3
+ {% endif %}
4
+ {{ env.wgsl.resourceDeclarations }}
5
+
6
+ const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
7
+
8
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
9
+ fn main(
10
+ @builtin(global_invocation_id) gid: vec3<u32>,
11
+ @builtin(num_workgroups) nwg: vec3<u32>
12
+ ) {
13
+ // Grid-stride loop: the dispatch is clamped to the maxComputeWorkgroupsPerDimension
14
+ // workgroups-per-dimension limit, so a thread may copy more than one
15
+ // element for very large tensors.
16
+ let stride = nwg.x * WG;
17
+ for (var i = gid.x; i < params.count; i += stride) {
18
+ y[i] = x[i];
19
+ }
20
+ }
build/webgpu/manifest.json ADDED
@@ -0,0 +1,477 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "domain": "ai.onnx",
3
+ "name": "Max",
4
+ "sinceVersion": 13,
5
+ "description": "Computes the elementwise maximum across one or more input tensors with NumPy-style multidirectional broadcasting. All inputs must share the same data type, and the output has the broadcasted shape.",
6
+ "inputs": [
7
+ { "role": "A", "dtype": "T", "description": "First input tensor." },
8
+ { "role": "B", "dtype": "T", "optional": true, "description": "Second input tensor, broadcast-compatible with A." },
9
+ {
10
+ "role": "C",
11
+ "dtype": "T",
12
+ "optional": true,
13
+ "description": "Third input tensor, broadcast-compatible with A and B."
14
+ },
15
+ {
16
+ "role": "D",
17
+ "dtype": "T",
18
+ "optional": true,
19
+ "description": "Fourth input tensor, broadcast-compatible with all other inputs."
20
+ },
21
+ {
22
+ "role": "E",
23
+ "dtype": "T",
24
+ "optional": true,
25
+ "description": "Fifth input tensor, broadcast-compatible with all other inputs."
26
+ }
27
+ ],
28
+ "outputs": [
29
+ {
30
+ "role": "max",
31
+ "dtype": "T",
32
+ "rank": "max(ranks.A, ranks.B if present.b else 0, ranks.C if present.b and present.c else 0, ranks.D if present.b and present.c and present.d else 0, ranks.E if present.b and present.c and present.d and present.e else 0)",
33
+ "shape": "variadicShape",
34
+ "description": "Elementwise maximum of all input tensors."
35
+ }
36
+ ],
37
+ "typeConstraints": { "T": ["float32", "float16", "int32", "uint32", "int16", "int8", "uint8"] },
38
+ "args": {
39
+ "a": { "kind": "tensor", "semantic": "A", "role": "input" },
40
+ "b": { "kind": "tensor", "semantic": "B", "role": "input", "required": false },
41
+ "c": { "kind": "tensor", "semantic": "C", "role": "input2", "required": false },
42
+ "y": { "kind": "tensor", "semantic": "max", "role": "output" },
43
+ "d": { "kind": "tensor", "semantic": "D", "role": "input3", "required": false },
44
+ "e": { "kind": "tensor", "semantic": "E", "role": "input4", "required": false }
45
+ },
46
+ "tunables": { "WORKGROUP_SIZE": 256 },
47
+ "derive": {
48
+ "variadicInputCount": "1 + (1 if present.b else 0) + (1 if (present.b and present.c) else 0) + (1 if (present.b and present.c and present.d) else 0) + (1 if (present.b and present.c and present.d and present.e) else 0)",
49
+ "variadicShape": "broadcastShape(broadcastShape(broadcastShape(broadcastShape(shapes.A, shapes.B), shapes.C), shapes.D), shapes.E) if present.b and present.c and present.d and present.e else (broadcastShape(broadcastShape(broadcastShape(shapes.A, shapes.B), shapes.C), shapes.D) if present.b and present.c and present.d else (broadcastShape(broadcastShape(shapes.A, shapes.B), shapes.C) if present.b and present.c else (broadcastShape(shapes.A, shapes.B) if present.b else shapes.A)))",
50
+ "flatVec4OutputOk": "numel(shapes.y) > 0 and numel(shapes.y) % 4 == 0 and f16Ok(dtypes.T)",
51
+ "broadcastOutputOk": "f16Ok(dtypes.T)"
52
+ },
53
+ "bindingSets": {
54
+ "identity": [
55
+ { "name": "x", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
56
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
57
+ {
58
+ "name": "params",
59
+ "semantic": "kernel.params",
60
+ "buffer": { "type": "uniform" },
61
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
62
+ }
63
+ ],
64
+ "vec4Two": [
65
+ {
66
+ "name": "a",
67
+ "arg": "a",
68
+ "semantic": "A",
69
+ "buffer": { "type": "read-only-storage" },
70
+ "elementType": "$vectorScalar"
71
+ },
72
+ {
73
+ "name": "b",
74
+ "arg": "b",
75
+ "semantic": "B",
76
+ "buffer": { "type": "read-only-storage" },
77
+ "elementType": "$vectorScalar"
78
+ },
79
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
80
+ {
81
+ "name": "params",
82
+ "semantic": "kernel.params",
83
+ "buffer": { "type": "uniform" },
84
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
85
+ }
86
+ ],
87
+ "vec4Three": [
88
+ {
89
+ "name": "a",
90
+ "arg": "a",
91
+ "semantic": "A",
92
+ "buffer": { "type": "read-only-storage" },
93
+ "elementType": "$vectorScalar"
94
+ },
95
+ {
96
+ "name": "b",
97
+ "arg": "b",
98
+ "semantic": "B",
99
+ "buffer": { "type": "read-only-storage" },
100
+ "elementType": "$vectorScalar"
101
+ },
102
+ {
103
+ "name": "c",
104
+ "arg": "c",
105
+ "semantic": "C",
106
+ "buffer": { "type": "read-only-storage" },
107
+ "elementType": "$vectorScalar"
108
+ },
109
+ { "name": "y", "arg": "y", "semantic": "max", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
110
+ {
111
+ "name": "params",
112
+ "semantic": "kernel.params",
113
+ "buffer": { "type": "uniform" },
114
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
115
+ }
116
+ ],
117
+ "vec4Four": [
118
+ {
119
+ "name": "a",
120
+ "arg": "a",
121
+ "semantic": "A",
122
+ "buffer": { "type": "read-only-storage" },
123
+ "elementType": "$vectorScalar"
124
+ },
125
+ {
126
+ "name": "b",
127
+ "arg": "b",
128
+ "semantic": "B",
129
+ "buffer": { "type": "read-only-storage" },
130
+ "elementType": "$vectorScalar"
131
+ },
132
+ {
133
+ "name": "c",
134
+ "arg": "c",
135
+ "semantic": "C",
136
+ "buffer": { "type": "read-only-storage" },
137
+ "elementType": "$vectorScalar"
138
+ },
139
+ {
140
+ "name": "d",
141
+ "arg": "d",
142
+ "semantic": "D",
143
+ "buffer": { "type": "read-only-storage" },
144
+ "elementType": "$vectorScalar"
145
+ },
146
+ { "name": "y", "arg": "y", "semantic": "max", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
147
+ {
148
+ "name": "params",
149
+ "semantic": "kernel.params",
150
+ "buffer": { "type": "uniform" },
151
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
152
+ }
153
+ ],
154
+ "scalarTwo": [
155
+ { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
156
+ { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
157
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
158
+ {
159
+ "name": "params",
160
+ "semantic": "kernel.params",
161
+ "buffer": { "type": "uniform" },
162
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
163
+ }
164
+ ],
165
+ "scalarThree": [
166
+ { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
167
+ { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
168
+ { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
169
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
170
+ {
171
+ "name": "params",
172
+ "semantic": "kernel.params",
173
+ "buffer": { "type": "uniform" },
174
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
175
+ }
176
+ ],
177
+ "scalarFour": [
178
+ { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
179
+ { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
180
+ { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
181
+ { "name": "d", "arg": "d", "semantic": "D", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
182
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
183
+ {
184
+ "name": "params",
185
+ "semantic": "kernel.params",
186
+ "buffer": { "type": "uniform" },
187
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
188
+ }
189
+ ],
190
+ "vec4Five": [
191
+ {
192
+ "name": "a",
193
+ "arg": "a",
194
+ "semantic": "A",
195
+ "buffer": { "type": "read-only-storage" },
196
+ "elementType": "$vectorScalar"
197
+ },
198
+ {
199
+ "name": "b",
200
+ "arg": "b",
201
+ "semantic": "B",
202
+ "buffer": { "type": "read-only-storage" },
203
+ "elementType": "$vectorScalar"
204
+ },
205
+ {
206
+ "name": "c",
207
+ "arg": "c",
208
+ "semantic": "C",
209
+ "buffer": { "type": "read-only-storage" },
210
+ "elementType": "$vectorScalar"
211
+ },
212
+ {
213
+ "name": "d",
214
+ "arg": "d",
215
+ "semantic": "D",
216
+ "buffer": { "type": "read-only-storage" },
217
+ "elementType": "$vectorScalar"
218
+ },
219
+ {
220
+ "name": "e",
221
+ "arg": "e",
222
+ "semantic": "E",
223
+ "buffer": { "type": "read-only-storage" },
224
+ "elementType": "$vectorScalar"
225
+ },
226
+ { "name": "y", "arg": "y", "semantic": "max", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
227
+ {
228
+ "name": "params",
229
+ "semantic": "kernel.params",
230
+ "buffer": { "type": "uniform" },
231
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
232
+ }
233
+ ],
234
+ "scalarFive": [
235
+ { "name": "a", "arg": "a", "semantic": "A", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
236
+ { "name": "b", "arg": "b", "semantic": "B", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
237
+ { "name": "c", "arg": "c", "semantic": "C", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
238
+ { "name": "d", "arg": "d", "semantic": "D", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
239
+ { "name": "e", "arg": "e", "semantic": "E", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
240
+ { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
241
+ {
242
+ "name": "params",
243
+ "semantic": "kernel.params",
244
+ "buffer": { "type": "uniform" },
245
+ "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
246
+ }
247
+ ]
248
+ },
249
+ "variants": [
250
+ {
251
+ "id": "single_input_identity",
252
+ "priority": 30,
253
+ "when": ["variadicInputCount == 1", "ranks.A == ranks.y", "numel(shapes.A) == numel(shapes.y)", "f16Ok(dtypes.T)"],
254
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
255
+ "passes": [
256
+ {
257
+ "id": "main",
258
+ "name": "Max.Identity",
259
+ "source": { "shader": "datamove-elementwise-copy.wgsl.jinja" },
260
+ "bindings": "identity",
261
+ "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
262
+ }
263
+ ]
264
+ },
265
+ {
266
+ "id": "same_shape_vec4_two_input",
267
+ "priority": 20,
268
+ "when": ["variadicInputCount == 2", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "flatVec4OutputOk"],
269
+ "constants": {
270
+ "scalar": "dtypes.T",
271
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
272
+ "usesF16": "dtypes.T == \"f16\""
273
+ },
274
+ "passes": [
275
+ {
276
+ "id": "main",
277
+ "name": "Max.vec4",
278
+ "source": { "shader": "minmax-vec4.wgsl.jinja", "inputs": { "op": "\"max\"", "hasC": "false" } },
279
+ "bindings": "vec4Two",
280
+ "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
281
+ }
282
+ ]
283
+ },
284
+ {
285
+ "id": "same_shape_vec4_three_input",
286
+ "priority": 25,
287
+ "when": ["variadicInputCount == 3", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "flatVec4OutputOk"],
288
+ "constants": {
289
+ "scalar": "dtypes.T",
290
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
291
+ "usesF16": "dtypes.T == \"f16\""
292
+ },
293
+ "passes": [
294
+ {
295
+ "id": "main",
296
+ "name": "Max.vec4_3",
297
+ "source": { "shader": "minmax-vec4.wgsl.jinja", "inputs": { "op": "\"max\"", "hasC": "true" } },
298
+ "bindings": "vec4Three",
299
+ "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
300
+ }
301
+ ]
302
+ },
303
+ {
304
+ "id": "same_shape_vec4_four_input",
305
+ "priority": 27,
306
+ "when": ["variadicInputCount == 4", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "sameShape(shapes.D, shapes.y)", "flatVec4OutputOk"],
307
+ "constants": {
308
+ "scalar": "dtypes.T",
309
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
310
+ "usesF16": "dtypes.T == \"f16\""
311
+ },
312
+ "passes": [
313
+ {
314
+ "id": "main",
315
+ "name": "Max.vec4_3",
316
+ "source": {
317
+ "shader": "minmax-vec4.wgsl.jinja",
318
+ "inputs": { "op": "\"max\"", "hasC": "true", "hasD": "true" }
319
+ },
320
+ "bindings": "vec4Four",
321
+ "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
322
+ }
323
+ ]
324
+ },
325
+ {
326
+ "id": "same_shape_vec4_five_input",
327
+ "priority": 28,
328
+ "when": ["variadicInputCount == 5", "sameShape(shapes.A, shapes.y)", "sameShape(shapes.B, shapes.y)", "sameShape(shapes.C, shapes.y)", "sameShape(shapes.D, shapes.y)", "sameShape(shapes.E, shapes.y)", "flatVec4OutputOk"],
329
+ "constants": {
330
+ "scalar": "dtypes.T",
331
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
332
+ "usesF16": "dtypes.T == \"f16\""
333
+ },
334
+ "passes": [
335
+ {
336
+ "id": "main",
337
+ "name": "Max.vec4_4",
338
+ "source": {
339
+ "shader": "minmax-vec4.wgsl.jinja",
340
+ "inputs": {
341
+ "op": "\"max\"",
342
+ "hasC": "true",
343
+ "hasD": "true",
344
+ "hasE": "true",
345
+ "extraInputs": "[\"c\", \"d\", \"e\"]"
346
+ }
347
+ },
348
+ "bindings": "vec4Five",
349
+ "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
350
+ }
351
+ ]
352
+ },
353
+ {
354
+ "id": "broadcast_two_input",
355
+ "when": ["variadicInputCount == 2", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "broadcastOutputOk"],
356
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
357
+ "passes": [
358
+ {
359
+ "id": "main",
360
+ "name": "Max",
361
+ "source": {
362
+ "shader": "minmax-broadcast.wgsl.jinja",
363
+ "inputs": {
364
+ "aShape": "shapes.A",
365
+ "bShape": "shapes.B",
366
+ "yShape": "shapes.y",
367
+ "aRank": "ranks.A",
368
+ "bRank": "ranks.B",
369
+ "yRank": "ranks.y",
370
+ "hasC": "false",
371
+ "op": "\"max\""
372
+ }
373
+ },
374
+ "bindings": "scalarTwo",
375
+ "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
376
+ }
377
+ ]
378
+ },
379
+ {
380
+ "id": "broadcast_three_input",
381
+ "priority": 20,
382
+ "when": ["variadicInputCount == 3", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "broadcastOutputOk"],
383
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
384
+ "passes": [
385
+ {
386
+ "id": "main",
387
+ "name": "Max",
388
+ "source": {
389
+ "shader": "minmax-broadcast.wgsl.jinja",
390
+ "inputs": {
391
+ "aShape": "shapes.A",
392
+ "bShape": "shapes.B",
393
+ "cShape": "shapes.C",
394
+ "yShape": "shapes.y",
395
+ "aRank": "ranks.A",
396
+ "bRank": "ranks.B",
397
+ "cRank": "ranks.C",
398
+ "yRank": "ranks.y",
399
+ "hasC": "true",
400
+ "op": "\"max\""
401
+ }
402
+ },
403
+ "bindings": "scalarThree",
404
+ "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
405
+ }
406
+ ]
407
+ },
408
+ {
409
+ "id": "broadcast_four_input",
410
+ "priority": 22,
411
+ "when": ["variadicInputCount == 4", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "ranks.D <= ranks.y", "broadcastOutputOk"],
412
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
413
+ "passes": [
414
+ {
415
+ "id": "main",
416
+ "name": "Max",
417
+ "source": {
418
+ "shader": "minmax-broadcast.wgsl.jinja",
419
+ "inputs": {
420
+ "aShape": "shapes.A",
421
+ "bShape": "shapes.B",
422
+ "cShape": "shapes.C",
423
+ "yShape": "shapes.y",
424
+ "aRank": "ranks.A",
425
+ "bRank": "ranks.B",
426
+ "cRank": "ranks.C",
427
+ "yRank": "ranks.y",
428
+ "hasC": "true",
429
+ "op": "\"max\"",
430
+ "dShape": "shapes.D",
431
+ "dRank": "ranks.D",
432
+ "hasD": "true"
433
+ }
434
+ },
435
+ "bindings": "scalarFour",
436
+ "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
437
+ }
438
+ ]
439
+ },
440
+ {
441
+ "id": "broadcast_five_input",
442
+ "priority": 23,
443
+ "when": ["variadicInputCount == 5", "ranks.A <= ranks.y", "ranks.B <= ranks.y", "ranks.C <= ranks.y", "ranks.D <= ranks.y", "ranks.E <= ranks.y", "broadcastOutputOk"],
444
+ "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
445
+ "passes": [
446
+ {
447
+ "id": "main",
448
+ "name": "Max",
449
+ "source": {
450
+ "shader": "minmax-broadcast.wgsl.jinja",
451
+ "inputs": {
452
+ "aShape": "shapes.A",
453
+ "bShape": "shapes.B",
454
+ "cShape": "shapes.C",
455
+ "yShape": "shapes.y",
456
+ "aRank": "ranks.A",
457
+ "bRank": "ranks.B",
458
+ "cRank": "ranks.C",
459
+ "yRank": "ranks.y",
460
+ "hasC": "true",
461
+ "op": "\"max\"",
462
+ "dShape": "shapes.D",
463
+ "dRank": "ranks.D",
464
+ "hasD": "true",
465
+ "hasE": "true",
466
+ "extraInputs": "[\"c\", \"d\", \"e\"]",
467
+ "eRank": "ranks.E",
468
+ "eShape": "shapes.E"
469
+ }
470
+ },
471
+ "bindings": "scalarFive",
472
+ "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
473
+ }
474
+ ]
475
+ }
476
+ ]
477
+ }
build/webgpu/metadata.json ADDED
@@ -0,0 +1,20 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "name": "ai.onnx.Max",
3
+ "id": "_ai_onnx_max_webgpu_07f8f49",
4
+ "version": 1,
5
+ "license": "Apache-2.0",
6
+ "backend": { "type": "webgpu" },
7
+ "digest": {
8
+ "algorithm": "sha256",
9
+ "files": {
10
+ "bench.json": "RNgdFMBNRNyK6rosPqaYVu6EZMBRczN8MaT1PXgx5qs=",
11
+ "datamove-elementwise-copy.wgsl.jinja": "J5yC2bAddPiP+odXLgVGS3TJ9jeNsfRTvedKrj/fhZg=",
12
+ "manifest.json": "5ogrMJ0Z1/7/SjIm5WqLTPzgTFa4TQ1h0VsARPggnLI=",
13
+ "minmax-broadcast.wgsl.jinja": "PfO5v9vE6fpTpa/V/4d6sfxmLxTnjx65wMuLfCLBxQI=",
14
+ "minmax-vec4.wgsl.jinja": "8EEl0UubHzIfW8gQD2Ra6SzRpRL4uv8ao7U3FgQxm+g=",
15
+ "test.json": "3Te1JTq0Cf0qFrtb5eAvL3A5SvDrmI+XmzWbln4BJ5s="
16
+ }
17
+ },
18
+ "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
19
+ "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Max" }
20
+ }
build/webgpu/minmax-broadcast.wgsl.jinja ADDED
@@ -0,0 +1,122 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% set extraInputs = source.extraInputs if source.extraInputs is defined else (["c"] if source.hasC else []) + (["d"] if source.hasD else []) %}
2
+ {% if usesF16 %}
3
+ enable f16;
4
+ {% endif %}
5
+ {{ env.wgsl.resourceDeclarations }}
6
+ {% if scalar != "i32" and scalar != "u32" %}
7
+
8
+ fn is_nan_f32(value: f32) -> bool {
9
+ let bits = bitcast<u32>(value);
10
+ return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
11
+ }
12
+
13
+ {% endif %}
14
+ {% macro offset_fn(fn_name, opShape, opRank, op_same, op_numel, outShape, outRank, out_numel) %}
15
+ fn {{ fn_name }}({% if out_numel != 0 and op_numel != 1 %}out_index: u32{% endif %}) -> u32 {
16
+ {% if out_numel == 0 %}
17
+ return 0u;
18
+ {% elif op_numel == 1 %}
19
+ return 0u;
20
+ {% elif op_same %}
21
+ return out_index;
22
+ {% else %}
23
+ var offset = 0u;
24
+ {% for axis in range(outRank) %}
25
+ {% set op_axis = axis - (outRank - opRank) %}
26
+ {% if op_axis >= 0 and opShape[op_axis] != 1 %}
27
+ {% set c_stride = namespace(value=1) %}
28
+ {% for j in range(axis + 1, outRank) %}
29
+ {% set c_stride.value = c_stride.value * outShape[j] %}
30
+ {% endfor %}
31
+ {% set op_stride = namespace(value=1) %}
32
+ {% for j in range(op_axis + 1, opRank) %}
33
+ {% set op_stride.value = op_stride.value * opShape[j] %}
34
+ {% endfor %}
35
+ {% if c_stride.value == 1 %}
36
+ let coord{{ axis }} = out_index % {{ outShape[axis] }}u;
37
+ {% else %}
38
+ let coord{{ axis }} = (out_index / {{ c_stride.value }}u) % {{ outShape[axis] }}u;
39
+ {% endif %}
40
+ {% if op_stride.value == 1 %}
41
+ offset = offset + coord{{ axis }};
42
+ {% else %}
43
+ offset = offset + coord{{ axis }} * {{ op_stride.value }}u;
44
+ {% endif %}
45
+ {% endif %}
46
+ {% endfor %}
47
+ return offset;
48
+ {% endif %}
49
+ }
50
+ {%- endmacro %}{% macro broadcast_offset_call(fn_name, opShape, outShape, out_index) %}
51
+ {% set op_numel = namespace(value=1) %}
52
+ {% for d in opShape %}{% set op_numel.value = op_numel.value * d %}{% endfor %}
53
+ {% set out_numel = namespace(value=1) %}
54
+ {% for d in outShape %}{% set out_numel.value = out_numel.value * d %}{% endfor %}
55
+ {{ fn_name }}({% if out_numel.value != 0 and op_numel.value != 1 %}{{ out_index }}{% endif %})
56
+ {%- endmacro %}{% macro broadcast_offset_fn(fn_name, opShape, opRank, outShape, outRank) %}
57
+ {% set op_numel = namespace(value=1) %}
58
+ {% for d in opShape %}
59
+ {% set op_numel.value = op_numel.value * d %}
60
+ {% endfor %}
61
+ {% set out_numel = namespace(value=1) %}
62
+ {% for d in outShape %}
63
+ {% set out_numel.value = out_numel.value * d %}
64
+ {% endfor %}
65
+ {% set op_same = namespace(value=(opRank == outRank)) %}
66
+ {% if op_same.value %}
67
+ {% for axis in range(outRank) %}
68
+ {% if opShape[axis] != outShape[axis] %}
69
+ {% set op_same.value = false %}
70
+ {% endif %}
71
+ {% endfor %}
72
+ {% endif %}
73
+ {{ offset_fn(fn_name, opShape, opRank, op_same.value, op_numel.value, outShape, outRank, out_numel.value) }}
74
+ {%- endmacro %}
75
+
76
+
77
+
78
+ {{ broadcast_offset_fn("a_offset", source.aShape, source.aRank, source.yShape, source.yRank) }}
79
+
80
+ {{ broadcast_offset_fn("b_offset", source.bShape, source.bRank, source.yShape, source.yRank) }}
81
+
82
+ {% for n in extraInputs %}
83
+ {{ broadcast_offset_fn(n ~ "_offset", source[n ~ "Shape"], source[n ~ "Rank"], source.yShape, source.yRank) }}
84
+
85
+ {% endfor %}
86
+
87
+
88
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
89
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
90
+ // 2D-folded flat index: gid.y carries the high bits past the
91
+ // maxComputeWorkgroupsPerDimension limit.
92
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
93
+ if (i >= params.count) {
94
+ return;
95
+ }
96
+ {% if scalar == "i32" or scalar == "u32" %}
97
+ let av = a[{{ broadcast_offset_call("a_offset", source.aShape, source.yShape, "i") }}];
98
+ let bv = b[{{ broadcast_offset_call("b_offset", source.bShape, source.yShape, "i") }}];
99
+ var out = select(bv, av, av > bv);
100
+ {% for n in extraInputs %}
101
+ let {{ n }}v = {{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "i") }}];
102
+ out = select({{ n }}v, out, out > {{ n }}v);
103
+ {% endfor %}
104
+ y[i] = out;
105
+ {% else %}
106
+ let av = f32(a[{{ broadcast_offset_call("a_offset", source.aShape, source.yShape, "i") }}]);
107
+ let bv = f32(b[{{ broadcast_offset_call("b_offset", source.bShape, source.yShape, "i") }}]);
108
+ var out = max(av, bv);
109
+ out = select(out, av, is_nan_f32(av));
110
+ out = select(out, bv, is_nan_f32(bv));
111
+ {% for n in extraInputs %}
112
+ let {{ n }}v = f32({{ n }}[{{ broadcast_offset_call(n ~ "_offset", source[n ~ "Shape"], source.yShape, "i") }}]);
113
+ let m{{ n }} = max(out, {{ n }}v);
114
+ {% if loop.first %}
115
+ // Built-in min/max may drop a NaN operand. Re-inject {{ n }}v's NaN and
116
+ // restore out's own NaN from an earlier input.
117
+ {% endif %}
118
+ out = select(select(m{{ n }}, {{ n }}v, is_nan_f32({{ n }}v)), out, is_nan_f32(out));
119
+ {% endfor %}
120
+ y[i] = {{ scalar }}(out);
121
+ {% endif %}
122
+ }
build/webgpu/minmax-vec4.wgsl.jinja ADDED
@@ -0,0 +1,68 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {% set extraInputs = source.extraInputs if source.extraInputs is defined else (["c"] if source.hasC else []) + (["d"] if source.hasD else []) %}
2
+ {% if usesF16 %}
3
+ enable f16;
4
+ {% endif %}
5
+ {{ env.wgsl.resourceDeclarations }}
6
+ {% if scalar != "i32" and scalar != "u32" %}
7
+
8
+ fn is_nan_f32(value: f32) -> bool {
9
+ let bits = bitcast<u32>(value);
10
+ return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
11
+ }
12
+
13
+
14
+ fn is_nan_vec4(value: vec4<f32>) -> vec4<bool> {
15
+ return vec4<bool>(
16
+ is_nan_f32(value.x),
17
+ is_nan_f32(value.y),
18
+ is_nan_f32(value.z),
19
+ is_nan_f32(value.w)
20
+ );
21
+ }
22
+
23
+ {% endif %}
24
+ @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
25
+ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
26
+ // 2D-folded flat index: gid.y carries the high bits past the
27
+ // maxComputeWorkgroupsPerDimension limit.
28
+ let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
29
+ if (i >= params.count) {
30
+ return;
31
+ }
32
+ let av = a[i];
33
+ let bv = b[i];
34
+ {% if scalar == "i32" or scalar == "u32" %}
35
+ var out = select(bv, av, av > bv);
36
+ {% for n in extraInputs %}
37
+ let {{ n }}v = {{ n }}[i];
38
+ out = select({{ n }}v, out, out > {{ n }}v);
39
+ {% endfor %}
40
+ y[i] = out;
41
+ {% elif scalar == "f16" %}
42
+ let avf = vec4<f32>(av);
43
+ let bvf = vec4<f32>(bv);
44
+ var out = max(avf, bvf);
45
+ out = select(out, avf, is_nan_vec4(avf));
46
+ out = select(out, bvf, is_nan_vec4(bvf));
47
+ {% for n in extraInputs %}
48
+ let {{ n }}vf = vec4<f32>({{ n }}[i]);
49
+ let m{{ n }} = max(out, {{ n }}vf);
50
+ {% if loop.first %}
51
+ // Built-in min/max may drop a NaN operand. Re-inject {{ n }}vf's NaN and
52
+ // restore out's own NaN from an earlier input.
53
+ {% endif %}
54
+ out = select(select(m{{ n }}, {{ n }}vf, is_nan_vec4({{ n }}vf)), out, is_nan_vec4(out));
55
+ {% endfor %}
56
+ y[i] = vec4<f16>(out);
57
+ {% else %}
58
+ var out = max(av, bv);
59
+ out = select(out, av, is_nan_vec4(av));
60
+ out = select(out, bv, is_nan_vec4(bv));
61
+ {% for n in extraInputs %}
62
+ let {{ n }}v = {{ n }}[i];
63
+ let m{{ n }} = max(out, {{ n }}v);
64
+ out = select(select(m{{ n }}, {{ n }}v, is_nan_vec4({{ n }}v)), out, is_nan_vec4(out));
65
+ {% endfor %}
66
+ y[i] = out;
67
+ {% endif %}
68
+ }
build/webgpu/test.json ADDED
@@ -0,0 +1,998 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+ {
2
+ "op": "ai.onnx.Max",
3
+ "cases": [
4
+ {
5
+ "name": "int16_max_arity_boundaries",
6
+ "inputs": {
7
+ "a": {
8
+ "dtype": "int16",
9
+ "shape": [5],
10
+ "data": { "kind": "values", "values": [32767, -32768, -32768, -32768, -32768] }
11
+ },
12
+ "b": {
13
+ "dtype": "int16",
14
+ "shape": [5],
15
+ "data": { "kind": "values", "values": [-32768, 32766, -32768, -32768, -32768] }
16
+ },
17
+ "c": {
18
+ "dtype": "int16",
19
+ "shape": [5],
20
+ "data": { "kind": "values", "values": [-32768, -32768, 1, -32768, -32768] }
21
+ },
22
+ "d": {
23
+ "dtype": "int16",
24
+ "shape": [5],
25
+ "data": { "kind": "values", "values": [-32768, -32768, -32768, 0, -32768] }
26
+ },
27
+ "e": {
28
+ "dtype": "int16",
29
+ "shape": [5],
30
+ "data": { "kind": "values", "values": [-32768, -32768, -32768, -32768, -1] }
31
+ }
32
+ },
33
+ "outputs": {
34
+ "y": {
35
+ "dtype": "int16",
36
+ "shape": [5],
37
+ "tolerance": 0,
38
+ "data": { "kind": "values", "values": [32767, 32766, 1, 0, -1] }
39
+ }
40
+ }
41
+ },
42
+ {
43
+ "name": "max_arity_float16_positions",
44
+ "provenance": {
45
+ "notes": "Synthetic five-input float16 Max contract fixture; each bounded input position uniquely wins one output lane."
46
+ },
47
+ "inputs": {
48
+ "a": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [9.0, 0.0, 0.0, 0.0, 0.0] } },
49
+ "b": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [0.0, 8.0, 0.0, 0.0, 0.0] } },
50
+ "c": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [0.0, 0.0, 7.0, 0.0, 0.0] } },
51
+ "d": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 6.0, 0.0] } },
52
+ "e": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 5.0] } }
53
+ },
54
+ "outputs": {
55
+ "y": {
56
+ "dtype": "float16",
57
+ "shape": [5],
58
+ "tolerance": 0,
59
+ "data": { "kind": "values", "values": [9.0, 8.0, 7.0, 6.0, 5.0] }
60
+ }
61
+ }
62
+ },
63
+ {
64
+ "name": "max_arity_int32_positions",
65
+ "provenance": {
66
+ "notes": "Synthetic five-input int32 Max contract fixture; each bounded input position uniquely wins one output lane."
67
+ },
68
+ "inputs": {
69
+ "a": { "dtype": "int32", "shape": [5], "data": { "kind": "values", "values": [9, 0, 0, 0, 0] } },
70
+ "b": { "dtype": "int32", "shape": [5], "data": { "kind": "values", "values": [0, 8, 0, 0, 0] } },
71
+ "c": { "dtype": "int32", "shape": [5], "data": { "kind": "values", "values": [0, 0, 7, 0, 0] } },
72
+ "d": { "dtype": "int32", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 6, 0] } },
73
+ "e": { "dtype": "int32", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 0, 5] } }
74
+ },
75
+ "outputs": {
76
+ "y": { "dtype": "int32", "shape": [5], "tolerance": 0, "data": { "kind": "values", "values": [9, 8, 7, 6, 5] } }
77
+ }
78
+ },
79
+ {
80
+ "name": "max_arity_int8_positions",
81
+ "provenance": {
82
+ "notes": "Synthetic five-input int8 Max contract fixture; each bounded input position uniquely wins one output lane."
83
+ },
84
+ "inputs": {
85
+ "a": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [9, 0, 0, 0, 0] } },
86
+ "b": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [0, 8, 0, 0, 0] } },
87
+ "c": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [0, 0, 7, 0, 0] } },
88
+ "d": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 6, 0] } },
89
+ "e": { "dtype": "int8", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 0, 5] } }
90
+ },
91
+ "outputs": {
92
+ "y": { "dtype": "int8", "shape": [5], "tolerance": 0, "data": { "kind": "values", "values": [9, 8, 7, 6, 5] } }
93
+ }
94
+ },
95
+ {
96
+ "name": "max_arity_uint32_positions",
97
+ "provenance": {
98
+ "notes": "Synthetic five-input uint32 Max contract fixture; each bounded input position uniquely wins one output lane."
99
+ },
100
+ "inputs": {
101
+ "a": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [9, 0, 0, 0, 0] } },
102
+ "b": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [0, 8, 0, 0, 0] } },
103
+ "c": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [0, 0, 7, 0, 0] } },
104
+ "d": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 6, 0] } },
105
+ "e": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 0, 5] } }
106
+ },
107
+ "outputs": {
108
+ "y": {
109
+ "dtype": "uint32",
110
+ "shape": [5],
111
+ "tolerance": 0,
112
+ "data": { "kind": "values", "values": [9, 8, 7, 6, 5] }
113
+ }
114
+ }
115
+ },
116
+ {
117
+ "name": "max_arity_uint8_positions",
118
+ "provenance": {
119
+ "notes": "Synthetic five-input uint8 Max contract fixture; each bounded input position uniquely wins one output lane."
120
+ },
121
+ "inputs": {
122
+ "a": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [9, 0, 0, 0, 0] } },
123
+ "b": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [0, 8, 0, 0, 0] } },
124
+ "c": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [0, 0, 7, 0, 0] } },
125
+ "d": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 6, 0] } },
126
+ "e": { "dtype": "uint8", "shape": [5], "data": { "kind": "values", "values": [0, 0, 0, 0, 5] } }
127
+ },
128
+ "outputs": {
129
+ "y": { "dtype": "uint8", "shape": [5], "tolerance": 0, "data": { "kind": "values", "values": [9, 8, 7, 6, 5] } }
130
+ }
131
+ },
132
+ {
133
+ "name": "f32_positive_subnormal_max_zero_gpu_gap",
134
+ "skipGpu": {
135
+ "category": "permanent",
136
+ "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes float32 subnormals to zero, so the positive subnormal becomes 0 and max(.,0) returns 0 instead of the subnormal."
137
+ },
138
+ "provenance": {
139
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
140
+ "test": "MathOpTest.Max_6",
141
+ "notes": "Positive subnormal inputs are greater than zero and should be selected by Max."
142
+ },
143
+ "inputs": {
144
+ "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 1e-39] } },
145
+ "b": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 0.0 } }
146
+ },
147
+ "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
148
+ },
149
+ {
150
+ "name": "f32_positive_subnormal_max_zero_scalar_gpu_gap",
151
+ "skipGpu": {
152
+ "category": "permanent",
153
+ "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: Metal flushes float32 subnormals to zero, so the positive subnormal becomes 0 and max(.,0) returns 0 instead of the subnormal."
154
+ },
155
+ "provenance": {
156
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
157
+ "test": "MathOpTest.Max_6",
158
+ "notes": "Scalar broadcast companion: positive subnormal inputs remain greater than zero."
159
+ },
160
+ "inputs": {
161
+ "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } },
162
+ "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
163
+ },
164
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
165
+ },
166
+ {
167
+ "name": "same_shape",
168
+ "inputs": {
169
+ "a": {
170
+ "dtype": "float32",
171
+ "shape": [2, 3],
172
+ "data": { "kind": "values", "values": [1.0, 5.0, -2.0, 4.0, 0.0, 6.0] }
173
+ },
174
+ "b": {
175
+ "dtype": "float32",
176
+ "shape": [2, 3],
177
+ "data": { "kind": "values", "values": [3.0, 2.0, -4.0, 8.0, 1.0, 1.0] }
178
+ }
179
+ },
180
+ "outputs": { "y": { "dtype": "float32", "shape": [2, 3] } }
181
+ },
182
+ {
183
+ "name": "row_broadcast_f16",
184
+ "inputs": {
185
+ "a": { "dtype": "float16", "shape": [2, 4] },
186
+ "b": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [0.1, 0.2, 0.3, 0.4] } }
187
+ },
188
+ "outputs": { "y": { "dtype": "float16", "shape": [2, 4] } },
189
+ "tolerance": 0.001
190
+ },
191
+ {
192
+ "name": "true_scalar_rhs_broadcast",
193
+ "inputs": {
194
+ "a": {
195
+ "dtype": "float32",
196
+ "shape": [2, 3],
197
+ "data": { "kind": "values", "values": [-3.0, -1.0, 0.0, 1.0, 2.0, 3.0] }
198
+ },
199
+ "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.5] } }
200
+ },
201
+ "outputs": { "y": { "dtype": "float32", "shape": [2, 3], "tolerance": 0.000001 } }
202
+ },
203
+ {
204
+ "name": "nan_propagates",
205
+ "inputs": {
206
+ "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": ["NaN", 1.0, "NaN", 2.0] } },
207
+ "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [3.0, "NaN", "NaN", 1.0] } }
208
+ },
209
+ "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001, "allowNaN": true } }
210
+ },
211
+ {
212
+ "name": "rank0_scalar_scalar_output",
213
+ "inputs": {
214
+ "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-7.0] } },
215
+ "b": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-3.0] } }
216
+ },
217
+ "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
218
+ },
219
+ {
220
+ "name": "int32_exact_above_float24",
221
+ "inputs": {
222
+ "a": {
223
+ "dtype": "int32",
224
+ "shape": [4],
225
+ "data": { "kind": "values", "values": [16777216, 16777217, -16777217, -16777216] }
226
+ },
227
+ "b": {
228
+ "dtype": "int32",
229
+ "shape": [4],
230
+ "data": { "kind": "values", "values": [16777217, 16777216, -16777216, -16777217] }
231
+ }
232
+ },
233
+ "outputs": { "y": { "dtype": "int32", "shape": [4], "tolerance": 0 } }
234
+ },
235
+ {
236
+ "name": "same_shape_vec4_three_input",
237
+ "inputs": {
238
+ "a": {
239
+ "dtype": "float32",
240
+ "shape": [8],
241
+ "data": { "kind": "values", "values": [1.0, 5.0, -2.0, 4.0, 0.0, 6.0, 10.0, -10.0] }
242
+ },
243
+ "b": {
244
+ "dtype": "float32",
245
+ "shape": [8],
246
+ "data": { "kind": "values", "values": [3.0, 2.0, -4.0, 8.0, 1.0, 1.0, 9.0, -9.0] }
247
+ },
248
+ "c": {
249
+ "dtype": "float32",
250
+ "shape": [8],
251
+ "data": { "kind": "values", "values": [0.0, 7.0, -3.0, 2.0, -1.0, 8.0, 11.0, -11.0] }
252
+ }
253
+ },
254
+ "outputs": {
255
+ "y": {
256
+ "dtype": "float32",
257
+ "shape": [8],
258
+ "tolerance": 0.000001,
259
+ "data": { "kind": "values", "values": [3.0, 7.0, -2.0, 8.0, 1.0, 8.0, 11.0, -9.0] }
260
+ }
261
+ }
262
+ },
263
+ {
264
+ "name": "ort_float_three_input_same_shape",
265
+ "provenance": {
266
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
267
+ "test": "MathOpTest.Max_6"
268
+ },
269
+ "inputs": {
270
+ "a": {
271
+ "dtype": "float32",
272
+ "shape": [3, 3],
273
+ "data": { "kind": "values", "values": [1.0, 0.0, 1.0, -1.0, 1.1, -100.0, -5.4, 0.01, -10000.0] }
274
+ },
275
+ "b": {
276
+ "dtype": "float32",
277
+ "shape": [3, 3],
278
+ "data": { "kind": "values", "values": [1.0, 0.0, 2.0, -2.0, 2.2, 64.0, -1.0, 0.02, 0.1] }
279
+ },
280
+ "c": {
281
+ "dtype": "float32",
282
+ "shape": [3, 3],
283
+ "data": { "kind": "values", "values": [1.0, 0.0, 3.0, -3.0, 3.3, 64.0, 5.4, 0.03, 10000.0] }
284
+ }
285
+ },
286
+ "outputs": {
287
+ "y": {
288
+ "dtype": "float32",
289
+ "shape": [3, 3],
290
+ "tolerance": 0.000001,
291
+ "data": { "kind": "values", "values": [1.0, 0.0, 3.0, -1.0, 3.3, 64.0, 5.4, 0.03, 10000.0] }
292
+ }
293
+ }
294
+ },
295
+ {
296
+ "name": "ort_validated_four_inputs_same_shape_variadic",
297
+ "provenance": {
298
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
299
+ "test": "MathOpTest.Max_6",
300
+ "notes": "Extends ORT's same-shape Max coverage to a valid four-input ONNX variadic node."
301
+ },
302
+ "inputs": {
303
+ "a": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": 1.0 } },
304
+ "b": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": 3.0 } },
305
+ "c": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": 5.0 } },
306
+ "d": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "constant", "value": 7.0 } }
307
+ },
308
+ "outputs": {
309
+ "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0, "data": { "kind": "constant", "value": 7.0 } }
310
+ }
311
+ },
312
+ {
313
+ "name": "four_input_same_shape_vec4_reference_generated",
314
+ "provenance": {
315
+ "notes": "Reference-generated four-operand coverage. Every other four-input Max case pins its outputs, so prepareCase never consulted the reference and it went unnoticed that the reference reduced only A..C. The four operands are interleaved sinusoids of equal amplitude, so D supplies the maximum on a substantial share of lanes and dropping it changes the answer far above tolerance."
316
+ },
317
+ "inputs": {
318
+ "a": {
319
+ "dtype": "float32",
320
+ "shape": [2, 3, 2, 8],
321
+ "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23, "scale": 0.25 }
322
+ },
323
+ "b": {
324
+ "dtype": "float32",
325
+ "shape": [2, 3, 2, 8],
326
+ "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.07, "scale": 0.25 }
327
+ },
328
+ "c": {
329
+ "dtype": "float32",
330
+ "shape": [2, 3, 2, 8],
331
+ "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.13, "scale": 0.25 }
332
+ },
333
+ "d": {
334
+ "dtype": "float32",
335
+ "shape": [2, 3, 2, 8],
336
+ "data": { "kind": "fillFloat32", "sinStep": 0.05, "cosStep": 0.37, "scale": 0.25 }
337
+ }
338
+ },
339
+ "outputs": { "y": { "dtype": "float32", "shape": [2, 3, 2, 8], "tolerance": 0.000001, "relTolerance": 0.000001 } }
340
+ },
341
+ {
342
+ "name": "ort_four_inputs_nan_propagates_variadic",
343
+ "provenance": {
344
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
345
+ "test": "MathOpTest.Max_6",
346
+ "notes": "Extends ORT's variadic Max coverage with a fourth input that carries NaNs; any NaN input should propagate at that element."
347
+ },
348
+ "inputs": {
349
+ "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, "NaN", -5.0, 4.0] } },
350
+ "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [2.0, 3.0, "NaN", 1.0] } },
351
+ "c": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 5.0, -2.0, 8.0] } },
352
+ "d": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [3.0, 4.0, 6.0, "NaN"] } }
353
+ },
354
+ "outputs": {
355
+ "y": {
356
+ "dtype": "float32",
357
+ "shape": [4],
358
+ "tolerance": 0,
359
+ "allowNaN": true,
360
+ "data": { "kind": "values", "values": [3.0, "NaN", "NaN", "NaN"] }
361
+ }
362
+ }
363
+ },
364
+ {
365
+ "name": "ort_float_three_input_broadcast",
366
+ "provenance": {
367
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
368
+ "test": "MathOpTest.Max_12_Float"
369
+ },
370
+ "inputs": {
371
+ "a": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
372
+ "b": {
373
+ "dtype": "float32",
374
+ "shape": [3, 3],
375
+ "data": { "kind": "values", "values": [10.0, 20.0, 30.0, 40.0, 50.0, 60.0, 70.0, 80.0, 90.0] }
376
+ },
377
+ "c": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [-1.0, -2.0, 300.0] } }
378
+ },
379
+ "outputs": {
380
+ "y": {
381
+ "dtype": "float32",
382
+ "shape": [3, 3],
383
+ "tolerance": 0.000001,
384
+ "data": { "kind": "values", "values": [10.0, 20.0, 30.0, 40.0, 50.0, 60.0, 300.0, 300.0, 300.0] }
385
+ }
386
+ }
387
+ },
388
+ {
389
+ "name": "ort_float_four_input_broadcast",
390
+ "provenance": {
391
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
392
+ "test": "MathOpTest.Max_12_Float",
393
+ "notes": "Extends ORT's multidirectional broadcast case to a valid four-input ONNX variadic Max node."
394
+ },
395
+ "inputs": {
396
+ "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
397
+ "b": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [10.0, 20.0, 30.0] } },
398
+ "c": { "dtype": "float32", "shape": [3, 1, 1], "data": { "kind": "values", "values": [100.0, 200.0, 300.0] } },
399
+ "d": {
400
+ "dtype": "float32",
401
+ "shape": [1, 1, 3],
402
+ "data": { "kind": "values", "values": [1000.0, 2000.0, 3000.0] }
403
+ }
404
+ },
405
+ "outputs": {
406
+ "y": {
407
+ "dtype": "float32",
408
+ "shape": [3, 3, 3],
409
+ "tolerance": 0,
410
+ "data": {
411
+ "kind": "values",
412
+ "values": [1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0, 1000.0, 2000.0, 3000.0]
413
+ }
414
+ }
415
+ }
416
+ },
417
+ {
418
+ "name": "ort_float_nan_broadcast",
419
+ "provenance": {
420
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
421
+ "test": "MathOpTest.Max_12_Float_Nan"
422
+ },
423
+ "inputs": {
424
+ "a": {
425
+ "dtype": "float32",
426
+ "shape": [3, 3],
427
+ "data": { "kind": "values", "values": ["NaN", "NaN", "NaN", -0.5, 0.0, -2.0, 0.5, 0.0, 2.0] }
428
+ },
429
+ "b": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [0.0, -1.0, 1.0] } }
430
+ },
431
+ "outputs": {
432
+ "y": {
433
+ "dtype": "float32",
434
+ "shape": [3, 3],
435
+ "tolerance": 0.000001,
436
+ "allowNaN": true,
437
+ "data": { "kind": "values", "values": ["NaN", "NaN", "NaN", -0.5, 0.0, -1.0, 1.0, 1.0, 2.0] }
438
+ }
439
+ }
440
+ },
441
+ {
442
+ "name": "ort_float_2input_broadcast",
443
+ "provenance": {
444
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
445
+ "test": "MathOpTest.Max_8_2inputbroadcast"
446
+ },
447
+ "inputs": {
448
+ "a": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
449
+ "b": {
450
+ "dtype": "float32",
451
+ "shape": [3, 3],
452
+ "data": { "kind": "values", "values": [10.0, 20.0, 30.0, 40.0, 50.0, 60.0, 70.0, 80.0, 90.0] }
453
+ }
454
+ },
455
+ "outputs": { "y": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
456
+ },
457
+ {
458
+ "name": "ort_float_nan_with_scalar",
459
+ "provenance": {
460
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
461
+ "test": "MathOpTest.Max_12_Float_Nan_with_scalar"
462
+ },
463
+ "inputs": {
464
+ "a": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": ["NaN", -0.5, 0.5] } },
465
+ "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
466
+ },
467
+ "outputs": { "y": { "dtype": "float32", "shape": [3, 1], "tolerance": 0.000001, "allowNaN": true } }
468
+ },
469
+ {
470
+ "name": "ort_float_scalar_nan_broadcast",
471
+ "provenance": {
472
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
473
+ "test": "MathOpTest.Max_12_Float_with_scalar_Nan"
474
+ },
475
+ "inputs": {
476
+ "a": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.25, -0.25, -0.5, 0.5] } },
477
+ "b": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": ["NaN"] } }
478
+ },
479
+ "outputs": { "y": { "dtype": "float32", "shape": [2, 2], "tolerance": 0.000001, "allowNaN": true } }
480
+ },
481
+ {
482
+ "name": "ort_f16_matrix_vector",
483
+ "provenance": {
484
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
485
+ "test": "MathOpTest.Max_13_Float16_MatrixVector"
486
+ },
487
+ "inputs": {
488
+ "a": {
489
+ "dtype": "float16",
490
+ "shape": [4, 3],
491
+ "data": { "kind": "values", "values": [1.0, 1.0, 1.0, -0.5, 0.0, -2.0, 0.0, 0.5, 0.75, 0.5, 0.0, 2.0] }
492
+ },
493
+ "b": { "dtype": "float16", "shape": [4, 1], "data": { "kind": "values", "values": [0.0, -1.0, 0.5, 1.0] } }
494
+ },
495
+ "outputs": { "y": { "dtype": "float16", "shape": [4, 3], "tolerance": 0.002 } }
496
+ },
497
+ {
498
+ "name": "ort_f16_vector_matrix",
499
+ "provenance": {
500
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
501
+ "test": "MathOpTest.Max_13_Float16_VectorMatrix"
502
+ },
503
+ "inputs": {
504
+ "a": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "values", "values": [0.0, -1.0, 1.0] } },
505
+ "b": {
506
+ "dtype": "float16",
507
+ "shape": [3, 3],
508
+ "data": { "kind": "values", "values": [1.0, 1.0, 1.0, -0.5, 0.0, -2.0, 0.5, 0.0, 2.0] }
509
+ }
510
+ },
511
+ "outputs": { "y": { "dtype": "float16", "shape": [3, 3], "tolerance": 0.002 } }
512
+ },
513
+ {
514
+ "name": "ort_f16_nan_pair",
515
+ "provenance": {
516
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
517
+ "test": "MathOpTest.Max_13_Float16_Nan"
518
+ },
519
+ "inputs": {
520
+ "a": { "dtype": "float16", "shape": [4, 1], "data": { "kind": "values", "values": [-1.0, "NaN", 1.0, 0.5] } },
521
+ "b": { "dtype": "float16", "shape": [4, 1], "data": { "kind": "values", "values": [0.5, 1.0, 0.25, "NaN"] } }
522
+ },
523
+ "outputs": { "y": { "dtype": "float16", "shape": [4, 1], "tolerance": 0.002, "allowNaN": true } }
524
+ },
525
+ {
526
+ "name": "ort_f16_nan_with_scalar",
527
+ "provenance": {
528
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
529
+ "test": "MathOpTest.Max_13_Float16_Nan_with_scalar"
530
+ },
531
+ "inputs": {
532
+ "a": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "values", "values": [-1.0, "NaN", 1.0] } },
533
+ "b": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.25] } }
534
+ },
535
+ "outputs": { "y": { "dtype": "float16", "shape": [3, 1], "tolerance": 0.002, "allowNaN": true } }
536
+ },
537
+ {
538
+ "name": "ort_f16_scalar_nan_broadcast",
539
+ "provenance": {
540
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
541
+ "test": "MathOpTest.Max_13_Float16_with_scalar_Nan"
542
+ },
543
+ "inputs": {
544
+ "a": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "values", "values": [-0.5, 1.0, 1.5] } },
545
+ "b": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": ["NaN"] } }
546
+ },
547
+ "outputs": { "y": { "dtype": "float16", "shape": [3, 1], "tolerance": 0.002, "allowNaN": true } }
548
+ },
549
+ {
550
+ "name": "ort_f16_three_input",
551
+ "provenance": {
552
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
553
+ "test": "MathOpTest.Max_12_MLFloat16"
554
+ },
555
+ "inputs": {
556
+ "a": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-1.0, -1.0, -1.0] } },
557
+ "b": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-2.0, -1.0, -2.0] } },
558
+ "c": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-3.0, -2.0, -3.0] } }
559
+ },
560
+ "outputs": {
561
+ "y": {
562
+ "dtype": "float16",
563
+ "shape": [1, 3],
564
+ "tolerance": 0.002,
565
+ "data": { "kind": "values", "values": [-1.0, -1.0, -1.0] }
566
+ }
567
+ }
568
+ },
569
+ {
570
+ "name": "ort_f16_scalar0_three_input",
571
+ "provenance": {
572
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
573
+ "test": "MathOpTest.Max_12_MLFloat16_Scalar0"
574
+ },
575
+ "inputs": {
576
+ "a": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-1.0] } },
577
+ "b": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-11.0, -12.0, -22.0] } },
578
+ "c": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-10.0, -11.0, -13.0] } }
579
+ },
580
+ "outputs": {
581
+ "y": {
582
+ "dtype": "float16",
583
+ "shape": [1, 3],
584
+ "tolerance": 0.002,
585
+ "data": { "kind": "values", "values": [-1.0, -1.0, -1.0] }
586
+ }
587
+ }
588
+ },
589
+ {
590
+ "name": "ort_f16_scalar1_three_input",
591
+ "provenance": {
592
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
593
+ "test": "MathOpTest.Max_12_MLFloat16_Scalar1"
594
+ },
595
+ "inputs": {
596
+ "a": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-1.0, -2.0, -3.0] } },
597
+ "b": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [2.0] } },
598
+ "c": { "dtype": "float16", "shape": [1, 3], "data": { "kind": "values", "values": [-2.0, -3.0, -4.0] } }
599
+ },
600
+ "outputs": {
601
+ "y": {
602
+ "dtype": "float16",
603
+ "shape": [1, 3],
604
+ "tolerance": 0.002,
605
+ "data": { "kind": "values", "values": [2.0, 2.0, 2.0] }
606
+ }
607
+ }
608
+ },
609
+ {
610
+ "name": "ort_int8_pair_broadcast",
611
+ "provenance": {
612
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
613
+ "test": "MathOpTest.Max_12_Int8",
614
+ "notes": "Two-input projection of ORT's broadcast case for logical int8 storage."
615
+ },
616
+ "inputs": {
617
+ "a": {
618
+ "dtype": "int8",
619
+ "shape": [3, 3],
620
+ "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90] }
621
+ },
622
+ "b": { "dtype": "int8", "shape": [3, 1], "data": { "kind": "values", "values": [-1, -2, 127] } }
623
+ },
624
+ "outputs": { "y": { "dtype": "int8", "shape": [3, 3], "tolerance": 0 } }
625
+ },
626
+ {
627
+ "name": "ort_uint8_pair_broadcast",
628
+ "provenance": {
629
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
630
+ "test": "MathOpTest.Max_12_UInt8",
631
+ "notes": "Two-input projection of ORT's broadcast case for logical uint8 storage."
632
+ },
633
+ "inputs": {
634
+ "a": {
635
+ "dtype": "uint8",
636
+ "shape": [3, 3],
637
+ "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90] }
638
+ },
639
+ "b": { "dtype": "uint8", "shape": [3, 1], "data": { "kind": "values", "values": [100, 20, 30] } }
640
+ },
641
+ "outputs": { "y": { "dtype": "uint8", "shape": [3, 3], "tolerance": 0 } }
642
+ },
643
+ {
644
+ "name": "ort_int8_three_input_broadcast",
645
+ "provenance": {
646
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
647
+ "test": "MathOpTest.Max_12_Int8"
648
+ },
649
+ "inputs": {
650
+ "a": { "dtype": "int8", "shape": [1, 3], "data": { "kind": "values", "values": [1, 2, 3] } },
651
+ "b": {
652
+ "dtype": "int8",
653
+ "shape": [3, 3],
654
+ "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90] }
655
+ },
656
+ "c": { "dtype": "int8", "shape": [3, 1], "data": { "kind": "values", "values": [-1, -2, 127] } }
657
+ },
658
+ "outputs": { "y": { "dtype": "int8", "shape": [3, 3], "tolerance": 0 } }
659
+ },
660
+ {
661
+ "name": "ort_uint8_three_input_broadcast",
662
+ "provenance": {
663
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
664
+ "test": "MathOpTest.Max_12_UInt8"
665
+ },
666
+ "inputs": {
667
+ "a": { "dtype": "uint8", "shape": [1, 3], "data": { "kind": "values", "values": [1, 20, 30] } },
668
+ "b": {
669
+ "dtype": "uint8",
670
+ "shape": [3, 3],
671
+ "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90] }
672
+ },
673
+ "c": { "dtype": "uint8", "shape": [3, 1], "data": { "kind": "values", "values": [100, 20, 30] } }
674
+ },
675
+ "outputs": { "y": { "dtype": "uint8", "shape": [3, 3], "tolerance": 0 } }
676
+ },
677
+ {
678
+ "name": "ort_int32_three_input_broadcast",
679
+ "provenance": {
680
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
681
+ "test": "MathOpTest.Max_12_Int32"
682
+ },
683
+ "inputs": {
684
+ "a": { "dtype": "int32", "shape": [1, 3], "data": { "kind": "values", "values": [1, 2, 3] } },
685
+ "b": {
686
+ "dtype": "int32",
687
+ "shape": [3, 3],
688
+ "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90] }
689
+ },
690
+ "c": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [-1, -2, 300] } }
691
+ },
692
+ "outputs": { "y": { "dtype": "int32", "shape": [3, 3], "tolerance": 0 } }
693
+ },
694
+ {
695
+ "name": "ort_uint32_three_input_broadcast",
696
+ "provenance": {
697
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
698
+ "test": "MathOpTest.Max_12_UInt32"
699
+ },
700
+ "inputs": {
701
+ "a": { "dtype": "uint32", "shape": [1, 3], "data": { "kind": "values", "values": [1, 2, 3] } },
702
+ "b": {
703
+ "dtype": "uint32",
704
+ "shape": [3, 3],
705
+ "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60, 70, 80, 90] }
706
+ },
707
+ "c": { "dtype": "uint32", "shape": [3, 1], "data": { "kind": "values", "values": [1, 2, 300] } }
708
+ },
709
+ "outputs": { "y": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
710
+ },
711
+ {
712
+ "name": "onnx_backend_max_float16",
713
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_float16" },
714
+ "inputs": {
715
+ "a": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } },
716
+ "b": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.0, 4.0, 4.0] } }
717
+ },
718
+ "outputs": { "y": { "dtype": "float16", "shape": [3], "tolerance": 0 } }
719
+ },
720
+ {
721
+ "name": "onnx_backend_max_int32",
722
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_int32" },
723
+ "inputs": {
724
+ "a": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [3, 2, 1] } },
725
+ "b": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 4, 4] } }
726
+ },
727
+ "outputs": { "y": { "dtype": "int32", "shape": [3], "tolerance": 0 } }
728
+ },
729
+ {
730
+ "name": "onnx_backend_max_int8",
731
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_int8" },
732
+ "inputs": {
733
+ "a": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [3, 2, 1] } },
734
+ "b": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [1, 4, 4] } }
735
+ },
736
+ "outputs": { "y": { "dtype": "int8", "shape": [3], "tolerance": 0 } }
737
+ },
738
+ {
739
+ "name": "onnx_backend_max_two_inputs",
740
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_two_inputs" },
741
+ "inputs": {
742
+ "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } },
743
+ "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 4.0, 4.0] } }
744
+ },
745
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
746
+ },
747
+ {
748
+ "name": "onnx_backend_max_uint32",
749
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_uint32" },
750
+ "inputs": {
751
+ "a": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [3, 2, 1] } },
752
+ "b": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 4, 4] } }
753
+ },
754
+ "outputs": { "y": { "dtype": "uint32", "shape": [3], "tolerance": 0 } }
755
+ },
756
+ {
757
+ "name": "onnx_backend_max_uint8",
758
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_uint8" },
759
+ "inputs": {
760
+ "a": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [3, 2, 1] } },
761
+ "b": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [1, 4, 4] } }
762
+ },
763
+ "outputs": { "y": { "dtype": "uint8", "shape": [3], "tolerance": 0 } }
764
+ },
765
+ {
766
+ "name": "onnx_backend_max_one_input_identity",
767
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_one_input" },
768
+ "inputs": { "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } } },
769
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
770
+ },
771
+ {
772
+ "name": "onnx_backend_max_example_three_inputs",
773
+ "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_max_example" },
774
+ "inputs": {
775
+ "a": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [3.0, 2.0, 1.0] } },
776
+ "b": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 4.0, 4.0] } },
777
+ "c": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [2.0, 5.0, 3.0] } }
778
+ },
779
+ "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
780
+ },
781
+ {
782
+ "name": "ort_dim_zero_equal_rank",
783
+ "provenance": {
784
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
785
+ "test": "MathOpTest.DimWithZeroHandling",
786
+ "notes": "Projected from ORT's binary elementwise zero-dimension Add coverage to generic ONNX multidirectional broadcasting."
787
+ },
788
+ "inputs": {
789
+ "a": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
790
+ "b": { "dtype": "float32", "shape": [3, 0], "data": { "kind": "values", "values": [] } }
791
+ },
792
+ "outputs": { "y": { "dtype": "float32", "shape": [3, 0], "tolerance": 0 } }
793
+ },
794
+ {
795
+ "name": "ort_dim_zero_scalar_broadcast",
796
+ "provenance": {
797
+ "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
798
+ "test": "MathOpTest.DimWithZeroHandling",
799
+ "notes": "Projected from ORT's binary elementwise zero-dimension Add coverage to generic ONNX multidirectional broadcasting."
800
+ },
801
+ "inputs": {
802
+ "a": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } },
803
+ "b": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
804
+ },
805
+ "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
806
+ },
807
+ {
808
+ "name": "single_input_identity_ignores_d",
809
+ "inputs": {
810
+ "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, -5.0, 3.0, 0.0] } },
811
+ "d": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [2.0, -3.0, 0.5, 7.0] } }
812
+ },
813
+ "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
814
+ },
815
+ {
816
+ "name": "same_shape_vec4_ignores_d_when_c_absent",
817
+ "inputs": {
818
+ "a": {
819
+ "dtype": "float32",
820
+ "shape": [8],
821
+ "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
822
+ },
823
+ "b": {
824
+ "dtype": "float32",
825
+ "shape": [8],
826
+ "data": { "kind": "values", "values": [0.5, 1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5] }
827
+ },
828
+ "d": {
829
+ "dtype": "float32",
830
+ "shape": [8],
831
+ "data": { "kind": "values", "values": [10.0, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 100.0] }
832
+ }
833
+ },
834
+ "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0 } }
835
+ },
836
+ {
837
+ "name": "f32_infinity_pair",
838
+ "inputs": {
839
+ "a": {
840
+ "dtype": "float32",
841
+ "shape": [6],
842
+ "data": { "kind": "values", "values": ["Infinity", "-Infinity", "Infinity", "-Infinity", 1.0, 2.0] }
843
+ },
844
+ "b": {
845
+ "dtype": "float32",
846
+ "shape": [6],
847
+ "data": { "kind": "values", "values": ["-Infinity", 0.0, 1.0, "Infinity", "Infinity", "-Infinity"] }
848
+ }
849
+ },
850
+ "outputs": {
851
+ "y": {
852
+ "dtype": "float32",
853
+ "shape": [6],
854
+ "tolerance": 0,
855
+ "data": { "kind": "values", "values": ["Infinity", 0.0, "Infinity", "Infinity", "Infinity", 2.0] }
856
+ }
857
+ }
858
+ },
859
+ {
860
+ "name": "int8_extreme_values_broadcast",
861
+ "inputs": {
862
+ "a": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, 127, -128, 0] } },
863
+ "b": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [127, -128, 0, -1] } }
864
+ },
865
+ "outputs": {
866
+ "y": { "dtype": "int8", "shape": [4], "tolerance": 0, "data": { "kind": "values", "values": [127, 127, 0, 0] } }
867
+ }
868
+ },
869
+ {
870
+ "name": "five_input_same_shape_variadic",
871
+ "inputs": {
872
+ "a": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 9.0, -3.0, 4.0] } },
873
+ "b": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [2.0, 8.0, -4.0, 3.0] } },
874
+ "c": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [3.0, 7.0, -5.0, 2.0] } },
875
+ "d": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [4.0, 6.0, -6.0, 1.0] } },
876
+ "e": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [5.0, 5.0, -7.0, 0.0] } }
877
+ },
878
+ "outputs": {
879
+ "y": {
880
+ "dtype": "float32",
881
+ "shape": [4],
882
+ "tolerance": 0,
883
+ "data": { "kind": "values", "values": [5.0, 9.0, -3.0, 4.0] }
884
+ }
885
+ }
886
+ },
887
+ {
888
+ "name": "rank8_broadcast_two_input",
889
+ "inputs": {
890
+ "a": {
891
+ "dtype": "float32",
892
+ "shape": [1, 2, 1, 2, 1, 2, 2, 3],
893
+ "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
894
+ },
895
+ "b": {
896
+ "dtype": "float32",
897
+ "shape": [2, 1, 2, 1, 2, 1, 2, 3],
898
+ "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 }
899
+ }
900
+ },
901
+ "outputs": { "y": { "dtype": "float32", "shape": [2, 2, 2, 2, 2, 2, 2, 3], "tolerance": 0.000001 } }
902
+ },
903
+ {
904
+ "name": "int32_three_input_same_shape_vec4",
905
+ "provenance": {
906
+ "notes": "int32 at arity 3 exercises the flat vec4 same-shape path and its extra-operand fold. Each operand supplies the maximum on at least two lanes, so dropping or double-counting any operand changes the answer. Valid ONNX Max-13: the op is variadic over any numeric T."
907
+ },
908
+ "inputs": {
909
+ "a": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [5, -3, 12, 0, 7, -20, 33, 4] } },
910
+ "b": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [9, -7, 2, 15, -1, -25, 11, 40] } },
911
+ "c": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [1, -1, 6, 8, 21, -30, 5, 12] } }
912
+ },
913
+ "outputs": { "y": { "dtype": "int32", "shape": [8], "tolerance": 0 } }
914
+ },
915
+ {
916
+ "name": "uint32_four_input_same_shape_vec4",
917
+ "provenance": {
918
+ "notes": "uint32 at arity 4 on the flat vec4 same-shape path: the u32 branch of the same integer lane of minmax-vec4, with TWO extra-operand folds so the second fold reads the value the first produced. Lanes 0-3 each put one operand above 2^31 against small peers, so an accidental signed compare would pick the wrong operand on every one of them; lanes 4-7 spread the win evenly over a, b, c and d. Valid ONNX Max-13: variadic over numeric T with all four operands the same shape."
919
+ },
920
+ "inputs": {
921
+ "a": {
922
+ "dtype": "uint32",
923
+ "shape": [8],
924
+ "data": { "kind": "values", "values": [4000000000, 10, 15, 40, 1500, 40, 50, 60] }
925
+ },
926
+ "b": {
927
+ "dtype": "uint32",
928
+ "shape": [8],
929
+ "data": { "kind": "values", "values": [100, 3500000000, 25, 45, 200, 1600, 55, 65] }
930
+ },
931
+ "c": {
932
+ "dtype": "uint32",
933
+ "shape": [8],
934
+ "data": { "kind": "values", "values": [200, 20, 2147483648, 50, 210, 70, 1700, 75] }
935
+ },
936
+ "d": {
937
+ "dtype": "uint32",
938
+ "shape": [8],
939
+ "data": { "kind": "values", "values": [300, 30, 35, 4294967295, 220, 80, 85, 1800] }
940
+ }
941
+ },
942
+ "outputs": { "y": { "dtype": "uint32", "shape": [8], "tolerance": 0 } }
943
+ },
944
+ {
945
+ "name": "f16_three_input_same_shape_vec4",
946
+ "provenance": {
947
+ "notes": "float16 at arity 3 exercises the flat vec4 same-shape path, including widening to f32, the extra-operand maximum fold, NaN re-injection, and narrowing. Every value is exactly representable in float16 and each operand wins at least two lanes."
948
+ },
949
+ "inputs": {
950
+ "a": {
951
+ "dtype": "float16",
952
+ "shape": [8],
953
+ "data": { "kind": "values", "values": [3.5, -2.5, 0.25, 6.0, -3.0, 3.5, -0.75, 0.5] }
954
+ },
955
+ "b": {
956
+ "dtype": "float16",
957
+ "shape": [8],
958
+ "data": { "kind": "values", "values": [2.5, -0.5, -0.5, 2.0, -1.0, 2.5, -8.0, 7.75] }
959
+ },
960
+ "c": {
961
+ "dtype": "float16",
962
+ "shape": [8],
963
+ "data": { "kind": "values", "values": [0.5, -1.5, 4.5, 3.0, -2.0, 5.5, -4.0, 1.5] }
964
+ }
965
+ },
966
+ "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.002 } }
967
+ },
968
+ {
969
+ "name": "f16_four_input_same_shape_vec4",
970
+ "provenance": {
971
+ "notes": "float16 at arity 4 on the flat vec4 same-shape path (same_shape_vec4_four_input, numel 8). Two extra-operand folds in the f16 lane of minmax-vec4, so the second fold consumes the first fold's result and the loop.first-guarded comment renders on the first pass only. Valid ONNX Max-13 for T = float16 with four same-shape operands. Every value is exactly representable in float16 and each of a, b, c, d supplies the maximum on exactly two lanes, so no operand can be dropped."
972
+ },
973
+ "inputs": {
974
+ "a": {
975
+ "dtype": "float16",
976
+ "shape": [2, 4],
977
+ "data": { "kind": "values", "values": [3.5, -2.5, 0.25, 6.0, -1.0, 3.5, -4.0, 0.5] }
978
+ },
979
+ "b": {
980
+ "dtype": "float16",
981
+ "shape": [2, 4],
982
+ "data": { "kind": "values", "values": [2.5, -0.5, -0.5, 2.0, -3.0, 6.5, -8.0, 5.5] }
983
+ },
984
+ "c": {
985
+ "dtype": "float16",
986
+ "shape": [2, 4],
987
+ "data": { "kind": "values", "values": [0.5, -1.5, 4.5, 3.0, -2.0, 5.5, -0.75, 1.5] }
988
+ },
989
+ "d": {
990
+ "dtype": "float16",
991
+ "shape": [2, 4],
992
+ "data": { "kind": "values", "values": [-0.5, -3.5, 1.25, 9.0, -6.0, 2.5, -1.0, 7.75] }
993
+ }
994
+ },
995
+ "outputs": { "y": { "dtype": "float16", "shape": [2, 4], "tolerance": 0.002 } }
996
+ }
997
+ ]
998
+ }