Skip to content

Commit eacac4c

Browse files
nicobossnicoboss
authored
Merge branch 'ggml-org:master' into mradermacher
2 parents abf158c + 8101786 commit eacac4c

File tree

7 files changed

+509
-24
lines changed

7 files changed

+509
-24
lines changed

‎common/chat.cpp‎

Lines changed: 152 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -622,6 +622,7 @@ const char * common_chat_format_name(common_chat_format format) {
622622
case COMMON_CHAT_FORMAT_COMMAND_R7B: return "Command R7B";
623623
case COMMON_CHAT_FORMAT_GRANITE: return "Granite";
624624
case COMMON_CHAT_FORMAT_GPT_OSS: return "GPT-OSS";
625+
case COMMON_CHAT_FORMAT_SEED_OSS: return "Seed-OSS";
625626
default:
626627
throw std::runtime_error("Unknown chat format");
627628
}
@@ -2059,6 +2060,94 @@ static void common_chat_parse_granite(common_chat_msg_parser & builder) {
20592060
}
20602061
}
20612062

2063+
static void common_chat_parse_seed_oss(common_chat_msg_parser & builder) {
2064+
// Parse thinking tags first - this handles the main reasoning content
2065+
builder.try_parse_reasoning("<seed:think>", "</seed:think>");
2066+
2067+
if (!builder.syntax().parse_tool_calls) {
2068+
builder.add_content(builder.consume_rest());
2069+
return;
2070+
}
2071+
2072+
// Parse tool calls - Seed-OSS uses <seed:tool_call> format
2073+
static const common_regex tool_call_begin_regex("<seed:tool_call>");
2074+
static const common_regex tool_call_end_regex("</seed:tool_call>");
2075+
static const common_regex function_regex("<function=([^>]+)>");
2076+
static const common_regex param_regex("<parameter=([^>]+)>");
2077+
2078+
while (auto tool_res = builder.try_find_regex(tool_call_begin_regex)) {
2079+
builder.consume_spaces(); // Consume whitespace after <seed:tool_call>
2080+
2081+
// Look for function call inside tool call, ignore any content before it
2082+
if (auto func_res = builder.try_find_regex(function_regex, std::string::npos, false)) {
2083+
auto function_name = builder.str(func_res->groups[1]);
2084+
2085+
// Parse Seed-OSS parameters <parameter=name>value</parameter>
2086+
json args = json::object();
2087+
// Parse all parameters
2088+
while (auto param_res = builder.try_find_regex(param_regex, std::string::npos, false)) {
2089+
// again, ignore noise around parameters
2090+
auto param_name = builder.str(param_res->groups[1]);
2091+
builder.move_to(param_res->groups[0].end);
2092+
builder.consume_spaces(); // Consume whitespace after parameter
2093+
auto savedPos = builder.pos();
2094+
if (auto param_parse = builder.try_find_literal("</parameter>")) {
2095+
auto param = param_parse->prelude;
2096+
builder.move_to(savedPos);
2097+
try {
2098+
if (auto param_res = builder.try_consume_json()) {
2099+
args[param_name] = param_res->json;
2100+
} else {
2101+
args[param_name] = param;
2102+
}
2103+
} catch (json::exception &) {
2104+
args[param_name] = param;
2105+
}
2106+
} else {
2107+
throw common_chat_msg_partial_exception("Incomplete tool parameter");
2108+
}
2109+
}
2110+
// Look for closing function tag
2111+
auto end_func = builder.try_find_literal("</function>");
2112+
if (end_func) {
2113+
builder.move_to(end_func->groups[0].end);
2114+
builder.consume_spaces(); // Consume whitespace after </function>
2115+
2116+
// Add the tool call with parsed arguments, but only if we REALLY got the literal
2117+
auto eaten_fragment = builder.input().substr(end_func->groups[0].begin, end_func->groups[0].end);
2118+
auto funlen = std::string("</function>").length();
2119+
if (eaten_fragment.length() >= funlen && eaten_fragment.substr(0, funlen) == std::string("</function>")) {
2120+
if (!builder.add_tool_call(function_name, "", args.dump())) {
2121+
throw common_chat_msg_partial_exception("Incomplete tool call");
2122+
}
2123+
} else {
2124+
throw common_chat_msg_partial_exception("Incomplete tool call");
2125+
}
2126+
} else {
2127+
throw common_chat_msg_partial_exception("Incomplete tool call");
2128+
}
2129+
// Look for closing tool call tag
2130+
if (auto end_tool = builder.try_find_regex(tool_call_end_regex, std::string::npos, false)) {
2131+
builder.move_to(end_tool->groups[0].end);
2132+
builder.consume_spaces(); // Consume trailing whitespace after tool call
2133+
} else {
2134+
throw common_chat_msg_partial_exception("Incomplete tool call");
2135+
}
2136+
} else {
2137+
// No function found - don't consume content here, let it be handled at the end
2138+
break;
2139+
}
2140+
}
2141+
2142+
// Consume any remaining whitespace after all tool call processing
2143+
builder.consume_spaces();
2144+
auto remaining = builder.consume_rest();
2145+
// If there's any non-whitespace content remaining, add it as content
2146+
if (!string_strip(remaining).empty()) {
2147+
builder.add_content(remaining);
2148+
}
2149+
}
2150+
20622151
static common_chat_params common_chat_params_init_without_tools(const common_chat_template & tmpl, const struct templates_params & inputs) {
20632152
common_chat_params data;
20642153
data.prompt = apply(tmpl, inputs);
@@ -2075,8 +2164,62 @@ static common_chat_params common_chat_params_init_without_tools(const common_cha
20752164
return data;
20762165
}
20772166

2167+
static common_chat_params common_chat_params_init_seed_oss(
2168+
const common_chat_template & tmpl,
2169+
templates_params & params,
2170+
const common_chat_templates_inputs & inputs)
2171+
{
2172+
common_chat_params data;
2173+
data.prompt = apply(tmpl, params);
2174+
data.format = COMMON_CHAT_FORMAT_SEED_OSS;
2175+
if (string_ends_with(data.prompt, "<seed:think>")) {
2176+
if (!inputs.enable_thinking) {
2177+
data.prompt += "</seed:think>";
2178+
} else {
2179+
data.thinking_forced_open = true;
2180+
}
2181+
}
2182+
2183+
if (params.tools.is_array() && !params.tools.empty()) {
2184+
data.grammar_lazy = inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_REQUIRED;
2185+
data.grammar = build_grammar([&](const common_grammar_builder & builder) {
2186+
std::vector<std::string> tool_rules;
2187+
foreach_function(params.tools, [&](const json & tool) {
2188+
const auto & function = tool.at("function");
2189+
std::string name = function.at("name");
2190+
auto parameters = function.at("parameters");
2191+
builder.resolve_refs(parameters);
2192+
2193+
// Create rule for Seed-OSS function call format
2194+
std::string param_rules;
2195+
if (parameters.contains("properties")) {
2196+
for (const auto & [key, value] : parameters.at("properties").items()) {
2197+
param_rules += "\"<parameter=" + key + ">\"" + builder.add_schema(name + "-arg-" + key, value) +
2198+
"\"</parameter>\"";
2199+
}
2200+
}
2201+
2202+
tool_rules.push_back(builder.add_rule(name + "-call",
2203+
"\"<seed:tool_call>\" space \"<function=" + name + ">\" space " +
2204+
param_rules +
2205+
" \"</function>\" space \"</seed:tool_call>\""));
2206+
});
2207+
2208+
data.grammar_triggers.push_back({ COMMON_GRAMMAR_TRIGGER_TYPE_WORD, "<seed:tool_call>" });
2209+
2210+
data.preserved_tokens = {
2211+
"<seed:think>", "</seed:think>", "<seed:tool_call>", "</seed:tool_call>",
2212+
"<function=", "</function>", "<parameter=", "</parameter>",
2213+
};
2214+
2215+
builder.add_rule("root", string_join(tool_rules, " | "));
2216+
});
2217+
}
2218+
return data;
2219+
}
2220+
20782221
static common_chat_params common_chat_templates_apply_jinja(
2079-
const struct common_chat_templates * tmpls,
2222+
const struct common_chat_templates * tmpls,
20802223
const struct common_chat_templates_inputs & inputs)
20812224
{
20822225
templates_params params;
@@ -2145,6 +2288,11 @@ static common_chat_params common_chat_templates_apply_jinja(
21452288
return common_chat_params_init_gpt_oss(tmpl, params);
21462289
}
21472290

2291+
// Seed-OSS
2292+
if (src.find("<seed:think>") != std::string::npos) {
2293+
return common_chat_params_init_seed_oss(tmpl, params, inputs);
2294+
}
2295+
21482296
// Use generic handler when mixing tools + JSON schema.
21492297
// TODO: support that mix in handlers below.
21502298
if ((params.tools.is_array() && params.json_schema.is_object())) {
@@ -2303,6 +2451,9 @@ static void common_chat_parse(common_chat_msg_parser & builder) {
23032451
case COMMON_CHAT_FORMAT_GPT_OSS:
23042452
common_chat_parse_gpt_oss(builder);
23052453
break;
2454+
case COMMON_CHAT_FORMAT_SEED_OSS:
2455+
common_chat_parse_seed_oss(builder);
2456+
break;
23062457
default:
23072458
throw std::runtime_error(std::string("Unsupported format: ") + common_chat_format_name(builder.syntax().format));
23082459
}

‎common/chat.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -111,6 +111,7 @@ enum common_chat_format {
111111
COMMON_CHAT_FORMAT_COMMAND_R7B,
112112
COMMON_CHAT_FORMAT_GRANITE,
113113
COMMON_CHAT_FORMAT_GPT_OSS,
114+
COMMON_CHAT_FORMAT_SEED_OSS,
114115

115116
COMMON_CHAT_FORMAT_COUNT, // Not a format, just the # formats
116117
};

‎ggml/src/ggml-cuda/binbcast.cu‎

Lines changed: 47 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -57,7 +57,11 @@ static __global__ void k_bin_bcast(const src0_t * src0, const src1_t * src1, dst
5757
const int i10 = i0 % ne10;
5858

5959
float result = src0_row ? (float) src0_row[i0] : 0.0f;
60-
result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10])));
60+
if constexpr (sizeof...(src1_ptrs) > 0) {
61+
result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10])));
62+
} else {
63+
result = bin_op(result, (float)src1[i_src1 + i10]);
64+
}
6165

6266
dst_row[i0] = (dst_t) result;
6367
}
@@ -96,7 +100,11 @@ static __global__ void k_bin_bcast_unravel(const src0_t * src0, const src1_t *
96100
const int i10 = i0 % ne10;
97101

98102
float result = src0_row ? (float) src0_row[i0] : 0.0f;
99-
result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10])));
103+
if constexpr (sizeof...(src1_ptrs) > 0) {
104+
result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10])));
105+
} else {
106+
result = bin_op(result, (float)src1[i_src1 + i10]);
107+
}
100108

101109
dst_row[i0] = (dst_t) result;
102110
}
@@ -231,23 +239,43 @@ static void launch_bin_bcast_pack(const ggml_tensor * src0, const ggml_tensor *
231239

232240
if (block_nums.z > 65535) {
233241
int block_num = (ne0 * ne1 * ne2 * ne3 + block_size - 1) / block_size;
234-
k_bin_bcast_unravel<bin_op, src0_t, src1_t, dst_t>
235-
<<<block_num, block_size, 0, stream>>>(src0_dd, src1_dd, dst_dd,
236-
ne0, ne1, ne2, ne3,
237-
ne10, ne11, ne12, ne13,
238-
/* s0, */ s1, s2, s3,
239-
/* s00,*/ s01, s02, s03,
240-
/* s10,*/ s11, s12,s13,
241-
(const src1_t *) dst->src[I + 1]->data...);
242+
if constexpr (sizeof...(I) > 0) {
243+
k_bin_bcast_unravel<bin_op, src0_t, src1_t, dst_t>
244+
<<<block_num, block_size, 0, stream>>>(src0_dd, src1_dd, dst_dd,
245+
ne0, ne1, ne2, ne3,
246+
ne10, ne11, ne12, ne13,
247+
/* s0, */ s1, s2, s3,
248+
/* s00,*/ s01, s02, s03,
249+
/* s10,*/ s11, s12,s13,
250+
(const src1_t *) dst->src[I + 1]->data...);
251+
} else {
252+
k_bin_bcast_unravel<bin_op, src0_t, src1_t, dst_t>
253+
<<<block_num, block_size, 0, stream>>>(src0_dd, src1_dd, dst_dd,
254+
ne0, ne1, ne2, ne3,
255+
ne10, ne11, ne12, ne13,
256+
/* s0, */ s1, s2, s3,
257+
/* s00,*/ s01, s02, s03,
258+
/* s10,*/ s11, s12,s13);
259+
}
242260
} else {
243-
k_bin_bcast<bin_op, src0_t, src1_t, dst_t>
244-
<<<block_nums, block_dims, 0, stream>>>(src0_dd, src1_dd, dst_dd,
245-
ne0, ne1, ne2, ne3,
246-
ne10, ne11, ne12, ne13,
247-
/* s0, */ s1, s2, s3,
248-
/* s00,*/ s01, s02, s03,
249-
/* s10,*/ s11, s12,s13,
250-
(const src1_t *) dst->src[I + 1]->data...);
261+
if constexpr (sizeof...(I) > 0) {
262+
k_bin_bcast<bin_op, src0_t, src1_t, dst_t>
263+
<<<block_nums, block_dims, 0, stream>>>(src0_dd, src1_dd, dst_dd,
264+
ne0, ne1, ne2, ne3,
265+
ne10, ne11, ne12, ne13,
266+
/* s0, */ s1, s2, s3,
267+
/* s00,*/ s01, s02, s03,
268+
/* s10,*/ s11, s12,s13,
269+
(const src1_t *) dst->src[I + 1]->data...);
270+
} else {
271+
k_bin_bcast<bin_op, src0_t, src1_t, dst_t>
272+
<<<block_nums, block_dims, 0, stream>>>(src0_dd, src1_dd, dst_dd,
273+
ne0, ne1, ne2, ne3,
274+
ne10, ne11, ne12, ne13,
275+
/* s0, */ s1, s2, s3,
276+
/* s00,*/ s01, s02, s03,
277+
/* s10,*/ s11, s12,s13);
278+
}
251279
}
252280
}
253281
}
@@ -327,7 +355,7 @@ static void ggml_cuda_op_bin_bcast(
327355
}
328356

329357
void ggml_cuda_op_repeat(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
330-
ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_repeat>>(dst, dst->src[0], dst, nullptr, dst->src[0]->data, dst->data, ctx.stream());
358+
ggml_cuda_op_bin_bcast<bin_bcast_cuda<op_repeat, 0>>(dst, dst->src[0], dst, nullptr, dst->src[0]->data, dst->data, ctx.stream());
331359
}
332360

333361
void ggml_cuda_op_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {

‎ggml/src/ggml-cuda/ggml-cuda.cu‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2842,7 +2842,7 @@ static bool ggml_cuda_can_fuse(const struct ggml_cgraph * cgraph, int node_idx,
28422842
const ggml_tensor *add = nullptr;
28432843

28442844
if (ops.size() == 3 && ops.begin()[2] == GGML_OP_ADD) {
2845-
add = cgraph->nodes[node_idx+1];
2845+
add = cgraph->nodes[node_idx+2];
28462846
}
28472847

28482848
GGML_ASSERT(rms_norm->src[0]->type == GGML_TYPE_F32);

‎ggml/src/ggml-cuda/norm.cu‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -127,6 +127,7 @@ static __global__ void rms_norm_f32(const float * x, float * dst,
127127
const int add_nrows = 0,
128128
const int add_nchannels = 0,
129129
const int add_nsamples = 0) {
130+
130131
const int nrows = gridDim.x;
131132
const int nchannels = gridDim.y;
132133

@@ -135,6 +136,8 @@ static __global__ void rms_norm_f32(const float * x, float * dst,
135136
const int sample = blockIdx.z;
136137
const int tid = threadIdx.x;
137138

139+
static_assert(!do_add || do_multiply, "fusing add is not supported without multiplying");
140+
138141
x += sample*stride_sample + channel*stride_channel + row*stride_row;
139142
dst += ((sample*nchannels + channel)*nrows + row)*ncols;
140143

@@ -185,9 +188,6 @@ static __global__ void rms_norm_f32(const float * x, float * dst,
185188
} else if constexpr (do_multiply) {
186189
const int mul_col = col % mul_ncols;
187190
dst[col] = scale * x[col] * mul[mul_col];
188-
} else if constexpr (do_add) {
189-
const int add_col = col % add_ncols;
190-
dst[col] += add[add_col];
191191
} else {
192192
dst[col] = scale * x[col];
193193
}

0 commit comments

Comments
 (0)