sync 6fdf6301e2bb

Browse files

Files changed (6) hide show

README.md +3 -3
build/webgpu/manifest.json +19 -41
build/webgpu/metadata.json +7 -7
build/webgpu/test.json +2 -2
build/webgpu/unary-scalar.wgsl.jinja +7 -12
build/webgpu/unary-vec4.wgsl.jinja +20 -13

README.md CHANGED Viewed

@@ -39,14 +39,14 @@ See the [ONNX `Tanh` spec](https://onnx.ai/onnx/operators/onnx__Tanh.html) for t
 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
-- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
 - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
 - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
-npm install --save-exact @huggingface/kernels@0.0.1-preview.2
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
@@ -60,5 +60,5 @@ Replace each `*Data` placeholder with a typed array containing the corresponding
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/ai.onnx.Tanh", { version: 1 });
-const { y } = await kernel({ x: { data: xData, shape: [] } });
 ```

 - [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, provenance)
 - [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
 - [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark cases
 - [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
 - [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/ai.onnx.Tanh", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [3] } });
 ```

build/webgpu/manifest.json CHANGED Viewed

@@ -11,20 +11,13 @@
       "id": "same_layout_vec4",
       "priority": 20,
       "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
-      "derive": {
-        "scalar": "dtypes.T",
-        "usesF16": "dtypes.T == \"f16\"",
-        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
-      },
       "passes": [
         {
           "id": "main",
-          "name": "Tanh.vec4",
           "shader": "unary-vec4.wgsl.jinja",
-          "derive": {
-            "op": "\"tanh\"",
-            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
-          },
           "bindings": ["x", "y", "params_unary"],
           "dispatch": {
             "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 @@
     {
       "id": "scalar",
       "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
-      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Tanh",
           "shader": "unary-scalar.wgsl.jinja",
-          "derive": { "op": "\"tanh\"", "itemsPerInvocation": 4 },
-          "bindings": ["x_2", "y_2", "params_2_unary"],
           "dispatch": {
             "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -56,34 +49,29 @@
     {
       "id": "same_layout_vec4_tail",
       "priority": 19,
-      "when": ["numel(shapes.x) > 4", "numel(shapes.x) % 4 != 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
-      "derive": {
-        "scalar": "dtypes.T",
-        "usesF16": "dtypes.T == \"f16\"",
-        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
-      },
       "passes": [
         {
           "id": "bulk",
-          "name": "Tanh.vec4Bulk",
           "shader": "unary-vec4.wgsl.jinja",
           "derive": {
-            "op": "\"tanh\"",
-            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
           },
           "bindings": ["x", "y", "params_unary_tail"],
           "dispatch": {
-            "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
-            "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
         },
         {
           "id": "tail",
-          "name": "Tanh.tail",
           "shader": "unary-scalar.wgsl.jinja",
-          "derive": { "op": "\"tanh\"", "itemsPerInvocation": 4, "tailOnly": true },
-          "bindings": ["x_2", "y_2", "params_2_unary_tail"],
           "dispatch": {
             "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
@@ -96,25 +84,15 @@
   "bindings": {
     "x": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
     "y": { "buffer": "storage", "elementType": "$vectorScalar" },
-    "params_unary": {
-      "buffer": "uniform",
-      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }],
-      "name": "params"
-    },
-    "x_2": { "buffer": "read-only-storage", "name": "x", "elementType": "$scalar" },
-    "y_2": { "buffer": "storage", "name": "y", "elementType": "$scalar" },
-    "params_2_unary": {
-      "buffer": "uniform",
-      "name": "params",
-      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
-    },
     "params_unary_tail": {
-      "buffer": "uniform",
       "struct": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }],
       "name": "params"
     },
     "params_2_unary_tail": {
-      "buffer": "uniform",
       "name": "params",
       "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
     }

       "id": "same_layout_vec4",
       "priority": 20,
       "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
+          "name": "Tanh.Vec4",
           "shader": "unary-vec4.wgsl.jinja",
+          "derive": { "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1" },
           "bindings": ["x", "y", "params_unary"],
           "dispatch": {
             "x": "min(ceilDiv((ceilDiv(numel(shapes.y) / 4, 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "scalar",
       "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Tanh",
           "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "itemsPerInvocation": 4 },
+          "bindings": ["x_scalar", "y_scalar", "params_2_unary"],
           "dispatch": {
             "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
     {
       "id": "same_layout_vec4_tail",
       "priority": 19,
+      "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)"],
+      "derive": { "scalar": "dtypes.T", "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "bulk",
+          "name": "Tanh.Vec4Bulk",
           "shader": "unary-vec4.wgsl.jinja",
           "derive": {
+            "vec4PerThread": "tunables.TAIL_VEC4_PER_THREAD | default(4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)"
           },
           "bindings": ["x", "y", "params_unary_tail"],
           "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), vec4PerThread)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), vec4PerThread)), (tunables.WORKGROUP_SIZE)), 65535)",
             "z": 1
           }
         },
         {
           "id": "tail",
+          "name": "Tanh.Tail",
           "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "itemsPerInvocation": 4, "tailOnly": true },
+          "bindings": ["x_scalar", "y_scalar", "params_2_unary_tail"],
           "dispatch": {
             "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
             "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
   "bindings": {
     "x": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
     "y": { "buffer": "storage", "elementType": "$vectorScalar" },
+    "params_unary": { "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }], "name": "params" },
+    "x_scalar": { "buffer": "read-only-storage", "name": "x", "elementType": "$scalar" },
+    "y_scalar": { "buffer": "storage", "name": "y", "elementType": "$scalar" },
+    "params_2_unary": { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] },
     "params_unary_tail": {
       "struct": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }],
       "name": "params"
     },
     "params_2_unary_tail": {
       "name": "params",
       "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,6 +1,6 @@
 {
   "name": "ai.onnx.Tanh",
-  "id": "_ai_onnx_tanh_webgpu_e799420",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
@@ -8,15 +8,15 @@
     "algorithm": "sha256",
     "files": {
       "bench.json": "RnXzzrufm42zv2CsF1m967BosaRKOF9QYe3VV219CEs=",
-      "manifest.json": "IXlIZTqkPYG1aflezhJhBZhuofy7SJTnUQ1Gqdrpkc8=",
-      "test.json": "2BDGX+xy3u+uSe/U8FVXTWYznJzeSBv8BczW/s06ixI=",
-      "unary-scalar.wgsl.jinja": "HaNuP3u6dDbIs+NJvsK1lzzK6QWnlWLHHc0T/qho7g4=",
-      "unary-vec4.wgsl.jinja": "pz7xSOlVXRtQ1i7OpGZzVYi+x0cY/DgX73TGGFNOXmM="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": {
       "same_layout_vec4": ["unary-vec4.wgsl.jinja"],
       "scalar": ["unary-scalar.wgsl.jinja"],

 {
   "name": "ai.onnx.Tanh",
+  "id": "_ai_onnx_tanh_webgpu_80179bf",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
     "algorithm": "sha256",
     "files": {
       "bench.json": "RnXzzrufm42zv2CsF1m967BosaRKOF9QYe3VV219CEs=",
+      "manifest.json": "aaVwij2m5ZcwSI0h75dtzCnu/AhdoTzwKbAUZZxZz9k=",
+      "test.json": "EM9ZVdzYCk7XIYovEuL1dePwuWNWOAZs8oaWeuA/2Rc=",
+      "unary-scalar.wgsl.jinja": "ltP+E80gaUwt38kWaUo2bwPm+Uk3l1K9e9U+DmxBx28=",
+      "unary-vec4.wgsl.jinja": "4ccLQCNyFZb8+HLGq1UKpUKV+gEv+6VmV/bWC8h58c0="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": {
       "same_layout_vec4": ["unary-vec4.wgsl.jinja"],
       "scalar": ["unary-scalar.wgsl.jinja"],

build/webgpu/test.json CHANGED Viewed

@@ -199,7 +199,7 @@
     {
       "name": "f16_scalar_fallback_saturation_and_tiny",
       "provenance": {
-        "source": "repository-authored",
         "notes": "Seven float16 elements exercise scalar processing. Exact float16 inputs cover saturation at magnitude 10 and the tiny-normal identity branch at 2^-13 without relying on subnormal preservation."
       },
       "inputs": {
@@ -214,7 +214,7 @@
     {
       "name": "f16_scalar_fallback_midrange",
       "provenance": {
-        "source": "repository-authored",
         "notes": "Second odd-numel (5%4!=0) f16 scalar-fallback case over the steep mid-range of tanh where output precision matters most; complements the saturation/tiny case. f16-exact inputs (quarter-fractions)."
       },
       "inputs": {

     {
       "name": "f16_scalar_fallback_saturation_and_tiny",
       "provenance": {
+        "source": "synthetic",
         "notes": "Seven float16 elements exercise scalar processing. Exact float16 inputs cover saturation at magnitude 10 and the tiny-normal identity branch at 2^-13 without relying on subnormal preservation."
       },
       "inputs": {
     {
       "name": "f16_scalar_fallback_midrange",
       "provenance": {
+        "source": "synthetic",
         "notes": "Second odd-numel (5%4!=0) f16 scalar-fallback case over the steep mid-range of tanh where output precision matters most; complements the saturation/tiny case. f16-exact inputs (quarter-fractions)."
       },
       "inputs": {

build/webgpu/unary-scalar.wgsl.jinja CHANGED Viewed

@@ -22,16 +22,11 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
-  }
-{% endmacro %}
-// Scalar unary elementwise implementation. Specialization emits only the
-// selected operation and any numerical helper it requires.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-{% macro emit_tanh_safe() %}
 fn tanh_safe(x: f32) -> f32 {
   // tanh rounds to its saturated value for these tails in f32. Return that
   // value directly, including for infinite input, before invoking the builtin.
@@ -41,9 +36,9 @@ fn tanh_safe(x: f32) -> f32 {
   // relying on backend-specific builtin behavior near zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
-}{% endmacro %}
-{{ emit_tanh_safe() }}
 {{ flat_tail_open() }}
   y[i] = {{ scalar }}(tanh_safe(f32(x[i])));
-{{ flat_tail_close() -}}
 }

   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
+  }{% endmacro %}
+// Scalar unary elementwise kernel: applies the selected operation, with the
+// numerical helpers it uses, to each element.
 {{ env.wgsl.resourceDeclarations }}
 fn tanh_safe(x: f32) -> f32 {
   // tanh rounds to its saturated value for these tails in f32. Return that
   // value directly, including for infinite input, before invoking the builtin.
   // relying on backend-specific builtin behavior near zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
+}
 {{ flat_tail_open() }}
   y[i] = {{ scalar }}(tanh_safe(f32(x[i])));
+{{ flat_tail_close() }}
 }

build/webgpu/unary-vec4.wgsl.jinja CHANGED Viewed

@@ -2,7 +2,6 @@
 // component.
 {{ env.wgsl.resourceDeclarations }}
-{% macro emit_tanh_safe() %}
 fn tanh_safe(x: f32) -> f32 {
   // tanh rounds to its saturated value for these tails in f32. Return that
   // value directly, including for infinite input, before invoking the builtin.
@@ -12,36 +11,44 @@ fn tanh_safe(x: f32) -> f32 {
   // relying on backend-specific builtin behavior near zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
-}{% endmacro %}
-{{ emit_tanh_safe() }}
-{% set vec4PerThread = vec4PerThread %}
 {% if vec4PerThread > 1 %}
 const ITEMS: u32 = {{ vec4PerThread }}u;
-{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
-  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
 {% if vec4PerThread > 1 %}
   // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
   // access consecutive words on every step, while each lane can keep several
   // independent loads in flight.
-  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   let span = (params.count + ITEMS - 1u) / ITEMS;
   for (var j = 0u; j < ITEMS; j = j + 1u) {
     let i = tid + j * span;
     if (i >= params.count) {
       break;
     }
 {% else %}
-  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
-  if (i >= params.count) {
-    return;
-  }
 {% endif %}
   let xv = x[i];
   let fv = vec4<f32>(xv);
   y[i] = {{ vectorScalar }}(vec4<f32>(tanh_safe(fv.x), tanh_safe(fv.y), tanh_safe(fv.z), tanh_safe(fv.w)));

 // component.
 {{ env.wgsl.resourceDeclarations }}
 fn tanh_safe(x: f32) -> f32 {
   // tanh rounds to its saturated value for these tails in f32. Return that
   // value directly, including for infinite input, before invoking the builtin.
   // relying on backend-specific builtin behavior near zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
+}
+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};{% else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }{% endif %}{% endmacro %}
 {% if vec4PerThread > 1 %}
 const ITEMS: u32 = {{ vec4PerThread }}u;
+{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {% if vec4PerThread > 1 %}
   // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
   // access consecutive words on every step, while each lane can keep several
   // independent loads in flight.
+{{ flat_index_2d(tunables.WORKGROUP_SIZE, "tid", "") }}
   let span = (params.count + ITEMS - 1u) / ITEMS;
+  // Lanes at or past span exist only because the dispatch rounds up to whole
+  // workgroups. Lane span + k would start on lane k's second group and rewrite
+  // up to ITEMS - 1 groups another lane already stored.
+  if (tid >= span) {
+    return;
+  }
   for (var j = 0u; j < ITEMS; j = j + 1u) {
     let i = tid + j * span;
     if (i >= params.count) {
       break;
     }
 {% else %}
+{{ flat_index_2d(tunables.WORKGROUP_SIZE) }}
 {% endif %}
   let xv = x[i];
   let fv = vec4<f32>(xv);
   y[i] = {{ vectorScalar }}(vec4<f32>(tanh_safe(fv.x), tanh_safe(fv.y), tanh_safe(fv.z), tanh_safe(fv.w)));