Constant CLOCK_PROBE_MSL 

pub(crate) const CLOCK_PROBE_MSL: &str = "#include <metal_stdlib>\nusing namespace metal;\n\n#line 1 \"device/clock.metal\"\n// Dependent-FMA clock probe used to expose GPU throttling.\n\nkernel void clock_probe(device float* out [[buffer(0)]],\n                        constant uint& iters [[buffer(1)]],\n                        uint gid [[thread_position_in_grid]]) {\n    float x = 1.0f + (float)gid * 1e-9f;\n    for (uint i = 0; i < iters; i++) {\n        x = fma(x, 1.0000001f, 1e-7f);\n        x = fma(x, 0.9999999f, -1e-7f);\n    }\n    out[gid] = x;\n}\n\n";

Device diagnostic, compiled separately from model kernels.