Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Expand

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
2e455da
·
verified ·
1 Parent(s): 8f90598

sync 6fdf6301e2bb

Browse files
Files changed (7) hide show
  1. README.md +3 -3
  2. build/webgpu/bench.json +13 -3
  3. build/webgpu/expand-vec4.wgsl.jinja +10 -7
  4. build/webgpu/expand.wgsl.jinja +9 -6
  5. build/webgpu/manifest.json +18 -34
  6. build/webgpu/metadata.json +8 -8
  7. build/webgpu/test.json +82 -0
README.md CHANGED
@@ -33,7 +33,7 @@ See the [ONNX `Expand` spec](https://onnx.ai/onnx/operators/onnx__Expand.html) f
33
 
34
  | Variable | Allowed dtypes |
35
  | --- | --- |
36
- | `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
37
  | `S` | `int64` |
38
 
39
  ## Files
@@ -41,14 +41,14 @@ See the [ONNX `Expand` spec](https://onnx.ai/onnx/operators/onnx__Expand.html) f
41
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
42
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
43
  - [`test.json`](build/webgpu/test.json) — correctness cases
44
- - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
45
  - [`expand-vec4.wgsl.jinja`](build/webgpu/expand-vec4.wgsl.jinja)
46
  - [`expand.wgsl.jinja`](build/webgpu/expand.wgsl.jinja)
47
 
48
  ## Use with `@huggingface/kernels`
49
 
50
  ```sh
51
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
52
  ```
53
 
54
  Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
 
33
 
34
  | Variable | Allowed dtypes |
35
  | --- | --- |
36
+ | `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool`, `int64` |
37
  | `S` | `int64` |
38
 
39
  ## Files
 
41
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
42
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
43
  - [`test.json`](build/webgpu/test.json) — correctness cases
44
+ - [`bench.json`](build/webgpu/bench.json) — benchmark cases
45
  - [`expand-vec4.wgsl.jinja`](build/webgpu/expand-vec4.wgsl.jinja)
46
  - [`expand.wgsl.jinja`](build/webgpu/expand.wgsl.jinja)
47
 
48
  ## Use with `@huggingface/kernels`
49
 
50
  ```sh
51
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
52
  ```
53
 
54
  Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
build/webgpu/bench.json CHANGED
@@ -1,6 +1,16 @@
1
  {
2
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
3
  "cases": [
 
 
 
 
 
 
 
 
 
 
4
  {
5
  "name": "expand-f32-1m",
6
  "preset": "smoke",
@@ -29,7 +39,7 @@
29
  {
30
  "name": "expand-f32-unaligned-row-1m",
31
  "provenance": {
32
- "notes": "A 1025-wide row: the output last extent is not a multiple of four, so one four-lane group per row crosses the row boundary. The vec4 arm resolves that group per lane and keeps the single-offset path for the rest, rather than sending the whole tensor to the scalar route."
33
  },
34
  "preset": "stress",
35
  "vars": { "dtype": "float32", "count": 1049600 },
@@ -44,7 +54,7 @@
44
  "name": "expand-f32-mask-8x1x1x197-to-8x8x197x197-vit-pathology",
45
  "preset": "stress",
46
  "provenance": {
47
- "notes": "Key-mask broadcast [B,1,1,S] -> [B,H,S,S] at ViT-B/16 token count: broadcast_vec4 needs the output last extent % 4 == 0, and 197 fails it, so the 2.5M-element broadcast runs the scalar broadcast route. Control: expand-f32-mask-8x1x1x200-to-8x8x200x200-control."
48
  },
49
  "vars": { "dtype": "float32" },
50
  "inputs": {
@@ -62,7 +72,7 @@
62
  "name": "expand-f32-mask-8x1x1x200-to-8x8x200x200-control",
63
  "preset": "stress",
64
  "provenance": {
65
- "notes": "Nearest multiple-of-4 token count for expand-f32-mask-8x1x1x197-to-8x8x197x197-vit-pathology (3% more bytes) on the broadcast_vec4 route."
66
  },
67
  "vars": { "dtype": "float32" },
68
  "inputs": {
 
1
  {
2
  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
3
  "cases": [
4
+ {
5
+ "name": "qwen-image21-vae-repeat-axis1-1152x1x128x128-to-1152x8x128x128",
6
+ "preset": "model",
7
+ "inputs": {
8
+ "input": { "dtype": "float32", "shape": [1152, 1, 128, 128], "dist": "normal", "seed": 20268, "scale": 1 },
9
+ "shape": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [1152, 8, 128, 128] } }
10
+ },
11
+ "outputs": { "output": { "dtype": "float32", "shape": [1152, 8, 128, 128], "dist": "empty" } },
12
+ "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.input) + numel(shapes.output)) * 4" }] }
13
+ },
14
  {
15
  "name": "expand-f32-1m",
16
  "preset": "smoke",
 
39
  {
40
  "name": "expand-f32-unaligned-row-1m",
41
  "provenance": {
42
+ "notes": "A 1,025-wide output row is not divisible by four; four-element storage groups at row boundaries must preserve each row's broadcast values."
43
  },
44
  "preset": "stress",
45
  "vars": { "dtype": "float32", "count": 1049600 },
 
54
  "name": "expand-f32-mask-8x1x1x197-to-8x8x197x197-vit-pathology",
55
  "preset": "stress",
56
  "provenance": {
57
+ "notes": "Broadcasts a [B,1,1,S] key mask to [B,H,S,S] at the ViT-B/16 token count of 197. The paired 200-token case checks an aligned width."
58
  },
59
  "vars": { "dtype": "float32" },
60
  "inputs": {
 
72
  "name": "expand-f32-mask-8x1x1x200-to-8x8x200x200-control",
73
  "preset": "stress",
74
  "provenance": {
75
+ "notes": "The next multiple-of-four token count above the paired 197-token case, with about 3% more bytes."
76
  },
77
  "vars": { "dtype": "float32" },
78
  "inputs": {
build/webgpu/expand-vec4.wgsl.jinja CHANGED
@@ -1,3 +1,11 @@
 
 
 
 
 
 
 
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
  {% set x_numel = namespace(value=1) %}
@@ -19,8 +27,8 @@
19
  {% endfor %}
20
  {% if not rowAligned and x_mode != "scalar" %}
21
  const INNER: u32 = {{ yShape[yRank - 1] }}u;
22
- {% endif %}
23
 
 
24
  {% if x_mode == "scalar" %}
25
  // Each invocation splats the single input value across four consecutive outputs.
26
  {% elif x_mode == "splat" %}
@@ -68,12 +76,7 @@ fn x_offset(out_index: u32) -> u32 {
68
 
69
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
70
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
71
- // 2D-folded flat vec4 index: gid.y carries the high bits past the
72
- // per-axis dispatch fold width (the dispatch caps x and spills into y).
73
- let i4 = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
74
- if (i4 >= params.count) {
75
- return;
76
- }
77
  {% if x_mode != "scalar" %}
78
  let base = i4 * 4u;
79
  {% endif %}
 
1
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
2
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
3
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
+ // per-axis workgroup fold width.
5
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
6
+ if ({{ name }} >= {{ bound }}) {
7
+ return;
8
+ }{% endmacro %}
9
  {{ env.wgsl.resourceDeclarations }}
10
 
11
  {% set x_numel = namespace(value=1) %}
 
27
  {% endfor %}
28
  {% if not rowAligned and x_mode != "scalar" %}
29
  const INNER: u32 = {{ yShape[yRank - 1] }}u;
 
30
 
31
+ {% endif %}
32
  {% if x_mode == "scalar" %}
33
  // Each invocation splats the single input value across four consecutive outputs.
34
  {% elif x_mode == "splat" %}
 
76
 
77
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
78
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
79
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE, "i4") }}
 
 
 
 
 
80
  {% if x_mode != "scalar" %}
81
  let base = i4 * 4u;
82
  {% endif %}
build/webgpu/expand.wgsl.jinja CHANGED
@@ -1,3 +1,11 @@
 
 
 
 
 
 
 
 
1
  {{ env.wgsl.resourceDeclarations }}
2
 
3
  {% set xUsesIndex = namespace(value=false) %}
@@ -37,11 +45,6 @@ fn x_offset({% if xUsesIndex.value %}out_index: u32{% endif %}) -> u32 {
37
 
38
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
39
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
40
- // 2D-folded flat index: gid.y carries the high bits past the
41
- // per-axis dispatch fold width (outputs > 16.7M elements).
42
- let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
43
- if (i >= params.count) {
44
- return;
45
- }
46
  y[i] = x[x_offset({% if xUsesIndex.value %}i{% endif %})];
47
  }
 
1
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
2
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
3
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
+ // per-axis workgroup fold width.
5
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
6
+ if ({{ name }} >= {{ bound }}) {
7
+ return;
8
+ }{% endmacro %}
9
  {{ env.wgsl.resourceDeclarations }}
10
 
11
  {% set xUsesIndex = namespace(value=false) %}
 
45
 
46
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
47
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
48
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE) }}
 
 
 
 
 
49
  y[i] = x[x_offset({% if xUsesIndex.value %}i{% endif %})];
50
  }
build/webgpu/manifest.json CHANGED
@@ -8,38 +8,35 @@
8
  },
9
  "outputs": { "output": { "dtype": "T", "rank": "dim(shapes.shape, 0)" } },
10
  "typeConstraints": {
11
- "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"],
12
  "S": ["int64"]
13
  },
14
  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
15
- "derive": { "scalar": "dtypes.T" },
 
 
 
 
 
 
 
16
  "when": ["ranks.input <= ranks.output", "f16Ok(dtypes.T)"],
17
  "bindings": {
18
- "x": { "arg": "input", "buffer": "read-only-storage", "elementType": "$scalar" },
19
- "y": { "arg": "output", "buffer": "storage", "elementType": "$vectorScalar" },
20
- "params": {
21
- "buffer": "uniform",
22
- "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
23
- }
24
  },
25
  "variants": [
26
  {
27
  "id": "broadcast_vec4",
28
  "priority": 10,
29
- "when": ["ranks.output >= 1", "dim(shapes.output, -1) % 4 == 0"],
30
- "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
31
  "passes": [
32
  {
33
  "id": "main",
34
- "name": "Expand.vec4",
35
  "shader": "expand-vec4.wgsl.jinja",
36
- "derive": {
37
- "xShape": "shapes.input",
38
- "yShape": "shapes.output",
39
- "xRank": "ranks.input",
40
- "yRank": "ranks.output",
41
- "rowAligned": "true"
42
- },
43
  "bindings": ["x", "y", "params"],
44
  "dispatch": {
45
  "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -52,20 +49,13 @@
52
  {
53
  "id": "broadcast_vec4_unaligned",
54
  "priority": 10,
55
- "when": ["ranks.output >= 1", "dim(shapes.output, -1) % 4 != 0", "numel(shapes.output) % 4 == 0"],
56
- "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
57
  "passes": [
58
  {
59
  "id": "main",
60
- "name": "Expand.vec4Unaligned",
61
  "shader": "expand-vec4.wgsl.jinja",
62
- "derive": {
63
- "xShape": "shapes.input",
64
- "yShape": "shapes.output",
65
- "xRank": "ranks.input",
66
- "yRank": "ranks.output",
67
- "rowAligned": "false"
68
- },
69
  "bindings": ["x", "y", "params"],
70
  "dispatch": {
71
  "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -82,12 +72,6 @@
82
  "id": "main",
83
  "name": "Expand",
84
  "shader": "expand.wgsl.jinja",
85
- "derive": {
86
- "xShape": "shapes.input",
87
- "yShape": "shapes.output",
88
- "xRank": "ranks.input",
89
- "yRank": "ranks.output"
90
- },
91
  "bindings": [
92
  "x",
93
  { "arg": "output", "name": "y", "elementType": "$scalar" },
 
8
  },
9
  "outputs": { "output": { "dtype": "T", "rank": "dim(shapes.shape, 0)" } },
10
  "typeConstraints": {
11
+ "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool", "int64"],
12
  "S": ["int64"]
13
  },
14
  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
15
+ "derive": {
16
+ "scalar": "dtypes.T",
17
+ "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
18
+ "xShape": "shapes.input",
19
+ "yShape": "shapes.output",
20
+ "xRank": "ranks.input",
21
+ "yRank": "ranks.output"
22
+ },
23
  "when": ["ranks.input <= ranks.output", "f16Ok(dtypes.T)"],
24
  "bindings": {
25
+ "x": { "arg": "input", "elementType": "$scalar" },
26
+ "y": { "arg": "output", "elementType": "$vectorScalar" },
27
+ "params": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }] }
 
 
 
28
  },
29
  "variants": [
30
  {
31
  "id": "broadcast_vec4",
32
  "priority": 10,
33
+ "when": ["dtypes.T != \"vec2<u32>\"", "ranks.output >= 1", "dim(shapes.output, -1) % 4 == 0"],
 
34
  "passes": [
35
  {
36
  "id": "main",
37
+ "name": "Expand.Vec4",
38
  "shader": "expand-vec4.wgsl.jinja",
39
+ "derive": { "rowAligned": "true" },
 
 
 
 
 
 
40
  "bindings": ["x", "y", "params"],
41
  "dispatch": {
42
  "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
 
49
  {
50
  "id": "broadcast_vec4_unaligned",
51
  "priority": 10,
52
+ "when": ["dtypes.T != \"vec2<u32>\"", "ranks.output >= 1", "dim(shapes.output, -1) % 4 != 0", "numel(shapes.output) % 4 == 0"],
 
53
  "passes": [
54
  {
55
  "id": "main",
56
+ "name": "Expand.Vec4Unaligned",
57
  "shader": "expand-vec4.wgsl.jinja",
58
+ "derive": { "rowAligned": "false" },
 
 
 
 
 
 
59
  "bindings": ["x", "y", "params"],
60
  "dispatch": {
61
  "x": "min(ceilDiv((numel(shapes.output) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
 
72
  "id": "main",
73
  "name": "Expand",
74
  "shader": "expand.wgsl.jinja",
 
 
 
 
 
 
75
  "bindings": [
76
  "x",
77
  { "arg": "output", "name": "y", "elementType": "$scalar" },
build/webgpu/metadata.json CHANGED
@@ -1,22 +1,22 @@
1
  {
2
  "name": "ai.onnx.Expand",
3
- "id": "_ai_onnx_expand_webgpu_88532df",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
- "bench.json": "9ZZh45r+mNj29neJarPavktd9cc25IL9TLMj6M6ei7Q=",
11
- "expand-vec4.wgsl.jinja": "yHWI7Plp/V+IZJzuSr9LbcylQVf8PXrqUQUBdk7pMWo=",
12
- "expand.wgsl.jinja": "/fbA5AqgtlnuHr7J5LhwY8s3le1Mopoxwnhoc+krT58=",
13
- "manifest.json": "GHy2hH3gY8kKKh8BNXzyUK1iMPJsZF0z3ffelnUabF8=",
14
- "test.json": "ATB5R2TC849vaeIc1aoSyB0D6GWQs1rURocCzn/gUmQ="
15
  }
16
  },
17
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
18
  "webgpu": {
19
- "manifestSpec": "2.0",
20
  "variants": {
21
  "broadcast_vec4": ["expand-vec4.wgsl.jinja"],
22
  "broadcast_vec4_unaligned": ["expand-vec4.wgsl.jinja"],
 
1
  {
2
  "name": "ai.onnx.Expand",
3
+ "id": "_ai_onnx_expand_webgpu_ed79f31",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
7
  "digest": {
8
  "algorithm": "sha256",
9
  "files": {
10
+ "bench.json": "dJIB22THU9XT7VROcflGk19V/pQuSU8qNezZv2pYLgg=",
11
+ "expand-vec4.wgsl.jinja": "JbOkE7xSBKBqk/gA/6HEqUlhzDYvZSCySbH/Hdf5y38=",
12
+ "expand.wgsl.jinja": "PtY82ihEZt18i1Tb0mIY100bJRawUl/X5bJLH58Ecbs=",
13
+ "manifest.json": "kF/0cg7bLBeJ6JwiTfYdcFMoTxw2D2HvTy2PRstmUCg=",
14
+ "test.json": "0y7LngCk7iCRRywYALOpc18/OpBow8m58BfgUeM+JNE="
15
  }
16
  },
17
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
18
  "webgpu": {
19
+ "manifestSpec": "2.1",
20
  "variants": {
21
  "broadcast_vec4": ["expand-vec4.wgsl.jinja"],
22
  "broadcast_vec4_unaligned": ["expand-vec4.wgsl.jinja"],
build/webgpu/test.json CHANGED
@@ -1,5 +1,17 @@
1
  {
2
  "cases": [
 
 
 
 
 
 
 
 
 
 
 
 
3
  {
4
  "name": "int16_broadcast_boundaries",
5
  "inputs": {
@@ -601,6 +613,76 @@
601
  "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } }
602
  },
603
  "outputs": { "output": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 } }
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
604
  }
605
  ]
606
  }
 
1
  {
2
  "cases": [
3
+ {
4
+ "name": "repeat_second_axis_with_contiguous_inner_block",
5
+ "inputs": {
6
+ "input": {
7
+ "dtype": "float32",
8
+ "shape": [2, 1, 3, 8],
9
+ "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 }
10
+ },
11
+ "shape": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [2, 3, 3, 8] } }
12
+ },
13
+ "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 3, 8], "tolerance": 0 } }
14
+ },
15
  {
16
  "name": "int16_broadcast_boundaries",
17
  "inputs": {
 
613
  "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } }
614
  },
615
  "outputs": { "output": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 } }
616
+ },
617
+ {
618
+ "name": "ort_uint8_scalar_to_3x3",
619
+ "provenance": {
620
+ "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
621
+ "test": "ExpandOpTest.Expand_3x3_uint8"
622
+ },
623
+ "inputs": {
624
+ "input": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [5] } },
625
+ "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
626
+ },
627
+ "outputs": { "output": { "dtype": "uint8", "shape": [3, 3], "tolerance": 0 } }
628
+ },
629
+ {
630
+ "name": "ort_uint8_row_vector_to_rows",
631
+ "provenance": {
632
+ "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
633
+ "test": "ExpandOpTest.Expand_3x1_uint8"
634
+ },
635
+ "inputs": {
636
+ "input": { "dtype": "uint8", "shape": [3], "data": { "kind": "values", "values": [11, 22, 33] } },
637
+ "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
638
+ },
639
+ "outputs": { "output": { "dtype": "uint8", "shape": [3, 3], "tolerance": 0 } }
640
+ },
641
+ {
642
+ "name": "ort_uint8_col_vector_splat_to_four_cols",
643
+ "provenance": {
644
+ "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
645
+ "test": "ExpandOpTest.Expand_1x4_uint8"
646
+ },
647
+ "inputs": {
648
+ "input": { "dtype": "uint8", "shape": [3, 1], "data": { "kind": "values", "values": [11, 22, 33] } },
649
+ "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 4] } }
650
+ },
651
+ "outputs": { "output": { "dtype": "uint8", "shape": [3, 4], "tolerance": 0 } }
652
+ },
653
+ {
654
+ "name": "int64_full_range_int16_broadcast_boundaries",
655
+ "inputs": {
656
+ "input": {
657
+ "dtype": "int64",
658
+ "shape": [1, 4],
659
+ "data": {
660
+ "kind": "cycle",
661
+ "values": ["0", "4294967296", "8589934593", "-4294967297", "9223372036854775807", "-9223372036854775808", "9223372036854775806", "-1"]
662
+ }
663
+ },
664
+ "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } }
665
+ },
666
+ "outputs": { "output": { "dtype": "int64", "shape": [2, 4], "tolerance": 0, "relTolerance": 0 } },
667
+ "tolerance": 0,
668
+ "relTolerance": 0
669
+ },
670
+ {
671
+ "name": "int64_full_range_scalar_to_3x3",
672
+ "inputs": {
673
+ "input": {
674
+ "dtype": "int64",
675
+ "shape": [1],
676
+ "data": {
677
+ "kind": "cycle",
678
+ "values": ["4294967296", "8589934593", "-4294967297", "9223372036854775807", "-9223372036854775808", "9223372036854775806", "-1", "0"]
679
+ }
680
+ },
681
+ "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
682
+ },
683
+ "outputs": { "output": { "dtype": "int64", "shape": [3, 3], "tolerance": 0, "relTolerance": 0 } },
684
+ "tolerance": 0,
685
+ "relTolerance": 0
686
  }
687
  ]
688
  }