Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Constant

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
6eee053
·
verified ·
1 Parent(s): f57a051

sync 6fdf6301e2bb

Browse files
Files changed (4) hide show
  1. README.md +2 -2
  2. build/webgpu/constant.wgsl.jinja +17 -43
  3. build/webgpu/manifest.json +3 -6
  4. build/webgpu/metadata.json +5 -5
README.md CHANGED
@@ -43,13 +43,13 @@ Attributes and default values (overridable per request):
43
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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
  - [`constant.wgsl.jinja`](build/webgpu/constant.wgsl.jinja)
48
 
49
  ## Use with `@huggingface/kernels`
50
 
51
  ```sh
52
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
53
  ```
54
 
55
  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.
 
43
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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 cases
47
  - [`constant.wgsl.jinja`](build/webgpu/constant.wgsl.jinja)
48
 
49
  ## Use with `@huggingface/kernels`
50
 
51
  ```sh
52
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
53
  ```
54
 
55
  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/constant.wgsl.jinja CHANGED
@@ -1,69 +1,43 @@
1
- {% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
2
- {% if note == "dispatch-limit" %}
3
- // 2D-folded flat index: gid.y carries the high bits past the dispatch's
4
- // per-axis workgroup fold width (outputs > 16.7M elements).
5
- {% elif note == "limit" %}
6
- // 2D-folded flat index: gid.y carries the high bits past the dispatch's
7
- // per-axis workgroup fold width.
8
- {% elif note == "device-axis" %}
9
- // The flat dispatch is folded across x/y at a fixed per-axis workgroup
10
- // width; gid.y carries the high portion of the output index.
11
- {% elif note == "vec4-limit" %}
12
- // 2D-folded flat vec4 index: gid.y carries the high bits past the dispatch's
13
- // per-axis workgroup fold width (the dispatch caps x and spills into y).
14
- {% elif note == "element-limit" %}
15
- // 2D-folded flat element index: gid.y carries the high bits past the
16
- // dispatch's per-axis workgroup fold width.
17
- {% elif note == "dispatch" %}
18
  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
19
  // per-axis workgroup fold width.
20
- {% endif %}
21
- {% if bound == "" %}
22
- let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
23
- {%- elif guardInline %}
24
- let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
25
- if ({{ name }} >= {{ bound }}) { return; }
26
- {%- else %}
27
- let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
28
- if ({{ name }} >= {{ bound }}) {
29
- return;
30
- }
31
- {%- endif %}
32
- {% endmacro %}
33
-
34
  {{ env.wgsl.resourceDeclarations }}
35
  {% set emitValueTable = hasValues and valueCount > 0 %}
36
  {% set special = namespace(value=false) %}
37
  {% for value in values %}
38
  {% if value == "NaN" or value != value or value == "Infinity" or value == "-Infinity" or (value == 0 and 1.0 / value < 0.0) %}{% set special.value = true %}{% endif %}
39
  {% endfor %}
40
- {% macro constant_bits(value) %}
41
- {%- if value == "NaN" or value != value -%}0x7fc00000u
42
- {%- elif value == "Infinity" -%}0x7f800000u
43
- {%- elif value == "-Infinity" -%}0xff800000u
44
- {%- elif value == 0 and 1.0 / value < 0.0 -%}0x80000000u
45
- {%- else -%}bitcast<u32>(f32({{ value }}))
46
- {%- endif -%}
47
- {% endmacro %}
48
  {% if emitValueTable and special.value %}
 
49
  var<private> constant_bits_table: array<u32, {{ valueCount }}> = array<u32, {{ valueCount }}>(
50
  {% for value in values %}
51
- {{ constant_bits(value) }}{% if not loop.last %},{% endif %}
52
  {% endfor %}
53
  );
54
  {% elif emitValueTable %}
 
55
  var<private> constant_values: array<{{ scalar }}, {{ valueCount }}> = array<{{ scalar }}, {{ valueCount }}>(
56
  {% for value in values %}
57
- {{ scalar }}({{ value }}){% if not loop.last %},{% endif %}
58
  {% endfor %}
59
  );
60
-
61
  {% endif %}
 
62
  const COUNT: u32 = {{ count }}u;
63
 
64
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
65
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
66
- {{ flat_index_2d("i", "COUNT", guardInline=true) }}
67
  {% if emitValueTable %}
68
  {% if special.value %}
69
  output[i] = {{ scalar }}(bitcast<f32>(constant_bits_table[i]));
 
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 }}) { return; }{% endmacro %}
 
 
 
 
 
 
 
 
 
 
 
 
7
  {{ env.wgsl.resourceDeclarations }}
8
  {% set emitValueTable = hasValues and valueCount > 0 %}
9
  {% set special = namespace(value=false) %}
10
  {% for value in values %}
11
  {% if value == "NaN" or value != value or value == "Infinity" or value == "-Infinity" or (value == 0 and 1.0 / value < 0.0) %}{% set special.value = true %}{% endif %}
12
  {% endfor %}
13
+ {% macro constant_bits(value) %}{%
14
+ if value == "NaN" or value != value %}0x7fc00000u{%
15
+ elif value == "Infinity" %}0x7f800000u{%
16
+ elif value == "-Infinity" %}0xff800000u{%
17
+ elif value == 0 and 1.0 / value < 0.0 %}0x80000000u{%
18
+ else %}bitcast<u32>(f32({{ value }})){%
19
+ endif %}{% endmacro %}
 
20
  {% if emitValueTable and special.value %}
21
+
22
  var<private> constant_bits_table: array<u32, {{ valueCount }}> = array<u32, {{ valueCount }}>(
23
  {% for value in values %}
24
+ {{ constant_bits(value) }}{{ "," if not loop.last else "" }}
25
  {% endfor %}
26
  );
27
  {% elif emitValueTable %}
28
+
29
  var<private> constant_values: array<{{ scalar }}, {{ valueCount }}> = array<{{ scalar }}, {{ valueCount }}>(
30
  {% for value in values %}
31
+ {{ scalar }}({{ value }}){{ "," if not loop.last else "" }}
32
  {% endfor %}
33
  );
 
34
  {% endif %}
35
+
36
  const COUNT: u32 = {{ count }}u;
37
 
38
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
39
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
40
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE, "i", "COUNT", guardInline=true) }}
41
  {% if emitValueTable %}
42
  {% if special.value %}
43
  output[i] = {{ scalar }}(bitcast<f32>(constant_bits_table[i]));
build/webgpu/manifest.json CHANGED
@@ -7,7 +7,7 @@
7
  "attributes": { "value": {}, "value_float": {}, "value_floats": {} },
8
  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "uint32", "int8", "uint8", "bool"] },
9
  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
10
- "derive": { "scalar": "dtypes.T" },
11
  "variants": [
12
  {
13
  "id": "literal_values",
@@ -21,10 +21,7 @@
21
  "id": "main",
22
  "name": "Constant",
23
  "shader": "constant.wgsl.jinja",
24
- "derive": {
25
- "count": "numel(shapes.output)",
26
- "values": "attrs.value.values if has(attrs, \"value\") else attrs.value_floats"
27
- },
28
  "bindings": ["output"],
29
  "dispatch": {
30
  "x": "min(ceilDiv((numel(shapes.output)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -43,7 +40,7 @@
43
  "id": "main",
44
  "name": "Constant",
45
  "shader": "constant.wgsl.jinja",
46
- "derive": { "count": "numel(shapes.output)", "values": [] },
47
  "bindings": [
48
  "output",
49
  { "name": "params", "struct": [{ "name": "value", "type": "f32", "value": "attrs.value_float" }] }
 
7
  "attributes": { "value": {}, "value_float": {}, "value_floats": {} },
8
  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "uint32", "int8", "uint8", "bool"] },
9
  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
10
+ "derive": { "scalar": "dtypes.T", "count": "numel(shapes.output)" },
11
  "variants": [
12
  {
13
  "id": "literal_values",
 
21
  "id": "main",
22
  "name": "Constant",
23
  "shader": "constant.wgsl.jinja",
24
+ "derive": { "values": "attrs.value.values if has(attrs, \"value\") else attrs.value_floats" },
 
 
 
25
  "bindings": ["output"],
26
  "dispatch": {
27
  "x": "min(ceilDiv((numel(shapes.output)), (tunables.WORKGROUP_SIZE)), 65535)",
 
40
  "id": "main",
41
  "name": "Constant",
42
  "shader": "constant.wgsl.jinja",
43
+ "derive": { "values": [] },
44
  "bindings": [
45
  "output",
46
  { "name": "params", "struct": [{ "name": "value", "type": "f32", "value": "attrs.value_float" }] }
build/webgpu/metadata.json CHANGED
@@ -1,6 +1,6 @@
1
  {
2
  "name": "ai.onnx.Constant",
3
- "id": "_ai_onnx_constant_webgpu_2a92cba",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
@@ -8,14 +8,14 @@
8
  "algorithm": "sha256",
9
  "files": {
10
  "bench.json": "HxS+gUldt+otuLxtLGKs9T9R75abm94dZaQF+3LqaeA=",
11
- "constant.wgsl.jinja": "86oxsLXP1xToF/ZrbjOqxCz2jURzVf7HOiOrzPF3Qbs=",
12
- "manifest.json": "6KPtd+xByY7jpUB7WgoCMrt0mfTWagKuR86+rKb9fPY=",
13
  "test.json": "Mb0u5uDh/gytPn6T4PNNs5Uk+7MTzodeTx5sfDjEnxs="
14
  }
15
  },
16
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
17
  "webgpu": {
18
- "manifestSpec": "2.0",
19
  "variants": { "literal_values": ["constant.wgsl.jinja"], "fill_float": ["constant.wgsl.jinja"] }
20
  }
21
  }
 
1
  {
2
  "name": "ai.onnx.Constant",
3
+ "id": "_ai_onnx_constant_webgpu_b50517a",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
 
8
  "algorithm": "sha256",
9
  "files": {
10
  "bench.json": "HxS+gUldt+otuLxtLGKs9T9R75abm94dZaQF+3LqaeA=",
11
+ "constant.wgsl.jinja": "Pejhwr5ah4yqyKrt9bnH3fkokw5b+HeR1/Y5br5xEDk=",
12
+ "manifest.json": "oIxGRp7XdlLRqt/1o6iupAwZrORXptCr/MKzd/CFS+c=",
13
  "test.json": "Mb0u5uDh/gytPn6T4PNNs5Uk+7MTzodeTx5sfDjEnxs="
14
  }
15
  },
16
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
17
  "webgpu": {
18
+ "manifestSpec": "2.1",
19
  "variants": { "literal_values": ["constant.wgsl.jinja"], "fill_float": ["constant.wgsl.jinja"] }
20
  }
21
  }