upstream-archive/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-21 13:37:29 -05:00
Code Issues Packages Projects Releases Wiki Activity

metal : support qwen4exp hc ops (#29000)

Browse Source 
Add support for the new DSV4 HC op variants used by qwen4exp:
- hc_pre with per-element sigmoid gate (gated variant)
- hc_post with identity mixing (comb == nullptr)

Assisted-by: pi:llama.cpp/Qwen3.8-27B
This commit is contained in:
Georgi Gerganov
2026-09-19 11:27:30 +03:00
committed by GitHub
parent b23701f77d
commit 59fc5a1ca3
6 changed files with 106 additions and 22 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-metal/ggml-metal-device.cpp
+21 -6
View File
@@ -496,14 +496,29 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_lightning_indexe
return res;
}
ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_dsv4_hc(ggml_metal_library_t lib, ggml_op op) {
ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_dsv4_hc(ggml_metal_library_t lib, const ggml_tensor * op) {
const char * name = nullptr;
switch (op) {
case GGML_OP_DSV4_HC_COMB: name = "kernel_dsv4_hc_comb_f32"; break;
case GGML_OP_DSV4_HC_PRE: name = "kernel_dsv4_hc_pre_f32"; break;
case GGML_OP_DSV4_HC_POST: name = "kernel_dsv4_hc_post_f32"; break;
default: GGML_ABORT("fatal error");
switch (op->op) {
case GGML_OP_DSV4_HC_COMB:
name = "kernel_dsv4_hc_comb_f32";
break;
case GGML_OP_DSV4_HC_PRE:
if (ggml_get_op_params_i32(op, 1) != 0) {
name = "kernel_dsv4_hc_pre_gated_f32";
} else {
name = "kernel_dsv4_hc_pre_f32";
}
break;
case GGML_OP_DSV4_HC_POST:
if (op->src[3]) {
name = "kernel_dsv4_hc_post_f32";
} else {
name = "kernel_dsv4_hc_post_nocomb_f32";
}
break;
default:
GGML_ABORT("fatal error");
}
ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
ggml/src/ggml-metal/ggml-metal-device.h
+1 -1
View File
@@ -126,7 +126,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_cumsum_ad
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_tri (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_soft_max (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_lightning_indexer (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_dsv4_hc (ggml_metal_library_t lib, enum ggml_op op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_dsv4_hc (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_ssm_conv (ggml_metal_library_t lib, const struct ggml_tensor * op);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_ssm_conv_batched (ggml_metal_library_t lib, const struct ggml_tensor * op, int ssm_conv_bs);
struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_ssm_scan (ggml_metal_library_t lib, const struct ggml_tensor * op, bool tail);
ggml/src/ggml-metal/ggml-metal-device.m
+3 -7
View File
@@ -1802,8 +1802,6 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
op->src[1]->type == GGML_TYPE_F32 &&
op->type == GGML_TYPE_F32 &&
op->src[0]->ne[1] == 4 &&
op->src[1]->ne[0] == 4 &&
op->src[1]->ne[2] == 1 &&
ggml_is_contiguous_rows(op->src[0]) &&
ggml_is_contiguous_rows(op->src[1]);
case GGML_OP_DSV4_HC_POST:
@@ -1811,17 +1809,15 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
op->src[0]->type == GGML_TYPE_F32 &&
op->src[1]->type == GGML_TYPE_F32 &&
op->src[2]->type == GGML_TYPE_F32 &&
op->src[3] != NULL &&
op->src[3]->type == GGML_TYPE_F32 &&
(op->src[3] == NULL || op->src[3]->type == GGML_TYPE_F32) &&
op->type == GGML_TYPE_F32 &&
op->src[1]->ne[1] == 4 &&
op->src[2]->ne[0] == 4 &&
op->src[3]->ne[0] == 4 &&
op->src[3]->ne[1] == 4 &&
(op->src[3] == NULL || (op->src[3]->ne[0] == 4 && op->src[3]->ne[1] == 4)) &&
ggml_is_contiguous_rows(op->src[0]) &&
ggml_is_contiguous_rows(op->src[1]) &&
ggml_is_contiguous_rows(op->src[2]) &&
ggml_is_contiguous_rows(op->src[3]);
(op->src[3] == NULL || ggml_is_contiguous_rows(op->src[3]));
case GGML_OP_SSM_SCAN:
return has_simdgroup_reduction;
case GGML_OP_SSM_CONV:
ggml/src/ggml-metal/ggml-metal-impl.h
+2
View File
@@ -1283,8 +1283,10 @@ typedef struct {
uint64_t nb_x2;
uint64_t nb_w0;
uint64_t nb_w1;
uint64_t nb_w2;
uint64_t nb_d0;
uint64_t nb_d1;
float scale;
} ggml_metal_kargs_dsv4_hc_pre;
typedef struct {
ggml/src/ggml-metal/ggml-metal-ops.cpp
+12 -7
View File
@@ -1405,7 +1405,7 @@ int ggml_metal_op_dsv4_hc(ggml_metal_op_t ctx, int idx) {
ggml_tensor * op = ctx->node(idx);
ggml_metal_encoder_t enc = ctx->enc;
auto pipeline = ggml_metal_library_get_pipeline_dsv4_hc(ctx->lib, op->op);
auto pipeline = ggml_metal_library_get_pipeline_dsv4_hc(ctx->lib, op);
ggml_metal_encoder_set_pipeline(enc, pipeline);
@@ -1467,8 +1467,10 @@ int ggml_metal_op_dsv4_hc(ggml_metal_op_t ctx, int idx) {
/*.nb_x2 =*/ x->nb[2],
/*.nb_w0 =*/ weights->nb[0],
/*.nb_w1 =*/ weights->nb[1],
/*.nb_w2 =*/ weights->nb[2],
/*.nb_d0 =*/ op->nb[0],
/*.nb_d1 =*/ op->nb[1],
/*.scale =*/ ggml_get_op_params_f32(op, 0),
};
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0);
@@ -1491,7 +1493,6 @@ int ggml_metal_op_dsv4_hc(ggml_metal_op_t ctx, int idx) {
GGML_ASSERT(x->type == GGML_TYPE_F32);
GGML_ASSERT(residual->type == GGML_TYPE_F32);
GGML_ASSERT(post->type == GGML_TYPE_F32);
GGML_ASSERT(comb->type == GGML_TYPE_F32);
GGML_ASSERT(op->type == GGML_TYPE_F32);
GGML_ASSERT(residual->ne[1] == 4);
@@ -1505,9 +1506,9 @@ int ggml_metal_op_dsv4_hc(ggml_metal_op_t ctx, int idx) {
/*.nb_r2 =*/ residual->nb[2],
/*.nb_p0 =*/ post->nb[0],
/*.nb_p1 =*/ post->nb[1],
/*.nb_c0 =*/ comb->nb[0],
/*.nb_c1 =*/ comb->nb[1],
/*.nb_c2 =*/ comb->nb[2],
/*.nb_c0 =*/ comb ? comb->nb[0] : 0,
/*.nb_c1 =*/ comb ? comb->nb[1] : 0,
/*.nb_c2 =*/ comb ? comb->nb[2] : 0,
/*.nb_d0 =*/ op->nb[0],
/*.nb_d1 =*/ op->nb[1],
/*.nb_d2 =*/ op->nb[2],
@@ -1517,8 +1518,12 @@ int ggml_metal_op_dsv4_hc(ggml_metal_op_t ctx, int idx) {
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(x), 1);
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(residual), 2);
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(post), 3);
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(comb), 4);
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(op), 5);
if (comb) {
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(comb), 4);
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(op), 5);
} else {
ggml_metal_encoder_set_buffer(enc, ggml_metal_get_buffer_id(op), 4);
}
const int n_tiles = (args.n_embd + 31)/32;
const int nsg = std::min(4, n_tiles);
ggml/src/ggml-metal/kernels/misc.metal
+67 -1
View File
@@ -531,7 +531,73 @@ kernel void kernel_dsv4_hc_pre_f32(
result = fma(*(device const float *) (xb + ih*args.nb_x1), w[ih], result);
}
*(device float *) (dst + i0*args.nb_d0 + it*args.nb_d1) = result;
*(device float *) (dst + i0*args.nb_d0 + it*args.nb_d1) = args.scale*result;
}
kernel void kernel_dsv4_hc_pre_gated_f32(
constant ggml_metal_kargs_dsv4_hc_pre & args,
device const char * x,
device const char * gate,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
constexpr ushort hc = 4;
const int it = tgpig.y;
const int i0 = ((int) tgpig.x*ntg.y + sgitg)*32 + tiisg;
if (i0 >= args.n_embd) {
return;
}
device const char * xb = x + i0*args.nb_x0 + it*args.nb_x2;
device const char * gb = gate + i0*args.nb_w0 + it*args.nb_w2;
float result = 0.0f;
FOR_UNROLL (ushort ih = 0; ih < hc; ++ih) {
const float g = 1.0f/(1.0f + exp(-*(device const float *) (gb + ih*args.nb_w1)));
result = fma(*(device const float *) (xb + ih*args.nb_x1), g, result);
}
*(device float *) (dst + i0*args.nb_d0 + it*args.nb_d1) = args.scale*result;
}
kernel void kernel_dsv4_hc_post_nocomb_f32(
constant ggml_metal_kargs_dsv4_hc_post & args,
device const char * x,
device const char * residual,
device const char * post,
device char * dst,
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]],
ushort3 ntg[[threads_per_threadgroup]]) {
constexpr ushort hc = 4;
const int it = tgpig.y;
const int i0 = ((int) tgpig.x*ntg.y + sgitg)*32 + tiisg;
float post_lane = 0.0f;
if (tiisg < hc) {
post_lane = *(device const float *) (post + tiisg*args.nb_p0 + it*args.nb_p1);
}
float post_reg[hc];
FOR_UNROLL (ushort idst = 0; idst < hc; ++idst) {
post_reg[idst] = simd_shuffle(post_lane, idst);
}
if (i0 >= args.n_embd) {
return;
}
const float xv = *(device const float *) (x + i0*args.nb_x0 + it*args.nb_x1);
device const char * rb = residual + i0*args.nb_r0 + it*args.nb_r2;
FOR_UNROLL (ushort idst = 0; idst < hc; ++idst) {
const float rv = *(device const float *) (rb + idst*args.nb_r1);
*(device float *) (dst + i0*args.nb_d0 + idst*args.nb_d1 + it*args.nb_d2) = xv*post_reg[idst] + rv;
}
}
kernel void kernel_dsv4_hc_post_f32(