From 1c10b364a5618aafe4656f6c936b1c382f47c69e Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 22 Oct 2026 06:55:59 -0400 Subject: [PATCH 1/2] jouleclaw-backend-apple: 84 clippy warnings -> 10 #414 left this crate as the reason the `-D warnings` gate covers loader-gguf on default features only: clippy reports it through the path dependency, so gating the `fabric-meter` configuration needs it clean too. I described its warnings in #414 as "unused in THAT configuration and needed in others, so they want cfg work rather than deletion". That was wrong on two counts, and measuring it directly settled both: - the crate has NO `[features]` at all, only a macOS-conditional dependency block, so there is one configuration and the unused items are unconditionally unused; - there are 84 of them, not 43. The 43 was what leaked through loader-gguf's lint, not the crate's own count. 82 of the 84 are in mps_matmul.rs. 41 came from clippy's own suggestions (manual_is_multiple_of 23, unnecessary_cast 8, needless_option_as_deref 3, excessive_precision, doc_lazy_continuation, double_parens, chunks_exact_to_as_chunks). The other 42 -- 27 `unused_unsafe` and 15 `unused_imports` -- carry no machine-applicable suggestion at all. Neither `clippy --fix` nor `cargo fix` touches them, with or without `--all-targets`. They are plain code, not macro expansions, so they were fixed line-guided from rustc's own line numbers: - `|| unsafe { .. }` -> `|| { .. }` where objc2 0.6 made the call safe. Dropping a block rustc has already proved redundant cannot change behaviour. - `use objc2_metal::{MTLComputePipelineState as _, MTLLibrary as _}` -> just the name the function needs. The unused one is a trait brought in as `_`; the type itself is referenced fully qualified on the lines above, which is why it reads as used at a glance. Checked before deleting, because this is the shape that broke #372: each unused import was matched against every other mention of the name in its file and against the file's six `cfg(not(target_os = "macos"))` branches. None appears in one, and objc2-metal is a macOS-only dependency, so no other configuration can want them. Verified: 0 errors, 37 backend-apple tests (32 pass, 5 ignored), and loader-gguf still builds and its fabric-meter `process_calibration` passes. Co-Authored-By: Claude Opus 5 --- apply-sug.py | 85 ++++++++++ bac.py | 22 +++ fix-unused.py | 63 ++++++++ .../examples/mpp_diag.rs | 2 +- .../examples/mpp_gemm_bench.rs | 2 +- .../jouleclaw-backend-apple/src/mps_matmul.rs | 150 +++++++++--------- 6 files changed, 247 insertions(+), 77 deletions(-) create mode 100644 apply-sug.py create mode 100644 bac.py create mode 100644 fix-unused.py diff --git a/apply-sug.py b/apply-sug.py new file mode 100644 index 000000000..584840846 --- /dev/null +++ b/apply-sug.py @@ -0,0 +1,85 @@ +#!/usr/bin/env python3 +"""Apply clippy's suggested replacements for a whitelist of lints. + + apply-sug.py [ ...] + +Only MachineApplicable suggestions reach `cargo clippy --fix`; the rest are +still exact spans with exact text. Rejects template suggestions, which are a +bare lowercase word in angle brackets that no `::` precedes -- `needless_range_loop` +emits `for (i, ) in ..`, leaving the body's `w[i]` uses for a human, and +applying one literally does not parse. A turbofish such as `as_chunks::<2>()` +is a real suggestion and must not be caught by that. +""" +import json, re, subprocess, sys, collections, pathlib + +pkg, feats = sys.argv[1], sys.argv[2] +want = set(sys.argv[3:]) +cmd = ["cargo","clippy","--manifest-path","jouleclaw/jouleclaw-rs/Cargo.toml", + "-p",pkg,"--all-targets","--no-deps","--keep-going","--message-format=json"] +if feats: + cmd += ["--features", feats] +proc = subprocess.run(cmd, capture_output=True, text=True) + +edits = collections.defaultdict(list) +seen, skipped = set(), collections.Counter() +for line in proc.stdout.splitlines(): + try: + m = json.loads(line) + except json.JSONDecodeError: + continue + if m.get("reason") != "compiler-message": + continue + d = m.get("message", {}) + if d.get("level") != "warning": + continue + code = ((d.get("code") or {}).get("code") or "").replace("clippy::", "") + if code not in want: + continue + def walk(msg): + for sp in msg.get("spans", []): + if sp.get("suggested_replacement") is not None: + yield sp + for ch in msg.get("children", []): + yield from walk(ch) + got = list(walk(d)) + if not got: + skipped[code] += 1 + continue + for sp in got: + repl = sp["suggested_replacement"] + if re.search(r"(?", repl): + skipped[code] += 1 + continue + key = (sp["file_name"], sp["byte_start"], sp["byte_end"]) + if key in seen: + continue + seen.add(key) + edits[sp["file_name"]].append((sp["byte_start"], sp["byte_end"], repl, code)) + +WS = pathlib.Path("jouleclaw/jouleclaw-rs") +applied = collections.Counter() +for fn, es in edits.items(): + p = pathlib.Path(fn) + if not p.exists(): + p = WS / fn + if not p.exists(): + print(f" MISSING {fn}") + continue + data = p.read_bytes() + es.sort(key=lambda e: e[0], reverse=True) + last = None + for start, end, repl, code in es: + if last is not None and end > last: + skipped[code] += 1 + continue + data = data[:start] + repl.encode() + data[end:] + last = start + applied[code] += 1 + p.write_bytes(data) +print(f"applied {sum(applied.values())} edits in {len(edits)} files") +for c, n in applied.most_common(): + print(f" {n:4d} {c}") +if skipped: + print("left alone (no suggestion / overlapping / template):") + for c, n in skipped.most_common(): + print(f" {n:4d} {c}") diff --git a/bac.py b/bac.py new file mode 100644 index 000000000..5632aff89 --- /dev/null +++ b/bac.py @@ -0,0 +1,22 @@ +import json, subprocess, collections +cmd=["cargo","clippy","--manifest-path","jouleclaw/jouleclaw-rs/Cargo.toml", + "-p","jouleclaw-backend-apple","--all-targets","--no-deps","--keep-going","--message-format=json"] +pr=subprocess.run(cmd,capture_output=True,text=True) +c=collections.Counter(); e=0; seen=set(); files=collections.Counter() +for l in pr.stdout.splitlines(): + try: m=json.loads(l) + except Exception: continue + if m.get("reason")!="compiler-message": continue + d=m.get("message",{}); lv=d.get("level") + if lv not in ("warning","error"): continue + code=(d.get("code") or {}).get("code") or "(uncoded)" + sp=next((s for s in d.get("spans",[]) if s.get("is_primary")),None) + fn=str(sp.get("file_name") if sp else "?") + k=(code,fn,sp.get("line_start") if sp else None) + if k in seen: continue + seen.add(k) + if lv=="error": e+=1; continue + c[code]+=1; files[fn.split("/")[-1]]+=1 +print(f"errors={e} warnings={sum(c.values())}") +for k,v in c.most_common(12): print(f" {v:4d} {k}") +print(" by file:", dict(files)) diff --git a/fix-unused.py b/fix-unused.py new file mode 100644 index 000000000..c667561ac --- /dev/null +++ b/fix-unused.py @@ -0,0 +1,63 @@ +import json, re, subprocess, pathlib, collections + +cmd=["cargo","clippy","--manifest-path","jouleclaw/jouleclaw-rs/Cargo.toml", + "-p","jouleclaw-backend-apple","--all-targets","--no-deps","--keep-going", + "--message-format=json"] +pr=subprocess.run(cmd,capture_output=True,text=True) +WS=pathlib.Path("jouleclaw/jouleclaw-rs") + +unsafe_sites=collections.defaultdict(set) # file -> {line} +import_sites=collections.defaultdict(list) # file -> [(line, name)] +for l in pr.stdout.splitlines(): + try: m=json.loads(l) + except Exception: continue + if m.get("reason")!="compiler-message": continue + d=m.get("message",{}) + if d.get("level")!="warning": continue + code=(d.get("code") or {}).get("code") or "" + sp=next((s for s in d.get("spans",[]) if s.get("is_primary")),None) + if not sp: continue + if code=="unused_unsafe": + unsafe_sites[sp["file_name"]].add(sp["line_start"]) + elif code=="unused_imports": + for nm in re.findall(r"`([^`]+)`", d.get("message","")): + import_sites[sp["file_name"]].append((sp["line_start"], nm.strip())) + +changed=collections.Counter() +for fn in set(list(unsafe_sites) + list(import_sites)): + p=pathlib.Path(fn) + if not p.exists(): p=WS/fn + lines=p.read_text().split("\n") + + # `|| unsafe {` -> `|| {`: rustc says the block is redundant, so dropping + # the keyword cannot change behaviour. + for ln in sorted(unsafe_sites.get(fn,())): + i=ln-1 + if i < len(lines) and "unsafe" in lines[i]: + new=re.sub(r"\bunsafe\s+\{", "{", lines[i], count=1) + if new!=lines[i]: + lines[i]=new; changed["unused_unsafe"]+=1 + + # Drop the one unused name from a brace list, or the whole `use` if it was + # the only name. + for ln, nm in sorted(set(import_sites.get(fn,[])), reverse=True): + i=ln-1 + if i >= len(lines): continue + s=lines[i] + base=nm.split("::")[-1] + # `Name as _, ` or `, Name as _` or bare `Name, ` / `, Name` + pat=rf"\b{re.escape(base)}(\s+as\s+_)?\s*,\s*|\s*,\s*\b{re.escape(base)}(\s+as\s+_)?\b" + new=re.sub(pat, "", s, count=1) + if new==s: + continue + # A brace list reduced to one entry reads better unbraced. + new=re.sub(r"\{\s*([A-Za-z_][\w:]*(\s+as\s+_)?)\s*\}", r"\1", new) + if re.search(r"\{\s*\}", new): + lines[i]=None # the import is now empty; drop the line + else: + lines[i]=new + changed["unused_imports"]+=1 + + p.write_text("\n".join(l for l in lines if l is not None)) + +print("applied:", dict(changed)) diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_diag.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_diag.rs index a8f786fdc..c70f1a1d5 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_diag.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_diag.rs @@ -13,7 +13,7 @@ fn main() { MTLComputeCommandEncoder, MTLComputePipelineState, MTLCreateSystemDefaultDevice, MTLDevice, MTLLibrary, MTLResourceOptions, MTLSize, }; - fn f32_to_bf16(v: f32) -> u16 { let b = v.to_bits(); (((b + (((b >> 16) & 1) + 0x7fff)) >> 16)) as u16 } + fn f32_to_bf16(v: f32) -> u16 { let b = v.to_bits(); ((b + (((b >> 16) & 1) + 0x7fff)) >> 16) as u16 } const KERNEL: &str = r#" #include diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_gemm_bench.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_gemm_bench.rs index c169daacc..f50041e82 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_gemm_bench.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/examples/mpp_gemm_bench.rs @@ -25,7 +25,7 @@ fn main() { MTLComputeCommandEncoder, MTLComputePipelineState, MTLCreateSystemDefaultDevice, MTLDevice, MTLLibrary, MTLResourceOptions, MTLSize, }; - fn f32_to_bf16(v: f32) -> u16 { let b = v.to_bits(); (((b + (((b >> 16) & 1) + 0x7fff)) >> 16)) as u16 } + fn f32_to_bf16(v: f32) -> u16 { let b = v.to_bits(); ((b + (((b >> 16) & 1) + 0x7fff)) >> 16) as u16 } fn bf16_to_f32(h: u16) -> f32 { f32::from_bits((h as u32) << 16) } const KERNEL: &str = r#" diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs index 43e3e9227..998e3cb1b 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs @@ -192,7 +192,7 @@ pub mod inner { fn context() -> Option<&'static MpsContext> { static CTX: OnceLock> = OnceLock::new(); - CTX.get_or_init(|| unsafe { + CTX.get_or_init(|| { let device = MTLCreateSystemDefaultDevice()?; let queue = device.newCommandQueue()?; Some(MpsContext { @@ -335,7 +335,7 @@ pub mod inner { assert_eq!(w_bf16.len(), n * k * 2); let ctx = context().ok_or("no Metal device")?; let mut half: Vec = Vec::with_capacity(n * k); - for c in w_bf16.chunks_exact(2) { + for c in w_bf16.as_chunks::<2>().0 { let v = f32::from_bits((u16::from_le_bytes([c[0], c[1]]) as u32) << 16); if v.is_finite() && v.abs() > 65504.0 { return Err(format!("weight {v} overflows f16")); @@ -451,11 +451,11 @@ kernel void silu_mul_f32_to_f16( fn silu_mul_pipeline() -> Option<&'static Retained>> { use objc2_foundation::NSString; - use objc2_metal::{MTLComputePipelineState as _, MTLLibrary as _}; + use objc2_metal::MTLLibrary as _; static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(SILU_MUL_MSL); let lib = ctx @@ -494,11 +494,11 @@ kernel void silu_mul_f32_to_f32( fn silu_mul_f32_pipeline() -> Option<&'static Retained>> { use objc2_foundation::NSString; - use objc2_metal::{MTLComputePipelineState as _, MTLLibrary as _}; + use objc2_metal::MTLLibrary as _; static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx.device .newLibraryWithSource_options_error(&NSString::from_str(SILU_MUL_F32_MSL), None) @@ -528,11 +528,11 @@ kernel void residual_add_f32( fn residual_add_pipeline() -> Option<&'static Retained>> { use objc2_foundation::NSString; - use objc2_metal::{MTLComputePipelineState as _, MTLLibrary as _}; + use objc2_metal::MTLLibrary as _; static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx.device .newLibraryWithSource_options_error(&NSString::from_str(RESIDUAL_ADD_MSL), None) @@ -636,7 +636,7 @@ kernel void silu_mul_rows( use objc2_metal::MTLLibrary as _; static P: OnceLock>>> = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx.device .newLibraryWithSource_options_error(&NSString::from_str(RMSNORM_ROWS_MSL), None).ok()?; @@ -651,7 +651,7 @@ kernel void silu_mul_rows( use objc2_metal::MTLLibrary as _; static P: OnceLock>>> = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx.device .newLibraryWithSource_options_error(&NSString::from_str(ROPE_ROWS_MSL), None).ok()?; @@ -666,7 +666,7 @@ kernel void silu_mul_rows( use objc2_metal::MTLLibrary as _; static P: OnceLock>>> = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx.device .newLibraryWithSource_options_error(&NSString::from_str(SILU_MUL_ROWS_MSL), None).ok()?; @@ -867,7 +867,7 @@ kernel void relu2_f32_to_f16( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(RELU2_MSL); let lib = ctx.device.newLibraryWithSource_options_error(&src, None).ok()?; @@ -1054,7 +1054,7 @@ kernel void rmsnorm_f32_to_f16( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(RMSNORM_MSL); let lib = ctx @@ -1105,7 +1105,7 @@ kernel void gemm_bf16_mpp(device bfloat* A_buf [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(MATMUL2D_MSL); // Compile FAILS on pre-Metal-4 (the MPP header is @@ -1149,7 +1149,7 @@ kernel void gemm_bf16_mpp(device bfloat* A_buf [[buffer(0)]], // C declared [N,M] (transposed) makes the example // grid/slice correct for NON-SQUARE too. Tile dims: // M%64, N%32, K%32 (≥ one tile each). - if m % 64 != 0 || m < 64 || n % 32 != 0 || n < 32 || k % 32 != 0 { + if !m.is_multiple_of(64) || m < 64 || !n.is_multiple_of(32) || n < 32 || !k.is_multiple_of(32) { return Err(format!("matmul2d shape unsupported: M={m} N={n} K={k}")); } assert_eq!(a.len(), m * k); @@ -1407,7 +1407,7 @@ kernel void qknorm_rope( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(QKNORM_ROPE_MSL); let lib = ctx @@ -1539,7 +1539,7 @@ kernel void rope_rotate_half( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(ROPE_ROTATE_HALF_MSL); let lib = ctx @@ -1584,7 +1584,7 @@ kernel void scale_qk( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(SCALE_QK_MSL); let lib = ctx @@ -1606,7 +1606,7 @@ kernel void scale_qk( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(ROPE_INTERLEAVED_MSL); let lib = ctx @@ -1898,7 +1898,7 @@ kernel void attn_decode( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(KV_APPEND_MSL); let lib = ctx.device.newLibraryWithSource_options_error(&src, None).ok()?; @@ -1915,7 +1915,7 @@ kernel void attn_decode( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(ATTN_DECODE_MSL); let lib = ctx.device.newLibraryWithSource_options_error(&src, None).ok()?; @@ -1957,7 +1957,7 @@ kernel void attn_decode( let cap = cap.min(4096); // threadgroup score scratch bound let ctx = context()?; let bytes = n_kv * cap * hd * 2; - unsafe { + { let k = ctx .device .newBufferWithLength_options(bytes, MTLResourceOptions::StorageModeShared)?; @@ -2062,7 +2062,7 @@ kernel void attn_decode( Some(scale_qk_pipeline().ok_or("scale_qk pipeline unavailable")?) } }; - let qk_pipe = qk_pipe.as_deref(); + let qk_pipe = qk_pipe; let append_pipe = kv_append_pipeline().ok_or("kv_append pipeline unavailable")?; let attn_pipe = attn_decode_pipeline().ok_or("attn pipeline unavailable")?; // Slot bookkeeping (CPU side, before encode). @@ -2612,7 +2612,7 @@ kernel void gemv_q8(device const uchar* W [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -2642,7 +2642,7 @@ kernel void gemv_q8(device const uchar* W [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 32 != 0 { + if !k.is_multiple_of(32) { return Err(format!("k ({k}) not a multiple of 32")); } let nb = k / 32; @@ -2671,7 +2671,7 @@ kernel void gemv_q8(device const uchar* W [[buffer(0)]], }; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let _g = ctx.encode_lock.lock().unwrap(); for i in 0..m { @@ -2819,7 +2819,7 @@ kernel void gemv_q4k(device const uchar* W [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -2939,7 +2939,7 @@ kernel void gemm_q4k_m4(device const uchar* W [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -3040,7 +3040,7 @@ kernel void sgemm_q4k(device const float* A [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -3059,10 +3059,10 @@ kernel void sgemm_q4k(device const float* A [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 256 != 0 { + if !k.is_multiple_of(256) { return Err(format!("k ({k}) not a multiple of 256")); } - if out_dim % 8 != 0 { + if !out_dim.is_multiple_of(8) { return Err(format!("out_dim ({out_dim}) not a multiple of 8")); } let nb = k / 256; @@ -3095,7 +3095,7 @@ kernel void sgemm_q4k(device const float* A [[buffer(0)]], let mp = m.div_ceil(8) * 8; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; // Zero-padded activation [mp, k]. let a_buf = ctx @@ -3212,7 +3212,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -3230,10 +3230,10 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 256 != 0 { + if !k.is_multiple_of(256) { return Err(format!("k ({k}) not a multiple of 256")); } - if out_dim % 8 != 0 { + if !out_dim.is_multiple_of(8) { return Err(format!("out_dim ({out_dim}) not a multiple of 8")); } let nb = k / 256; @@ -3266,7 +3266,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], let mp = m.div_ceil(8) * 8; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let a_buf = ctx .device @@ -3314,7 +3314,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 256 != 0 { + if !k.is_multiple_of(256) { return Err(format!("k ({k}) not a multiple of 256")); } let nb = k / 256; @@ -3343,7 +3343,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], }; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let _g = ctx.encode_lock.lock().unwrap(); for i in 0..m { @@ -3400,7 +3400,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 256 != 0 { + if !k.is_multiple_of(256) { return Err(format!("k ({k}) not a multiple of 256")); } let nb = k / 256; @@ -3436,7 +3436,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], let n_tiles = m.div_ceil(TILE); unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let set_u32 = |enc: &ProtocolObject, v: u32, idx: usize| { enc.setBytes_length_atIndex( @@ -3511,7 +3511,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], } }; unsafe { - use objc2_metal::{MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize}; + use objc2_metal::{MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize}; let buf_a = ctx.device.newBufferWithBytes_length_options( std::ptr::NonNull::new(a.as_ptr() as *mut _).unwrap(), k * 4, MTLResourceOptions::StorageModeShared).ok_or("a")?; @@ -3631,7 +3631,7 @@ kernel void gemv_q6k(device const uchar* W [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -3740,7 +3740,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let lib = ctx .device @@ -3760,7 +3760,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 256 != 0 { + if !k.is_multiple_of(256) { return Err(format!("k ({k}) not a multiple of 256")); } let nb = k / 256; @@ -3796,7 +3796,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], let n_tiles = m.div_ceil(TILE); unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let set_u32 = |enc: &ProtocolObject, v: u32, idx: usize| { enc.setBytes_length_atIndex( @@ -3853,7 +3853,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), m * k); assert_eq!(c.len(), m * out_dim); - if k % 256 != 0 { + if !k.is_multiple_of(256) { return Err(format!("k ({k}) not a multiple of 256")); } let nb = k / 256; @@ -3882,7 +3882,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], }; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let _g = ctx.encode_lock.lock().unwrap(); for i in 0..m { @@ -3935,15 +3935,15 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], /// (gate/up Q4_K, down Q6_K). The gate/up/act intermediates stay /// in GPU buffers across stages (no CPU round-trip); the 3 GEMVs /// + activation collapse from 3 command buffers (3 commit+waits) - /// into one. `a` is post-norm `[h]`; gate/up Q4_K `[f·(h/256)·144]`, - /// down Q6_K `[h·(f/256)·210]`; `out` `[h]`. h,f multiples of 256. + /// into one. `a` is post-norm `[h]`; gate/up Q4_K `[f·(h/256)·144]`, + /// down Q6_K `[h·(f/256)·210]`; `out` `[h]`. h,f multiples of 256. pub fn mlp_block_q4k_q6k( a: &[f32], gate_w: &[u8], up_w: &[u8], down_w: &[u8], h: usize, f: usize, out: &mut [f32], ) -> Result<(), String> { assert_eq!(a.len(), h); assert_eq!(out.len(), h); - if h % 256 != 0 || f % 256 != 0 { + if !h.is_multiple_of(256) || !f.is_multiple_of(256) { return Err(format!("h ({h}) and f ({f}) must be multiples of 256")); } let nb_g = h / 256; // gate/up contraction blocks @@ -4063,7 +4063,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], ) -> Result<(), String> { assert_eq!(a.len(), h); assert_eq!(out.len(), q_out + 2 * kv_out); - if h % 256 != 0 { + if !h.is_multiple_of(256) { return Err(format!("h ({h}) must be a multiple of 256")); } let nb = h / 256; @@ -4091,7 +4091,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], let v_buf = upload(v_w)?; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let buf_a = ctx.device.newBufferWithBytes_length_options( std::ptr::NonNull::new(a.as_ptr() as *mut _).unwrap(), h * 4, @@ -4255,7 +4255,7 @@ kernel void attn_decode_f32( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(RMSNORM_F32F32_MSL); let lib = ctx.device.newLibraryWithSource_options_error(&src, None).ok()?; @@ -4272,7 +4272,7 @@ kernel void attn_decode_f32( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(ATTN_DECODE_F32_MSL); let lib = ctx.device.newLibraryWithSource_options_error(&src, None).ok()?; @@ -4361,7 +4361,7 @@ kernel void batched_causal_attn_f32( static P: OnceLock< Option>>, > = OnceLock::new(); - P.get_or_init(|| unsafe { + P.get_or_init(|| { let ctx = context()?; let src = NSString::from_str(BATCHED_CAUSAL_ATTN_MSL); let lib = ctx.device.newLibraryWithSource_options_error(&src, None).ok()?; @@ -4391,7 +4391,7 @@ kernel void batched_causal_attn_f32( let pipe = batched_causal_attn_pipeline().ok_or("batched_causal_attn pipeline")?; unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let up = |data: &[f32]| { mctx.device.newBufferWithBytes_length_options( @@ -4466,11 +4466,11 @@ kernel void batched_causal_attn_f32( assert_eq!(k_out.len(), n_kv * hd); assert_eq!(v_out.len(), n_kv * hd); assert_eq!(o_out.len(), k_dim); - if k_dim % 256 != 0 { + if !k_dim.is_multiple_of(256) { return Err(format!("k_dim ({k_dim}) must be a multiple of 256")); } let o_in = n_q * hd; - if o_in % 256 != 0 { + if !o_in.is_multiple_of(256) { return Err(format!("n_q*hd ({o_in}) must be a multiple of 256")); } let nb_in = k_dim / 256; // 256-blocks per input row for q/k/v @@ -4519,7 +4519,7 @@ kernel void batched_causal_attn_f32( Some(scale_qk_pipeline().ok_or("scale_qk pipeline unavailable")?) } }; - let qk_pipe = qk_pipe.as_deref(); + let qk_pipe = qk_pipe; let append_pipe = kv_append_pipeline().ok_or("kv_append pipeline unavailable")?; let attn_pipe = attn_decode_f32_pipeline().ok_or("attn pipeline unavailable")?; @@ -4562,7 +4562,7 @@ kernel void batched_causal_attn_f32( unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let upload = |data: *const u8, len: usize, what: &str| { ctx.device @@ -4843,11 +4843,11 @@ kernel void batched_causal_attn_f32( assert_eq!(ffn_norm.len(), k_dim); assert_eq!(hd, kv.hd); assert_eq!(n_kv, kv.n_kv); - if k_dim % 256 != 0 || ffn % 256 != 0 { + if !k_dim.is_multiple_of(256) || !ffn.is_multiple_of(256) { return Err(format!("k_dim ({k_dim}) and ffn ({ffn}) must be multiples of 256")); } let o_in = n_q * hd; - if o_in % 256 != 0 { + if !o_in.is_multiple_of(256) { return Err(format!("n_q*hd ({o_in}) must be a multiple of 256")); } let nb_in = k_dim / 256; // input blocks for q/k/v + gate/up @@ -4895,7 +4895,7 @@ kernel void batched_causal_attn_f32( } QkStage::ScaleQk { .. } => Some(scale_qk_pipeline().ok_or("scale_qk pipeline")?), }; - let qk_pipe = qk_pipe.as_deref(); + let qk_pipe = qk_pipe; let append_pipe = kv_append_pipeline().ok_or("kv_append pipeline")?; let attn_pipe = attn_decode_f32_pipeline().ok_or("attn pipeline")?; @@ -5238,7 +5238,7 @@ kernel void batched_causal_attn_f32( let kd = n_kv * hd; assert_eq!(k_out.len(), m * kd); assert_eq!(v_out.len(), m * kd); - if hidden % 256 != 0 || ffn % 256 != 0 || qd % 256 != 0 { + if !hidden.is_multiple_of(256) || !ffn.is_multiple_of(256) || !qd.is_multiple_of(256) { return Err("hidden, ffn, n_q*hd must be multiples of 256".into()); } if m == 0 || m > 4096 { @@ -5293,7 +5293,7 @@ kernel void batched_causal_attn_f32( let mp = m.div_ceil(8) * 8; // padded row count for the 8×8 tiled GEMM unsafe { use objc2_metal::{ - MTLCommandEncoder, MTLComputeCommandEncoder, MTLComputePipelineState as _, MTLSize, + MTLCommandEncoder, MTLComputeCommandEncoder, MTLSize, }; let alloc = |elems: usize| { ctx.device.newBufferWithLength_options(elems * 4, MTLResourceOptions::StorageModeShared) @@ -6457,7 +6457,7 @@ mod f16_tests { }; let mut got = vec![0f32; m * n]; super::gemm_bt_resident_f16(&a_f, &weight, m, &mut got).unwrap(); - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); for i in 0..m { for j in 0..n { let mut acc = 0f32; @@ -6534,8 +6534,8 @@ mod f16_tests { (-2.0, 0xc000), (65504.0, 0x7bff), // Subnormals: 2⁻¹⁵ → man 512; 2⁻²⁴ → man 1. - (3.0517578125e-5, 0x0200), - (5.9604644775390625e-8, 0x0001), + (3.051_757_8e-5, 0x0200), + (5.960_464_5e-8, 0x0001), // Mid-subnormal: 1e-6 ≈ 16.78 ULP → 17. (1e-6, 0x0011), // Below half the smallest subnormal → zero. @@ -6553,7 +6553,7 @@ mod f16_tests { #[test] fn attn_block_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (kd, hd, n_qh, n_kvh, cap, steps) = (32usize, 8usize, 4usize, 2usize, 3usize, 5usize); let (nq, nkv) = (n_qh * hd, n_kvh * hd); let nw: Vec = (0..kd).map(|i| 0.8 + ((i % 5) as f32) * 0.1).collect(); @@ -6688,7 +6688,7 @@ mod f16_tests { #[test] fn attn_block_interleaved_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (kd, hd, n_qh, n_kvh, cap, steps) = (32usize, 8usize, 4usize, 2usize, 3usize, 5usize); let (nq, nkv) = (n_qh * hd, n_kvh * hd); let nw: Vec = (0..kd).map(|i| 0.8 + ((i % 5) as f32) * 0.1).collect(); @@ -6814,7 +6814,7 @@ mod f16_tests { #[test] fn attn_block_rotate_half_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (kd, hd, n_qh, n_kvh, cap, steps) = (32usize, 8usize, 4usize, 2usize, 3usize, 5usize); let (nq, nkv) = (n_qh * hd, n_kvh * hd); let nw: Vec = (0..kd).map(|i| 0.8 + ((i % 5) as f32) * 0.1).collect(); @@ -6981,7 +6981,7 @@ mod f16_tests { #[test] fn attn_block_scale_qk_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (kd, hd, n_qh, n_kvh, cap, steps) = (32usize, 8usize, 4usize, 2usize, 3usize, 5usize); let (nq, nkv) = (n_qh * hd, n_kvh * hd); let nw: Vec = (0..kd).map(|i| 0.8 + ((i % 5) as f32) * 0.1).collect(); @@ -7083,7 +7083,7 @@ mod f16_tests { #[test] fn norm_qkv_rope_window_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (kd, hd, n_qh, n_kvh) = (64usize, 8usize, 4usize, 2usize); let (nq, nkv) = (n_qh * hd, n_kvh * hd); let h: Vec = (0..kd).map(|i| ((i % 13) as f32 - 6.0) * 0.21).collect(); @@ -7171,7 +7171,7 @@ mod f16_tests { #[test] fn norm_gemm_window_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (k, n1, n2) = (64usize, 24usize, 16usize); let h: Vec = (0..k).map(|i| ((i % 13) as f32 - 6.0) * 0.21).collect(); let nw: Vec = (0..k).map(|i| 0.8 + ((i % 5) as f32) * 0.1).collect(); @@ -7221,7 +7221,7 @@ mod f16_tests { #[test] fn moe_window_matches_cpu_reference() { let to_bf16 = |v: f32| -> u16 { (v.to_bits() >> 16) as u16 }; - let bf = |v: f32| f32::from_bits(((v.to_bits() >> 16) as u32) << 16); + let bf = |v: f32| f32::from_bits((v.to_bits() >> 16) << 16); let (k, f, h, n_p) = (32usize, 24usize, 16usize, 3usize); let a: Vec = (0..k).map(|i| ((i % 11) as f32 - 5.0) * 0.09).collect(); let mk_w = |seed: usize, rows: usize, cols: usize| -> (Vec, Vec) { -- 2.54.0 (Apple Git-157) From 19904251b7baaaa6e568e0d33953a7694e51d7fa Mon Sep 17 00:00:00 2001 From: Claude Date: Thu, 22 Oct 2026 06:58:23 -0400 Subject: [PATCH 2/2] jouleclaw-backend-apple: 0 warnings, and the gate covers both configurations 10 -> 0, and the `-D warnings` step now lints loader-gguf under `rayon fabric-meter` as well as on default, plus backend-apple on its own. That closes the per-configuration lint hole #373 opened and #414 halved. The last ten, all in mps_matmul.rs: - `struct PipelineCell` with its `unsafe impl Send`/`Sync` was never constructed and mentioned nowhere else -- scaffolding for a OnceLock wrapper the pipelines ended up doing directly. Removed, and the space reused for a `KeepAlive` alias, which is what the two `type_complexity` warnings wanted: the inline form is three generic layers deep and appeared twice. - `DeviceKvCache` had a public `len` and no `is_empty`. - three `let qk_pipe = qk_pipe;` no-op rebindings from a refactor. - `too_many_arguments` and `needless_range_loop` allowed crate-wide, the same two decisions omni and loader-gguf already record and for the same reasons: a GPU dispatch signature is its kernel's parameter list, and these kernels walk several buffers at one index. The comment this replaces in standards-build.yml claimed the fabric-meter configuration could not be gated because backend-apple's warnings "are unused in THAT configuration and needed in others, so they want cfg work rather than deletion". Measuring the crate directly refuted both halves: it has no `[features]` at all, so there is one configuration, and the count was 84 rather than the 43 that leaked through loader-gguf's lint. Verified: all ten lint specs in the step pass, 37 backend-apple tests (32 pass, 5 ignored), 343 loader-gguf tests, 77 omni tests. Co-Authored-By: Claude Opus 5 --- .github/workflows/standards-build.yml | 16 ++-- apply-sug.py | 85 ------------------- bac.py | 22 ----- fix-unused.py | 63 -------------- .../crates/jouleclaw-backend-apple/src/lib.rs | 8 ++ .../jouleclaw-backend-apple/src/mps_matmul.rs | 20 ++--- 6 files changed, 26 insertions(+), 188 deletions(-) delete mode 100644 apply-sug.py delete mode 100644 bac.py delete mode 100644 fix-unused.py diff --git a/.github/workflows/standards-build.yml b/.github/workflows/standards-build.yml index 51310f697..6773b75ff 100644 --- a/.github/workflows/standards-build.yml +++ b/.github/workflows/standards-build.yml @@ -597,12 +597,14 @@ jobs: # And the crate omni loads its weights through, which was linted # by nothing at all: 161 warnings when first measured, now 0. # - # Default features only. Under `rayon fabric-meter` this crate is - # still clean, but clippy also reports jouleclaw-backend-apple - # through the path dependency (`--no-deps` does not exclude - # workspace members), and that crate has 43 warnings of its own -- - # mostly `unused_unsafe` and `unused_imports` that are unused in - # THAT configuration and needed in others, so they want cfg work - # rather than deletion. + # Both configurations. `--no-deps` does not exclude workspace path + # dependencies, so the fabric-meter line also covers + # jouleclaw-backend-apple, which is why that crate is linted on its + # own below rather than only incidentally. lint jouleclaw/jouleclaw-rs -p jouleclaw-loader-gguf + lint jouleclaw/jouleclaw-rs -p jouleclaw-loader-gguf --features "rayon fabric-meter" + # The Apple backend: 84 warnings when first measured on its own, now + # 0. It has no features, only a macOS-conditional dependency block, + # so there is one configuration to keep clean. + lint jouleclaw/jouleclaw-rs -p jouleclaw-backend-apple lint sandbox/sandbox-rs -p sandbox-meter -p sandbox-proof diff --git a/apply-sug.py b/apply-sug.py deleted file mode 100644 index 584840846..000000000 --- a/apply-sug.py +++ /dev/null @@ -1,85 +0,0 @@ -#!/usr/bin/env python3 -"""Apply clippy's suggested replacements for a whitelist of lints. - - apply-sug.py [ ...] - -Only MachineApplicable suggestions reach `cargo clippy --fix`; the rest are -still exact spans with exact text. Rejects template suggestions, which are a -bare lowercase word in angle brackets that no `::` precedes -- `needless_range_loop` -emits `for (i, ) in ..`, leaving the body's `w[i]` uses for a human, and -applying one literally does not parse. A turbofish such as `as_chunks::<2>()` -is a real suggestion and must not be caught by that. -""" -import json, re, subprocess, sys, collections, pathlib - -pkg, feats = sys.argv[1], sys.argv[2] -want = set(sys.argv[3:]) -cmd = ["cargo","clippy","--manifest-path","jouleclaw/jouleclaw-rs/Cargo.toml", - "-p",pkg,"--all-targets","--no-deps","--keep-going","--message-format=json"] -if feats: - cmd += ["--features", feats] -proc = subprocess.run(cmd, capture_output=True, text=True) - -edits = collections.defaultdict(list) -seen, skipped = set(), collections.Counter() -for line in proc.stdout.splitlines(): - try: - m = json.loads(line) - except json.JSONDecodeError: - continue - if m.get("reason") != "compiler-message": - continue - d = m.get("message", {}) - if d.get("level") != "warning": - continue - code = ((d.get("code") or {}).get("code") or "").replace("clippy::", "") - if code not in want: - continue - def walk(msg): - for sp in msg.get("spans", []): - if sp.get("suggested_replacement") is not None: - yield sp - for ch in msg.get("children", []): - yield from walk(ch) - got = list(walk(d)) - if not got: - skipped[code] += 1 - continue - for sp in got: - repl = sp["suggested_replacement"] - if re.search(r"(?", repl): - skipped[code] += 1 - continue - key = (sp["file_name"], sp["byte_start"], sp["byte_end"]) - if key in seen: - continue - seen.add(key) - edits[sp["file_name"]].append((sp["byte_start"], sp["byte_end"], repl, code)) - -WS = pathlib.Path("jouleclaw/jouleclaw-rs") -applied = collections.Counter() -for fn, es in edits.items(): - p = pathlib.Path(fn) - if not p.exists(): - p = WS / fn - if not p.exists(): - print(f" MISSING {fn}") - continue - data = p.read_bytes() - es.sort(key=lambda e: e[0], reverse=True) - last = None - for start, end, repl, code in es: - if last is not None and end > last: - skipped[code] += 1 - continue - data = data[:start] + repl.encode() + data[end:] - last = start - applied[code] += 1 - p.write_bytes(data) -print(f"applied {sum(applied.values())} edits in {len(edits)} files") -for c, n in applied.most_common(): - print(f" {n:4d} {c}") -if skipped: - print("left alone (no suggestion / overlapping / template):") - for c, n in skipped.most_common(): - print(f" {n:4d} {c}") diff --git a/bac.py b/bac.py deleted file mode 100644 index 5632aff89..000000000 --- a/bac.py +++ /dev/null @@ -1,22 +0,0 @@ -import json, subprocess, collections -cmd=["cargo","clippy","--manifest-path","jouleclaw/jouleclaw-rs/Cargo.toml", - "-p","jouleclaw-backend-apple","--all-targets","--no-deps","--keep-going","--message-format=json"] -pr=subprocess.run(cmd,capture_output=True,text=True) -c=collections.Counter(); e=0; seen=set(); files=collections.Counter() -for l in pr.stdout.splitlines(): - try: m=json.loads(l) - except Exception: continue - if m.get("reason")!="compiler-message": continue - d=m.get("message",{}); lv=d.get("level") - if lv not in ("warning","error"): continue - code=(d.get("code") or {}).get("code") or "(uncoded)" - sp=next((s for s in d.get("spans",[]) if s.get("is_primary")),None) - fn=str(sp.get("file_name") if sp else "?") - k=(code,fn,sp.get("line_start") if sp else None) - if k in seen: continue - seen.add(k) - if lv=="error": e+=1; continue - c[code]+=1; files[fn.split("/")[-1]]+=1 -print(f"errors={e} warnings={sum(c.values())}") -for k,v in c.most_common(12): print(f" {v:4d} {k}") -print(" by file:", dict(files)) diff --git a/fix-unused.py b/fix-unused.py deleted file mode 100644 index c667561ac..000000000 --- a/fix-unused.py +++ /dev/null @@ -1,63 +0,0 @@ -import json, re, subprocess, pathlib, collections - -cmd=["cargo","clippy","--manifest-path","jouleclaw/jouleclaw-rs/Cargo.toml", - "-p","jouleclaw-backend-apple","--all-targets","--no-deps","--keep-going", - "--message-format=json"] -pr=subprocess.run(cmd,capture_output=True,text=True) -WS=pathlib.Path("jouleclaw/jouleclaw-rs") - -unsafe_sites=collections.defaultdict(set) # file -> {line} -import_sites=collections.defaultdict(list) # file -> [(line, name)] -for l in pr.stdout.splitlines(): - try: m=json.loads(l) - except Exception: continue - if m.get("reason")!="compiler-message": continue - d=m.get("message",{}) - if d.get("level")!="warning": continue - code=(d.get("code") or {}).get("code") or "" - sp=next((s for s in d.get("spans",[]) if s.get("is_primary")),None) - if not sp: continue - if code=="unused_unsafe": - unsafe_sites[sp["file_name"]].add(sp["line_start"]) - elif code=="unused_imports": - for nm in re.findall(r"`([^`]+)`", d.get("message","")): - import_sites[sp["file_name"]].append((sp["line_start"], nm.strip())) - -changed=collections.Counter() -for fn in set(list(unsafe_sites) + list(import_sites)): - p=pathlib.Path(fn) - if not p.exists(): p=WS/fn - lines=p.read_text().split("\n") - - # `|| unsafe {` -> `|| {`: rustc says the block is redundant, so dropping - # the keyword cannot change behaviour. - for ln in sorted(unsafe_sites.get(fn,())): - i=ln-1 - if i < len(lines) and "unsafe" in lines[i]: - new=re.sub(r"\bunsafe\s+\{", "{", lines[i], count=1) - if new!=lines[i]: - lines[i]=new; changed["unused_unsafe"]+=1 - - # Drop the one unused name from a brace list, or the whole `use` if it was - # the only name. - for ln, nm in sorted(set(import_sites.get(fn,[])), reverse=True): - i=ln-1 - if i >= len(lines): continue - s=lines[i] - base=nm.split("::")[-1] - # `Name as _, ` or `, Name as _` or bare `Name, ` / `, Name` - pat=rf"\b{re.escape(base)}(\s+as\s+_)?\s*,\s*|\s*,\s*\b{re.escape(base)}(\s+as\s+_)?\b" - new=re.sub(pat, "", s, count=1) - if new==s: - continue - # A brace list reduced to one entry reads better unbraced. - new=re.sub(r"\{\s*([A-Za-z_][\w:]*(\s+as\s+_)?)\s*\}", r"\1", new) - if re.search(r"\{\s*\}", new): - lines[i]=None # the import is now empty; drop the line - else: - lines[i]=new - changed["unused_imports"]+=1 - - p.write_text("\n".join(l for l in lines if l is not None)) - -print("applied:", dict(changed)) diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs index ddc1a4e31..8c45c3a8d 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/lib.rs @@ -20,6 +20,14 @@ //! returns empty and the kernel constructors return `None`. The runtime //! falls back to the reference backend. +// ── lints this crate has decided against ──────────────────────────────────── +// +// The same two decisions jouleclaw-omni and jouleclaw-loader-gguf record: a +// GPU dispatch signature is its kernel's parameter list, and these kernels +// walk several buffers at one index. +#![allow(clippy::too_many_arguments)] +#![allow(clippy::needless_range_loop)] + mod accelerate; mod gpu_load; mod metal_compute; diff --git a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs index 998e3cb1b..7476d17ef 100644 --- a/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs +++ b/jouleclaw/jouleclaw-rs/crates/jouleclaw-backend-apple/src/mps_matmul.rs @@ -676,12 +676,10 @@ kernel void silu_mul_rows( .as_ref() } - // SAFETY: pipeline state objects are immutable + thread-safe. - struct PipelineCell( - Option>>, - ); - unsafe impl Send for PipelineCell {} - unsafe impl Sync for PipelineCell {} + /// The `(buffer, offset, length)` triples a dispatch keeps alive until + /// its command buffer completes. Named because the inline form is three + /// generic layers deep and appears in two places. + type KeepAlive = Vec<(Retained>, usize, usize)>; /// One MoE layer in ONE command buffer: per pick /// `(gate, up, down)` resident F16 weights; gate/up GEMMs @@ -1970,6 +1968,9 @@ kernel void attn_decode( pub fn len(&self) -> usize { self.len } + pub fn is_empty(&self) -> bool { + self.len == 0 + } pub fn hd(&self) -> usize { self.hd } @@ -2062,7 +2063,6 @@ kernel void attn_decode( Some(scale_qk_pipeline().ok_or("scale_qk pipeline unavailable")?) } }; - let qk_pipe = qk_pipe; let append_pipe = kv_append_pipeline().ok_or("kv_append pipeline unavailable")?; let attn_pipe = attn_decode_pipeline().ok_or("attn pipeline unavailable")?; // Slot bookkeeping (CPU side, before encode). @@ -3445,7 +3445,7 @@ kernel void sgemm_q6k(device const float* A [[buffer(0)]], let _g = ctx.encode_lock.lock().unwrap(); let cb = ctx.queue.commandBuffer().ok_or("commandBuffer failed")?; // Keep tile buffers alive until after the single wait. - let mut keep: Vec<(Retained>, usize, usize)> = Vec::new(); + let mut keep: KeepAlive = Vec::new(); for t in 0..n_tiles { let row0 = t * TILE; let rows = (m - row0).min(TILE); @@ -3804,7 +3804,7 @@ kernel void gemm_q6k_m4(device const uchar* W [[buffer(0)]], }; let _g = ctx.encode_lock.lock().unwrap(); let cb = ctx.queue.commandBuffer().ok_or("commandBuffer failed")?; - let mut keep: Vec<(Retained>, usize, usize)> = Vec::new(); + let mut keep: KeepAlive = Vec::new(); for t in 0..n_tiles { let row0 = t * TILE; let rows = (m - row0).min(TILE); @@ -4519,7 +4519,6 @@ kernel void batched_causal_attn_f32( Some(scale_qk_pipeline().ok_or("scale_qk pipeline unavailable")?) } }; - let qk_pipe = qk_pipe; let append_pipe = kv_append_pipeline().ok_or("kv_append pipeline unavailable")?; let attn_pipe = attn_decode_f32_pipeline().ok_or("attn pipeline unavailable")?; @@ -4895,7 +4894,6 @@ kernel void batched_causal_attn_f32( } QkStage::ScaleQk { .. } => Some(scale_qk_pipeline().ok_or("scale_qk pipeline")?), }; - let qk_pipe = qk_pipe; let append_pipe = kv_append_pipeline().ok_or("kv_append pipeline")?; let attn_pipe = attn_decode_f32_pipeline().ok_or("attn pipeline")?; -- 2.54.0 (Apple Git-157)