Hugging Face's logo

Kernels:
webgpu-kernels
/
ai.onnx.Sinh

Kernel card Files
xet
 Community
Xenova HF Staff commited on
Commit
ca8a3a6
·
verified ·
1 Parent(s): 618a3b5

sync 6fdf6301e2bb

Browse files
Files changed (5) hide show
  1. README.md +2 -2
  2. build/webgpu/manifest.json +19 -41
  3. build/webgpu/metadata.json +6 -6
  4. build/webgpu/unary-scalar.wgsl.jinja +7 -10
  5. build/webgpu/unary-vec4.wgsl.jinja +21 -12
README.md CHANGED
@@ -39,14 +39,14 @@ See the [ONNX `Sinh` spec](https://onnx.ai/onnx/operators/onnx__Sinh.html) for t
39
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
40
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
41
  - [`test.json`](build/webgpu/test.json) — correctness cases
42
- - [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
43
  - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
44
  - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
45
 
46
  ## Use with `@huggingface/kernels`
47
 
48
  ```sh
49
- npm install --save-exact @huggingface/kernels@0.0.1-preview.2
50
  ```
51
 
52
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 
39
  - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
40
  - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
41
  - [`test.json`](build/webgpu/test.json) — correctness cases
42
+ - [`bench.json`](build/webgpu/bench.json) — benchmark cases
43
  - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
44
  - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
45
 
46
  ## Use with `@huggingface/kernels`
47
 
48
  ```sh
49
+ npm install --save-exact @huggingface/kernels@0.0.1-preview.3
50
  ```
51
 
52
  Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
build/webgpu/manifest.json CHANGED
@@ -11,20 +11,13 @@
11
  "id": "same_layout_vec4",
12
  "priority": 20,
13
  "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
14
- "derive": {
15
- "scalar": "dtypes.T",
16
- "usesF16": "dtypes.T == \"f16\"",
17
- "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
18
- },
19
  "passes": [
20
  {
21
  "id": "main",
22
- "name": "Sinh.vec4",
23
  "shader": "unary-vec4.wgsl.jinja",
24
- "derive": {
25
- "op": "\"sinh\"",
26
- "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
27
- },
28
  "bindings": ["x", "y", "params_unary"],
29
  "dispatch": {
30
  "x": "min(ceilDiv((ceilDiv(numel(shapes.y) / 4, 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -37,14 +30,14 @@
37
  {
38
  "id": "elementwise",
39
  "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
40
- "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
41
  "passes": [
42
  {
43
  "id": "main",
44
  "name": "Sinh",
45
  "shader": "unary-scalar.wgsl.jinja",
46
- "derive": { "op": "\"sinh\"", "itemsPerInvocation": 4 },
47
- "bindings": ["x_2", "y_2", "params_2_unary"],
48
  "dispatch": {
49
  "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
50
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -56,34 +49,29 @@
56
  {
57
  "id": "same_layout_vec4_tail",
58
  "priority": 19,
59
- "when": ["numel(shapes.x) > 4", "numel(shapes.x) % 4 != 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
60
- "derive": {
61
- "scalar": "dtypes.T",
62
- "usesF16": "dtypes.T == \"f16\"",
63
- "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
64
- },
65
  "passes": [
66
  {
67
  "id": "bulk",
68
- "name": "Sinh.vec4Bulk",
69
  "shader": "unary-vec4.wgsl.jinja",
70
  "derive": {
71
- "op": "\"sinh\"",
72
- "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
73
  },
74
  "bindings": ["x", "y", "params_unary_tail"],
75
  "dispatch": {
76
- "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
77
- "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
78
  "z": 1
79
  }
80
  },
81
  {
82
  "id": "tail",
83
- "name": "Sinh.tail",
84
  "shader": "unary-scalar.wgsl.jinja",
85
- "derive": { "op": "\"sinh\"", "itemsPerInvocation": 4, "tailOnly": true },
86
- "bindings": ["x_2", "y_2", "params_2_unary_tail"],
87
  "dispatch": {
88
  "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
89
  "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -96,25 +84,15 @@
96
  "bindings": {
97
  "x": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
98
  "y": { "buffer": "storage", "elementType": "$vectorScalar" },
99
- "params_unary": {
100
- "buffer": "uniform",
101
- "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }],
102
- "name": "params"
103
- },
104
- "x_2": { "buffer": "read-only-storage", "name": "x", "elementType": "$scalar" },
105
- "y_2": { "buffer": "storage", "name": "y", "elementType": "$scalar" },
106
- "params_2_unary": {
107
- "buffer": "uniform",
108
- "name": "params",
109
- "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
110
- },
111
  "params_unary_tail": {
112
- "buffer": "uniform",
113
  "struct": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }],
114
  "name": "params"
115
  },
116
  "params_2_unary_tail": {
117
- "buffer": "uniform",
118
  "name": "params",
119
  "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
120
  }
 
11
  "id": "same_layout_vec4",
12
  "priority": 20,
13
  "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
14
+ "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
 
 
 
 
15
  "passes": [
16
  {
17
  "id": "main",
18
+ "name": "Sinh.Vec4",
19
  "shader": "unary-vec4.wgsl.jinja",
20
+ "derive": { "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1" },
 
 
 
21
  "bindings": ["x", "y", "params_unary"],
22
  "dispatch": {
23
  "x": "min(ceilDiv((ceilDiv(numel(shapes.y) / 4, 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
 
30
  {
31
  "id": "elementwise",
32
  "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
33
+ "derive": { "scalar": "dtypes.T" },
34
  "passes": [
35
  {
36
  "id": "main",
37
  "name": "Sinh",
38
  "shader": "unary-scalar.wgsl.jinja",
39
+ "derive": { "itemsPerInvocation": 4 },
40
+ "bindings": ["x_scalar", "y_scalar", "params_2_unary"],
41
  "dispatch": {
42
  "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
43
  "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
 
49
  {
50
  "id": "same_layout_vec4_tail",
51
  "priority": 19,
52
+ "when": ["numel(shapes.x) > 4", "numel(shapes.x) % 4 != 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)", "ceilDiv(numel(shapes.y), 4) >= ceilDiv(tunables.MIN_TAIL_ELEMENTS | default(0), 4)"],
53
+ "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
 
 
 
 
54
  "passes": [
55
  {
56
  "id": "bulk",
57
+ "name": "Sinh.Vec4Bulk",
58
  "shader": "unary-vec4.wgsl.jinja",
59
  "derive": {
60
+ "vec4PerThread": "tunables.TAIL_VEC4_PER_THREAD | default(4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)"
 
61
  },
62
  "bindings": ["x", "y", "params_unary_tail"],
63
  "dispatch": {
64
+ "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), vec4PerThread)), (tunables.WORKGROUP_SIZE)), 65535)",
65
+ "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), vec4PerThread)), (tunables.WORKGROUP_SIZE)), 65535)",
66
  "z": 1
67
  }
68
  },
69
  {
70
  "id": "tail",
71
+ "name": "Sinh.Tail",
72
  "shader": "unary-scalar.wgsl.jinja",
73
+ "derive": { "itemsPerInvocation": 4, "tailOnly": true },
74
+ "bindings": ["x_scalar", "y_scalar", "params_2_unary_tail"],
75
  "dispatch": {
76
  "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
77
  "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
 
84
  "bindings": {
85
  "x": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
86
  "y": { "buffer": "storage", "elementType": "$vectorScalar" },
87
+ "params_unary": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }], "name": "params" },
88
+ "x_scalar": { "buffer": "read-only-storage", "name": "x", "elementType": "$scalar" },
89
+ "y_scalar": { "buffer": "storage", "name": "y", "elementType": "$scalar" },
90
+ "params_2_unary": { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] },
 
 
 
 
 
 
 
 
91
  "params_unary_tail": {
 
92
  "struct": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }],
93
  "name": "params"
94
  },
95
  "params_2_unary_tail": {
 
96
  "name": "params",
97
  "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
98
  }
build/webgpu/metadata.json CHANGED
@@ -1,6 +1,6 @@
1
  {
2
  "name": "ai.onnx.Sinh",
3
- "id": "_ai_onnx_sinh_webgpu_b96213f",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
@@ -8,15 +8,15 @@
8
  "algorithm": "sha256",
9
  "files": {
10
  "bench.json": "3w+U8M8rGtbqIqW6WnRLBu62utNSad1GJA4BxXRVLQI=",
11
- "manifest.json": "aHgpF84gGKJOYEGx+b/5LoYGD6Y3AjLS7fpHIPHRyHk=",
12
  "test.json": "tftTmQffHsMEjVVdLSM2Vhf9OWNywgma4+Ddv8g4t+k=",
13
- "unary-scalar.wgsl.jinja": "KcPlNLhcmGz6ub7cHgHUbQmKjvDO9c7lI7lplxt+ytM=",
14
- "unary-vec4.wgsl.jinja": "GJq1g0QZFyAB3YUoOOwZrsvc/stvs7kMoZDo8KW2Wxo="
15
  }
16
  },
17
- "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
18
  "webgpu": {
19
- "manifestSpec": "2.0",
20
  "variants": {
21
  "same_layout_vec4": ["unary-vec4.wgsl.jinja"],
22
  "elementwise": ["unary-scalar.wgsl.jinja"],
 
1
  {
2
  "name": "ai.onnx.Sinh",
3
+ "id": "_ai_onnx_sinh_webgpu_9d85402",
4
  "version": 1,
5
  "license": "Apache-2.0",
6
  "backend": { "type": "webgpu" },
 
8
  "algorithm": "sha256",
9
  "files": {
10
  "bench.json": "3w+U8M8rGtbqIqW6WnRLBu62utNSad1GJA4BxXRVLQI=",
11
+ "manifest.json": "ucdixchthLp8GigedAYpc8MjLccK0iyikqWu9OBSEX0=",
12
  "test.json": "tftTmQffHsMEjVVdLSM2Vhf9OWNywgma4+Ddv8g4t+k=",
13
+ "unary-scalar.wgsl.jinja": "q80Bk3YnC/RTGDEKsD931n1+yr7d+mYjpgCH8xluDdA=",
14
+ "unary-vec4.wgsl.jinja": "bFAekisCfkkfXspsIECH0ShKGaP/+YQDcHmk29BGfsQ="
15
  }
16
  },
17
+ "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
18
  "webgpu": {
19
+ "manifestSpec": "2.1",
20
  "variants": {
21
  "same_layout_vec4": ["unary-vec4.wgsl.jinja"],
22
  "elementwise": ["unary-scalar.wgsl.jinja"],
build/webgpu/unary-scalar.wgsl.jinja CHANGED
@@ -22,15 +22,11 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
22
  for (var i = begin; i < end; i = i + 1u) {
23
  {%- endmacro %}
24
  {% macro flat_tail_close() %}
25
- }
26
- {% endmacro %}
27
-
28
- // Scalar unary elementwise implementation. Specialization emits only the
29
- // selected operation and any numerical helper it requires.
30
- {% if usesF16 %}
31
- enable f16;
32
- {% endif %}
33
  {{ env.wgsl.resourceDeclarations }}
 
34
  {% macro emit_sinh_safe() %}
35
  fn sinh_safe(x: f32) -> f32 {
36
  // Beyond the finite boundary, sinh is signed infinity. Just below it,
@@ -44,10 +40,11 @@ fn sinh_safe(x: f32) -> f32 {
44
  return select(magnitude, -magnitude, x < 0.0);
45
  }
46
  return sinh(x);
47
- }{% endmacro %}
 
48
  {{ emit_sinh_safe() }}
49
  {{ flat_tail_open() }}
50
  let v = f32(x[i]);
51
  y[i] = {{ scalar }}(sinh_safe(v));
52
- {{ flat_tail_close() -}}
53
  }
 
22
  for (var i = begin; i < end; i = i + 1u) {
23
  {%- endmacro %}
24
  {% macro flat_tail_close() %}
25
+ }{% endmacro %}
26
+ // Scalar unary elementwise kernel: applies the selected operation, with the
27
+ // numerical helpers it uses, to each element.
 
 
 
 
 
28
  {{ env.wgsl.resourceDeclarations }}
29
+
30
  {% macro emit_sinh_safe() %}
31
  fn sinh_safe(x: f32) -> f32 {
32
  // Beyond the finite boundary, sinh is signed infinity. Just below it,
 
40
  return select(magnitude, -magnitude, x < 0.0);
41
  }
42
  return sinh(x);
43
+ }
44
+ {% endmacro %}
45
  {{ emit_sinh_safe() }}
46
  {{ flat_tail_open() }}
47
  let v = f32(x[i]);
48
  y[i] = {{ scalar }}(sinh_safe(v));
49
+ {{ flat_tail_close() }}
50
  }
build/webgpu/unary-vec4.wgsl.jinja CHANGED
@@ -15,36 +15,45 @@ fn sinh_safe(x: f32) -> f32 {
15
  return select(magnitude, -magnitude, x < 0.0);
16
  }
17
  return sinh(x);
18
- }{% endmacro %}
 
19
  {{ emit_sinh_safe() }}
20
-
21
- {% set vec4PerThread = vec4PerThread %}
 
 
 
 
 
 
 
 
22
  {% if vec4PerThread > 1 %}
23
  const ITEMS: u32 = {{ vec4PerThread }}u;
24
- {% endif %}
25
 
 
26
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
27
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
28
- // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
29
- // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
30
  {% if vec4PerThread > 1 %}
31
  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
32
  // access consecutive words on every step, while each lane can keep several
33
  // independent loads in flight.
34
- let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
35
  let span = (params.count + ITEMS - 1u) / ITEMS;
 
 
 
 
 
 
36
  for (var j = 0u; j < ITEMS; j = j + 1u) {
37
  let i = tid + j * span;
38
  if (i >= params.count) {
39
  break;
40
  }
41
  {% else %}
42
- let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
43
- if (i >= params.count) {
44
- return;
45
- }
46
  {% endif %}
47
-
48
  let xv = x[i];
49
  let fv = vec4<f32>(xv);
50
  y[i] = {{ vectorScalar }}(vec4<f32>(sinh_safe(fv.x), sinh_safe(fv.y), sinh_safe(fv.z), sinh_safe(fv.w)));
 
15
  return select(magnitude, -magnitude, x < 0.0);
16
  }
17
  return sinh(x);
18
+ }
19
+ {% endmacro %}
20
  {{ emit_sinh_safe() }}
21
+ {% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
22
+ {% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
23
+ // 2D-folded flat index: gid.y carries the high bits past the dispatch's
24
+ // per-axis workgroup fold width.
25
+ {% if bound == "" %}
26
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% else %}
27
+ let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
28
+ if ({{ name }} >= {{ bound }}) {
29
+ return;
30
+ }{% endif %}{% endmacro %}
31
  {% if vec4PerThread > 1 %}
32
  const ITEMS: u32 = {{ vec4PerThread }}u;
 
33
 
34
+ {% endif %}
35
  @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
36
  fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 
 
37
  {% if vec4PerThread > 1 %}
38
  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
39
  // access consecutive words on every step, while each lane can keep several
40
  // independent loads in flight.
41
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE, "tid", "") }}
42
  let span = (params.count + ITEMS - 1u) / ITEMS;
43
+ // Lanes at or past span exist only because the dispatch rounds up to whole
44
+ // workgroups. Lane span + k would start on lane k's second group and rewrite
45
+ // up to ITEMS - 1 groups another lane already stored.
46
+ if (tid >= span) {
47
+ return;
48
+ }
49
  for (var j = 0u; j < ITEMS; j = j + 1u) {
50
  let i = tid + j * span;
51
  if (i >= params.count) {
52
  break;
53
  }
54
  {% else %}
55
+ {{ flat_index_2d(tunables.WORKGROUP_SIZE) }}
 
 
 
56
  {% endif %}
 
57
  let xv = x[i];
58
  let fv = vec4<f32>(xv);
59
  y[i] = {{ vectorScalar }}(vec4<f32>(sinh_safe(fv.x), sinh_safe(fv.y), sinh_safe(fv.z), sinh_safe(fv.w)));