feat(s2): elementwise unary-math kernels (exp/log/sqrt/erf/…) on gfx1151 + AVX-512 - #135
Merged
Merged
Conversation
…151 + AVX-512 The S2 scalar-math / stability family gets a real optimized device kernel on both devices we have hardware for, pairing with the reduction family. ROCm (gfx1151): `tessera_rocm.unary` directive + `generate-rocm-unary-kernel` pass → a flat per-element kernel (one thread/element), the unary sibling of the activation lane. Covers exp/log/sqrt/rsqrt/reciprocal/abs(absolute)/sign/erf/ tanh/sigmoid/log1p/expm1/softplus (softplus stable: log1p(exp(-|x|))+max(x,0)); transcendentals lower through math → ROCDL. New runtime.launch() lane `rocm_unary_compiled`, dispatched by op name; f16/bf16/f32 storage, f32 compute. Validated on gfx1151 vs numpy across kind × dtype × shape (143 tests). x86 (AVX-512): `avx512_unary_f32.cpp` — the algebraic subset (sqrt/rsqrt/ reciprocal/abs/neg/sign) with direct intrinsics (no polynomial approx) + scalar reference + on-device test (31/31). The transcendentals stay numpy-reference on CPU (no fused x86 claim). Wiring: executor table + execution_matrix KNOWN_EXECUTORS/_MATRIX row + 13 _ROCM_COMPILED manifest entries (status `compiled`) + fixtures + ROCM_AUDIT Landed entry + regenerated dashboards (runtime_abi/execution_matrix/test_cov). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 9102444e59
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
Comment on lines
+94
to
+98
| Value pos = b.create<arith::CmpFOp>(loc, ogt, x, zero); | ||
| Value neg = b.create<arith::CmpFOp>(loc, olt, x, zero); | ||
| Value negOne = cst(b, loc, f32, -1.0f); | ||
| Value s = b.create<arith::SelectOp>(loc, neg, negOne, zero); | ||
| y = b.create<arith::SelectOp>(loc, pos, one, s); |
There was a problem hiding this comment.
When tessera.sign is run on tensors containing NaNs, both ordered comparisons here are false, so the compiled ROCm lane writes 0 instead of propagating the NaN. The existing public/reference behavior uses np.sign (python/tessera/__init__.py:3647 and runtime._apple_gpu_unary_numpy), which preserves NaNs, and the new AVX-512 sign path uses the same ordered-compare logic. Please handle unordered inputs before falling back to zero for ±0, otherwise NaN-containing inputs silently diverge from the reference.
Useful? React with 👍 / 👎.
Merged
4 tasks
gstoner
pushed a commit
that referenced
this pull request
Jun 26, 2026
… on gfx1151 + AVX-512 Adds the binary-arithmetic compiled lane — the 2-operand sibling of the elementwise unary-math lane (#135) — closing the S2 binary gaps sub/div/pow/ maximum/minimum on both devices we can run on. ROCm (gfx1151): - `tessera_rocm.binary` ODS op + `generate-rocm-binary-kernel` pass (flat 2-operand per-element kernel, mirrors the silu_mul 2-operand template); pow lowers via math→ROCDL; maximum/minimum are IEEE NaN-propagating. - runtime `rocm_binary_compiled` lane (_build/_execute + executor table + execution_matrix catalog/row); `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_binary_compiled.py: validated on gfx1151 vs numpy across kind × dtype × shape incl. rank-3 + NaN-propagation, + a GPU-free codegen gate. x86 (AVX-512): - avx512_binary_f32.cpp: direct-intrinsic subset sub/div/maximum/minimum (NaN-propagating via unordered-compare blend); pow stays numpy-reference (no fused x86 claim, as unary leaves the transcendentals). Validated standalone (test_binary.cpp, ALL PASSED on the AVX-512 host). No `x86:fused` manifest slot claimed — the Graph-IR→x86 lowering isn't wired yet (kernel-exists != end-to-end-wired, Decision #25). Regenerated drift dashboards (runtime_abi/runtime_execution_matrix/test_coverage) + ROCM_AUDIT.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jun 26, 2026
… on gfx1151 + AVX-512 (#136) * feat(s2): elementwise binary-arithmetic kernels (sub/div/pow/max/min) on gfx1151 + AVX-512 Adds the binary-arithmetic compiled lane — the 2-operand sibling of the elementwise unary-math lane (#135) — closing the S2 binary gaps sub/div/pow/ maximum/minimum on both devices we can run on. ROCm (gfx1151): - `tessera_rocm.binary` ODS op + `generate-rocm-binary-kernel` pass (flat 2-operand per-element kernel, mirrors the silu_mul 2-operand template); pow lowers via math→ROCDL; maximum/minimum are IEEE NaN-propagating. - runtime `rocm_binary_compiled` lane (_build/_execute + executor table + execution_matrix catalog/row); `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_binary_compiled.py: validated on gfx1151 vs numpy across kind × dtype × shape incl. rank-3 + NaN-propagation, + a GPU-free codegen gate. x86 (AVX-512): - avx512_binary_f32.cpp: direct-intrinsic subset sub/div/maximum/minimum (NaN-propagating via unordered-compare blend); pow stays numpy-reference (no fused x86 claim, as unary leaves the transcendentals). Validated standalone (test_binary.cpp, ALL PASSED on the AVX-512 host). No `x86:fused` manifest slot claimed — the Graph-IR→x86 lowering isn't wired yet (kernel-exists != end-to-end-wired, Decision #25). Regenerated drift dashboards (runtime_abi/runtime_execution_matrix/test_coverage) + ROCM_AUDIT.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(tests): refresh suite-by-suite test counts (fast ~11,686 / full ~12,464) The perf-doc drift gate (test_perf_doc_fast_and_full_counts_are_current) went red: the suite grew past the doc's stale ~9,776 fast / ~10,554 full figures (±15% gate). Measured 2026-06-26 via `pytest tests/unit --collect-only`: 11,686 fast / 778 deselected / 12,464 full. Pre-existing drift surfaced by this PR's added tests; updates the table + prose to current counts. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * docs(audit): resync docs_freshness dashboard after rebase onto main consolidation The rebase onto main (5b06bb1 docs consolidation 34→22) shifted the freshness audit to 89 catalogued docs; regenerate so the drift gate matches the merged state. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
4 tasks
gstoner
added a commit
that referenced
this pull request
Jun 26, 2026
…loor/round) on gfx1151 + AVX-512 (#138) Extends the existing elementwise unary-math lane (#135) with 12 more ops, closing the S2 scalar-math trig/special + numeric-helper rounding gaps on both devices. ROCm (gfx1151) — generate-rocm-unary-kernel gains: - trig/special: cos, tan, sinh, cosh, asin, acos, atan, erfc (math.{Cos,Tan,Sinh,Cosh,Asin,Acos,Atan,Erfc} -> ocml) - rounding: floor, ceil, round, trunc; `round` = math.RoundEvenOp to match numpy's round-half-to-even. All 12 lower through the same ROCDL path; runtime _ROCM_UNARY_OPS + manifest + fixtures extended. test_rocm_unary_compiled.py validates all on gfx1151 vs numpy (tan domain bounded inside (-pi/2, pi/2); asin/acos to |x|<=1; erfc ref added) + the GPU-free codegen gate covers the new kinds. x86 (AVX-512): - avx512_unary_f32.cpp gains the rounding subset floor/ceil/trunc/round as direct `_mm512_roundscale_ps` intrinsics (round = ties-to-even / std::nearbyint). Trig/erfc are transcendental and stay numpy-reference on CPU (no fused x86 claim), as the lane already leaves exp/log/erf. Validated standalone (test_unary.cpp, ALL PASSED). No x86:fused slot claimed (Decision #25). Regenerated test_coverage dashboard + ROCM_AUDIT.md. Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
The S2 scalar-math / stability family gets a real optimized device kernel on both devices we have hardware for (gfx1151 + AVX-512 x86), pairing with the reduction family (#133/#134). This is the per-op multi-device completeness bar: a genuine optimized kernel + an on-device numerical test per device, not just the numpy reference.
ROCm (gfx1151)
tessera_rocm.unaryODS op +generate-rocm-unary-kernelpass → a flat per-element kernel (one thread/element), the unary sibling of the activation lane.exp/log/sqrt/rsqrt/reciprocal/abs(absolute)/sign/erf/tanh/sigmoid/log1p/expm1/softplus. Softplus uses the stable formlog1p(exp(-|x|)) + max(x,0); transcendentals lower through themath→ ROCDL path.runtime.launch()lanerocm_unary_compiled, dispatched by op name; f16/bf16/f32 storage, f32 compute.test_rocm_unary_compiled.py, 143 passed).x86 (AVX-512)
avx512_unary_f32.cpp— the algebraic subset (sqrt/rsqrt/reciprocal/abs/neg/sign) with direct AVX-512 intrinsics (no polynomial approximation) + scalar reference + on-device test (test_unary.cpp, 31/31 on the AVX-512 host).Wiring
Executor table +
execution_matrixKNOWN_EXECUTORS/_MATRIXrow + 13_ROCM_COMPILEDmanifest entries (statuscompiled) +_NUMERICAL_FIXTURES+ ROCM_AUDIT Landed entry + regenerated dashboards (runtime_abi/runtime_execution_matrix/test_coverage). mypy + ruff clean; 17 generated docs in sync; gate/drift tests green.🤖 Generated with Claude Code