feat(s2): elementwise comparison kernels (eq/ne/lt/le/gt/ge) on gfx1151 + AVX-512 - #137
Merged
Merged
Conversation
…51 + AVX-512 Adds the S2 comparison family — the first compiled lane with a non-float (boolean) result — on both devices this box runs. eq/ne/lt/le/gt/ge over float inputs, i8/bool output, NaN semantics matching numpy (ordered everywhere except ne, which is unordered-not-equal). ROCm (gfx1151): - `tessera_rocm.compare` ODS op + `generate-rocm-compare-kernel` pass: flat 2-operand kernel, arith.cmpf (OEQ/UNE/OLT/OLE/OGT/OGE) -> extui i1->i8, memref<?xi8> output. - runtime `rocm_compare_compiled` lane (_build/_execute with asymmetric in/out element sizes: float in, 1-byte bool out) + executor table + execution_matrix catalog/row; `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_compare_compiled.py: validated on gfx1151 vs numpy across kind × dtype × shape incl. rank-3 + NaN semantics, + a GPU-free codegen gate. x86 (AVX-512): - avx512_compare_f32.cpp: `_mm512_cmp_ps_mask` (matching _CMP_*_OQ / _CMP_NEQ_UQ predicates) + `_mm_maskz_set1_epi8` to expand the mask to 0/1 bytes. C's native float operators already match numpy's NaN rule, so the scalar reference uses them directly. Validated standalone (test_compare.cpp, ALL PASSED incl. NaN). No `x86:fused` manifest slot claimed — Graph-IR->x86 lowering still unwired (Decision #25). Regenerated drift dashboards + ROCM_AUDIT.md. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
4 tasks
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.
Adds the S2 comparison family — the first compiled lane with a non-float (boolean) result — on both devices this box runs (gfx1151 + AVX-512).
eq/ne/lt/le/gt/geover float inputs →i8/bool output, with NaN semantics matching numpy. Follows the per-op-family quality bar: a real optimized device kernel and an on-device numerical test vs numpy.ROCm (gfx1151)
tessera_rocm.compareODS op +generate-rocm-compare-kernelpass — flat 2-operand per-element kernel:arith.cmpf→extui i1→i8,memref<?xi8>output. NaN: ordered predicates (OEQ/OLT/OLE/OGT/OGE, NaN→false) exceptne=UNE(NaN→true), matchingnp.not_equal/np.equal.rocm_compare_compiledlane with asymmetric in/out element sizes (float in, 1-byte bool out) + executor table + execution-matrix catalog/row +_ROCM_COMPILED/_NUMERICAL_FIXTURES.test_rocm_compare_compiled.py: 64 tests pass on gfx1151 across kind × dtype × shape incl. rank-3 + dedicated NaN-semantics test, plus a GPU-free codegen gate (6 kinds codegen → ROCDL,memref<?xi8>+cmpfasserted, bad kind rejected).x86 (AVX-512)
avx512_compare_f32.cpp:_mm512_cmp_ps_mask(predicates_CMP_*_OQ/_CMP_NEQ_UQ) +_mm_maskz_set1_epi8(mask, 1)to expand the 16-lane mask to 0/1 bytes (AVX512BW+VL). C's native float operators already follow numpy's NaN rule, so the scalar reference uses them directly.test_compare.cpp: validated standalone on the AVX-512 host — ALL PASSED (aligned/tail/small/wide/degenerate + NaN across all 6 kinds).Honesty / scope
x86:fusedmanifest slot claimed — Graph-IR→x86 lowering isn't wired yet (Decision Apple GPU MLA: weight-absorption decode (the bandwidth win) + KV-cache shape #25), matching the reduce/unary/binary lanes.comparisonfamily (eq/ne/lt/le/gt/ge, all previously_planned). Regenerated drift dashboards (runtime_abi,runtime_execution_matrix,test_coverage) +ROCM_AUDIT.md. mypy + ruff clean.Test plan
tessera-optrebuilds; 6 kinds codegen + lower to ROCDL; bad kind rejected;f32,f32,i8,indexsignaturetest_rocm_compare_compiled.py— 64 passed on gfx1151test_compare.cpp— ALL PASSED on AVX-512 host🤖 Generated with Claude Code