Skip to content

Metal: in-place ops on the owning backend, ROPE_APPLY kernel, group norm batch fix, no FP8/INT8 mul_mat abort #30

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
danielhanchen merged 5 commits into from
Oct 12, 2026
Merged

Metal: in-place ops on the owning backend, ROPE_APPLY kernel, group norm batch fix, no FP8/INT8 mul_mat abort #30

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
a179391
ggml-metal: native ROPE_APPLY kernel (precomputed RoPE used by Flux, …
danielhanchen Oct 12, 2026
dd51513
ggml-metal: report FP8 and INT8 mul_mat as unsupported instead of abo…
danielhanchen Oct 12, 2026
4381979
ggml-metal: group norm handles batch > 1, sizes its threadgroup to th…
danielhanchen Oct 12, 2026
ad2f736
ggml-backend: run an in-place op on the backend that owns its destina…
danielhanchen Oct 12, 2026
55575cf
Trim comments
danielhanchen Oct 12, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
205 changes: 205 additions & 0 deletions scripts/unsloth/ggml-patches/0010-ggml-metal-rope-apply.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,205 @@
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..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) {
{
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,47 @@ 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);
+
+ 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..b47563c6 100644
--- a/src/ggml-metal/kernels/rope.metal
+++ b/src/ggml-metal/kernels/rope.metal
@@ -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<float>;
template [[host_name("kernel_rope_vision_f16")]] kernel kernel_rope_vision_t kernel_rope_vision<half>;
+
+// 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,
+ 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;
+ }
+ }
+}
18 changes: 18 additions & 0 deletions scripts/unsloth/ggml-patches/0011-ggml-metal-no-fp8-int8-mul-mat.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,18 @@
diff --git a/src/ggml-metal/ggml-metal-device.m b/src/ggml-metal/ggml-metal-device.m
index 498ee7ca..80aaf0ec 100644
--- a/src/ggml-metal/ggml-metal-device.m
+++ b/src/ggml-metal/ggml-metal-device.m
@@ -1507,6 +1507,13 @@ static bool ggml_metal_supports_mul_mat_op(
return false;
}

+ // 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) {
+ return false;
+ }
+
if (op->src[1]->type != GGML_TYPE_F16) {
return true;
}
61 changes: 61 additions & 0 deletions scripts/unsloth/ggml-patches/0012-ggml-metal-group-norm-threads.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,61 @@
diff --git a/src/ggml-metal/ggml-metal-ops.cpp b/src/ggml-metal/ggml-metal-ops.cpp
index 031adb8a..aacf12c0 100644
--- a/src/ggml-metal/ggml-metal-ops.cpp
+++ b/src/ggml-metal/ggml-metal-ops.cpp
@@ -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);

- int nth = 32; // SIMD width
- //while (nth < ne00/4 && nth < ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)) {
- // nth *= 2;
- //}
+ // 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));
- //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 +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);
+ 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..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(
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;
+ // 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;
}
32 changes: 32 additions & 0 deletions scripts/unsloth/ggml-patches/0013-ggml-sched-inplace-op-owner-backend.patch
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,32 @@
diff --git a/src/ggml-backend.cpp b/src/ggml-backend.cpp
index 8c348734..cc597430 100644
--- a/src/ggml-backend.cpp
+++ b/src/ggml-backend.cpp
@@ -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 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)) {
+ 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;