Add Metal capability probe for GitHub-hosted macOS runners

.github/runner/metal-probe.swift

@@ -0,0 +1,132 @@
+// Probe whether Metal is usable where this runs: on a GitHub-hosted macOS
+// runner (macOS VM) and inside the iOS simulator on that runner. It creates
+// the default device, runs a compute kernel and a render pass, and checks the
+// results on the CPU. Exit code 0 means Metal works end to end.
+import Foundation
+import Metal
+
+func fail(_ message: String) -> Never {
+    print("FAIL: \(message)")
+    exit(1)
+}
+
+guard let device = MTLCreateSystemDefaultDevice() else {
+    fail("MTLCreateSystemDefaultDevice() returned nil: no Metal device")
+}
+
+print("device: \(device.name)")
+#if os(macOS)
+print("all devices: \(MTLCopyAllDevices().map { $0.name })")
+print("lowPower: \(device.isLowPower) headless: \(device.isHeadless) removable: \(device.isRemovable)")
+#endif
+print("families: apple7=\(device.supportsFamily(.apple7)) apple8=\(device.supportsFamily(.apple8)) metal3=\(device.supportsFamily(.metal3)) common3=\(device.supportsFamily(.common3))")
+print("maxThreadsPerThreadgroup: \(device.maxThreadsPerThreadgroup)")
+print("maxBufferLength: \(device.maxBufferLength)")
+print("argumentBuffersSupport tier: \(device.argumentBuffersSupport.rawValue)")
+print("unifiedMemory: \(device.hasUnifiedMemory)")
+
+let source = """
+#include <metal_stdlib>
+using namespace metal;
+
+kernel void square(device const float* in [[buffer(0)]],
+                   device float* out [[buffer(1)]],
+                   uint id [[thread_position_in_grid]]) {
+  out[id] = in[id] * in[id];
+}
+
+struct VSOut {
+  float4 position [[position]];
+};
+
+vertex VSOut vs(uint vid [[vertex_id]]) {
+  // One triangle that covers the whole clip space.
+  float2 p[3] = { float2(-1, -1), float2(3, -1), float2(-1, 3) };
+  VSOut o;
+  o.position = float4(p[vid], 0, 1);
+  return o;
+}
+
+fragment float4 fs() {
+  return float4(0, 1, 0, 1);
+}
+"""
+
+do {
+    let library = try device.makeLibrary(source: source, options: nil)
+    guard let queue = device.makeCommandQueue() else { fail("no command queue") }
+
+    // Compute
+    guard let squareFn = library.makeFunction(name: "square") else { fail("no kernel") }
+    let pipeline = try device.makeComputePipelineState(function: squareFn)
+    let count = 4096
+    var input = (0..<count).map { Float($0) }
+    guard let inBuf = device.makeBuffer(bytes: &input, length: count * 4, options: .storageModeShared),
+          let outBuf = device.makeBuffer(length: count * 4, options: .storageModeShared) else {
+        fail("no buffers")
+    }
+    let cmd = queue.makeCommandBuffer()!
+    let enc = cmd.makeComputeCommandEncoder()!
+    enc.setComputePipelineState(pipeline)
+    enc.setBuffer(inBuf, offset: 0, index: 0)
+    enc.setBuffer(outBuf, offset: 0, index: 1)
+    let tg = min(pipeline.maxTotalThreadsPerThreadgroup, count)
+    enc.dispatchThreads(MTLSize(width: count, height: 1, depth: 1),
+                        threadsPerThreadgroup: MTLSize(width: tg, height: 1, depth: 1))
+    enc.endEncoding()
+    cmd.commit()
+    cmd.waitUntilCompleted()
+    if let error = cmd.error { fail("compute command buffer: \(error)") }
+    let out = outBuf.contents().bindMemory(to: Float.self, capacity: count)
+    let computeOK = (0..<count).allSatisfy { out[$0] == Float($0) * Float($0) }
+    print("compute kernel: \(computeOK ? "correct" : "WRONG RESULT")")
+    if !computeOK { exit(1) }
+
+    // Render
+    let width = 64, height = 64
+    let desc = MTLTextureDescriptor.texture2DDescriptor(pixelFormat: .bgra8Unorm,
+                                                        width: width, height: height,
+                                                        mipmapped: false)
+    desc.usage = [.renderTarget, .shaderRead]
+    desc.storageMode = .private
+    guard let target = device.makeTexture(descriptor: desc) else { fail("no render target") }
+    let pipeDesc = MTLRenderPipelineDescriptor()
+    pipeDesc.vertexFunction = library.makeFunction(name: "vs")
+    pipeDesc.fragmentFunction = library.makeFunction(name: "fs")
+    pipeDesc.colorAttachments[0].pixelFormat = .bgra8Unorm
+    let renderPipeline = try device.makeRenderPipelineState(descriptor: pipeDesc)
+    let pass = MTLRenderPassDescriptor()
+    pass.colorAttachments[0].texture = target
+    pass.colorAttachments[0].loadAction = .clear
+    pass.colorAttachments[0].storeAction = .store
+    pass.colorAttachments[0].clearColor = MTLClearColor(red: 1, green: 0, blue: 0, alpha: 1)
+    let bytesPerRow = width * 4
+    guard let readback = device.makeBuffer(length: bytesPerRow * height, options: .storageModeShared) else {
+        fail("no readback buffer")
+    }
+    let rcmd = queue.makeCommandBuffer()!
+    let renc = rcmd.makeRenderCommandEncoder(descriptor: pass)!
+    renc.setRenderPipelineState(renderPipeline)
+    renc.drawPrimitives(type: .triangle, vertexStart: 0, vertexCount: 3)
+    renc.endEncoding()
+    let blit = rcmd.makeBlitCommandEncoder()!
+    blit.copy(from: target, sourceSlice: 0, sourceLevel: 0,
+              sourceOrigin: MTLOrigin(x: 0, y: 0, z: 0),
+              sourceSize: MTLSize(width: width, height: height, depth: 1),
+              to: readback, destinationOffset: 0,
+              destinationBytesPerRow: bytesPerRow, destinationBytesPerImage: bytesPerRow * height)
+    blit.endEncoding()
+    rcmd.commit()
+    rcmd.waitUntilCompleted()
+    if let error = rcmd.error { fail("render command buffer: \(error)") }
+    let px = readback.contents().bindMemory(to: UInt8.self, capacity: bytesPerRow * height)
+    // BGRA: a fully covered green target reads as 0, 255, 0, 255.
+    let renderOK = (0..<(width * height)).allSatisfy { i in
+        px[i * 4] == 0 && px[i * 4 + 1] == 255 && px[i * 4 + 2] == 0 && px[i * 4 + 3] == 255
+    }
+    print("render pass: \(renderOK ? "correct" : "WRONG RESULT") (pixel 0 = \(px[0]),\(px[1]),\(px[2]),\(px[3]))")
+    if !renderOK { exit(1) }
+    print("OK: Metal compute and render work here")
+} catch {
+    fail("\(error)")
+}

.github/workflows/metal-probe.yml

@@ -0,0 +1,131 @@
+# Experiment: does a GitHub-hosted macOS runner expose a usable Metal device,
+# on the host and inside the iOS simulator? That is the one thing the iOS e2e
+# job needs that a self-hosted Mac was assumed to be required for.
+name: Metal probe
+on:
+  push:
+    branches:
+      - claude/trusting-dirac-20bbq7
+  workflow_dispatch:
+
+jobs:
+  probe:
+    runs-on: macos-latest
+    timeout-minutes: 20
+    steps:
+      - name: Checkout
+        uses: actions/checkout@v4
+
+      - name: Runner info
+        run: |
+          sw_vers
+          uname -m
+          sysctl -n hw.model machdep.cpu.brand_string hw.ncpu hw.memsize || true
+          xcodebuild -version
+          system_profiler SPDisplaysDataType
+
+      - name: Metal on the host
+        run: swift .github/runner/metal-probe.swift
+
+      - name: Boot an iOS simulator
+        run: |
+          # Same selection as the e2e job in ci.yml: newest iOS runtime and the
+          # newest "iPhone N Pro" device type.
+          RUNTIME=$(xcrun simctl list runtimes available --json | jq -r '
+            .runtimes | map(select(.platform == "iOS"))
+            | sort_by(.version | split(".") | map(tonumber)) | last | .identifier')
+          DEVICE_TYPE=$(xcrun simctl list devicetypes --json | jq -r '
+            .devicetypes | map(select(.name | test("^iPhone [0-9]+ Pro$")))
+            | sort_by(.name | capture("^iPhone (?<n>[0-9]+) Pro$").n | tonumber) | last | .identifier')
+          echo "runtime=$RUNTIME device=$DEVICE_TYPE"
+          UDID=$(xcrun simctl create metal-probe "$DEVICE_TYPE" "$RUNTIME")
+          echo "UDID=$UDID" >> $GITHUB_ENV
+          xcrun simctl boot "$UDID"
+          xcrun simctl bootstatus "$UDID" -b
+
+      - name: Metal in the iOS simulator
+        run: |
+          SDK=$(xcrun -sdk iphonesimulator --show-sdk-path)
+          SDK_VERSION=$(xcrun -sdk iphonesimulator --show-sdk-version)
+          xcrun -sdk iphonesimulator swiftc \
+            -target "arm64-apple-ios${SDK_VERSION}-simulator" -sdk "$SDK" \
+            .github/runner/metal-probe.swift -o "$RUNNER_TEMP/metal-probe-sim"
+          xcrun simctl spawn "$UDID" "$RUNNER_TEMP/metal-probe-sim"
+
+  # Second half of the experiment: the iOS e2e job from ci.yml, unchanged
+  # except for runs-on, on a fresh GitHub-hosted Mac. Lint, docs and the
+  # release dry run are left out: they are not what the self-hosted machine
+  # was assumed to be needed for.
+  e2e-hosted:
+    runs-on: macos-latest
+    timeout-minutes: 60
+    env:
+      METRO_PORT: 8083
+      E2E_PORT: 4244
+      IOS_SIMULATOR: webgpu-ci
+    steps:
+      - name: Checkout
+        uses: actions/checkout@v4
+
+      - name: Setup
+        uses: ./.github/actions/setup
+        with:
+          github_token: ${{ secrets.GITHUB_TOKEN }}
+
+      - name: Build packages
+        run: yarn turbo run build --filter='./packages/*'
+
+      - name: Install CocoaPods
+        working-directory: apps/example/ios
+        run: pod install
+
+      - name: Start Package Manager
+        working-directory: apps/example
+        run: CI=true yarn start --port $METRO_PORT --reset-cache &
+
+      - name: Build example for iOS
+        working-directory: apps/example/ios
+        run: |
+          which xcpretty >/dev/null || sudo gem install xcpretty
+          set -o pipefail && xcodebuild \
+            -workspace example.xcworkspace \
+            -scheme Example \
+            -configuration Debug \
+            -sdk iphonesimulator \
+            -destination "generic/platform=iOS Simulator" \
+            -derivedDataPath build \
+            COMPILER_INDEX_STORE_ENABLE=NO \
+            DISABLE_MANUAL_TARGET_ORDER_BUILD_WARNING=YES \
+            RCT_METRO_PORT=$METRO_PORT \
+            -disableAutomaticPackageResolution \
+            build | xcpretty
+
+      - name: Prepare iOS simulator
+        run: |
+          RUNTIME=$(xcrun simctl list runtimes available --json | jq -r '
+            .runtimes | map(select(.platform == "iOS"))
+            | sort_by(.version | split(".") | map(tonumber)) | last | .identifier')
+          DEVICE_TYPE=$(xcrun simctl list devicetypes --json | jq -r '
+            .devicetypes | map(select(.name | test("^iPhone [0-9]+ Pro$")))
+            | sort_by(.name | capture("^iPhone (?<n>[0-9]+) Pro$").n | tonumber) | last | .identifier')
+          echo "Creating $IOS_SIMULATOR ($DEVICE_TYPE, $RUNTIME)"
+          UDID=$(xcrun simctl create "$IOS_SIMULATOR" "$DEVICE_TYPE" "$RUNTIME")
+          echo "IOS_SIMULATOR_UDID=$UDID" >> $GITHUB_ENV
+          xcrun simctl boot "$UDID"
+          xcrun simctl bootstatus "$UDID" -b
+
+      - name: Install and launch app on Simulator
+        run: |
+          xcrun simctl install "$IOS_SIMULATOR_UDID" apps/example/ios/build/Build/Products/Debug-iphonesimulator/ReactTestApp.app
+          xcrun simctl launch "$IOS_SIMULATOR_UDID" com.microsoft.ReactTestApp
+
+      - name: Run e2e tests
+        working-directory: packages/webgpu
+        run: yarn test
+
+      - name: Upload test snapshots
+        if: always()
+        uses: actions/upload-artifact@v4
+        with:
+          name: webgpu-test-snapshots-hosted
+          path: packages/webgpu/src/__tests__/snapshots/