Skip to content

Commit ac7fced

Browse files
nicobossnicoboss
authored
Merge branch 'ggml-org:master' into mradermacher
2 parents 7fa25bc + f9fb33f commit ac7fced

File tree

19 files changed

+848
-49
lines changed

19 files changed

+848
-49
lines changed

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

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3640,9 +3640,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
36403640
case GGML_OP_CONV_2D_DW:
36413641
case GGML_OP_CONV_TRANSPOSE_2D:
36423642
case GGML_OP_POOL_2D:
3643-
case GGML_OP_SUM:
36443643
case GGML_OP_ACC:
36453644
return true;
3645+
case GGML_OP_SUM:
3646+
return ggml_is_contiguous_rows(op->src[0]);
36463647
case GGML_OP_ARGSORT:
36473648
// TODO: Support arbitrary column width
36483649
return op->src[0]->ne[0] <= 1024;

‎ggml/src/ggml-metal/ggml-metal-device.m‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -662,6 +662,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
662662
case GGML_OP_LOG:
663663
return ggml_is_contiguous(op->src[0]) && op->src[0]->type == GGML_TYPE_F32;
664664
case GGML_OP_SUM:
665+
return has_simdgroup_reduction && ggml_is_contiguous(op->src[0]);
665666
case GGML_OP_SUM_ROWS:
666667
case GGML_OP_MEAN:
667668
case GGML_OP_SOFT_MAX:

‎ggml/src/ggml-metal/ggml-metal-ops.cpp‎

Lines changed: 14 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -866,12 +866,25 @@ int ggml_metal_op_sum(ggml_metal_op_t ctx, int idx) {
866866

867867
ggml_metal_pipeline_t pipeline = ggml_metal_library_get_pipeline_sum(lib, op);
868868

869+
int nth = 32; // SIMD width
870+
871+
while (nth < (int) n && nth < ggml_metal_pipeline_max_theads_per_threadgroup(pipeline)) {
872+
nth *= 2;
873+
}
874+
875+
nth = std::min(nth, ggml_metal_pipeline_max_theads_per_threadgroup(pipeline));
876+
nth = std::min(nth, (int) n);
877+
878+
const int nsg = (nth + 31) / 32;
879+
869880
ggml_metal_encoder_set_pipeline(enc, pipeline);
870881
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), 0);
871882
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), 1);
872883
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), 2);
873884

874-
ggml_metal_encoder_dispatch_threadgroups(enc, 1, 1, 1, 1, 1, 1);
885+
ggml_metal_encoder_set_threadgroup_memory_size(enc, nsg * sizeof(float), 0);
886+
887+
ggml_metal_encoder_dispatch_threadgroups(enc, 1, 1, 1, nth, 1, 1);
875888

876889
return 1;
877890
}

‎ggml/src/ggml-metal/ggml-metal.metal‎

Lines changed: 36 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1727,18 +1727,48 @@ kernel void kernel_op_sum_f32(
17271727
constant ggml_metal_kargs_sum & args,
17281728
device const float * src0,
17291729
device float * dst,
1730-
ushort tiitg[[thread_index_in_threadgroup]]) {
1730+
threadgroup float * shmem_f32 [[threadgroup(0)]],
1731+
uint3 tgpig[[threadgroup_position_in_grid]],
1732+
ushort3 tpitg[[thread_position_in_threadgroup]],
1733+
ushort sgitg[[simdgroup_index_in_threadgroup]],
1734+
ushort tiisg[[thread_index_in_simdgroup]],
1735+
ushort3 ntg[[threads_per_threadgroup]]) {
17311736

1732-
if (tiitg != 0) {
1737+
if (args.np == 0) {
17331738
return;
17341739
}
17351740

1736-
float acc = 0.0f;
1737-
for (ulong i = 0; i < args.np; ++i) {
1738-
acc += src0[i];
1741+
const uint nsg = (ntg.x + 31) / 32;
1742+
1743+
float sumf = 0;
1744+
1745+
for (int64_t i0 = tpitg.x; i0 < args.np; i0 += ntg.x) {
1746+
sumf += src0[i0];
17391747
}
17401748

1741-
dst[0] = acc;
1749+
sumf = simd_sum(sumf);
1750+
1751+
if (tiisg == 0) {
1752+
shmem_f32[sgitg] = sumf;
1753+
}
1754+
1755+
threadgroup_barrier(mem_flags::mem_threadgroup);
1756+
1757+
float total = 0;
1758+
1759+
if (sgitg == 0) {
1760+
float v = 0;
1761+
1762+
if (tpitg.x < nsg) {
1763+
v = shmem_f32[tpitg.x];
1764+
}
1765+
1766+
total = simd_sum(v);
1767+
1768+
if (tpitg.x == 0) {
1769+
dst[0] = total;
1770+
}
1771+
}
17421772
}
17431773

17441774
template <bool norm>

‎tests/test-backend-ops.cpp‎

Lines changed: 18 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -4588,20 +4588,31 @@ struct test_topk_moe: public test_case {
45884588
struct test_sum : public test_case {
45894589
const ggml_type type;
45904590
const std::array<int64_t, 4> ne;
4591+
const std::array<int64_t, 4> permute;
4592+
bool _use_permute;
45914593

45924594
std::string vars() override {
4593-
return VARS_TO_STR2(type, ne);
4595+
std::string v = VARS_TO_STR2(type, ne);
4596+
if (_use_permute) v += "," + VAR_TO_STR(permute);
4597+
return v;
45944598
}
45954599

45964600
test_sum(ggml_type type = GGML_TYPE_F32,
4597-
std::array<int64_t, 4> ne = {10, 5, 4, 3})
4598-
: type(type), ne(ne) {}
4601+
std::array<int64_t, 4> ne = {10, 5, 4, 3},
4602+
std::array<int64_t, 4> permute = {0, 0, 0, 0})
4603+
: type(type), ne(ne), permute(permute),
4604+
_use_permute(permute[0] + permute[1] + permute[2] + permute[3] > 0) {}
45994605

46004606
ggml_tensor * build_graph(ggml_context * ctx) override {
46014607
ggml_tensor * a = ggml_new_tensor(ctx, type, 4, ne.data());
46024608
ggml_set_param(a);
46034609
ggml_set_name(a, "a");
46044610

4611+
if (_use_permute) {
4612+
a = ggml_permute(ctx, a, permute[0], permute[1], permute[2], permute[3]);
4613+
ggml_set_name(a, "a_permuted");
4614+
}
4615+
46054616
ggml_tensor * out = ggml_sum(ctx, a);
46064617
ggml_set_name(out, "out");
46074618

@@ -6724,6 +6735,9 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
67246735

67256736
test_cases.emplace_back(new test_sum());
67266737
test_cases.emplace_back(new test_sum_rows());
6738+
test_cases.emplace_back(new test_sum(GGML_TYPE_F32, {11, 5, 6, 3}, {0, 2, 1, 3})); // row-contiguous but non-contiguous
6739+
test_cases.emplace_back(new test_sum(GGML_TYPE_F32, {11, 5, 6, 3}, {0, 3, 2, 1}));
6740+
test_cases.emplace_back(new test_sum(GGML_TYPE_F32, {11, 5, 6, 3}, {0, 1, 3, 2}));
67276741
test_cases.emplace_back(new test_sum_rows(GGML_TYPE_F32, { 11, 5, 6, 3 }, true, false));
67286742
test_cases.emplace_back(new test_sum_rows(GGML_TYPE_F32, { 11, 5, 6, 3 }, false, true));
67296743
test_cases.emplace_back(new test_sum_rows(GGML_TYPE_F32, { 11, 5, 6, 3 }, true, true));
@@ -6734,6 +6748,7 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
67346748
test_cases.emplace_back(new test_sum(GGML_TYPE_F32, { 33, 1024, 1, 1 }));
67356749
test_cases.emplace_back(new test_sum_rows(GGML_TYPE_F32, { 33, 1024, 1, 1 }));
67366750
test_cases.emplace_back(new test_sum(GGML_TYPE_F32, { 33, 256, 1, 1 }));
6751+
test_cases.emplace_back(new test_sum(GGML_TYPE_F32, { 33, 256, 1, 1 }, { 1, 0, 2, 3 })); // sum dst not-contiguous
67376752
test_cases.emplace_back(new test_sum_rows(GGML_TYPE_F32, { 33, 256, 1, 1 }));
67386753
test_cases.emplace_back(new test_mean(GGML_TYPE_F32, { 33, 256, 1, 1 }));
67396754
test_cases.emplace_back(new test_mean(GGML_TYPE_F32, { 32769, 1, 1, 1 }));

‎tools/server/public/index.html.gz‎

2.75 KB
Binary file not shown.

‎tools/server/webui/src/lib/components/app/chat/ChatSettings/ChatSettingsDialog.svelte‎

Lines changed: 4 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -14,8 +14,7 @@
1414
import { ChatSettingsFooter, ChatSettingsFields } from '$lib/components/app';
1515
import * as Dialog from '$lib/components/ui/dialog';
1616
import { ScrollArea } from '$lib/components/ui/scroll-area';
17-
import { SETTING_CONFIG_DEFAULT } from '$lib/constants/settings-config';
18-
import { config, updateMultipleConfig, resetConfig } from '$lib/stores/settings.svelte';
17+
import { config, updateMultipleConfig } from '$lib/stores/settings.svelte';
1918
import { setMode } from 'mode-watcher';
2019
import type { Component } from 'svelte';
2120
@@ -267,16 +266,13 @@
267266
}
268267
269268
function handleReset() {
270-
resetConfig();
269+
localConfig = { ...config() };
271270
272-
localConfig = { ...SETTING_CONFIG_DEFAULT };
273-
274-
setMode(SETTING_CONFIG_DEFAULT.theme as 'light' | 'dark' | 'system');
275-
originalTheme = SETTING_CONFIG_DEFAULT.theme as string;
271+
setMode(localConfig.theme as 'light' | 'dark' | 'system');
272+
originalTheme = localConfig.theme as string;
276273
}
277274
278275
function handleSave() {
279-
// Validate custom JSON if provided
280276
if (localConfig.custom && typeof localConfig.custom === 'string' && localConfig.custom.trim()) {
281277
try {
282278
JSON.parse(localConfig.custom);

‎tools/server/webui/src/lib/components/app/chat/ChatSettings/ChatSettingsFields.svelte‎

Lines changed: 115 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,15 @@
11
<script lang="ts">
2+
import { RotateCcw } from '@lucide/svelte';
23
import { Checkbox } from '$lib/components/ui/checkbox';
34
import { Input } from '$lib/components/ui/input';
45
import Label from '$lib/components/ui/label/label.svelte';
56
import * as Select from '$lib/components/ui/select';
67
import { Textarea } from '$lib/components/ui/textarea';
78
import { SETTING_CONFIG_DEFAULT, SETTING_CONFIG_INFO } from '$lib/constants/settings-config';
89
import { supportsVision } from '$lib/stores/server.svelte';
10+
import { getParameterInfo, resetParameterToServerDefault } from '$lib/stores/settings.svelte';
11+
import { ParameterSyncService } from '$lib/services/parameter-sync';
12+
import ParameterSourceIndicator from './ParameterSourceIndicator.svelte';
913
import type { Component } from 'svelte';
1014
1115
interface Props {
@@ -16,22 +20,77 @@
1620
}
1721
1822
let { fields, localConfig, onConfigChange, onThemeChange }: Props = $props();
23+
24+
// Helper function to get parameter source info for syncable parameters
25+
function getParameterSourceInfo(key: string) {
26+
if (!ParameterSyncService.canSyncParameter(key)) {
27+
return null;
28+
}
29+
30+
return getParameterInfo(key);
31+
}
1932
</script>
2033

2134
{#each fields as field (field.key)}
2235
<div class="space-y-2">
2336
{#if field.type === 'input'}
24-
<Label for={field.key} class="block text-sm font-medium">
25-
{field.label}
26-
</Label>
37+
{@const paramInfo = getParameterSourceInfo(field.key)}
38+
{@const currentValue = String(localConfig[field.key] ?? '')}
39+
{@const propsDefault = paramInfo?.serverDefault}
40+
{@const isCustomRealTime = (() => {
41+
if (!paramInfo || propsDefault === undefined) return false;
2742

28-
<Input
29-
id={field.key}
30-
value={String(localConfig[field.key] ?? '')}
31-
onchange={(e) => onConfigChange(field.key, e.currentTarget.value)}
32-
placeholder={`Default: ${SETTING_CONFIG_DEFAULT[field.key] ?? 'none'}`}
33-
class="w-full md:max-w-md"
34-
/>
43+
// Apply same rounding logic for real-time comparison
44+
const inputValue = currentValue;
45+
const numericInput = parseFloat(inputValue);
46+
const normalizedInput = !isNaN(numericInput)
47+
? Math.round(numericInput * 1000000) / 1000000
48+
: inputValue;
49+
const normalizedDefault =
50+
typeof propsDefault === 'number'
51+
? Math.round(propsDefault * 1000000) / 1000000
52+
: propsDefault;
53+
54+
return normalizedInput !== normalizedDefault;
55+
})()}
56+
57+
<div class="flex items-center gap-2">
58+
<Label for={field.key} class="text-sm font-medium">
59+
{field.label}
60+
</Label>
61+
{#if isCustomRealTime}
62+
<ParameterSourceIndicator />
63+
{/if}
64+
</div>
65+
66+
<div class="relative w-full md:max-w-md">
67+
<Input
68+
id={field.key}
69+
value={currentValue}
70+
oninput={(e) => {
71+
// Update local config immediately for real-time badge feedback
72+
onConfigChange(field.key, e.currentTarget.value);
73+
}}
74+
placeholder={`Default: ${SETTING_CONFIG_DEFAULT[field.key] ?? 'none'}`}
75+
class="w-full {isCustomRealTime ? 'pr-8' : ''}"
76+
/>
77+
{#if isCustomRealTime}
78+
<button
79+
type="button"
80+
onclick={() => {
81+
resetParameterToServerDefault(field.key);
82+
// Trigger UI update by calling onConfigChange with the default value
83+
const defaultValue = propsDefault ?? SETTING_CONFIG_DEFAULT[field.key];
84+
onConfigChange(field.key, String(defaultValue));
85+
}}
86+
class="absolute top-1/2 right-2 inline-flex h-5 w-5 -translate-y-1/2 items-center justify-center rounded transition-colors hover:bg-muted"
87+
aria-label="Reset to default"
88+
title="Reset to default"
89+
>
90+
<RotateCcw class="h-3 w-3" />
91+
</button>
92+
{/if}
93+
</div>
3594
{#if field.help || SETTING_CONFIG_INFO[field.key]}
3695
<p class="mt-1 text-xs text-muted-foreground">
3796
{field.help || SETTING_CONFIG_INFO[field.key]}
@@ -59,14 +118,28 @@
59118
(opt: { value: string; label: string; icon?: Component }) =>
60119
opt.value === localConfig[field.key]
61120
)}
121+
{@const paramInfo = getParameterSourceInfo(field.key)}
122+
{@const currentValue = localConfig[field.key]}
123+
{@const propsDefault = paramInfo?.serverDefault}
124+
{@const isCustomRealTime = (() => {
125+
if (!paramInfo || propsDefault === undefined) return false;
62126

63-
<Label for={field.key} class="block text-sm font-medium">
64-
{field.label}
65-
</Label>
127+
// For select fields, do direct comparison (no rounding needed)
128+
return currentValue !== propsDefault;
129+
})()}
130+
131+
<div class="flex items-center gap-2">
132+
<Label for={field.key} class="text-sm font-medium">
133+
{field.label}
134+
</Label>
135+
{#if isCustomRealTime}
136+
<ParameterSourceIndicator />
137+
{/if}
138+
</div>
66139

67140
<Select.Root
68141
type="single"
69-
value={localConfig[field.key]}
142+
value={currentValue}
70143
onValueChange={(value) => {
71144
if (field.key === 'theme' && value && onThemeChange) {
72145
onThemeChange(value);
@@ -75,16 +148,34 @@
75148
}
76149
}}
77150
>
78-
<Select.Trigger class="w-full md:w-auto md:max-w-md">
79-
<div class="flex items-center gap-2">
80-
{#if selectedOption?.icon}
81-
{@const IconComponent = selectedOption.icon}
82-
<IconComponent class="h-4 w-4" />
83-
{/if}
84-
85-
{selectedOption?.label || `Select ${field.label.toLowerCase()}`}
86-
</div>
87-
</Select.Trigger>
151+
<div class="relative w-full md:w-auto md:max-w-md">
152+
<Select.Trigger class="w-full">
153+
<div class="flex items-center gap-2">
154+
{#if selectedOption?.icon}
155+
{@const IconComponent = selectedOption.icon}
156+
<IconComponent class="h-4 w-4" />
157+
{/if}
158+
159+
{selectedOption?.label || `Select ${field.label.toLowerCase()}`}
160+
</div>
161+
</Select.Trigger>
162+
{#if isCustomRealTime}
163+
<button
164+
type="button"
165+
onclick={() => {
166+
resetParameterToServerDefault(field.key);
167+
// Trigger UI update by calling onConfigChange with the default value
168+
const defaultValue = propsDefault ?? SETTING_CONFIG_DEFAULT[field.key];
169+
onConfigChange(field.key, String(defaultValue));
170+
}}
171+
class="absolute top-1/2 right-8 inline-flex h-5 w-5 -translate-y-1/2 items-center justify-center rounded transition-colors hover:bg-muted"
172+
aria-label="Reset to default"
173+
title="Reset to default"
174+
>
175+
<RotateCcw class="h-3 w-3" />
176+
</button>
177+
{/if}
178+
</div>
88179
<Select.Content>
89180
{#if field.options}
90181
{#each field.options as option (option.value)}

0 commit comments

Comments
 (0)