From a17939116df300e843dc1962cb99f8a4a4c8d6bf Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 11 Oct 2026 21:05:13 -0700 Subject: [PATCH 1/5] ggml-metal: native ROPE_APPLY kernel (precomputed RoPE used by Flux, Z-Image, Qwen-Image, Wan, LTXV) --- .../0010-ggml-metal-rope-apply.patch | 207 ++++++++++++++++++ 1 file changed, 207 insertions(+) create mode 100644 scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch diff --git a/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch b/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch new file mode 100644 index 0000000000..bf4918202f --- /dev/null +++ b/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch @@ -0,0 +1,207 @@ +diff --git a/src/ggml-metal/ggml-metal-device.cpp b/src/ggml-metal/ggml-metal-device.cpp +index dc6b695e..d1a624ad 100644 +--- a/src/ggml-metal/ggml-metal-device.cpp ++++ b/src/ggml-metal/ggml-metal-device.cpp +@@ -2391,6 +2391,20 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_roll(ggml_metal_ + return res; + } + ++ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_rope_apply(ggml_metal_library_t lib, const ggml_tensor * op) { ++ assert(op->op == GGML_OP_ROPE_APPLY); ++ GGML_UNUSED(op); ++ ++ const char * name = "kernel_rope_apply_f32"; ++ ++ ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); ++ if (!res.pipeline) { ++ res = ggml_metal_library_compile_pipeline(lib, name, name, nullptr); ++ } ++ ++ return res; ++} ++ + ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad(ggml_metal_library_t lib, const ggml_tensor * op) { + assert(op->op == GGML_OP_PAD); + +diff --git a/src/ggml-metal/ggml-metal-device.h b/src/ggml-metal/ggml-metal-device.h +index 1bdaecc7..b41813a9 100644 +--- a/src/ggml-metal/ggml-metal-device.h ++++ b/src/ggml-metal/ggml-metal-device.h +@@ -171,6 +171,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_upscale + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad (ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pad_reflect_1d (ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_roll (ggml_metal_library_t lib, const struct ggml_tensor * op); ++struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_rope_apply (ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_arange (ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_timestep_embedding(ggml_metal_library_t lib, const struct ggml_tensor * op); + struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_opt_step_adamw (ggml_metal_library_t lib, const struct ggml_tensor * op); +diff --git a/src/ggml-metal/ggml-metal-device.m b/src/ggml-metal/ggml-metal-device.m +index a81d3a8c..498ee7ca 100644 +--- a/src/ggml-metal/ggml-metal-device.m ++++ b/src/ggml-metal/ggml-metal-device.m +@@ -1681,6 +1681,10 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te + case GGML_OP_ROPE: + case GGML_OP_ROPE_BACK: + return true; ++ case GGML_OP_ROPE_APPLY: ++ return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && ++ op->type == GGML_TYPE_F32 && ggml_is_contiguous(op) && ++ op->src[0]->ne[1] <= 65535 && op->src[0]->ne[3] <= 65535; + case GGML_OP_IM2COL: + return ggml_is_contiguous(op->src[1]) && op->src[1]->type == GGML_TYPE_F32 && (op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_F32); + case GGML_OP_CONV_2D: +diff --git a/src/ggml-metal/ggml-metal-impl.h b/src/ggml-metal/ggml-metal-impl.h +index 490dd83a..d61e0fcb 100644 +--- a/src/ggml-metal/ggml-metal-impl.h ++++ b/src/ggml-metal/ggml-metal-impl.h +@@ -1166,6 +1166,22 @@ typedef struct { + int32_t s3; + } ggml_metal_kargs_roll; + ++typedef struct { ++ int32_t ne00; ++ int32_t ne01; ++ int32_t ne02; ++ int32_t ne03; ++ uint64_t nb00; ++ uint64_t nb01; ++ uint64_t nb02; ++ uint64_t nb03; ++ uint64_t nb10; ++ uint64_t nb11; ++ uint64_t nb12; ++ uint64_t nb13; ++ int32_t interleaved; ++} ggml_metal_kargs_rope_apply; ++ + typedef struct { + uint64_t nb1; + int dim; +diff --git a/src/ggml-metal/ggml-metal-ops.cpp b/src/ggml-metal/ggml-metal-ops.cpp +index 708703e5..e11d7c25 100644 +--- a/src/ggml-metal/ggml-metal-ops.cpp ++++ b/src/ggml-metal/ggml-metal-ops.cpp +@@ -410,6 +410,10 @@ static int ggml_metal_op_encode_impl(ggml_metal_op_t ctx, int idx) { + { + n_fuse = ggml_metal_op_rope(ctx, idx); + } break; ++ case GGML_OP_ROPE_APPLY: ++ { ++ n_fuse = ggml_metal_op_rope_apply(ctx, idx); ++ } break; + case GGML_OP_IM2COL: + { + n_fuse = ggml_metal_op_im2col(ctx, idx); +@@ -4977,6 +4981,48 @@ int ggml_metal_op_roll(ggml_metal_op_t ctx, int idx) { + return 1; + } + ++int ggml_metal_op_rope_apply(ggml_metal_op_t ctx, int idx) { ++ ggml_tensor * op = ctx->node(idx); ++ ++ ggml_metal_library_t lib = ctx->lib; ++ ggml_metal_encoder_t enc = ctx->enc; ++ ++ GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne); ++ GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb); ++ GGML_TENSOR_LOCALS(uint64_t, nb1, op->src[1], nb); ++ ++ ggml_metal_kargs_rope_apply args = { ++ /*.ne00 =*/ ne00, ++ /*.ne01 =*/ ne01, ++ /*.ne02 =*/ ne02, ++ /*.ne03 =*/ ne03, ++ /*.nb00 =*/ nb00, ++ /*.nb01 =*/ nb01, ++ /*.nb02 =*/ nb02, ++ /*.nb03 =*/ nb03, ++ /*.nb10 =*/ nb10, ++ /*.nb11 =*/ nb11, ++ /*.nb12 =*/ nb12, ++ /*.nb13 =*/ nb13, ++ /*.interleaved =*/ ggml_get_op_params_i32(op, 0) != 0 ? 1 : 0, ++ }; ++ ++ auto pipeline = ggml_metal_library_get_pipeline_rope_apply(lib, op); ++ ++ // one threadgroup per (token, head, batch) row, threads over the d/2 rotation pairs ++ const int nth = std::max(1, std::min(ne00/2, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline))); ++ ++ ggml_metal_encoder_set_pipeline(enc, pipeline); ++ ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0); ++ ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1); ++ ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), 2); ++ ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 3); ++ ++ ggml_metal_encoder_dispatch_threadgroups(enc, ne02, ne01, ne03, nth, 1, 1); ++ ++ return 1; ++} ++ + int ggml_metal_op_pad(ggml_metal_op_t ctx, int idx) { + ggml_tensor * op = ctx->node(idx); + +diff --git a/src/ggml-metal/ggml-metal-ops.h b/src/ggml-metal/ggml-metal-ops.h +index 583d1156..ef6be30f 100644 +--- a/src/ggml-metal/ggml-metal-ops.h ++++ b/src/ggml-metal/ggml-metal-ops.h +@@ -93,6 +93,7 @@ int ggml_metal_op_upscale (ggml_metal_op_t ctx, int idx); + int ggml_metal_op_pad (ggml_metal_op_t ctx, int idx); + int ggml_metal_op_pad_reflect_1d (ggml_metal_op_t ctx, int idx); + int ggml_metal_op_roll (ggml_metal_op_t ctx, int idx); ++int ggml_metal_op_rope_apply (ggml_metal_op_t ctx, int idx); + int ggml_metal_op_arange (ggml_metal_op_t ctx, int idx); + int ggml_metal_op_timestep_embedding(ggml_metal_op_t ctx, int idx); + int ggml_metal_op_argmax (ggml_metal_op_t ctx, int idx); +diff --git a/src/ggml-metal/kernels/rope.metal b/src/ggml-metal/kernels/rope.metal +index 401ceacb..bdcc1283 100644 +--- a/src/ggml-metal/kernels/rope.metal ++++ b/src/ggml-metal/kernels/rope.metal +@@ -331,3 +331,47 @@ template [[host_name("kernel_rope_multi_f16")]] kernel kernel_rope_multi_t kerne + + template [[host_name("kernel_rope_vision_f32")]] kernel kernel_rope_vision_t kernel_rope_vision; + template [[host_name("kernel_rope_vision_f16")]] kernel kernel_rope_vision_t kernel_rope_vision; ++ ++// precomputed RoPE: x [d, n_head, L, N] (any strides), pe [2, 2, d/2, L] rotation matrices ++// [[cos, -sin], [sin, cos]]; dst is contiguous [d, L, n_head*N] ++kernel void kernel_rope_apply_f32( ++ constant ggml_metal_kargs_rope_apply & args, ++ device const char * src0, ++ device const char * src1, ++ device char * dst, ++ uint3 tgpig[[threadgroup_position_in_grid]], ++ uint3 tpitg[[thread_position_in_threadgroup]], ++ uint3 ntg[[threads_per_threadgroup]]) { ++ const int64_t token = tgpig.x; ++ const int64_t head = tgpig.y; ++ const int64_t batch = tgpig.z; ++ ++ const int64_t half_dim = args.ne00/2; ++ ++ device const char * x_row = src0 + head*args.nb01 + token*args.nb02 + batch*args.nb03; ++ device const char * pe_row = src1 + token*args.nb13; ++ device float * out = (device float *) dst + (token + args.ne02*(head + args.ne01*batch))*args.ne00; ++ ++ for (int64_t j = tpitg.x; j < half_dim; j += ntg.x) { ++ const int64_t i0 = args.interleaved ? 2*j : j; ++ const int64_t i1 = args.interleaved ? 2*j + 1 : j + half_dim; ++ ++ const float a = *(device const float *) (x_row + i0*args.nb00); ++ const float b = *(device const float *) (x_row + i1*args.nb00); ++ ++ device const char * m = pe_row + j*args.nb12; ++ ++ const float c = *(device const float *) (m); ++ const float s = *(device const float *) (m + args.nb11); ++ ++ if (args.interleaved) { ++ const float neg_s = *(device const float *) (m + args.nb10); ++ const float c1 = *(device const float *) (m + args.nb10 + args.nb11); ++ out[i0] = a*c + b*neg_s; ++ out[i1] = a*s + b*c1; ++ } else { ++ out[i0] = a*c - b*s; ++ out[i1] = a*s + b*c; ++ } ++ } ++} From dd51513832e6df641fbccf218d6d2cf076cfb943 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 11 Oct 2026 21:05:13 -0700 Subject: [PATCH 2/5] ggml-metal: report FP8 and INT8 mul_mat as unsupported instead of aborting at pipeline lookup --- .../0011-ggml-metal-no-fp8-int8-mul-mat.patch | 20 +++++++++++++++++++ 1 file changed, 20 insertions(+) create mode 100644 scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch diff --git a/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch b/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch new file mode 100644 index 0000000000..b85993458c --- /dev/null +++ b/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch @@ -0,0 +1,20 @@ +diff --git a/src/ggml-metal/ggml-metal-device.m b/src/ggml-metal/ggml-metal-device.m +index 498ee7ca..418e6f04 100644 +--- a/src/ggml-metal/ggml-metal-device.m ++++ b/src/ggml-metal/ggml-metal-device.m +@@ -1507,6 +1507,15 @@ static bool ggml_metal_supports_mul_mat_op( + return false; + } + ++ // no mul_mv / mul_mm kernels exist for FP8 or INT8 weights: claiming support here makes the ++ // pipeline lookup fail and abort at encode time instead of letting the caller cast the weight ++ // (sd.cpp casts FP8 to BF16 when this returns false) or the scheduler run the op elsewhere ++ if (op->src[0]->type == GGML_TYPE_F8_E4M3 || ++ op->src[0]->type == GGML_TYPE_F8_E5M2 || ++ op->src[0]->type == GGML_TYPE_I8) { ++ return false; ++ } ++ + if (op->src[1]->type != GGML_TYPE_F16) { + return true; + } From 4381979e798244e94736425215b8045faf30c540 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 11 Oct 2026 21:05:13 -0700 Subject: [PATCH 3/5] ggml-metal: group norm handles batch > 1, sizes its threadgroup to the group, and fixes the reduction buffer reuse --- .../0012-ggml-metal-group-norm-threads.patch | 66 +++++++++++++++++++ 1 file changed, 66 insertions(+) create mode 100644 scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch diff --git a/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch b/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch new file mode 100644 index 0000000000..081d41aef9 --- /dev/null +++ b/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch @@ -0,0 +1,66 @@ +diff --git a/src/ggml-metal/ggml-metal-ops.cpp b/src/ggml-metal/ggml-metal-ops.cpp +index e11d7c25..33121471 100644 +--- a/src/ggml-metal/ggml-metal-ops.cpp ++++ b/src/ggml-metal/ggml-metal-ops.cpp +@@ -4103,13 +4103,16 @@ int ggml_metal_op_group_norm(ggml_metal_op_t ctx, int idx) { + + auto pipeline = ggml_metal_library_get_pipeline_group_norm(lib, op); + +- int nth = 32; // SIMD width +- //while (nth < ne00/4 && nth < ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)) { +- // nth *= 2; +- //} ++ // one threadgroup per group: with a single simdgroup per group, a VAE / UNet group norm ++ // (32 groups of up to millions of elements) occupies only 32 simdgroups of the whole GPU. ++ // Scale the threadgroup with the group size (the kernel reduces across simdgroups through ++ // the 32-float threadgroup buffer, so up to 32 simdgroups = 1024 threads). ++ const int64_t gs = (int64_t) ne00*ne01*((ne02 + ngrp - 1)/ngrp); + +- //nth = std::min(nth, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)); +- //nth = std::min(nth, ne00/4); ++ int nth = 32; // SIMD width ++ while (nth < 1024 && (int64_t) nth*8 <= gs && nth < ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)) { ++ nth *= 2; ++ } + + const size_t smem = pipeline.smem; + +@@ -4120,7 +4123,9 @@ int ggml_metal_op_group_norm(ggml_metal_op_t ctx, int idx) { + + ggml_metal_encoder_set_threadgroup_memory_size(enc, smem, 0); + +- ggml_metal_encoder_dispatch_threadgroups(enc, ngrp, 1, 1, nth, 1, 1); ++ // one threadgroup per (group, batch): the kernel used to ignore ne03, leaving every batch ++ // after the first unnormalized ++ ggml_metal_encoder_dispatch_threadgroups(enc, ngrp*ne03, 1, 1, nth, 1, 1); + + return 1; + } +diff --git a/src/ggml-metal/kernels/norm.metal b/src/ggml-metal/kernels/norm.metal +index c76d0c7f..613a8ca6 100644 +--- a/src/ggml-metal/kernels/norm.metal ++++ b/src/ggml-metal/kernels/norm.metal +@@ -249,7 +249,12 @@ kernel void kernel_group_norm_f32( + const int64_t ne = args.ne00*args.ne01*args.ne02; + const int64_t gs = args.ne00*args.ne01*((args.ne02 + args.ngrp - 1) / args.ngrp); + +- int start = tgpig * gs; ++ // threadgroup tgpig normalizes group (tgpig % ngrp) of batch (tgpig / ngrp); src0 is contiguous ++ const int64_t i03 = tgpig / args.ngrp; ++ src0 += i03*ne; ++ dst += i03*ne; ++ ++ int start = (tgpig % args.ngrp) * gs; + int end = start + gs; + + start += tpitg; +@@ -294,6 +299,9 @@ kernel void kernel_group_norm_f32( + + tmp = simd_sum(tmp); + if (ntg > N_SIMDWIDTH) { ++ // every simdgroup must have read the mean reduction from buf before it is reused ++ threadgroup_barrier(mem_flags::mem_threadgroup); ++ + if (sgitg == 0) { + buf[tiisg] = 0.0f; + } From ad2f736291d4123639c99005529cca81d68dd920 Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 11 Oct 2026 21:05:14 -0700 Subject: [PATCH 4/5] ggml-backend: run an in-place op on the backend that owns its destination memory --- ...-ggml-sched-inplace-op-owner-backend.patch | 36 +++++++++++++++++++ 1 file changed, 36 insertions(+) create mode 100644 scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch diff --git a/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch b/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch new file mode 100644 index 0000000000..e428e95b4b --- /dev/null +++ b/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch @@ -0,0 +1,36 @@ +diff --git a/src/ggml-backend.cpp b/src/ggml-backend.cpp +index 8c348734..067c2376 100644 +--- a/src/ggml-backend.cpp ++++ b/src/ggml-backend.cpp +@@ -1294,6 +1294,31 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra + GGML_ASSERT(*cur_backend_id != -1); + } + ++ // pass 4.5: an in-place op (a non-view op whose result is a view of another tensor, e.g. ++ // ggml_mul_inplace / ggml_silu_inplace) writes into the memory of its view_src. If the op was ++ // assigned (by expansion or by a user pin) to a backend that cannot use that memory, because ++ // view_src is produced by an op another backend had to run, the kernel receives a pointer into ++ // a foreign buffer (Metal: "buffer is nil" and garbage / a crash; CUDA: an illegal address). ++ // Run such an op on the backend that owns the memory instead; pass 5 copies its other inputs. ++ for (int i = 0; i < graph->n_nodes; i++) { ++ struct ggml_tensor * node = graph->nodes[i]; ++ if (node->view_src == NULL || ggml_is_view_op(node->op)) { ++ continue; ++ } ++ int * cur_backend_id = &tensor_backend_id(node); ++ const int owner_id = node->view_src->buffer != NULL ++ ? ggml_backend_sched_backend_from_buffer(sched, node->view_src, node) ++ : tensor_backend_id(node->view_src); ++ if (owner_id == -1 || owner_id == *cur_backend_id || ++ ggml_backend_sched_buffer_supported(sched, node->view_src, *cur_backend_id)) { ++ continue; ++ } ++ if (ggml_backend_supports_op(sched->backends[owner_id], node)) { ++ *cur_backend_id = owner_id; ++ SET_CAUSE(node, "4.5.inplace"); ++ } ++ } ++ + // pass 5: split graph, find tensors that need to be copied + { + int i_split = 0; From 55575cfedad080d5c28ae09324a9cbbe8a6da3fd Mon Sep 17 00:00:00 2001 From: Daniel Han Date: Sun, 11 Oct 2026 21:47:33 -0700 Subject: [PATCH 5/5] Trim comments --- .../0010-ggml-metal-rope-apply.patch | 12 +++++------- .../0011-ggml-metal-no-fp8-int8-mul-mat.patch | 8 +++----- .../0012-ggml-metal-group-norm-threads.patch | 17 ++++++----------- ...13-ggml-sched-inplace-op-owner-backend.patch | 12 ++++-------- 4 files changed, 18 insertions(+), 31 deletions(-) diff --git a/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch b/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch index bf4918202f..f1ed135438 100644 --- a/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch +++ b/scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch @@ -78,7 +78,7 @@ index 490dd83a..d61e0fcb 100644 uint64_t nb1; int dim; diff --git a/src/ggml-metal/ggml-metal-ops.cpp b/src/ggml-metal/ggml-metal-ops.cpp -index 708703e5..e11d7c25 100644 +index 708703e5..031adb8a 100644 --- a/src/ggml-metal/ggml-metal-ops.cpp +++ b/src/ggml-metal/ggml-metal-ops.cpp @@ -410,6 +410,10 @@ static int ggml_metal_op_encode_impl(ggml_metal_op_t ctx, int idx) { @@ -92,7 +92,7 @@ index 708703e5..e11d7c25 100644 case GGML_OP_IM2COL: { n_fuse = ggml_metal_op_im2col(ctx, idx); -@@ -4977,6 +4981,48 @@ int ggml_metal_op_roll(ggml_metal_op_t ctx, int idx) { +@@ -4977,6 +4981,47 @@ int ggml_metal_op_roll(ggml_metal_op_t ctx, int idx) { return 1; } @@ -124,7 +124,6 @@ index 708703e5..e11d7c25 100644 + + auto pipeline = ggml_metal_library_get_pipeline_rope_apply(lib, op); + -+ // one threadgroup per (token, head, batch) row, threads over the d/2 rotation pairs + const int nth = std::max(1, std::min(ne00/2, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline))); + + ggml_metal_encoder_set_pipeline(enc, pipeline); @@ -154,16 +153,15 @@ index 583d1156..ef6be30f 100644 int ggml_metal_op_timestep_embedding(ggml_metal_op_t ctx, int idx); int ggml_metal_op_argmax (ggml_metal_op_t ctx, int idx); diff --git a/src/ggml-metal/kernels/rope.metal b/src/ggml-metal/kernels/rope.metal -index 401ceacb..bdcc1283 100644 +index 401ceacb..b47563c6 100644 --- a/src/ggml-metal/kernels/rope.metal +++ b/src/ggml-metal/kernels/rope.metal -@@ -331,3 +331,47 @@ template [[host_name("kernel_rope_multi_f16")]] kernel kernel_rope_multi_t kerne +@@ -331,3 +331,46 @@ template [[host_name("kernel_rope_multi_f16")]] kernel kernel_rope_multi_t kerne template [[host_name("kernel_rope_vision_f32")]] kernel kernel_rope_vision_t kernel_rope_vision; template [[host_name("kernel_rope_vision_f16")]] kernel kernel_rope_vision_t kernel_rope_vision; + -+// precomputed RoPE: x [d, n_head, L, N] (any strides), pe [2, 2, d/2, L] rotation matrices -+// [[cos, -sin], [sin, cos]]; dst is contiguous [d, L, n_head*N] ++// x [d, n_head, L, N] (any strides), pe [2, 2, d/2, L] of [[cos, -sin], [sin, cos]]; dst contiguous [d, L, n_head*N] +kernel void kernel_rope_apply_f32( + constant ggml_metal_kargs_rope_apply & args, + device const char * src0, diff --git a/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch b/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch index b85993458c..69eb343e59 100644 --- a/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch +++ b/scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch @@ -1,14 +1,12 @@ diff --git a/src/ggml-metal/ggml-metal-device.m b/src/ggml-metal/ggml-metal-device.m -index 498ee7ca..418e6f04 100644 +index 498ee7ca..80aaf0ec 100644 --- a/src/ggml-metal/ggml-metal-device.m +++ b/src/ggml-metal/ggml-metal-device.m -@@ -1507,6 +1507,15 @@ static bool ggml_metal_supports_mul_mat_op( +@@ -1507,6 +1507,13 @@ static bool ggml_metal_supports_mul_mat_op( return false; } -+ // no mul_mv / mul_mm kernels exist for FP8 or INT8 weights: claiming support here makes the -+ // pipeline lookup fail and abort at encode time instead of letting the caller cast the weight -+ // (sd.cpp casts FP8 to BF16 when this returns false) or the scheduler run the op elsewhere ++ // no FP8/INT8 mul_mv/mul_mm kernels: claiming support aborts at encode instead of letting sd.cpp cast or the op fall back + if (op->src[0]->type == GGML_TYPE_F8_E4M3 || + op->src[0]->type == GGML_TYPE_F8_E5M2 || + op->src[0]->type == GGML_TYPE_I8) { diff --git a/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch b/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch index 081d41aef9..2d7051b70d 100644 --- a/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch +++ b/scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch @@ -1,8 +1,8 @@ diff --git a/src/ggml-metal/ggml-metal-ops.cpp b/src/ggml-metal/ggml-metal-ops.cpp -index e11d7c25..33121471 100644 +index 031adb8a..aacf12c0 100644 --- a/src/ggml-metal/ggml-metal-ops.cpp +++ b/src/ggml-metal/ggml-metal-ops.cpp -@@ -4103,13 +4103,16 @@ int ggml_metal_op_group_norm(ggml_metal_op_t ctx, int idx) { +@@ -4103,13 +4103,13 @@ int ggml_metal_op_group_norm(ggml_metal_op_t ctx, int idx) { auto pipeline = ggml_metal_library_get_pipeline_group_norm(lib, op); @@ -10,10 +10,7 @@ index e11d7c25..33121471 100644 - //while (nth < ne00/4 && nth < ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)) { - // nth *= 2; - //} -+ // one threadgroup per group: with a single simdgroup per group, a VAE / UNet group norm -+ // (32 groups of up to millions of elements) occupies only 32 simdgroups of the whole GPU. -+ // Scale the threadgroup with the group size (the kernel reduces across simdgroups through -+ // the 32-float threadgroup buffer, so up to 32 simdgroups = 1024 threads). ++ // one simdgroup per group starves the GPU on VAE group norms; the 32-float reduction buffer caps nth at 1024 + const int64_t gs = (int64_t) ne00*ne01*((ne02 + ngrp - 1)/ngrp); - //nth = std::min(nth, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)); @@ -25,19 +22,17 @@ index e11d7c25..33121471 100644 const size_t smem = pipeline.smem; -@@ -4120,7 +4123,9 @@ int ggml_metal_op_group_norm(ggml_metal_op_t ctx, int idx) { +@@ -4120,7 +4120,7 @@ int ggml_metal_op_group_norm(ggml_metal_op_t ctx, int idx) { ggml_metal_encoder_set_threadgroup_memory_size(enc, smem, 0); - ggml_metal_encoder_dispatch_threadgroups(enc, ngrp, 1, 1, nth, 1, 1); -+ // one threadgroup per (group, batch): the kernel used to ignore ne03, leaving every batch -+ // after the first unnormalized + ggml_metal_encoder_dispatch_threadgroups(enc, ngrp*ne03, 1, 1, nth, 1, 1); return 1; } diff --git a/src/ggml-metal/kernels/norm.metal b/src/ggml-metal/kernels/norm.metal -index c76d0c7f..613a8ca6 100644 +index c76d0c7f..a301bbf0 100644 --- a/src/ggml-metal/kernels/norm.metal +++ b/src/ggml-metal/kernels/norm.metal @@ -249,7 +249,12 @@ kernel void kernel_group_norm_f32( @@ -45,7 +40,7 @@ index c76d0c7f..613a8ca6 100644 const int64_t gs = args.ne00*args.ne01*((args.ne02 + args.ngrp - 1) / args.ngrp); - int start = tgpig * gs; -+ // threadgroup tgpig normalizes group (tgpig % ngrp) of batch (tgpig / ngrp); src0 is contiguous ++ // src0 is contiguous + const int64_t i03 = tgpig / args.ngrp; + src0 += i03*ne; + dst += i03*ne; diff --git a/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch b/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch index e428e95b4b..c3d3b8579e 100644 --- a/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch +++ b/scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch @@ -1,17 +1,13 @@ diff --git a/src/ggml-backend.cpp b/src/ggml-backend.cpp -index 8c348734..067c2376 100644 +index 8c348734..cc597430 100644 --- a/src/ggml-backend.cpp +++ b/src/ggml-backend.cpp -@@ -1294,6 +1294,31 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra +@@ -1294,6 +1294,27 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra GGML_ASSERT(*cur_backend_id != -1); } -+ // pass 4.5: an in-place op (a non-view op whose result is a view of another tensor, e.g. -+ // ggml_mul_inplace / ggml_silu_inplace) writes into the memory of its view_src. If the op was -+ // assigned (by expansion or by a user pin) to a backend that cannot use that memory, because -+ // view_src is produced by an op another backend had to run, the kernel receives a pointer into -+ // a foreign buffer (Metal: "buffer is nil" and garbage / a crash; CUDA: an illegal address). -+ // Run such an op on the backend that owns the memory instead; pass 5 copies its other inputs. ++ // pass 4.5: an in-place op writes into its view_src's memory, so run it on the backend owning ++ // that buffer; elsewhere it gets a foreign pointer (Metal "buffer is nil", CUDA illegal address) + for (int i = 0; i < graph->n_nodes; i++) { + struct ggml_tensor * node = graph->nodes[i]; + if (node->view_src == NULL || ggml_is_view_op(node->op)) {