minimal_repro.swift

minimal_repro.swift

// Minimal reproduction: the Metal runtime compiler miscompiles integer
// division of a mulhi() result by a constant.
//
//     uint hi = mulhi(x, 500u);   // high 32 bits of x * 500, always < 500
//     out      = hi / 20u;        // expected < 25
//
// On Apple M1 Max / macOS 26.5.1 the observed result exactly matches division
// of the LOW 32 bits of the product instead of the high 32:
//
//     observed == uint((x &* 500) & 0xFFFF_FFFF) / 20    for every x tested
//
// The same failure pattern reproduces across tested non-power-of-two constant
// divisors and for modulo. Power-of-two divisors, runtime divisors, and tested
// non-division uses of the mulhi result are compiled correctly.
//
// Run:  swift minimal_repro.swift
import Metal

let source = """
#include <metal_stdlib>
using namespace metal;
kernel void divtest(device const uint *inp [[buffer(0)]],
                    device uint *outp      [[buffer(1)]],
                    uint tid [[thread_position_in_grid]]) {
    uint x  = inp[tid];
    uint hi = mulhi(x, 500u);  // = (ulong(x) * 500) >> 32, so hi < 500
    outp[tid] = hi / 20u;      // must be < 25
}
"""

guard let dev = MTLCreateSystemDefaultDevice() else { fatalError("no Metal device") }
print("Metal integer-division miscompile reproduction\n")
print("Metal device: \(dev.name)")
print("Shader language: Metal Shading Language (MSL)")
print("Compilation API: MTLDevice.makeLibrary(source:options: nil)")
print("Kernel expression: mulhi(x, 500u) / 20u")
print("CPU reference: Swift UInt64 arithmetic (not Metal)\n")
let lib = try dev.makeLibrary(source: source, options: nil)  // runtime compiler
let pso = try dev.makeComputePipelineState(function: lib.makeFunction(name: "divtest")!)

let xs: [UInt32] = [1, 1_000_000, 2_068_421_821, 3_538_566_664, 4_294_967_295,
                    1_408_978_359, 1_553_329_576]
let inBuf = dev.makeBuffer(bytes: xs, length: xs.count * 4, options: .storageModeShared)!
let outBuf = dev.makeBuffer(length: xs.count * 4, options: .storageModeShared)!

let cb = dev.makeCommandQueue()!.makeCommandBuffer()!
let enc = cb.makeComputeCommandEncoder()!
enc.setComputePipelineState(pso)
enc.setBuffer(inBuf, offset: 0, index: 0)
enc.setBuffer(outBuf, offset: 0, index: 1)
enc.dispatchThreads(MTLSize(width: xs.count, height: 1, depth: 1),
                    threadsPerThreadgroup: MTLSize(width: xs.count, height: 1, depth: 1))
enc.endEncoding()
cb.commit()
cb.waitUntilCompleted()

guard cb.status == .completed else {
    fatalError("command buffer failed: \(String(describing: cb.error))")
}
print("Command buffer status: completed\n")

let gpu = outBuf.contents().assumingMemoryBound(to: UInt32.self)
var wrong = 0
var wrongLowHalfMatches = 0
print("Expected: high32(x * 500) / 20")
print("Diagnostic prediction: low32(x * 500) / 20\n")
print("         x   CPU expected   Metal GPU   CPU low32 prediction   verdict")
for (i, x) in xs.enumerated() {
    // Independent CPU reference. These calculations do not run through Metal.
    let product = UInt64(x) * 500
    let expected = UInt32(product >> 32) / 20
    let lowHalfPrediction = UInt32(truncatingIfNeeded: product) / 20
    if gpu[i] != expected { wrong += 1 }
    if gpu[i] != expected && gpu[i] == lowHalfPrediction { wrongLowHalfMatches += 1 }
    let verdict = gpu[i] == expected
        ? "correct"
        : (gpu[i] == lowHalfPrediction ? "WRONG (matches low32)" : "WRONG (other result)")
    print(String(format: "%10u %14u %11u %22u   %@",
                 x, expected, gpu[i], lowHalfPrediction, verdict))
}

print()
if wrong == 0 {
    print("RESULT: no miscompile observed")
} else {
    print("RESULT: METAL MISCOMPILE REPRODUCED")
    print("\(wrong)/\(xs.count) Metal GPU results differ from the expected MSL result.")
    print("\(wrongLowHalfMatches)/\(xs.count) exactly match division of the low product half.")
}