feat(s2): elementwise bitwise kernels (and/or/xor/not) on gfx1151 + AVX-512 - #140
Conversation
…VX-512 Adds the S2 bitwise family over i32 integers — the integer sibling of the logical lane (#139) — on both devices. and/or/xor (binary) + not (unary), acting on the full bit pattern (no normalization, unlike logical). ROCm (gfx1151): - `tessera_rocm.bitwise` ODS op + `generate-rocm-bitwise-kernel` pass: arith.{andi,ori,xori} on i32; not = a ^ -1. Binary kinds emit a 4-arg signature; not emits a 3-arg (unary) signature. - runtime `rocm_bitwise_compiled` lane (arity-aware, i32 in/out) + executor table + execution_matrix catalog/row; `_ROCM_COMPILED` + `_NUMERICAL_FIXTURES`. - test_rocm_bitwise_compiled.py: validated on gfx1151 vs numpy across kind × shape incl. rank-3 + signed inputs, + a GPU-free codegen gate. x86 (AVX-512): - avx512_bitwise_i32.cpp: `_mm512_{and,or,xor}_si512` (16 i32/vector); not = xor all-ones. Validated standalone (test_bitwise.cpp, ALL PASSED). No x86:fused slot claimed (Decision #25). Completes the S2 logical-family registry gaps (logical_* in #139 + bitwise_* here). Regenerated drift dashboards + ROCM_AUDIT.md. 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: 61f709c941
ℹ️ 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".
| ctypes.c_int64(0), ctypes.c_int64(size), ctypes.c_int64(1)] | ||
|
|
||
| for arr_in in inputs: | ||
| host = np.ascontiguousarray(arr_in, dtype=np.int32).reshape(-1) |
There was a problem hiding this comment.
When a rocm_bitwise_compiled artifact is launched with non-int32 arrays, this copy silently casts each operand to int32 before feeding the memref<?xi32> kernel. That truncates or reinterprets unsupported inputs such as int64/uint64/float values and returns wrapped 32-bit bitwise results instead of failing, while the manifest and execution row advertise only i32 support and neighboring ROCm lanes reject unsupported dtypes. Please validate all operands are np.int32 before copying them to the device (or return a clear error).
Useful? React with 👍 / 👎.
1 participant
Adds the S2 bitwise family over i32 integers — the integer sibling of the logical lane (#139) — on both devices (gfx1151 + AVX-512).
and/or/xor(binary) +not(unary), acting on the full bit pattern (no normalization, unlike logical). Together with #139 this closes the registry's entirelogicalcategory (logical_*+bitwise_*).ROCm (gfx1151)
tessera_rocm.bitwiseODS op +generate-rocm-bitwise-kernelpass:arith.{andi,ori,xori}on i32;not=a ^ -1. Binary kinds emit a 4-arg signature;notemits a 3-arg (unary) signature.rocm_bitwise_compiledlane (arity-aware, i32 in/out) + executor table + execution-matrix catalog/row +_ROCM_COMPILED/_NUMERICAL_FIXTURES.test_rocm_bitwise_compiled.py: 19 tests pass on gfx1151 vs numpy across kind × shape incl. rank-3 + signed inputs, plus a GPU-free codegen gate (4 kinds → ROCDL, 3-arg/4-arg +memref<?xi32>asserted, bad kind rejected).x86 (AVX-512)
avx512_bitwise_i32.cpp:_mm512_{and,or,xor}_si512(16 i32/vector);not= xor all-ones.test_bitwise.cpp: validated standalone on the AVX-512 host — ALL PASSED.Honesty / scope
x86:fusedslot claimed (Decision Apple GPU MLA: weight-absorption decode (the bandwidth win) + KV-cache shape #25). Regenerated drift dashboards (runtime_abi,runtime_execution_matrix,test_coverage) +ROCM_AUDIT.md. mypy + ruff clean.Test plan
not=3-arg, rest=4-arg, i32 memrefstest_rocm_bitwise_compiled.py— 19 passed on gfx1151test_bitwise.cpp— ALL PASSED on AVX-512 host🤖 Generated with Claude Code