sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +64 -0
build/webgpu/bench.json +31 -0
build/webgpu/manifest.json +121 -0
build/webgpu/metadata.json +18 -0
build/webgpu/quick-gelu.wgsl.jinja +77 -0
build/webgpu/test.json +248 -0

README.md CHANGED Viewed

@@ -1,3 +1,67 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# com.microsoft.QuickGelu
+`com.microsoft`  ·  ONNX Runtime contrib operator  ·  contrib since_version 1
+## Description
+Computes `x * sigmoid(alpha * x)` elementwise, a fast approximation of GELU activation. The output has the same shape as the input. This WebGPU package implements float16 and float32; the schema-allowed double and bfloat16 types are not supported.
+See the [ONNX Runtime `QuickGelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.QuickGelu) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `X` | `T` | — | — | Input tensor of any shape. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `Y` | `T` | same as `X` | same as `X` | Output tensor; same shape as the input. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `alpha` | `1.702` | Scalar multiplier applied to `x` inside the sigmoid; defaults to 1.702, which approximates GELU. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, 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
+- [`quick-gelu.wgsl.jinja`](build/webgpu/quick-gelu.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/com.microsoft.QuickGelu", { version: 1 });
+const { Y } = await kernel({ X: { data: XData, shape: [2, 4] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,31 @@

+{
+  "op": "com.microsoft.QuickGelu",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "quickgelu-f32-4096x3072",
+      "preset": "smoke",
+      "vars": { "dtype": "float32" },
+      "inputs": { "X": { "shape": [4096, 3072], "dtype": "float32", "dist": "normal", "seed": 720, "scale": 2 } },
+      "outputs": { "Y": { "shape": [4096, 3072], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4096 * 3072 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "quickgelu-f16-4096x3072",
+      "preset": "model",
+      "vars": { "dtype": "float16" },
+      "inputs": { "X": { "shape": [4096, 3072], "dtype": "float16", "dist": "normal", "seed": 721, "scale": 2 } },
+      "outputs": { "Y": { "shape": [4096, 3072], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "4096 * 3072 * 2 * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "quickgelu-f32-scalar-cliff-odd-numel",
+      "preset": "stress",
+      "attrs": { "alpha": 1.702 },
+      "vars": { "dtype": "float32" },
+      "inputs": { "X": { "shape": [2049, 2047], "dtype": "float32", "dist": "normal", "seed": 7201, "scale": 2 } },
+      "outputs": { "Y": { "shape": [2049, 2047], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "2049 * 2047 * 2 * dtypeBytes(args.dtype)" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,121 @@

+{
+  "domain": "com.microsoft",
+  "name": "QuickGelu",
+  "sinceVersion": 1,
+  "description": "Computes `x * sigmoid(alpha * x)` elementwise, a fast approximation of GELU activation. The output has the same shape as the input. This WebGPU package implements float16 and float32; the schema-allowed double and bfloat16 types are not supported.",
+  "inputs": [{ "role": "X", "dtype": "T", "description": "Input tensor of any shape." }],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.X",
+      "shape": "shapes.X",
+      "description": "Output tensor; same shape as the input."
+    }
+  ],
+  "attributes": { "alpha": 1.702 },
+  "attributeDescriptions": {
+    "alpha": "Scalar multiplier applied to `x` inside the sigmoid; defaults to 1.702, which approximates GELU."
+  },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "X": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
+    "workgroupOk": "tunables.WORKGROUP_SIZE > 0 and tunables.WORKGROUP_SIZE <= deviceWorkgroupCap",
+    "baseOk": "workgroupOk and numel(shapes.X) == numel(shapes.Y) and f16Ok(dtypes.T)",
+    "vec4Ok": "numel(shapes.X) > 0 and numel(shapes.X) % 4 == 0"
+  },
+  "bindingSets": {
+    "scalarTail": [
+      { "name": "x", "arg": "X", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
+      { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "vec4",
+      "priority": 20,
+      "when": ["baseOk", "vec4Ok"],
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+        "vec4": true,
+        "vec4Tail": false
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "QuickGelu.vec4",
+          "source": { "shader": "quick-gelu.wgsl.jinja", "inputs": { "alpha": "attrs.alpha" } },
+          "bindings": [
+            {
+              "name": "x",
+              "arg": "X",
+              "semantic": "X",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "y",
+              "arg": "Y",
+              "semantic": "Y",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "vec4_tail",
+      "priority": 10,
+      "when": ["baseOk", "numel(shapes.X) > 0"],
+      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"", "vec4": false, "vec4Tail": true },
+      "passes": [
+        {
+          "id": "main",
+          "name": "QuickGelu.vec4Tail",
+          "source": { "shader": "quick-gelu.wgsl.jinja", "inputs": { "alpha": "attrs.alpha" } },
+          "bindings": "scalarTail",
+          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "scalar",
+      "priority": 0,
+      "when": ["baseOk", "true"],
+      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"", "vec4": false, "vec4Tail": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "QuickGelu.scalar",
+          "source": { "shader": "quick-gelu.wgsl.jinja", "inputs": { "alpha": "attrs.alpha" } },
+          "bindings": "scalarTail",
+          "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,18 @@

+{
+  "name": "com.microsoft.QuickGelu",
+  "id": "_com_microsoft_quickgelu_webgpu_75eb0ab",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "g3U1AfptsaN7o7RQkF70XpQkOLCsBrZIr0WbKXyeZF4=",
+      "manifest.json": "DEC3xwPqmogreS46virRjZUF4PzWGkJwGMQ0ku4ytok=",
+      "quick-gelu.wgsl.jinja": "vBvqUo8P4UGy3DVJFBNcoTPinSfCxWFoIX77f9s6Zro=",
+      "test.json": "ErgPevOReyIn0tBPjmUMYo73FIwanLWeFzQmyXeSOak="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.QuickGelu" }
+}

build/webgpu/quick-gelu.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,77 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+// com.microsoft.QuickGelu : Y = X * sigmoid(alpha * X)
+//   sigmoid here uses the stable two-branch form so the gate never overflows
+//   for extreme magnitudes (alpha*x = +/-1702 for x = -/+1000 with the default
+//   alpha): the naive 1/(1+exp(-alpha*x)) computes exp(+1702) = Inf and yields
+//   Inf/Inf = NaN, while exp(z)/(1+exp(z)) (z <= 0) and 1/(1+exp(-z)) (z >= 0)
+//   each only ever evaluate exp of a non-positive argument. ALPHA is compiled
+//   as a constant; alpha == 0 collapses to sigmoid(0) = 0.5 exactly (the
+//   z >= 0 branch: 1/(1+exp(0))).
+const ALPHA: f32 = f32({{ source.alpha }});
+fn sigmoid_stable(z: f32) -> f32 {
+  if (z >= 0.0) {
+    return 1.0 / (1.0 + exp(-z));
+  }
+  let e = exp(z);
+  return e / (1.0 + e);
+}
+fn quick_gelu(v: f32) -> f32 {
+  return v * sigmoid_stable(ALPHA * v);
+}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d() }}
+{% if vec4Tail %}
+  let base = i * 4u;
+{% for lane in range(4) %}
+  if (base + {{ lane }}u < params.count) {
+    y[base + {{ lane }}u] = {{ scalar }}(quick_gelu(f32(x[base + {{ lane }}u])));
+  }
+{% endfor %}
+{% elif vec4 %}
+  let fv = vec4<f32>(x[i]);
+  y[i] = vec4<{{ scalar }}>(vec4<f32>(quick_gelu(fv.x), quick_gelu(fv.y), quick_gelu(fv.z), quick_gelu(fv.w)));
+{% else %}
+  y[i] = {{ scalar }}(quick_gelu(f32(x[i])));
+{% endif %}
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,248 @@

+{
+  "op": "com.microsoft.QuickGelu",
+  "cases": [
+    {
+      "name": "dispatch_cliff_scalar_over_16M",
+      "attrs": { "alpha": 1.702 },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [16776961], "data": { "kind": "linspace", "start": -4.0, "end": 4.0 } }
+      },
+      "outputs": { "Y": { "dtype": "float32", "shape": [16776961], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "ort_default_alpha_extreme_safe_sigmoid",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu",
+        "notes": "Default alpha path over very large magnitudes; the safe sigmoid should not overflow."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "data": { "kind": "values", "values": [-1000.0, -100.0, -10.0, -1.0, 0.0, 1.0, 10.0, 100.0, 1000.0] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [0.0, 0.0, -2.980232238769531e-7, -0.15420421957969666, 0.0, 0.845795750617981, 10.0, 100.0, 1000.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_alpha_one_matches_silu_edges",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu"
+      },
+      "attrs": { "alpha": 1 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": { "kind": "values", "values": [-100.0, -10.0, -1.0, 0.0, 1.0, 10.0, 100.0] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [7],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [0.0, -0.0004538893699645996, -0.2689414322376251, 0.0, 0.7310585975646973, 9.99954605102539, 100.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_negative_alpha_flips_gate",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu"
+      },
+      "attrs": { "alpha": -1.702 },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-3.0, -1.0, 0.0, 1.0, 3.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [5],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [-2.981928825378418, -0.845795750617981, 0.0, 0.15420421957969666, 0.018071293830871582]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_negative_alpha_extreme_safe_sigmoid",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu",
+        "notes": "Upstream negative-alpha branch extended over the same large magnitudes as the positive-alpha ORT vector to guard the stable sigmoid reformulation."
+      },
+      "attrs": { "alpha": -1.702 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "data": { "kind": "values", "values": [-1000.0, -100.0, -10.0, -1.0, 0.0, 1.0, 10.0, 100.0, 1000.0] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [1, 9],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [-1000.0, -100.0, -10.0, -0.845795750617981, 0.0, 0.15420423448085785, 4.0579612914370955e-7, 0.0, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_empty_rank4",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "notes": "Empty tensors should preserve shape and produce no values."
+      },
+      "inputs": { "X": { "dtype": "float32", "shape": [1, 0, 2, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 2, 3], "data": { "kind": "values", "values": [] } } }
+    },
+    {
+      "name": "vec4_default_alpha_mixed_magnitudes",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu",
+        "notes": "numel divisible by 4 so the vec4 variant is exercised; default alpha over a spread of signs/magnitudes."
+      },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [-2.0, -0.5, 0.0, 0.5, 2.0, 1.0, -1.0, 4.0] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [-0.06434137374162674, -0.14961156249046326, 0.0, 0.35038843750953674, 1.935658574104309, 0.845795750617981, -0.15420423448085785, 3.9955852031707764]
+          }
+        }
+      }
+    },
+    {
+      "name": "alpha_zero_halves_input",
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu",
+        "notes": "Additional edge: alpha=0 makes sigmoid(alpha*x)=0.5 for every finite x."
+      },
+      "attrs": { "alpha": 0 },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-3.0, -1.0, 0.0, 1.0, 3.0] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-1.5, -0.5, 0.0, 0.5, 1.5] }
+        }
+      }
+    },
+    {
+      "name": "alpha_zero_subnormal_half_input_vec4_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu",
+        "notes": "With alpha=0, QuickGelu is exactly x/2, so signed subnormal inputs should not flush to zero in the vec4 path."
+      },
+      "attrs": { "alpha": 0 },
+      "inputs": {
+        "X": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38, -1e-38] }
+        }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39, -4.999999675228202e-39]
+          }
+        }
+      }
+    },
+    {
+      "name": "alpha_zero_subnormal_half_input_scalar_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/contrib_ops/activation_op_test.cc",
+        "test": "ActivationOpTest.QuickGelu",
+        "notes": "Scalar-path companion for alpha=0 exact half-input subnormal behavior."
+      },
+      "attrs": { "alpha": 0 },
+      "inputs": {
+        "X": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } }
+      },
+      "outputs": {
+        "Y": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 2e-45,
+          "data": { "kind": "values", "values": [4.99997305055738e-41, -4.99997305055738e-41, 4.999999675228202e-39] }
+        }
+      }
+    },
+    {
+      "name": "f16_values",
+      "attrs": { "alpha": 1.702 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-4.0, -2.0, -0.5, 0.0, 0.5, 1.0, 2.0, 4.0] }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [8], "tolerance": 0.005 } }
+    },
+    {
+      "name": "f16_default_alpha_extreme_stable_sigmoid",
+      "attrs": { "alpha": 1.702 },
+      "inputs": {
+        "X": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-1000.0, -100.0, -10.0, -1.0, 1.0, 10.0, 100.0, 1000.0] }
+        }
+      },
+      "outputs": { "Y": { "dtype": "float16", "shape": [8], "tolerance": 0.02 } }
+    }
+  ]
+}