Repository navigation
Retain 16-bit conversions in Metal round-trip arithmetic values #2238
Description
Problem
Generated Metal widens source short/ushort arithmetic locals to int/uint without preserving the 16-bit conversion at storage boundaries. Strict compilation succeeds, but round-trip execution differs from the original source.
#include <metal_stdlib>
using namespace metal;
kernel void select_value(const device uint* input [[buffer(0)]],
device uint* output [[buffer(1)]]) {
int groups = 65537;
short narrowed = groups;
const long at = long(narrowed);
const device uint* selected = input + at;
output[0] = selected[0];
}At cb7597e3, the Metal generator maps this local to int narrowed = groups;. The original selects input element 1; the generated artifact selects element 65537. The existing byte conversion helper explicitly excludes short and ushort.
Evidence
The public project/package/native-loader tests use 257 input words (17 + 13 * i), 32 poisoned output guards and input observations. Original Metal returns 30; generated Metal returns 0 on the same device. Eight short-specific controls fail: signed/unsigned storage, assignment, addition/subtraction updates, vector construction and component writes. The same controls pass native OpenGL after the independently tracked #2236/#2237 corrections. Wider scalar controls and byte scalar controls pass.
Required Behavior
- Preserve signed and unsigned 16-bit conversions at initialization, assignment, argument, return and explicit-constructor boundaries while retaining any deliberate widened arithmetic representation.
- Preserve short vector component types and compound-update conversions without repeating index or operand effects.
- Keep physical resource and aggregate storage layouts unchanged.
- Require native original/generated Metal parity, independent integer references, input observations and output guards, plus existing narrow-conversion regressions.
This extends the storage-conversion work in #2023 beyond bytes. It is separate from #2074's aggregate physical-field initializers. Mixed-width locals and vectors occur throughout the pinned MLX kernels, but the fix must apply to ordinary translated source without repository-specific rules.
Activity
Preserve source integer conversions in pointer bounds
NripeshN commentedon Oct 13, 2026
Implemented in 412f8e0f, tracked by #1974.
The change preserves lexical source types and fixed-width storage conversions in bounds, including helper-local narrow vectors, and corrects the corresponding narrow component/Metal conversion emission. Unsafe offset cases remain structured rejections. Native Metal/OpenGL controls retain input observations and output guards; 36 reduced HLSL controls pass strict DXC.
The PR records the full-run and focused verification separately. Windows numerical execution and complete MLX/runtime qualification are not inferred from these compiler checks. The separate private aggregate-state requirement remains open under #2239.
NripeshN commentedon Oct 13, 2026
Strict Compiler Regression in Discarded Postfix Updates
The published
f73dd3f5c71d40e79535780931e1220da023e2b5macOS porting job fails GEMM compilation because a narrowed postfix result is emitted even where the source discards it. A reduced kernel reproduces the same error locally:Metal round-trip generation emits:
xcrun -sdk macosx metal -Werror -crejects the final comma operand withexpression result unused [-Werror,-Wunused-value]. The wide host representation still needs the explicit short conversion; dropping that conversion or weakening compiler warnings would lose the original correctness requirement.The correction needs to distinguish discarded statement/loop-update results from postfix values used by another expression. Preserve exactly one update, source storage-width conversion, prefix/postfix value semantics and evaluation counts. Cover narrow scalar and vector-component updates, used postfix results, loop clauses and standalone statements, and keep strict native compilation and numerical controls.
The full GEMM failure is recorded in the macOS porting job. This is a follow-up to the storage-conversion change in this issue, not evidence that existing GEMM numerical cases passed. Separately, the copy and unary source-export jobs report generated identity mismatches and will require reviewed contract refreshes after compiler fixes are complete.
NripeshN commentedon Oct 13, 2026
An isolated follow-up confirms that discarded postfix values are responsible for the reduced strict-Metal failure. Wrapping only the top-level loop-update result in an explicit
voidconversion preserves the stored narrow update while making its discarded value explicit.Four reduced cases now compile with warnings treated as errors and match their unchanged original Metal sources: short and ushort loop counters, two comma-separated updates, and a postfix result consumed by an output assignment. Each checks 32 untouched output guards. The consumed-result case retains values
[0, 1, 2], so suppressing the unused result must not propagate into nested expressions.This is a process-local prototype, not a repository fix. Production files are frozen during the separate pointer/half-reduction regression. A complete change should cover ordinary discarded expression statements, loop update lists, indexed operands, value-consuming contexts, single evaluation and source-width conversion before publication. The original compiler failure remains retained separately.
NripeshN commentedon Oct 13, 2026
Expanded the isolated strict-Metal reproduction without changing production generation. All 16 cases compile with warnings treated as errors and match the original Metal kernel, with 32 output guards:
forinitializers and update lists.continueandbreak, private struct fields, and a narrow vector update.The prototype scopes discarded-result markers to the exact expression nodes and restores the prior scope afterward. It does not apply a discarded context to nested operands or non-void tail expressions. No production fix is included yet; the pending pointer/binary16 candidate is frozen for its full regression. A permanent change still needs repository tests and integration review. This extends the earlier four-case evidence rather than replacing it.
Preserve discarded narrow-integer updates in Metal
NripeshN commentedon Oct 13, 2026
The discarded-result compiler failure is fixed in
42fb93d2d185dac277023df40f3c04b0dc5e4529on #1974.The implementation uses explicit statement/loop boundaries rather than global discarded-expression state. It retains the narrow storage conversion and postfix computation, discarding only the unused result. Conditions, arguments, assignment operands, explicit returns and non-void tail expressions still consume their values normally.
Permanent tests cover 18 original/generated Metal cases with 32 output guards: byte and short boundaries, prefix/postfix and compound updates, loop initializers, ordered update lists,
continue/break, fields, vectors, returned/argument values and nested shadowing. Eight tail-expression tests check both used and discarded results. The complete required native module passed 381 tests.The unchanged pinned MLX float32 NN GEMM entry now passes strict Metal compilation and original-source/independent-reference parity across six specialization variants and 15 numerical cases. This directly clears the previously recorded GEMM loop-update compiler failure; it is not a full GEMM-family or upstream-suite claim.
The full local regression passed 45,014 tests with 47,727 skipped and no failures. All-file hooks, strict docs and support checks passed. The issue remains open for review of its broader conversion scope and hosted qualification.
Reconcile narrow-integer Metal corpus references
Add conformance and differential test harness (#2256)