mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-08-18 02:32:40 +02:00
Compare commits
10 Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Georgi Gerganov
|
04a134c70b | ||
|
Georgi Gerganov
|
4197155add | ||
|
Georgi Gerganov
|
cea66f4c5a | ||
 Fathi Boudra
|
4695f001fe | ||
 Titaniumtown
|
f275595dd1 | ||
 Neo Zhang
|
37a215c9e9 | ||
 Daniel Bevenius
|
4df29be4f4 | ||
 fairydreaming
|
3cb7ffb1a1 | ||
 Xuan-Son Nguyen
|
b94041a98e | ||
 Oğuzhan Akkaya
|
10bf611e53 |
@@ -22,6 +22,8 @@ jobs:
|
||||
steps:
|
||||
- name: Checkout
|
||||
uses: actions/checkout@v6
|
||||
with:
|
||||
ssh-key: ${{ secrets.DEPLOY_KEY_RELEASE }}
|
||||
|
||||
- name: Run release checks
|
||||
id: checks
|
||||
@@ -42,5 +44,10 @@ jobs:
|
||||
- name: Dry run summary
|
||||
if: ${{ github.event.inputs.dry_run == 'true' }}
|
||||
run: |
|
||||
echo "Dry run complete - all checks passed."
|
||||
echo "Would have created tag: ${{ steps.checks.outputs.version }}"
|
||||
if [[ "${{ steps.checks.outputs.checks_passed }}" == "true" ]]; then
|
||||
echo "Dry run complete - all checks passed."
|
||||
echo "Would have created tag: ${{ steps.checks.outputs.version }}"
|
||||
else
|
||||
echo "::error::Dry run found release check failures. A release tag would not be created."
|
||||
exit 1
|
||||
fi
|
||||
|
||||
@@ -1598,6 +1598,7 @@ jobs:
|
||||
uses: actions/checkout@v6
|
||||
with:
|
||||
fetch-depth: 0
|
||||
ssh-key: ${{ secrets.DEPLOY_KEY_RELEASE }}
|
||||
|
||||
- name: Determine tag name
|
||||
id: tag
|
||||
|
||||
+72
-27
@@ -470,36 +470,80 @@ std::vector<common_chat_msg> common_chat_msgs_parse_oaicompat(const json & messa
|
||||
return msgs;
|
||||
}
|
||||
|
||||
struct messages_inp_normalizer {
|
||||
const jinja::caps & caps;
|
||||
|
||||
messages_inp_normalizer(const jinja::caps & c) : caps(c) {}
|
||||
|
||||
// handle supports_string_content / supports_typed_content
|
||||
// if string=true and array=false, convert array to string
|
||||
// if string=false and array=true, convert string to array
|
||||
// if both are true, do nothing
|
||||
json normalize(const json & messages) {
|
||||
bool only_string = caps.supports_string_content && !caps.supports_typed_content;
|
||||
bool only_typed = !caps.supports_string_content && caps.supports_typed_content;
|
||||
if ((!only_string && !only_typed) || !messages.is_array()) {
|
||||
return messages;
|
||||
}
|
||||
json normalized = json::array();
|
||||
for (const auto & msg : messages) {
|
||||
json copy = msg;
|
||||
auto it = copy.find("content");
|
||||
if (it != copy.end()) {
|
||||
if (only_typed && it->is_string()) {
|
||||
*it = json::array({
|
||||
json{
|
||||
{"type", "text"},
|
||||
{"text", it->get<std::string>()},
|
||||
}
|
||||
});
|
||||
} else if (only_string && it->is_array()) {
|
||||
*it = concat_content_parts(*it);
|
||||
}
|
||||
}
|
||||
normalized.push_back(std::move(copy));
|
||||
}
|
||||
return normalized;
|
||||
}
|
||||
|
||||
// join parts with newline, do not add newline before or after media markers
|
||||
static std::string concat_content_parts(const json & parts) {
|
||||
std::string text;
|
||||
bool last_was_media_marker = false;
|
||||
for (const auto & part : parts) {
|
||||
std::string type = part.value("type", "");
|
||||
bool add_new_line = true;
|
||||
if (type == "text") {
|
||||
add_new_line = !last_was_media_marker && !text.empty();
|
||||
last_was_media_marker = false;
|
||||
} else if (type == "media_marker") {
|
||||
add_new_line = false;
|
||||
last_was_media_marker = true;
|
||||
} else {
|
||||
LOG_WRN("Ignoring content part type: %s\n", type.c_str());
|
||||
continue;
|
||||
}
|
||||
|
||||
if (add_new_line) {
|
||||
text += '\n';
|
||||
}
|
||||
|
||||
text += part.value("text", "");
|
||||
}
|
||||
return text;
|
||||
}
|
||||
};
|
||||
|
||||
static json render_message_to_json(const std::vector<common_chat_msg> & msgs, const jinja::caps & c) {
|
||||
if (!c.supports_string_content && !c.supports_typed_content) {
|
||||
LOG_WRN("%s: Neither string content nor typed content is supported by the template. This is unexpected and may lead to issues.\n", __func__);
|
||||
}
|
||||
|
||||
bool only_string_accepted = c.supports_string_content && !c.supports_typed_content;
|
||||
bool only_typed_accepted = !c.supports_string_content && c.supports_typed_content;
|
||||
|
||||
json messages = json::array();
|
||||
for (const auto & msg : msgs) {
|
||||
if (only_string_accepted) {
|
||||
json jmsg = msg.to_json_oaicompat(/* concat_typed_text= */ true);
|
||||
messages.push_back(jmsg);
|
||||
} else if (only_typed_accepted) {
|
||||
json jmsg = msg.to_json_oaicompat(/* concat_typed_text= */ false);
|
||||
if (jmsg.at("content").is_string()) {
|
||||
jmsg["content"] = json::array({
|
||||
json{
|
||||
{"type", "text"},
|
||||
{"text", jmsg.at("content").get<std::string>()},
|
||||
}
|
||||
});
|
||||
}
|
||||
messages.push_back(jmsg);
|
||||
} else {
|
||||
json jmsg = msg.to_json_oaicompat(/* concat_typed_text= */ false);
|
||||
messages.push_back(jmsg);
|
||||
}
|
||||
messages.push_back(msg.to_json_oaicompat(/* concat_typed_text= */ false));
|
||||
}
|
||||
return messages;
|
||||
return messages_inp_normalizer(c).normalize(messages);
|
||||
}
|
||||
|
||||
// DEPRECATED: only used in tests
|
||||
@@ -892,8 +936,11 @@ static std::string common_chat_template_direct_apply_impl(
|
||||
const std::optional<json> & additional_context = std::nullopt) {
|
||||
jinja::context ctx(tmpl.source());
|
||||
|
||||
// messages_override is already built for this template, do not touch its content parts
|
||||
nlohmann::ordered_json inp = nlohmann::ordered_json{
|
||||
{"messages", messages_override.has_value() ? *messages_override : inputs.messages},
|
||||
{"messages", messages_override.has_value()
|
||||
? *messages_override
|
||||
: messages_inp_normalizer(tmpl.original_caps()).normalize(inputs.messages)},
|
||||
{"bos_token", tmpl.bos_token()},
|
||||
{"eos_token", tmpl.eos_token()},
|
||||
{"enable_thinking", inputs.enable_thinking},
|
||||
@@ -957,14 +1004,12 @@ static std::string common_chat_template_generation_prompt_impl(
|
||||
const std::optional<json> & tools_override = std::nullopt,
|
||||
const std::optional<json> & additional_context = std::nullopt) {
|
||||
|
||||
auto adjusted_messages = messages_override ? *messages_override : inputs.messages;
|
||||
|
||||
autoparser::generation_params params = inputs;
|
||||
params.add_generation_prompt = false;
|
||||
params.continue_final_message = COMMON_CHAT_CONTINUATION_NONE;
|
||||
std::string no_gen_prompt = common_chat_template_direct_apply_impl(tmpl, params, adjusted_messages, tools_override, additional_context);
|
||||
std::string no_gen_prompt = common_chat_template_direct_apply_impl(tmpl, params, messages_override, tools_override, additional_context);
|
||||
params.add_generation_prompt = true;
|
||||
std::string gen_prompt = common_chat_template_direct_apply_impl(tmpl, params, adjusted_messages, tools_override, additional_context);
|
||||
std::string gen_prompt = common_chat_template_direct_apply_impl(tmpl, params, messages_override, tools_override, additional_context);
|
||||
|
||||
size_t prefix_len = 0;
|
||||
size_t min_size = std::min(no_gen_prompt.size(), gen_prompt.size());
|
||||
|
||||
+10
-4
@@ -23,7 +23,7 @@ void caps_apply_preserve_reasoning(jinja::context & ctx, bool enabled) {
|
||||
ctx.set_val("preserve_thinking", mk_val<value_bool>(enabled));
|
||||
ctx.set_val("clear_thinking", mk_val<value_bool>(!enabled));
|
||||
ctx.set_val("truncate_history_thinking", mk_val<value_bool>(!enabled));
|
||||
ctx.set_val("drop_thinking", mk_val<value_bool>(!enabled));
|
||||
ctx.set_val("drop_thinking", mk_val<value_bool>(!enabled));
|
||||
}
|
||||
|
||||
void caps_apply_reasoning_effort(jinja::context & ctx, const std::string & effort) {
|
||||
@@ -117,6 +117,8 @@ caps caps_get(jinja::program & prog) {
|
||||
|
||||
JJ_DEBUG("%s\n", ">>> Running capability check: typed content");
|
||||
|
||||
static const std::string content_marker = "STRING_MARKER";
|
||||
|
||||
// case: typed content support
|
||||
caps_try_execute(
|
||||
prog,
|
||||
@@ -125,22 +127,26 @@ caps caps_get(jinja::program & prog) {
|
||||
return json::array({
|
||||
{
|
||||
{"role", "user"},
|
||||
{"content", "content"}
|
||||
{"content", content_marker}
|
||||
}
|
||||
});
|
||||
},
|
||||
nullptr, // ctx_fn
|
||||
nullptr, // tools_fn
|
||||
[&](context &, bool success, value & messages, value &, const std::string &) {
|
||||
[&](context &, bool success, value & messages, value &, const std::string & rendered) {
|
||||
auto & content = messages->at(0)->at("content");
|
||||
caps_print_stats(content, "messages[0].content");
|
||||
if (has_op(content, "selectattr") || has_op(content, "array_access")) {
|
||||
bool used_as_array = has_op(content, "selectattr") || has_op(content, "array_access");
|
||||
if (used_as_array) {
|
||||
// accessed as an array
|
||||
result.supports_typed_content = true;
|
||||
}
|
||||
if (!success) {
|
||||
// failed to execute with content as string
|
||||
result.supports_string_content = false;
|
||||
} else if (used_as_array && rendered.find(content_marker) == std::string::npos) {
|
||||
// edge case: string may be accessed for checking, but does not appear in the output
|
||||
result.supports_string_content = false;
|
||||
}
|
||||
}
|
||||
);
|
||||
|
||||
+3
-3
@@ -77,8 +77,8 @@ Legend:
|
||||
| MUL_MAT_ID | ❌ | 🟡 | ✅ | ✅ | 🟡 | 🟡 | 🟡 | ✅ | ✅ | 🟡 | 🟡 | ❌ |
|
||||
| NEG | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| NORM | ❌ | ✅ | ✅ | ✅ | 🟡 | ✅ | ✅ | ✅ | ✅ | 🟡 | ❌ | ❌ |
|
||||
| OPT_STEP_ADAMW | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ |
|
||||
| OPT_STEP_SGD | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ❌ | ❌ | ✅ | ❌ | ❌ | ❌ |
|
||||
| OPT_STEP_ADAMW | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
|
||||
| OPT_STEP_SGD | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
|
||||
| OUT_PROD | 🟡 | 🟡 | 🟡 | 🟡 | ❌ | ❌ | ❌ | 🟡 | 🟡 | ❌ | ❌ | 🟡 |
|
||||
| PAD | ❌ | 🟡 | ✅ | 🟡 | ❌ | 🟡 | 🟡 | 🟡 | ✅ | ✅ | ❌ | ❌ |
|
||||
| PAD_REFLECT_1D | ❌ | ✅ | ✅ | ✅ | ❌ | ✅ | ❌ | ✅ | ❌ | ❌ | ❌ | ❌ |
|
||||
@@ -98,7 +98,7 @@ Legend:
|
||||
| RWKV_WKV7 | ❌ | ❌ | ✅ | ✅ | ✅ | ✅ | ❌ | ✅ | ✅ | ❌ | ❌ | ❌ |
|
||||
| SCALE | ❌ | 🟡 | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| SET | ❌ | ❌ | ✅ | ✅ | ❌ | ✅ | ❌ | 🟡 | ✅ | ✅ | ❌ | ❌ |
|
||||
| SET_ROWS | ❌ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ❌ | ❌ |
|
||||
| SET_ROWS | ❌ | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | 🟡 | ✅ | 🟡 | 🟡 | ❌ | ❌ |
|
||||
| SGN | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | ❌ | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| SIGMOID | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
| SILU | ❌ | ✅ | ✅ | 🟡 | 🟡 | ✅ | 🟡 | ✅ | ✅ | ✅ | ❌ | ❌ |
|
||||
|
||||
+640
-20006
File diff suppressed because it is too large
Load Diff
@@ -4,6 +4,6 @@
|
||||
# Copyright (C) 2026 Intel Corporation
|
||||
# SPDX-License-Identifier: MIT
|
||||
|
||||
./build/bin/test-backend-ops support --output csv > docs/ops/SYCL.csv
|
||||
./build/bin/test-backend-ops -b SYCL0 support --output csv > docs/ops/SYCL.csv
|
||||
./scripts/create_ops_docs.py
|
||||
|
||||
|
||||
+1
-1
@@ -5,7 +5,7 @@ project("ggml" C CXX ASM)
|
||||
### GGML Version
|
||||
set(GGML_VERSION_MAJOR 0)
|
||||
set(GGML_VERSION_MINOR 20)
|
||||
set(GGML_VERSION_PATCH 0)
|
||||
set(GGML_VERSION_PATCH 1)
|
||||
set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
|
||||
|
||||
list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/")
|
||||
|
||||
+102
-56
@@ -349,8 +349,9 @@ static void ggml_cpy_f32_q8_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK8_0 == 0);
|
||||
const int num_blocks = ne / QK8_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK8_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q8_0, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -361,8 +362,10 @@ static void ggml_cpy_q8_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
GGML_ASSERT(ne % QK8_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK8_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q8_0_f32, QK8_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -373,9 +376,11 @@ static void ggml_cpy_q2_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK2_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK2_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_q_f32<cpy_blck_q2_0_f32, QK2_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11,
|
||||
ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -387,8 +392,9 @@ static void ggml_cpy_f32_q4_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ne / QK4_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -399,9 +405,11 @@ static void ggml_cpy_q4_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -414,8 +422,9 @@ static void ggml_cpy_f32_q4_1_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ne / QK4_1;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -426,9 +435,11 @@ static void ggml_cpy_q4_1_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -441,8 +452,9 @@ static void ggml_cpy_f32_q5_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ne / QK5_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_f32_q<cpy_blck_f32_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -453,9 +465,11 @@ static void ggml_cpy_q5_0_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -468,8 +482,9 @@ static void ggml_cpy_f32_q5_1_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK5_1 == 0);
|
||||
const int num_blocks = ne / QK5_1;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK5_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f32_q5_1, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03,
|
||||
ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -480,9 +495,11 @@ static void ggml_cpy_q5_1_f32_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK5_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02,
|
||||
nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13,
|
||||
@@ -494,9 +511,11 @@ static void ggml_cpy_mxfp4_f32_sycl(const char * cx, char * cdst, const int ne,
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ne;
|
||||
GGML_ASSERT(ne % QK_MXFP4 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_MXFP4, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_q_f32<cpy_blck_q_f32<dequantize_mxfp4, QK_MXFP4>, QK_MXFP4>(cx, cdst, ne, ne00, ne01, ne02, nb00,
|
||||
nb01, nb02, nb03, ne10, ne11, ne12,
|
||||
@@ -509,9 +528,10 @@ static void ggml_cpy_f32_iq4_nl_sycl(const char * cx, char * cdst, const int ne,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_NL == 0);
|
||||
const int num_blocks = ne / QK4_NL;
|
||||
const int num_blocks = ceil_div(ne / QK4_NL, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
cpy_f32_q<cpy_blck_f32_iq4_nl, QK4_NL>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11,
|
||||
ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -556,8 +576,9 @@ static void ggml_cpy_f16_q4_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ne / QK4_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f16_q4_0, QK4_0>(cx, cdst, ne, ne00, ne01, ne02,
|
||||
nb00, nb01, nb02, nb03,
|
||||
@@ -570,8 +591,9 @@ static void ggml_cpy_f16_q4_1_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ne / QK4_1;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f16_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02,
|
||||
nb00, nb01, nb02, nb03,
|
||||
@@ -584,8 +606,9 @@ static void ggml_cpy_f16_q5_0_sycl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ne / QK5_0;
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks), sycl::range<3>(1, 1, 1)),
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_f32_q<cpy_blck_f16_q5_0, QK5_0>(cx, cdst, ne, ne00, ne01, ne02,
|
||||
nb00, nb01, nb02, nb03,
|
||||
@@ -849,7 +872,8 @@ static void ggml_cpy_q8_0_q8_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK8_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK8_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
@@ -863,7 +887,8 @@ static void ggml_cpy_q5_0_q5_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK5_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
@@ -877,7 +902,8 @@ static void ggml_cpy_q5_1_q5_1(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK5_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK5_1, SYCL_CPY_BLOCK_SIZE);
|
||||
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE),
|
||||
@@ -892,7 +918,8 @@ static void ggml_cpy_q4_0_q4_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK4_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -906,8 +933,9 @@ static void ggml_cpy_q4_1_q4_1(const char * cx, char * cdst, const int ne, const
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
GGML_ASSERT(ne % QK4_1 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_1, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
cpy_q_q<block_q4_1, QK4_1>(cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb13, item_ct1);
|
||||
@@ -918,7 +946,8 @@ static void ggml_cpy_q1_0_q1_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK1_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK1_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
@@ -930,7 +959,8 @@ static void ggml_cpy_q2_0_q2_0(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK2_0 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK2_0, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -942,7 +972,8 @@ static void ggml_cpy_mxfp4_mxfp4(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_MXFP4 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_MXFP4, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
@@ -954,7 +985,8 @@ static void ggml_cpy_nvfp4_nvfp4(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_NVFP4 == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_NVFP4, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -966,7 +998,8 @@ static void ggml_cpy_q2_K_q2_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -978,7 +1011,8 @@ static void ggml_cpy_q3_K_q3_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -990,7 +1024,8 @@ static void ggml_cpy_q4_K_q4_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1002,7 +1037,8 @@ static void ggml_cpy_q5_K_q5_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1014,7 +1050,8 @@ static void ggml_cpy_q6_K_q6_K(const char * cx, char * cdst, const int ne, const
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1026,7 +1063,8 @@ static void ggml_cpy_iq2_xxs_iq2_xxs(const char * cx, char * cdst, const int ne,
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1038,7 +1076,8 @@ static void ggml_cpy_iq2_xs_iq2_xs(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1050,7 +1089,8 @@ static void ggml_cpy_iq2_s_iq2_s(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1062,7 +1102,8 @@ static void ggml_cpy_iq3_xxs_iq3_xxs(const char * cx, char * cdst, const int ne,
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1074,7 +1115,8 @@ static void ggml_cpy_iq1_s_iq1_s(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1086,7 +1128,8 @@ static void ggml_cpy_iq1_m_iq1_m(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1098,7 +1141,8 @@ static void ggml_cpy_iq4_nl_iq4_nl(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK4_NL == 0);
|
||||
const int num_blocks = ceil_div(ne / QK4_NL, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1110,7 +1154,8 @@ static void ggml_cpy_iq3_s_iq3_s(const char * cx, char * cdst, const int ne, con
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
@@ -1122,7 +1167,8 @@ static void ggml_cpy_iq4_xs_iq4_xs(const char * cx, char * cdst, const int ne, c
|
||||
const int ne02, const int nb00, const int nb01, const int nb02, const int nb03,
|
||||
const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
|
||||
const int nb12, const int nb13, queue_ptr stream) {
|
||||
const int num_blocks = ceil_div(ne, SYCL_CPY_BLOCK_SIZE);
|
||||
GGML_ASSERT(ne % QK_K == 0);
|
||||
const int num_blocks = ceil_div(ne / QK_K, SYCL_CPY_BLOCK_SIZE);
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<3>(sycl::range<3>(1, 1, num_blocks) * sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE), sycl::range<3>(1, 1, SYCL_CPY_BLOCK_SIZE)),
|
||||
[=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]]{
|
||||
|
||||
@@ -77,6 +77,7 @@
|
||||
#include "ggml-sycl/fill.hpp"
|
||||
#include "ggml-sycl/cumsum.hpp"
|
||||
#include "ggml-sycl/diag.hpp"
|
||||
#include "ggml-sycl/opt-step.hpp"
|
||||
#include "ggml-sycl/solve_tri.hpp"
|
||||
#include "ggml-sycl/gated_delta_net.hpp"
|
||||
#include "ggml-sycl/pool.hpp"
|
||||
@@ -5355,6 +5356,12 @@ static bool ggml_sycl_compute_forward(ggml_backend_sycl_context & ctx, struct gg
|
||||
case GGML_OP_GATED_DELTA_NET:
|
||||
ggml_sycl_gated_delta_net(ctx, dst);
|
||||
break;
|
||||
case GGML_OP_OPT_STEP_ADAMW:
|
||||
ggml_sycl_opt_step_adamw(ctx, dst);
|
||||
break;
|
||||
case GGML_OP_OPT_STEP_SGD:
|
||||
ggml_sycl_opt_step_sgd(ctx, dst);
|
||||
break;
|
||||
case GGML_OP_SSM_CONV:
|
||||
ggml_sycl_ssm_conv(ctx, dst);
|
||||
break;
|
||||
@@ -6263,6 +6270,8 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
|
||||
case GGML_OP_RWKV_WKV7:
|
||||
case GGML_OP_GATED_LINEAR_ATTN:
|
||||
case GGML_OP_GATED_DELTA_NET:
|
||||
case GGML_OP_OPT_STEP_ADAMW:
|
||||
case GGML_OP_OPT_STEP_SGD:
|
||||
return true;
|
||||
case GGML_OP_SSM_CONV:
|
||||
return op->type == GGML_TYPE_F32 &&
|
||||
|
||||
@@ -0,0 +1,131 @@
|
||||
#include "opt-step.hpp"
|
||||
|
||||
#define SYCL_OPT_STEP_BLOCK_SIZE 256
|
||||
|
||||
template <typename T>
|
||||
static void opt_step_adamw_f32_kernel(
|
||||
T * __restrict__ x,
|
||||
const T * __restrict__ g,
|
||||
T * __restrict__ g_m,
|
||||
T * __restrict__ g_v,
|
||||
const T * __restrict__ pars,
|
||||
const int64_t k,
|
||||
const sycl::nd_item<1> & item) {
|
||||
|
||||
const int64_t i = (int64_t) item.get_global_id(0);
|
||||
if (i >= k) {
|
||||
return;
|
||||
}
|
||||
|
||||
const float alpha = pars[0];
|
||||
const float beta1 = pars[1];
|
||||
const float beta2 = pars[2];
|
||||
const float eps = pars[3];
|
||||
const float wd = pars[4];
|
||||
const float beta1h = pars[5];
|
||||
const float beta2h = pars[6];
|
||||
|
||||
const float gi = g[i];
|
||||
const float gmi = g_m[i] * beta1 + gi * (1.0f - beta1);
|
||||
const float gvi = g_v[i] * beta2 + gi * gi * (1.0f - beta2);
|
||||
|
||||
g_m[i] = gmi;
|
||||
g_v[i] = gvi;
|
||||
|
||||
const float mh = gmi * beta1h;
|
||||
const float vh = sycl::sqrt(gvi * beta2h) + eps;
|
||||
|
||||
x[i] = x[i] * (1.0f - alpha * wd) - alpha * mh / vh;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
static void opt_step_sgd_f32_kernel(
|
||||
T * __restrict__ x,
|
||||
const T * __restrict__ g,
|
||||
const T * __restrict__ pars,
|
||||
const int64_t k,
|
||||
const sycl::nd_item<1> & item) {
|
||||
|
||||
const int64_t i = (int64_t) item.get_global_id(0);
|
||||
if (i >= k) {
|
||||
return;
|
||||
}
|
||||
|
||||
x[i] = x[i] * (1.0f - pars[0] * pars[1]) - pars[0] * g[i];
|
||||
}
|
||||
|
||||
void ggml_sycl_opt_step_adamw(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
||||
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/5);
|
||||
|
||||
const ggml_tensor * src0 = dst->src[0];
|
||||
const ggml_tensor * src0_grad = dst->src[1];
|
||||
const ggml_tensor * src0_grad_m = dst->src[2];
|
||||
const ggml_tensor * src0_grad_v = dst->src[3];
|
||||
const ggml_tensor * adamw_params = dst->src[4];
|
||||
|
||||
GGML_ASSERT(src0->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(src0_grad->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(src0_grad_m->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(src0_grad_v->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(adamw_params->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(ggml_is_contiguous(src0));
|
||||
GGML_ASSERT(ggml_is_contiguous(src0_grad));
|
||||
GGML_ASSERT(ggml_is_contiguous(src0_grad_m));
|
||||
GGML_ASSERT(ggml_is_contiguous(src0_grad_v));
|
||||
GGML_ASSERT(ggml_is_contiguous(adamw_params));
|
||||
GGML_ASSERT(ggml_are_same_shape(src0, src0_grad));
|
||||
GGML_ASSERT(ggml_are_same_shape(src0, src0_grad_m));
|
||||
GGML_ASSERT(ggml_are_same_shape(src0, src0_grad_v));
|
||||
GGML_ASSERT(ggml_nelements(adamw_params) == 7);
|
||||
|
||||
dpct::queue_ptr stream = ctx.stream();
|
||||
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
|
||||
|
||||
float * src0_d = (float *) src0->data;
|
||||
const float * src0_grad_d = (const float *) src0_grad->data;
|
||||
float * src0_grad_m_d = (float *) src0_grad_m->data;
|
||||
float * src0_grad_v_d = (float *) src0_grad_v->data;
|
||||
const float * adamw_params_d = (const float *) adamw_params->data;
|
||||
|
||||
const int64_t ne = ggml_nelements(src0);
|
||||
const int64_t num_blocks = (ne + SYCL_OPT_STEP_BLOCK_SIZE - 1) / SYCL_OPT_STEP_BLOCK_SIZE;
|
||||
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<1>(num_blocks * SYCL_OPT_STEP_BLOCK_SIZE, SYCL_OPT_STEP_BLOCK_SIZE),
|
||||
[=](sycl::nd_item<1> item) {
|
||||
opt_step_adamw_f32_kernel(src0_d, src0_grad_d, src0_grad_m_d, src0_grad_v_d, adamw_params_d, ne, item);
|
||||
});
|
||||
}
|
||||
|
||||
void ggml_sycl_opt_step_sgd(ggml_backend_sycl_context & ctx, ggml_tensor * dst) {
|
||||
scope_op_debug_print scope_dbg_print(__func__, dst, /*num_src=*/3);
|
||||
|
||||
const ggml_tensor * src0 = dst->src[0];
|
||||
const ggml_tensor * src0_grad = dst->src[1];
|
||||
const ggml_tensor * sgd_params = dst->src[2];
|
||||
|
||||
GGML_ASSERT(src0->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(src0_grad->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(sgd_params->type == GGML_TYPE_F32);
|
||||
GGML_ASSERT(ggml_is_contiguous(src0));
|
||||
GGML_ASSERT(ggml_is_contiguous(src0_grad));
|
||||
GGML_ASSERT(ggml_is_contiguous(sgd_params));
|
||||
GGML_ASSERT(ggml_are_same_shape(src0, src0_grad));
|
||||
GGML_ASSERT(ggml_nelements(sgd_params) == 2);
|
||||
|
||||
dpct::queue_ptr stream = ctx.stream();
|
||||
SYCL_CHECK(ggml_sycl_set_device(ctx.device));
|
||||
|
||||
float * src0_d = (float *) src0->data;
|
||||
const float * src0_grad_d = (const float *) src0_grad->data;
|
||||
const float * sgd_params_d = (const float *) sgd_params->data;
|
||||
|
||||
const int64_t ne = ggml_nelements(src0);
|
||||
const int64_t num_blocks = (ne + SYCL_OPT_STEP_BLOCK_SIZE - 1) / SYCL_OPT_STEP_BLOCK_SIZE;
|
||||
|
||||
stream->parallel_for(
|
||||
sycl::nd_range<1>(num_blocks * SYCL_OPT_STEP_BLOCK_SIZE, SYCL_OPT_STEP_BLOCK_SIZE),
|
||||
[=](sycl::nd_item<1> item) {
|
||||
opt_step_sgd_f32_kernel(src0_d, src0_grad_d, sgd_params_d, ne, item);
|
||||
});
|
||||
}
|
||||
@@ -0,0 +1,6 @@
|
||||
#pragma once
|
||||
|
||||
#include "common.hpp"
|
||||
|
||||
void ggml_sycl_opt_step_adamw(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
|
||||
void ggml_sycl_opt_step_sgd(ggml_backend_sycl_context & ctx, ggml_tensor * dst);
|
||||
@@ -11,6 +11,7 @@ SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
|
||||
REPO_ROOT="$(cd "$SCRIPT_DIR/.." && pwd)"
|
||||
|
||||
DRY_RUN=false
|
||||
CHECKS_PASSED=true
|
||||
for arg in "$@"; do
|
||||
case "$arg" in
|
||||
--dry-run) DRY_RUN=true ;;
|
||||
@@ -44,6 +45,7 @@ else
|
||||
if [[ "$RUNS" -eq 0 ]]; then
|
||||
if [[ "$DRY_RUN" == "true" ]]; then
|
||||
echo "Warning: no successful release.yml run found for HEAD (${SHA}) (dry run, continuing)."
|
||||
CHECKS_PASSED=false
|
||||
else
|
||||
echo "Error: no successful release.yml run found for HEAD (${SHA})"
|
||||
echo "The nightly build must complete successfully before making a release."
|
||||
@@ -73,6 +75,7 @@ else
|
||||
echo "$DIFF"
|
||||
if [[ "$DRY_RUN" == "true" ]]; then
|
||||
echo "Warning: would abort release due to ggml mismatch (dry run, continuing)."
|
||||
CHECKS_PASSED=false
|
||||
else
|
||||
echo "Error: ggml must match upstream before making a release."
|
||||
exit 1
|
||||
@@ -81,3 +84,7 @@ else
|
||||
echo "local ggml/ matches upstream ${GGML_VERSION}"
|
||||
fi
|
||||
fi
|
||||
|
||||
if [[ -n "${GITHUB_OUTPUT:-}" ]]; then
|
||||
echo "checks_passed=${CHECKS_PASSED}" >> "$GITHUB_OUTPUT"
|
||||
fi
|
||||
|
||||
@@ -1 +1 @@
|
||||
2d191b5dee1a591c41ee8a653ce42bfcd9c8716d
|
||||
3834fd814e74e8af277939dabd69ecc780affd21
|
||||
|
||||
@@ -396,8 +396,11 @@ static void llama_adapter_lora_init_impl(llama_model & model, const char * path_
|
||||
llama_file gguf_file(path_lora, "rb");
|
||||
std::vector<uint8_t> read_buf;
|
||||
auto set_tensor = [&](ggml_tensor * orig, ggml_tensor * dev) {
|
||||
size_t offs = gguf_get_data_offset(ctx_gguf.get()) + gguf_get_tensor_offset(ctx_gguf.get(), gguf_find_tensor(ctx_gguf.get(), orig->name));
|
||||
size_t size = ggml_nbytes(orig);
|
||||
const size_t offs = gguf_get_data_offset(ctx_gguf.get()) + gguf_get_tensor_offset(ctx_gguf.get(), gguf_find_tensor(ctx_gguf.get(), orig->name));
|
||||
const size_t size = ggml_nbytes(orig);
|
||||
if (offs + size < offs || offs + size > gguf_file.size()) {
|
||||
throw std::runtime_error(format("LoRA tensor '%s' data is not within the file bounds, file is corrupted or incomplete", orig->name));
|
||||
}
|
||||
read_buf.resize(size);
|
||||
gguf_file.seek(offs, SEEK_SET);
|
||||
gguf_file.read_raw(read_buf.data(), size);
|
||||
|
||||
@@ -180,10 +180,11 @@ llama_model_deepseek32::graph::graph(const llama_model & model, const llm_graph_
|
||||
|
||||
const int64_t n_indexer_head = hparams.indexer_n_head;
|
||||
const int64_t n_embd_indexer_head = hparams.indexer_head_size;
|
||||
const int64_t n_embd_indexer_head_rope = hparams.n_rot();
|
||||
const int64_t n_embd_indexer_head_nope = n_embd_indexer_head - n_embd_indexer_head_rope;
|
||||
const uint32_t n_indexer_top_k = hparams.indexer_top_k;
|
||||
|
||||
// the indexer head layous is [rope | nope]
|
||||
GGML_ASSERT(hparams.n_rot() <= n_embd_indexer_head);
|
||||
|
||||
const uint32_t kv_lora_rank = hparams.n_lora_kv;
|
||||
|
||||
// We have to pre-scale kq_scale and attn_factor to make the YaRN RoPE work correctly.
|
||||
@@ -233,28 +234,11 @@ llama_model_deepseek32::graph::graph(const llama_model & model, const llm_graph_
|
||||
ggml_tensor * indexer_q = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_q_b, qr);
|
||||
cb(indexer_q, "indexer_q", il);
|
||||
|
||||
// split into {n_embd_indexer_head_rope, n_indexer_head, n_tokens}
|
||||
ggml_tensor * indexer_q_pe =
|
||||
ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_rope, n_indexer_head, n_tokens,
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head, 0);
|
||||
cb(indexer_q_pe, "indexer_q_pe", il);
|
||||
|
||||
// and {n_embd_indexer_head_nope, n_indexer_head, n_tokens}
|
||||
ggml_tensor * indexer_q_nope =
|
||||
ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_nope, n_indexer_head, n_tokens,
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head,
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head_nope));
|
||||
cb(indexer_q_nope, "indexer_q_nope", il);
|
||||
|
||||
indexer_q_pe = ggml_rope_ext(ctx0, indexer_q_pe, inp_pos, nullptr, n_rot,
|
||||
// {n_embd_indexer_head, n_indexer_head, n_tokens}
|
||||
indexer_q = ggml_reshape_3d(ctx0, indexer_q, n_embd_indexer_head, n_indexer_head, n_tokens);
|
||||
indexer_q = ggml_rope_ext(ctx0, indexer_q, inp_pos, nullptr, n_rot,
|
||||
LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
cb(indexer_q_pe, "indexer_q_pe", il);
|
||||
|
||||
// {n_embd_indexer_head_rope + n_embd_indexer_head_nope, n_head, n_tokens}
|
||||
indexer_q = ggml_concat(ctx0, indexer_q_pe, indexer_q_nope, 0);
|
||||
cb(indexer_q, "indexer_q", il);
|
||||
|
||||
ggml_tensor * indexer_k = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_k, cur);
|
||||
@@ -263,28 +247,11 @@ llama_model_deepseek32::graph::graph(const llama_model & model, const llm_graph_
|
||||
indexer_k = build_norm(indexer_k, model.layers[il].indexer_k_norm, model.layers[il].indexer_k_norm_b, LLM_NORM, il);
|
||||
cb(indexer_k, "indexer_k", il);
|
||||
|
||||
// split into {n_embd_indexer_head_rope, 1, n_tokens}
|
||||
ggml_tensor * indexer_k_pe =
|
||||
ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_rope, 1, n_tokens,
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1, 0);
|
||||
cb(indexer_k_pe, "indexer_k_pe", il);
|
||||
|
||||
// and {n_embd_indexer_head_nope, 1, n_tokens}
|
||||
ggml_tensor * indexer_k_nope =
|
||||
ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_nope, 1, n_tokens,
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1,
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head_nope));
|
||||
cb(indexer_k_nope, "indexer_k_nope", il);
|
||||
|
||||
indexer_k_pe = ggml_rope_ext(ctx0, indexer_k_pe, inp_pos, nullptr, n_rot,
|
||||
// {n_embd_indexer_head, 1, n_tokens}
|
||||
indexer_k = ggml_reshape_3d(ctx0, indexer_k, n_embd_indexer_head, 1, n_tokens);
|
||||
indexer_k = ggml_rope_ext(ctx0, indexer_k, inp_pos, nullptr, n_rot,
|
||||
LLAMA_ROPE_TYPE_NEOX, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
cb(indexer_k_pe, "indexer_k_pe", il);
|
||||
|
||||
// {n_embd_indexer_head_rope + n_embd_indexer_head_nope, 1, n_tokens}
|
||||
indexer_k = ggml_concat(ctx0, indexer_k_pe, indexer_k_nope, 0);
|
||||
cb(indexer_k, "indexer_k", il);
|
||||
|
||||
// perform Hadamard transform on indexer q and k
|
||||
|
||||
+9
-42
@@ -216,10 +216,11 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par
|
||||
|
||||
const int64_t n_indexer_head = hparams.indexer_n_head;
|
||||
const int64_t n_embd_indexer_head = hparams.indexer_head_size;
|
||||
const int64_t n_embd_indexer_head_rope = hparams.n_rot();
|
||||
const int64_t n_embd_indexer_head_nope = n_embd_indexer_head - n_embd_indexer_head_rope;
|
||||
const uint32_t n_indexer_top_k = hparams.indexer_top_k;
|
||||
|
||||
// the indexer head layout is [rope | nope]
|
||||
GGML_ASSERT(hparams.n_rot() <= n_embd_indexer_head);
|
||||
|
||||
const uint32_t kv_lora_rank = hparams.n_lora_kv;
|
||||
|
||||
// We have to pre-scale kq_scale and attn_factor to make the YaRN RoPE work correctly.
|
||||
@@ -273,28 +274,11 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par
|
||||
ggml_tensor * indexer_q = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_q_b, qr);
|
||||
cb(indexer_q, "indexer_q", il);
|
||||
|
||||
// split into {n_embd_indexer_head_rope, n_indexer_head, n_tokens}
|
||||
ggml_tensor * indexer_q_pe =
|
||||
ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_rope, n_indexer_head, n_tokens,
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head, 0);
|
||||
cb(indexer_q_pe, "indexer_q_pe", il);
|
||||
|
||||
// and {n_embd_indexer_head_nope, n_indexer_head, n_tokens}
|
||||
ggml_tensor * indexer_q_nope =
|
||||
ggml_view_3d(ctx0, indexer_q, n_embd_indexer_head_nope, n_indexer_head, n_tokens,
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head) * n_indexer_head,
|
||||
ggml_row_size(indexer_q->type, n_embd_indexer_head_nope));
|
||||
cb(indexer_q_nope, "indexer_q_nope", il);
|
||||
|
||||
indexer_q_pe = ggml_rope_ext(ctx0, indexer_q_pe, inp_pos, nullptr, n_rot,
|
||||
// {n_embd_indexer_head, n_indexer_head, n_tokens}
|
||||
indexer_q = ggml_reshape_3d(ctx0, indexer_q, n_embd_indexer_head, n_indexer_head, n_tokens);
|
||||
indexer_q = ggml_rope_ext(ctx0, indexer_q, inp_pos, nullptr, n_rot,
|
||||
LLAMA_ROPE_TYPE_NORM, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
cb(indexer_q_pe, "indexer_q_pe", il);
|
||||
|
||||
// {n_embd_indexer_head_rope + n_embd_indexer_head_nope, n_head, n_tokens}
|
||||
indexer_q = ggml_concat(ctx0, indexer_q_pe, indexer_q_nope, 0);
|
||||
cb(indexer_q, "indexer_q", il);
|
||||
|
||||
ggml_tensor * indexer_k = ggml_mul_mat(ctx0, model.layers[il].indexer_attn_k, cur);
|
||||
@@ -303,28 +287,11 @@ llama_model_glm_dsa::graph::graph(const llama_model & model, const llm_graph_par
|
||||
indexer_k = build_norm(indexer_k, model.layers[il].indexer_k_norm, model.layers[il].indexer_k_norm_b, LLM_NORM, il);
|
||||
cb(indexer_k, "indexer_k", il);
|
||||
|
||||
// split into {n_embd_indexer_head_rope, 1, n_tokens}
|
||||
ggml_tensor * indexer_k_pe =
|
||||
ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_rope, 1, n_tokens,
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1, 0);
|
||||
cb(indexer_k_pe, "indexer_k_pe", il);
|
||||
|
||||
// and {n_embd_indexer_head_nope, 1, n_tokens}
|
||||
ggml_tensor * indexer_k_nope =
|
||||
ggml_view_3d(ctx0, indexer_k, n_embd_indexer_head_nope, 1, n_tokens,
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head),
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head) * 1,
|
||||
ggml_row_size(indexer_k->type, n_embd_indexer_head_nope));
|
||||
cb(indexer_k_nope, "indexer_k_nope", il);
|
||||
|
||||
indexer_k_pe = ggml_rope_ext(ctx0, indexer_k_pe, inp_pos, nullptr, n_rot,
|
||||
// {n_embd_indexer_head, 1, n_tokens}
|
||||
indexer_k = ggml_reshape_3d(ctx0, indexer_k, n_embd_indexer_head, 1, n_tokens);
|
||||
indexer_k = ggml_rope_ext(ctx0, indexer_k, inp_pos, nullptr, n_rot,
|
||||
LLAMA_ROPE_TYPE_NORM, n_ctx_orig, freq_base, freq_scale,
|
||||
ext_factor, attn_factor, beta_fast, beta_slow);
|
||||
cb(indexer_k_pe, "indexer_k_pe", il);
|
||||
|
||||
// {n_embd_indexer_head_rope + n_embd_indexer_head_nope, 1, n_tokens}
|
||||
indexer_k = ggml_concat(ctx0, indexer_k_pe, indexer_k_nope, 0);
|
||||
cb(indexer_k, "indexer_k", il);
|
||||
|
||||
// perform Hadamard transform on indexer q and k
|
||||
|
||||
+49
-2
@@ -10,6 +10,7 @@
|
||||
#include "jinja/parser.h"
|
||||
#include "jinja/lexer.h"
|
||||
#include "jinja/utils.h"
|
||||
#include "jinja/caps.h"
|
||||
|
||||
#include "testing.h"
|
||||
|
||||
@@ -33,6 +34,7 @@ static void test_array_methods(testing & t);
|
||||
static void test_object_methods(testing & t);
|
||||
static void test_hasher(testing & t);
|
||||
static void test_stats(testing & t);
|
||||
static void test_caps(testing & t);
|
||||
static void test_string_parts(testing & t);
|
||||
static void test_fuzzing(testing & t);
|
||||
|
||||
@@ -73,6 +75,7 @@ int main(int argc, char *argv[]) {
|
||||
if (!g_python_mode) {
|
||||
t.test("hasher", test_hasher);
|
||||
t.test("stats", test_stats);
|
||||
t.test("caps", test_caps);
|
||||
t.test("string parts", test_string_parts);
|
||||
t.test("fuzzing", test_fuzzing);
|
||||
}
|
||||
@@ -2059,6 +2062,51 @@ static void test_stats(testing & t) {
|
||||
});
|
||||
}
|
||||
|
||||
static void test_caps(testing & t) {
|
||||
static auto get_caps = [](const std::string & tmpl) -> jinja::caps {
|
||||
jinja::lexer lexer;
|
||||
auto lexer_res = lexer.tokenize(tmpl);
|
||||
|
||||
jinja::program prog = jinja::parse_from_tokens(lexer_res);
|
||||
|
||||
return jinja::caps_get(prog);
|
||||
};
|
||||
|
||||
t.test("string content", [](testing & t) {
|
||||
auto caps = get_caps(
|
||||
"{% for message in messages %}"
|
||||
"{{ message['role'] + ': ' + message['content'] }}"
|
||||
"{% endfor %}"
|
||||
);
|
||||
t.assert_true("supports string content", caps.supports_string_content);
|
||||
t.assert_true("does not support typed content", !caps.supports_typed_content);
|
||||
});
|
||||
|
||||
t.test("typed content, raises on string", [](testing & t) {
|
||||
// 'selectattr' is not a String filter, so it throws
|
||||
auto caps = get_caps(
|
||||
"{% for message in messages %}"
|
||||
"{% for content in message['content'] | selectattr('type', 'equalto', 'text') %}"
|
||||
"{{ content['text'] }}"
|
||||
"{% endfor %}"
|
||||
"{% endfor %}"
|
||||
);
|
||||
t.assert_true("does not support string content", !caps.supports_string_content);
|
||||
t.assert_true("supports typed content", caps.supports_typed_content);
|
||||
});
|
||||
|
||||
t.test("typed content, silently drops string", [](testing & t) {
|
||||
// no throw here, but content[0]['text'] is undefined for a string (MiniMax-M1 case)
|
||||
auto caps = get_caps(
|
||||
"{% for message in messages %}"
|
||||
"{{ message['content'][0]['text'] }}"
|
||||
"{% endfor %}"
|
||||
);
|
||||
t.assert_true("does not support string content", !caps.supports_string_content);
|
||||
t.assert_true("supports typed content", caps.supports_typed_content);
|
||||
});
|
||||
}
|
||||
|
||||
static void test_string_parts(testing & t) {
|
||||
static auto render = [](const std::string & tmpl, const json & vars) -> jinja::string {
|
||||
jinja::lexer lexer;
|
||||
@@ -2116,8 +2164,7 @@ static void test_template_cpp(testing & t, const std::string & name, const std::
|
||||
t.log("Actual : " + json(rendered).dump());
|
||||
}
|
||||
} catch (const jinja::not_implemented_exception & e) {
|
||||
// TODO @ngxson : remove this when the test framework supports skipping tests
|
||||
t.log("Skipped: " + std::string(e.what()));
|
||||
t.skip(e.what());
|
||||
}
|
||||
});
|
||||
}
|
||||
|
||||
+28
-3
@@ -21,6 +21,11 @@ struct testing {
|
||||
int failures = 0;
|
||||
int unnamed = 0;
|
||||
int exceptions = 0;
|
||||
int skipped = 0;
|
||||
|
||||
// set by skip(), read by the innermost test()
|
||||
bool skip_current = false;
|
||||
std::string skip_reason;
|
||||
|
||||
static constexpr std::size_t status_column = 80;
|
||||
|
||||
@@ -78,7 +83,12 @@ struct testing {
|
||||
}
|
||||
}
|
||||
|
||||
void print_result(const std::string &label, int new_failures, int new_assertions, const std::string &extra = "") const {
|
||||
void skip(const std::string &reason = "") {
|
||||
skip_current = true;
|
||||
skip_reason = reason;
|
||||
}
|
||||
|
||||
void print_result(const std::string &label, int new_failures, int new_assertions, const std::string &extra = "", bool was_skipped = false) const {
|
||||
std::string line = indent() + label;
|
||||
|
||||
std::string details;
|
||||
@@ -101,7 +111,7 @@ struct testing {
|
||||
line += " (" + details + ")";
|
||||
}
|
||||
|
||||
std::string status = (new_failures == 0) ? "[PASS]" : "[FAIL]";
|
||||
std::string status = new_failures != 0 ? "[FAIL]" : (was_skipped ? "[SKIP]" : "[PASS]");
|
||||
|
||||
if (line.size() + 1 < status_column) {
|
||||
line.append(status_column - line.size(), ' ');
|
||||
@@ -126,12 +136,26 @@ struct testing {
|
||||
int before_failures = failures;
|
||||
int before_assertions = assertions;
|
||||
|
||||
// do not let a skipped subtest also mark its parent as skipped
|
||||
bool outer_skip = skip_current;
|
||||
std::string outer_skip_reason = skip_reason;
|
||||
skip_current = false;
|
||||
skip_reason.clear();
|
||||
|
||||
run_with_exceptions([&] { f(*this); }, "test");
|
||||
|
||||
int new_failures = failures - before_failures;
|
||||
int new_assertions = assertions - before_assertions;
|
||||
|
||||
print_result(name, new_failures, new_assertions);
|
||||
bool was_skipped = skip_current && new_failures == 0;
|
||||
if (was_skipped) {
|
||||
++skipped;
|
||||
}
|
||||
|
||||
print_result(name, new_failures, new_assertions, was_skipped ? skip_reason : "", was_skipped);
|
||||
|
||||
skip_current = outer_skip;
|
||||
skip_reason = outer_skip_reason;
|
||||
|
||||
stack.pop_back();
|
||||
}
|
||||
@@ -238,6 +262,7 @@ struct testing {
|
||||
out << "assertions : " << assertions << "\n";
|
||||
out << "failures : " << failures << "\n";
|
||||
out << "exceptions : " << exceptions << "\n";
|
||||
out << "skipped : " << skipped << "\n";
|
||||
return failures == 0 ? 0 : 1;
|
||||
}
|
||||
};
|
||||
|
||||
@@ -846,7 +846,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
|
||||
invalid_param = true;
|
||||
break;
|
||||
}
|
||||
LOG_WRN("DEPRECATED: -mmp and --mmap are deprecated in favour of --load-mode. Please use --load-mode mmap instead.");
|
||||
LOG_WRN("DEPRECATED: -mmp and --mmap are deprecated in favour of --load-mode. Please use --load-mode mmap instead.\n");
|
||||
auto p = string_split<bool>(argv[i], split_delim);
|
||||
|
||||
std::vector<llama_load_mode> modes;
|
||||
@@ -865,7 +865,7 @@ static cmd_params parse_cmd_params(int argc, char ** argv) {
|
||||
invalid_param = true;
|
||||
break;
|
||||
}
|
||||
LOG_WRN("DEPRECATED: -dio and --direct-io are deprecated in favour of --load-mode. Please use --load-mode dio instead.");
|
||||
LOG_WRN("DEPRECATED: -dio and --direct-io are deprecated in favour of --load-mode. Please use --load-mode dio instead.\n");
|
||||
auto p = string_split<bool>(argv[i], split_delim);
|
||||
|
||||
std::vector<llama_load_mode> modes;
|
||||
|
||||
@@ -8,8 +8,7 @@
|
||||
SettingsChatImportExportTab,
|
||||
SettingsChatMobileHeader,
|
||||
SettingsChatToolsTab,
|
||||
SettingsFooter,
|
||||
SettingsRemoteAccess
|
||||
SettingsFooter
|
||||
} from '$lib/components/app/settings';
|
||||
import { Button } from '$lib/components/ui/button';
|
||||
import {
|
||||
@@ -153,8 +152,6 @@
|
||||
<SettingsChatToolsTab />
|
||||
{:else if currentSection.title === SETTINGS_SECTION_TITLES.IMPORT_EXPORT}
|
||||
<SettingsChatImportExportTab />
|
||||
{:else if currentSection.title === SETTINGS_SECTION_TITLES.REMOTE_ACCESS}
|
||||
<SettingsRemoteAccess />
|
||||
{:else if currentSection.fields}
|
||||
<div class="space-y-6">
|
||||
<SettingsChatFields
|
||||
|
||||
@@ -1,291 +0,0 @@
|
||||
<script lang="ts">
|
||||
import {
|
||||
AlertCircle,
|
||||
Check,
|
||||
Copy,
|
||||
Loader2,
|
||||
RefreshCw,
|
||||
Users,
|
||||
Wifi,
|
||||
WifiOff
|
||||
} from '@lucide/svelte';
|
||||
import { SettingsGroup } from '$lib/components/app/settings';
|
||||
import * as AlertDialog from '$lib/components/ui/alert-dialog';
|
||||
import { Badge } from '$lib/components/ui/badge';
|
||||
import { Button } from '$lib/components/ui/button';
|
||||
import { Input } from '$lib/components/ui/input';
|
||||
import { webrtcStore } from '$lib/stores/webrtc.svelte';
|
||||
import { fade } from 'svelte/transition';
|
||||
|
||||
// -- host state
|
||||
let codeCopied = $state(false);
|
||||
let showRegenerateDialog = $state(false);
|
||||
let regenerating = $state(false);
|
||||
|
||||
// -- client state
|
||||
let joinInput = $state('');
|
||||
let joinError = $state('');
|
||||
let joining = $state(false);
|
||||
|
||||
async function handleStartHost() {
|
||||
await webrtcStore.startHost();
|
||||
}
|
||||
|
||||
function handleStopHost() {
|
||||
webrtcStore.stopHost();
|
||||
}
|
||||
|
||||
function copyCode() {
|
||||
navigator.clipboard.writeText(webrtcStore.shareCode).then(() => {
|
||||
codeCopied = true;
|
||||
setTimeout(() => (codeCopied = false), 2000);
|
||||
});
|
||||
}
|
||||
|
||||
async function handleRegenerateConfirm() {
|
||||
showRegenerateDialog = false;
|
||||
regenerating = true;
|
||||
try {
|
||||
await webrtcStore.regenerateCodes();
|
||||
} finally {
|
||||
regenerating = false;
|
||||
}
|
||||
}
|
||||
|
||||
async function handleJoin() {
|
||||
joinError = '';
|
||||
const code = joinInput.trim().replace(/\s/g, '');
|
||||
|
||||
if (code.length < 40) {
|
||||
joinError = 'Code must be 40 characters';
|
||||
|
||||
return;
|
||||
}
|
||||
|
||||
joining = true;
|
||||
try {
|
||||
await webrtcStore.joinAsClient(code);
|
||||
} catch (e) {
|
||||
joinError = e instanceof Error ? e.message : String(e);
|
||||
} finally {
|
||||
joining = false;
|
||||
}
|
||||
}
|
||||
|
||||
function handleLeave() {
|
||||
webrtcStore.leaveAsClient();
|
||||
joinInput = '';
|
||||
joinError = '';
|
||||
}
|
||||
|
||||
// Display the share code broken into 8-char blocks for readability
|
||||
function formatCode(code: string): string {
|
||||
return code.match(/.{1,8}/g)?.join(' ') ?? code;
|
||||
}
|
||||
</script>
|
||||
|
||||
<div class="space-y-12" in:fade={{ duration: 150 }}>
|
||||
<!-- ------------------------------------------------------------------ -->
|
||||
<!-- HOST -->
|
||||
<!-- ------------------------------------------------------------------ -->
|
||||
<SettingsGroup title="Host">
|
||||
<div class="space-y-4">
|
||||
<p class="text-sm text-muted-foreground">
|
||||
Generate a code and share it so remote devices can connect to this instance.
|
||||
</p>
|
||||
|
||||
<!-- Share code block: shown whenever codes exist, even when host is inactive -->
|
||||
{#if webrtcStore.hasHostCodes}
|
||||
<div class="rounded-lg border border-border bg-muted/40 p-4">
|
||||
<p class="mb-2 text-xs font-medium uppercase tracking-wide text-muted-foreground">
|
||||
Share code
|
||||
</p>
|
||||
<p class="break-all font-mono text-sm tracking-widest select-all">
|
||||
{formatCode(webrtcStore.shareCode)}
|
||||
</p>
|
||||
<p class="mt-2 text-xs text-muted-foreground">
|
||||
Share this 40-character code with remote devices. It includes both the room ID and the
|
||||
passcode.
|
||||
</p>
|
||||
</div>
|
||||
|
||||
<div class="flex flex-wrap gap-2">
|
||||
<Button variant="outline" size="sm" onclick={copyCode} class="gap-1.5">
|
||||
{#if codeCopied}
|
||||
<Check class="h-3.5 w-3.5" />
|
||||
Copied
|
||||
{:else}
|
||||
<Copy class="h-3.5 w-3.5" />
|
||||
Copy code
|
||||
{/if}
|
||||
</Button>
|
||||
|
||||
<Button
|
||||
variant="outline"
|
||||
size="sm"
|
||||
onclick={() => (showRegenerateDialog = true)}
|
||||
disabled={regenerating}
|
||||
class="gap-1.5"
|
||||
>
|
||||
{#if regenerating}
|
||||
<Loader2 class="h-3.5 w-3.5 animate-spin" />
|
||||
Regenerating...
|
||||
{:else}
|
||||
<RefreshCw class="h-3.5 w-3.5" />
|
||||
Regenerate code
|
||||
{/if}
|
||||
</Button>
|
||||
</div>
|
||||
{/if}
|
||||
|
||||
<!-- Status row (only when host is active) -->
|
||||
{#if webrtcStore.mode === 'host'}
|
||||
<div class="flex items-center gap-3">
|
||||
{#if webrtcStore.status === 'connecting'}
|
||||
<Badge variant="secondary" class="gap-1.5">
|
||||
<Loader2 class="h-3 w-3 animate-spin" />
|
||||
Connecting to trackers...
|
||||
</Badge>
|
||||
{:else if webrtcStore.status === 'connected'}
|
||||
<Badge variant="default" class="gap-1.5">
|
||||
<Wifi class="h-3 w-3" />
|
||||
Active
|
||||
</Badge>
|
||||
<span class="flex items-center gap-1 text-sm text-muted-foreground">
|
||||
<Users class="h-3.5 w-3.5" />
|
||||
{webrtcStore.peerCount}
|
||||
{webrtcStore.peerCount === 1 ? 'client' : 'clients'} connected
|
||||
</span>
|
||||
{:else if webrtcStore.status === 'error'}
|
||||
<Badge variant="destructive" class="gap-1.5">
|
||||
<AlertCircle class="h-3 w-3" />
|
||||
Error
|
||||
</Badge>
|
||||
<span class="text-sm text-destructive">{webrtcStore.errorMessage}</span>
|
||||
{/if}
|
||||
</div>
|
||||
{/if}
|
||||
|
||||
<!-- Enable / Disable button -->
|
||||
{#if webrtcStore.mode === 'off' || webrtcStore.mode === 'client'}
|
||||
<Button
|
||||
variant="outline"
|
||||
onclick={handleStartHost}
|
||||
disabled={webrtcStore.mode === 'client'}
|
||||
>
|
||||
<Wifi class="h-4 w-4" />
|
||||
Enable remote access
|
||||
</Button>
|
||||
{:else}
|
||||
<Button variant="outline" onclick={handleStopHost}>
|
||||
<WifiOff class="h-4 w-4" />
|
||||
Disable remote access
|
||||
</Button>
|
||||
{/if}
|
||||
|
||||
{#if webrtcStore.mode === 'client'}
|
||||
<p class="text-xs text-muted-foreground">
|
||||
Disable join mode first before enabling host mode.
|
||||
</p>
|
||||
{/if}
|
||||
</div>
|
||||
</SettingsGroup>
|
||||
|
||||
<!-- ------------------------------------------------------------------ -->
|
||||
<!-- CLIENT / JOIN -->
|
||||
<!-- ------------------------------------------------------------------ -->
|
||||
<SettingsGroup title="Join">
|
||||
<div class="space-y-4">
|
||||
<p class="text-sm text-muted-foreground">
|
||||
Connect to a remote llama.cpp instance. All requests will be routed through the peer-to-peer
|
||||
tunnel.
|
||||
</p>
|
||||
|
||||
{#if webrtcStore.mode === 'client'}
|
||||
<div class="space-y-3">
|
||||
<div class="flex items-center gap-3">
|
||||
{#if webrtcStore.status === 'connecting'}
|
||||
<Badge variant="secondary" class="gap-1.5">
|
||||
<Loader2 class="h-3 w-3 animate-spin" />
|
||||
Connecting...
|
||||
</Badge>
|
||||
{:else if webrtcStore.status === 'connected'}
|
||||
<Badge variant="default" class="gap-1.5">
|
||||
<Wifi class="h-3 w-3" />
|
||||
Connected to host
|
||||
</Badge>
|
||||
{:else if webrtcStore.status === 'error'}
|
||||
<Badge variant="destructive" class="gap-1.5">
|
||||
<AlertCircle class="h-3 w-3" />
|
||||
Disconnected
|
||||
</Badge>
|
||||
<span class="text-sm text-destructive">{webrtcStore.errorMessage}</span>
|
||||
{/if}
|
||||
</div>
|
||||
|
||||
<Button variant="outline" onclick={handleLeave}>
|
||||
<WifiOff class="h-4 w-4" />
|
||||
Leave
|
||||
</Button>
|
||||
</div>
|
||||
{:else}
|
||||
<div class="space-y-3">
|
||||
<div class="space-y-1.5">
|
||||
<label for="join-code" class="text-sm font-medium">Access code</label>
|
||||
<Input
|
||||
id="join-code"
|
||||
placeholder="Paste the 40-character code from the host"
|
||||
bind:value={joinInput}
|
||||
disabled={joining || webrtcStore.mode === 'host'}
|
||||
class="font-mono"
|
||||
/>
|
||||
{#if joinError}
|
||||
<p class="text-sm text-destructive">{joinError}</p>
|
||||
{/if}
|
||||
</div>
|
||||
|
||||
<Button
|
||||
onclick={handleJoin}
|
||||
disabled={joining || joinInput.trim().length < 40 || webrtcStore.mode === 'host'}
|
||||
>
|
||||
{#if joining}
|
||||
<Loader2 class="h-4 w-4 animate-spin" />
|
||||
Connecting...
|
||||
{:else}
|
||||
<Wifi class="h-4 w-4" />
|
||||
Connect
|
||||
{/if}
|
||||
</Button>
|
||||
|
||||
{#if webrtcStore.mode === 'host'}
|
||||
<p class="text-xs text-muted-foreground">
|
||||
Disable host mode first before joining a remote server.
|
||||
</p>
|
||||
{/if}
|
||||
</div>
|
||||
{/if}
|
||||
</div>
|
||||
</SettingsGroup>
|
||||
|
||||
<p class="text-xs text-muted-foreground">
|
||||
Uses WebRTC with Google STUN servers for NAT traversal. Signaling via public WebTorrent
|
||||
trackers. No data is routed through any relay server.
|
||||
</p>
|
||||
</div>
|
||||
|
||||
<!-- Regenerate code confirmation dialog -->
|
||||
<AlertDialog.Root bind:open={showRegenerateDialog}>
|
||||
<AlertDialog.Content>
|
||||
<AlertDialog.Header>
|
||||
<AlertDialog.Title>Regenerate access code?</AlertDialog.Title>
|
||||
<AlertDialog.Description>
|
||||
This will create a new room and passcode. Any devices using the current code will be
|
||||
disconnected and will need to be updated with the new code.
|
||||
</AlertDialog.Description>
|
||||
</AlertDialog.Header>
|
||||
<AlertDialog.Footer>
|
||||
<AlertDialog.Cancel>Cancel</AlertDialog.Cancel>
|
||||
<AlertDialog.Action onclick={handleRegenerateConfirm}>Regenerate</AlertDialog.Action>
|
||||
</AlertDialog.Footer>
|
||||
</AlertDialog.Content>
|
||||
</AlertDialog.Root>
|
||||
@@ -74,10 +74,3 @@ export { default as SettingsChatFields } from './SettingsChat/SettingsChatFields
|
||||
* server favicons and permission management controls.
|
||||
*/
|
||||
export { default as SettingsChatToolsTab } from './SettingsChat/SettingsChatToolsTab.svelte';
|
||||
|
||||
/**
|
||||
* Remote Access configuration panel.
|
||||
* Host mode: generates a share code for remote clients.
|
||||
* Client mode: accepts a share code and routes all same-origin requests through a WebRTC tunnel.
|
||||
*/
|
||||
export { default as SettingsRemoteAccess } from './SettingsRemoteAccess.svelte';
|
||||
|
||||
@@ -18,7 +18,6 @@ export const SETTINGS_SECTION_SLUGS = {
|
||||
GENERAL: 'general',
|
||||
IMPORT_EXPORT: 'import-export',
|
||||
PENALTIES: 'penalties',
|
||||
REMOTE_ACCESS: 'remote-access',
|
||||
SAMPLING: 'sampling',
|
||||
TOOLS: 'tools'
|
||||
} as const;
|
||||
|
||||
@@ -13,7 +13,6 @@ import {
|
||||
Monitor,
|
||||
Moon,
|
||||
PencilRuler,
|
||||
Radio,
|
||||
Sliders,
|
||||
Sun
|
||||
} from '@lucide/svelte';
|
||||
@@ -37,7 +36,6 @@ export const SETTINGS_SECTION_TITLES = {
|
||||
GENERAL: 'General',
|
||||
IMPORT_EXPORT: 'Import/Export',
|
||||
PENALTIES: 'Penalties',
|
||||
REMOTE_ACCESS: 'Remote Access',
|
||||
SAMPLING: 'Sampling',
|
||||
TOOLS: 'Tools'
|
||||
} as const;
|
||||
@@ -48,11 +46,6 @@ const STANDALONE_SECTIONS: { title: SettingsSectionTitle; slug: string; icon: Co
|
||||
icon: Database,
|
||||
slug: SETTINGS_SECTION_SLUGS.IMPORT_EXPORT,
|
||||
title: SETTINGS_SECTION_TITLES.IMPORT_EXPORT
|
||||
},
|
||||
{
|
||||
icon: Radio,
|
||||
slug: SETTINGS_SECTION_SLUGS.REMOTE_ACCESS,
|
||||
title: SETTINGS_SECTION_TITLES.REMOTE_ACCESS
|
||||
}
|
||||
];
|
||||
const COLOR_MODE_OPTIONS: Array<{ value: string; label: string; icon: Component }> = [
|
||||
|
||||
@@ -1,304 +0,0 @@
|
||||
import { browser } from '$app/environment';
|
||||
import {
|
||||
ClientTunnel,
|
||||
generatePassCode,
|
||||
generateRoomCode,
|
||||
HostTunnel
|
||||
} from '$lib/utils/webrtc-tunnel';
|
||||
|
||||
// Stores the generated host codes; persists until explicitly regenerated.
|
||||
const HOST_CODES_KEY = 'llama_webrtc_host_codes';
|
||||
// Stores the active session (mode + codes) for auto-reconnect on reload.
|
||||
const SESSION_KEY = 'llama_webrtc_session';
|
||||
|
||||
type HostCodes = { roomCode: string; passCode: string };
|
||||
type SessionData = { mode: 'host' | 'client'; roomCode: string; passCode: string };
|
||||
type ConnectionStatus = 'idle' | 'connecting' | 'connected' | 'error';
|
||||
|
||||
class WebRTCStore {
|
||||
mode = $state<'off' | 'host' | 'client'>('off');
|
||||
status = $state<ConnectionStatus>('idle');
|
||||
peerCount = $state(0);
|
||||
errorMessage = $state('');
|
||||
|
||||
// Reflect the saved host codes; populated on init even when mode is 'off'.
|
||||
private _roomCode = $state('');
|
||||
private _passCode = $state('');
|
||||
|
||||
private hostTunnel: HostTunnel | null = null;
|
||||
private clientTunnel: ClientTunnel | null = null;
|
||||
// Requests that arrive while mode='client' but tunnel not yet open are held here.
|
||||
private connectionWaiters: Array<{ resolve: () => void; reject: (e: Error) => void }> = [];
|
||||
// The original window.fetch saved before the interceptor is installed.
|
||||
private originalFetch: typeof window.fetch | null = null;
|
||||
|
||||
constructor() {
|
||||
if (browser) {
|
||||
// Load persisted host codes so the UI can show them before host is enabled.
|
||||
const saved = this.readHostCodes();
|
||||
|
||||
if (saved) {
|
||||
this._roomCode = saved.roomCode;
|
||||
this._passCode = saved.passCode;
|
||||
}
|
||||
|
||||
this.restoreSession();
|
||||
}
|
||||
}
|
||||
|
||||
get roomCode(): string {
|
||||
return this._roomCode;
|
||||
}
|
||||
|
||||
get passCode(): string {
|
||||
return this._passCode;
|
||||
}
|
||||
|
||||
// Full 40-char code shared with remote clients.
|
||||
get shareCode(): string {
|
||||
return this._roomCode + this._passCode;
|
||||
}
|
||||
|
||||
get isConnected(): boolean {
|
||||
return this.status === 'connected';
|
||||
}
|
||||
|
||||
get hasHostCodes(): boolean {
|
||||
return this._roomCode !== '' && this._passCode !== '';
|
||||
}
|
||||
|
||||
// -------------------------------------------------------------------------
|
||||
// Host
|
||||
// -------------------------------------------------------------------------
|
||||
|
||||
async startHost(): Promise<void> {
|
||||
if (this.mode !== 'off') return;
|
||||
|
||||
// Reuse the persisted codes; generate once if none exist yet.
|
||||
let roomCode = this._roomCode;
|
||||
let passCode = this._passCode;
|
||||
|
||||
if (!roomCode || !passCode) {
|
||||
roomCode = generateRoomCode();
|
||||
passCode = generatePassCode();
|
||||
this._roomCode = roomCode;
|
||||
this._passCode = passCode;
|
||||
this.writeHostCodes({ passCode, roomCode });
|
||||
}
|
||||
|
||||
await this.activateHost(roomCode, passCode);
|
||||
}
|
||||
|
||||
/** Generate a fresh room + pass code. Restarts the tunnel if currently active. */
|
||||
async regenerateCodes(): Promise<void> {
|
||||
const roomCode = generateRoomCode();
|
||||
const passCode = generatePassCode();
|
||||
|
||||
this._roomCode = roomCode;
|
||||
this._passCode = passCode;
|
||||
this.writeHostCodes({ passCode, roomCode });
|
||||
|
||||
if (this.mode === 'host') {
|
||||
this.hostTunnel?.stop();
|
||||
this.hostTunnel = null;
|
||||
await this.activateHost(roomCode, passCode);
|
||||
}
|
||||
}
|
||||
|
||||
private async activateHost(roomCode: string, passCode: string): Promise<void> {
|
||||
this.mode = 'host';
|
||||
this.status = 'connecting';
|
||||
this.errorMessage = '';
|
||||
this.peerCount = 0;
|
||||
|
||||
const tunnel = new HostTunnel(roomCode, passCode, {
|
||||
onPeerCountChange: (count) => {
|
||||
this.peerCount = count;
|
||||
}
|
||||
});
|
||||
|
||||
try {
|
||||
await tunnel.start();
|
||||
this.hostTunnel = tunnel;
|
||||
this.status = 'connected';
|
||||
this.writeSession({ mode: 'host', passCode, roomCode });
|
||||
} catch (e) {
|
||||
this.hostTunnel = null;
|
||||
this.status = 'error';
|
||||
this.errorMessage = e instanceof Error ? e.message : String(e);
|
||||
}
|
||||
}
|
||||
|
||||
stopHost(): void {
|
||||
this.hostTunnel?.stop();
|
||||
this.hostTunnel = null;
|
||||
this.mode = 'off';
|
||||
this.status = 'idle';
|
||||
this.peerCount = 0;
|
||||
// Codes are intentionally kept: _roomCode/_passCode and HOST_CODES_KEY
|
||||
// remain so the user can re-enable without a new code.
|
||||
this.clearSession();
|
||||
}
|
||||
|
||||
// -------------------------------------------------------------------------
|
||||
// Client
|
||||
// -------------------------------------------------------------------------
|
||||
|
||||
async joinAsClient(shareCode: string): Promise<void> {
|
||||
if (shareCode.length < 40) throw new Error('Invalid code: must be 40 characters');
|
||||
|
||||
const roomCode = shareCode.slice(0, 8);
|
||||
const passCode = shareCode.slice(8);
|
||||
|
||||
await this.activateClient(roomCode, passCode);
|
||||
}
|
||||
|
||||
private async activateClient(roomCode: string, passCode: string): Promise<void> {
|
||||
this.mode = 'client';
|
||||
this.status = 'connecting';
|
||||
this.errorMessage = '';
|
||||
// Install the fetch interceptor synchronously (before any await) so that
|
||||
// requests fired by layout effects on the same tick are already captured.
|
||||
this.installInterceptor();
|
||||
|
||||
const tunnel = new ClientTunnel(roomCode, passCode, {
|
||||
onConnected: () => {
|
||||
this.status = 'connected';
|
||||
},
|
||||
onDisconnected: () => {
|
||||
this.status = 'error';
|
||||
this.errorMessage = 'Disconnected from host';
|
||||
}
|
||||
});
|
||||
|
||||
try {
|
||||
await tunnel.connect();
|
||||
this.clientTunnel = tunnel;
|
||||
this.writeSession({ mode: 'client', passCode, roomCode });
|
||||
// Release any requests that were queued while connecting.
|
||||
const waiters = this.connectionWaiters.splice(0);
|
||||
|
||||
for (const w of waiters) w.resolve();
|
||||
} catch (e) {
|
||||
this.clientTunnel = null;
|
||||
this.mode = 'off';
|
||||
this.status = 'error';
|
||||
this.errorMessage = e instanceof Error ? e.message : String(e);
|
||||
this.uninstallInterceptor();
|
||||
// Reject queued requests.
|
||||
const waiters = this.connectionWaiters.splice(0);
|
||||
const err = e instanceof Error ? e : new Error(String(e));
|
||||
|
||||
for (const w of waiters) w.reject(err);
|
||||
|
||||
throw e;
|
||||
}
|
||||
}
|
||||
|
||||
leaveAsClient(): void {
|
||||
this.uninstallInterceptor();
|
||||
this.clientTunnel?.disconnect();
|
||||
this.clientTunnel = null;
|
||||
this.mode = 'off';
|
||||
this.status = 'idle';
|
||||
this.clearSession();
|
||||
}
|
||||
|
||||
// -------------------------------------------------------------------------
|
||||
// Fetch interceptor (installed synchronously when client mode activates)
|
||||
// -------------------------------------------------------------------------
|
||||
|
||||
private installInterceptor(): void {
|
||||
if (this.originalFetch) return; // already installed
|
||||
|
||||
this.originalFetch = window.fetch.bind(window);
|
||||
window.fetch = (input: RequestInfo | URL, init?: RequestInit) => {
|
||||
try {
|
||||
const url =
|
||||
input instanceof Request ? input.url : input instanceof URL ? input.href : String(input);
|
||||
const parsed = new URL(url, window.location.href);
|
||||
|
||||
if (parsed.origin === window.location.origin) {
|
||||
return this.tunnelFetch(input, init);
|
||||
}
|
||||
} catch {
|
||||
// not a parseable URL — fall through
|
||||
}
|
||||
|
||||
return this.originalFetch!(input, init);
|
||||
};
|
||||
}
|
||||
|
||||
private uninstallInterceptor(): void {
|
||||
if (!this.originalFetch) return;
|
||||
|
||||
window.fetch = this.originalFetch;
|
||||
this.originalFetch = null;
|
||||
}
|
||||
|
||||
// -------------------------------------------------------------------------
|
||||
// Fetch proxy
|
||||
// -------------------------------------------------------------------------
|
||||
|
||||
tunnelFetch(input: RequestInfo | URL, init?: RequestInit): Promise<Response> {
|
||||
// If the tunnel is open, forward immediately.
|
||||
if (this.clientTunnel?.isConnected) {
|
||||
return this.clientTunnel.fetch(input, init);
|
||||
}
|
||||
|
||||
// If we are still connecting, queue the request until the tunnel opens.
|
||||
if (this.mode === 'client' && this.status === 'connecting') {
|
||||
return new Promise<void>((resolve, reject) => {
|
||||
this.connectionWaiters.push({ reject, resolve });
|
||||
}).then(() => this.clientTunnel!.fetch(input, init));
|
||||
}
|
||||
|
||||
throw new Error('tunnel not connected');
|
||||
}
|
||||
|
||||
// -------------------------------------------------------------------------
|
||||
// Persistence helpers
|
||||
// -------------------------------------------------------------------------
|
||||
|
||||
private readHostCodes(): HostCodes | null {
|
||||
try {
|
||||
const raw = localStorage.getItem(HOST_CODES_KEY);
|
||||
|
||||
return raw ? (JSON.parse(raw) as HostCodes) : null;
|
||||
} catch {
|
||||
return null;
|
||||
}
|
||||
}
|
||||
|
||||
private writeHostCodes(codes: HostCodes): void {
|
||||
localStorage.setItem(HOST_CODES_KEY, JSON.stringify(codes));
|
||||
}
|
||||
|
||||
private restoreSession(): void {
|
||||
try {
|
||||
const raw = localStorage.getItem(SESSION_KEY);
|
||||
|
||||
if (!raw) return;
|
||||
|
||||
const session = JSON.parse(raw) as SessionData;
|
||||
|
||||
if (session.mode === 'host') {
|
||||
void this.activateHost(session.roomCode, session.passCode);
|
||||
} else if (session.mode === 'client') {
|
||||
void this.activateClient(session.roomCode, session.passCode);
|
||||
}
|
||||
} catch {
|
||||
// ignore corrupt storage
|
||||
}
|
||||
}
|
||||
|
||||
private writeSession(data: SessionData): void {
|
||||
localStorage.setItem(SESSION_KEY, JSON.stringify(data));
|
||||
}
|
||||
|
||||
private clearSession(): void {
|
||||
localStorage.removeItem(SESSION_KEY);
|
||||
}
|
||||
}
|
||||
|
||||
export const webrtcStore = new WebRTCStore();
|
||||
@@ -1,729 +0,0 @@
|
||||
/**
|
||||
* WebRTC tunnel for remote llama.cpp access.
|
||||
*
|
||||
* Signaling uses the WebTorrent tracker WebSocket protocol (no external deps).
|
||||
* The room code is used as the info_hash rendezvous key; the pass code
|
||||
* authenticates the client on the data channel after the WebRTC handshake.
|
||||
*
|
||||
* Host: announces periodically, accepts incoming offers, relays HTTP requests
|
||||
* made by connected clients back to its own local server.
|
||||
* Client: sends an offer, authenticates via pass code, then all same-origin
|
||||
* fetch calls are transparently forwarded through the data channel.
|
||||
*/
|
||||
|
||||
const STUN_CONFIG: RTCConfiguration = {
|
||||
iceServers: [{ urls: 'stun:stun.l.google.com:19302' }, { urls: 'stun:stun1.l.google.com:19302' }]
|
||||
};
|
||||
const TRACKER_URLS = ['wss://tracker.openwebtorrent.com', 'wss://tracker.btorrent.xyz'];
|
||||
const ICE_GATHER_TIMEOUT_MS = 10_000;
|
||||
const ANNOUNCE_INTERVAL_MS = 30_000;
|
||||
const CONNECT_TIMEOUT_MS = 30_000;
|
||||
const TRACKER_CONNECT_TIMEOUT_MS = 10_000;
|
||||
// Characters that are unambiguous to read aloud or type
|
||||
const CODE_CHARS = 'ABCDEFGHJKMNPQRSTUVWXYZabcdefghjkmnpqrstuvwxyz23456789';
|
||||
|
||||
function randomStr(len: number): string {
|
||||
const bytes = new Uint8Array(len);
|
||||
|
||||
crypto.getRandomValues(bytes);
|
||||
let result = '';
|
||||
|
||||
for (const b of bytes) result += CODE_CHARS[b % CODE_CHARS.length];
|
||||
|
||||
return result;
|
||||
}
|
||||
|
||||
export function generateRoomCode(): string {
|
||||
return randomStr(8);
|
||||
}
|
||||
|
||||
export function generatePassCode(): string {
|
||||
return randomStr(32);
|
||||
}
|
||||
|
||||
// info_hash must be exactly 20 chars for WebTorrent trackers
|
||||
function roomToInfoHash(roomCode: string): string {
|
||||
return roomCode.padEnd(20, '0').slice(0, 20);
|
||||
}
|
||||
|
||||
function waitForIceComplete(pc: RTCPeerConnection): Promise<void> {
|
||||
return new Promise((resolve, reject) => {
|
||||
if (pc.iceGatheringState === 'complete') {
|
||||
resolve();
|
||||
|
||||
return;
|
||||
}
|
||||
|
||||
const timer = setTimeout(
|
||||
() => reject(new Error('ICE gathering timed out')),
|
||||
ICE_GATHER_TIMEOUT_MS
|
||||
);
|
||||
|
||||
pc.addEventListener('icegatheringstatechange', () => {
|
||||
if (pc.iceGatheringState === 'complete') {
|
||||
clearTimeout(timer);
|
||||
resolve();
|
||||
}
|
||||
});
|
||||
});
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Tracker signaling (WebTorrent WS tracker protocol)
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
type TrackerMsg = Record<string, unknown>;
|
||||
|
||||
class Tracker {
|
||||
private ws: WebSocket | null = null;
|
||||
private readonly infoHash: string;
|
||||
private readonly peerId: string;
|
||||
|
||||
onOffer?: (fromPeerId: string, offerId: string, offer: RTCSessionDescriptionInit) => void;
|
||||
onAnswer?: (offerId: string, answer: RTCSessionDescriptionInit) => void;
|
||||
onClose?: () => void;
|
||||
|
||||
constructor(infoHash: string, peerId: string) {
|
||||
this.infoHash = infoHash;
|
||||
this.peerId = peerId;
|
||||
}
|
||||
|
||||
connect(url: string): Promise<void> {
|
||||
return new Promise((resolve, reject) => {
|
||||
const ws = new WebSocket(url);
|
||||
|
||||
this.ws = ws;
|
||||
const timer = setTimeout(
|
||||
() => reject(new Error('tracker connect timeout')),
|
||||
TRACKER_CONNECT_TIMEOUT_MS
|
||||
);
|
||||
|
||||
ws.onopen = () => {
|
||||
clearTimeout(timer);
|
||||
resolve();
|
||||
};
|
||||
ws.onerror = () => {
|
||||
clearTimeout(timer);
|
||||
reject(new Error('tracker WebSocket error'));
|
||||
};
|
||||
ws.onclose = () => {
|
||||
this.onClose?.();
|
||||
};
|
||||
ws.onmessage = (event) => {
|
||||
try {
|
||||
const msg = JSON.parse(event.data as string) as TrackerMsg;
|
||||
|
||||
if (msg.offer && msg.peer_id && msg.offer_id) {
|
||||
this.onOffer?.(
|
||||
msg.peer_id as string,
|
||||
msg.offer_id as string,
|
||||
msg.offer as RTCSessionDescriptionInit
|
||||
);
|
||||
} else if (msg.answer && msg.offer_id) {
|
||||
this.onAnswer?.(msg.offer_id as string, msg.answer as RTCSessionDescriptionInit);
|
||||
}
|
||||
} catch {
|
||||
// ignore malformed tracker messages
|
||||
}
|
||||
};
|
||||
});
|
||||
}
|
||||
|
||||
private send(msg: TrackerMsg): void {
|
||||
if (this.ws?.readyState === WebSocket.OPEN) {
|
||||
this.ws.send(JSON.stringify(msg));
|
||||
}
|
||||
}
|
||||
|
||||
announce(
|
||||
opts: {
|
||||
numwant?: number;
|
||||
offers?: Array<{ offer_id: string; offer: RTCSessionDescriptionInit }>;
|
||||
} = {}
|
||||
): void {
|
||||
const msg: TrackerMsg = {
|
||||
action: 'announce',
|
||||
info_hash: this.infoHash,
|
||||
numwant: opts.numwant ?? 0,
|
||||
peer_id: this.peerId
|
||||
};
|
||||
|
||||
if (opts.offers) msg.offers = opts.offers;
|
||||
|
||||
this.send(msg);
|
||||
}
|
||||
|
||||
sendAnswer(toPeerId: string, offerId: string, answer: RTCSessionDescriptionInit): void {
|
||||
this.send({
|
||||
action: 'announce',
|
||||
answer,
|
||||
info_hash: this.infoHash,
|
||||
offer_id: offerId,
|
||||
peer_id: this.peerId,
|
||||
to_peer_id: toPeerId
|
||||
});
|
||||
}
|
||||
|
||||
close(): void {
|
||||
this.ws?.close();
|
||||
this.ws = null;
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Tunnel message types (JSON, sent over RTCDataChannel)
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
interface ReqMsg {
|
||||
type: 'req';
|
||||
id: string;
|
||||
method: string;
|
||||
path: string;
|
||||
headers: Record<string, string>;
|
||||
body: string | null; // base64 or null
|
||||
}
|
||||
interface ResStartMsg {
|
||||
type: 'res_start';
|
||||
id: string;
|
||||
status: number;
|
||||
headers: Record<string, string>;
|
||||
}
|
||||
interface ResChunkMsg {
|
||||
type: 'res_chunk';
|
||||
id: string;
|
||||
data: string; // base64
|
||||
}
|
||||
interface ResEndMsg {
|
||||
type: 'res_end';
|
||||
id: string;
|
||||
}
|
||||
interface ResErrMsg {
|
||||
type: 'res_err';
|
||||
id: string;
|
||||
message: string;
|
||||
}
|
||||
interface CancelMsg {
|
||||
type: 'cancel';
|
||||
id: string;
|
||||
}
|
||||
interface AuthMsg {
|
||||
type: 'auth';
|
||||
pass: string;
|
||||
}
|
||||
interface AuthOkMsg {
|
||||
type: 'auth_ok';
|
||||
}
|
||||
interface AuthFailMsg {
|
||||
type: 'auth_fail';
|
||||
}
|
||||
|
||||
type TunnelMsg =
|
||||
| ReqMsg
|
||||
| ResStartMsg
|
||||
| ResChunkMsg
|
||||
| ResEndMsg
|
||||
| ResErrMsg
|
||||
| CancelMsg
|
||||
| AuthMsg
|
||||
| AuthOkMsg
|
||||
| AuthFailMsg;
|
||||
|
||||
// Max bytes per res_chunk message (keeps data channel messages well below limits)
|
||||
const CHUNK_BYTES = 8192;
|
||||
|
||||
function uint8ToBase64(bytes: Uint8Array): string {
|
||||
let binary = '';
|
||||
|
||||
for (let i = 0; i < bytes.length; i++) binary += String.fromCharCode(bytes[i]);
|
||||
|
||||
return btoa(binary);
|
||||
}
|
||||
|
||||
function base64ToUint8(b64: string): Uint8Array {
|
||||
return Uint8Array.from(atob(b64), (c) => c.charCodeAt(0));
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// HostTunnel
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
export type HostCallbacks = {
|
||||
onPeerCountChange?: (count: number) => void;
|
||||
};
|
||||
|
||||
export class HostTunnel {
|
||||
private readonly passCode: string;
|
||||
private readonly infoHash: string;
|
||||
private readonly peerId: string;
|
||||
private readonly callbacks: HostCallbacks;
|
||||
|
||||
private trackers: Tracker[] = [];
|
||||
private peers = new Map<string, { pc: RTCPeerConnection; channel: RTCDataChannel }>();
|
||||
// AbortControllers for in-flight host-side fetch() calls, keyed by request id.
|
||||
private activeRequests = new Map<string, AbortController>();
|
||||
private announceTimer: ReturnType<typeof setInterval> | null = null;
|
||||
private stopped = false;
|
||||
|
||||
constructor(roomCode: string, passCode: string, callbacks: HostCallbacks = {}) {
|
||||
this.passCode = passCode;
|
||||
this.infoHash = roomToInfoHash(roomCode);
|
||||
this.peerId = randomStr(20);
|
||||
this.callbacks = callbacks;
|
||||
}
|
||||
|
||||
get peerCount(): number {
|
||||
return this.peers.size;
|
||||
}
|
||||
|
||||
async start(): Promise<void> {
|
||||
await this.connectTrackers();
|
||||
this.announceTimer = setInterval(() => {
|
||||
for (const t of this.trackers) t.announce({ numwant: 0 });
|
||||
}, ANNOUNCE_INTERVAL_MS);
|
||||
}
|
||||
|
||||
private async connectTrackers(): Promise<void> {
|
||||
for (const url of TRACKER_URLS) {
|
||||
try {
|
||||
await this.connectOneTracker(url);
|
||||
} catch {
|
||||
// try next
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
private async connectOneTracker(url: string): Promise<void> {
|
||||
const tracker = new Tracker(this.infoHash, this.peerId);
|
||||
|
||||
tracker.onOffer = (fromPeerId, offerId, offer) => {
|
||||
void this.handleOffer(tracker, fromPeerId, offerId, offer);
|
||||
};
|
||||
tracker.onClose = () => {
|
||||
this.trackers = this.trackers.filter((t) => t !== tracker);
|
||||
|
||||
if (!this.stopped) {
|
||||
setTimeout(() => void this.connectOneTracker(url), 5000);
|
||||
}
|
||||
};
|
||||
await tracker.connect(url);
|
||||
tracker.announce({ numwant: 0 });
|
||||
this.trackers.push(tracker);
|
||||
}
|
||||
|
||||
private async handleOffer(
|
||||
tracker: Tracker,
|
||||
fromPeerId: string,
|
||||
offerId: string,
|
||||
offer: RTCSessionDescriptionInit
|
||||
): Promise<void> {
|
||||
try {
|
||||
const pc = new RTCPeerConnection(STUN_CONFIG);
|
||||
|
||||
await pc.setRemoteDescription(new RTCSessionDescription(offer));
|
||||
const answer = await pc.createAnswer();
|
||||
|
||||
await pc.setLocalDescription(answer);
|
||||
await waitForIceComplete(pc);
|
||||
tracker.sendAnswer(fromPeerId, offerId, pc.localDescription!);
|
||||
|
||||
pc.ondatachannel = (event) => this.setupChannel(pc, fromPeerId, event.channel);
|
||||
} catch {
|
||||
// ignore failed handshakes
|
||||
}
|
||||
}
|
||||
|
||||
private setupChannel(pc: RTCPeerConnection, peerId: string, channel: RTCDataChannel): void {
|
||||
let authenticated = false;
|
||||
|
||||
channel.onclose = () => {
|
||||
if (this.peers.has(peerId)) {
|
||||
this.peers.delete(peerId);
|
||||
this.callbacks.onPeerCountChange?.(this.peers.size);
|
||||
}
|
||||
|
||||
pc.close();
|
||||
};
|
||||
|
||||
channel.onmessage = (event) => {
|
||||
try {
|
||||
const msg = JSON.parse(event.data as string) as TunnelMsg;
|
||||
|
||||
if (!authenticated) {
|
||||
if (msg.type === 'auth') {
|
||||
if (msg.pass === this.passCode) {
|
||||
authenticated = true;
|
||||
channel.send(JSON.stringify({ type: 'auth_ok' } satisfies AuthOkMsg));
|
||||
this.peers.set(peerId, { channel, pc });
|
||||
this.callbacks.onPeerCountChange?.(this.peers.size);
|
||||
} else {
|
||||
channel.send(JSON.stringify({ type: 'auth_fail' } satisfies AuthFailMsg));
|
||||
channel.close();
|
||||
}
|
||||
}
|
||||
|
||||
return;
|
||||
}
|
||||
|
||||
if (msg.type === 'req') {
|
||||
void this.handleRequest(channel, msg);
|
||||
} else if (msg.type === 'cancel') {
|
||||
this.activeRequests.get(msg.id)?.abort();
|
||||
}
|
||||
} catch {
|
||||
// ignore malformed messages
|
||||
}
|
||||
};
|
||||
}
|
||||
|
||||
private async handleRequest(channel: RTCDataChannel, msg: ReqMsg): Promise<void> {
|
||||
const { body, headers, id, method, path } = msg;
|
||||
const t0 = performance.now();
|
||||
const ac = new AbortController();
|
||||
|
||||
this.activeRequests.set(id, ac);
|
||||
|
||||
try {
|
||||
const init: RequestInit = { headers, method, signal: ac.signal };
|
||||
|
||||
if (body !== null) init.body = base64ToUint8(body).buffer as ArrayBuffer;
|
||||
|
||||
const response = await fetch(path, init);
|
||||
const resHeaders: Record<string, string> = {};
|
||||
|
||||
response.headers.forEach((v, k) => {
|
||||
resHeaders[k] = v;
|
||||
});
|
||||
|
||||
console.log(
|
||||
`[rtc] ${method} ${path} -> ${response.status} (${Math.round(performance.now() - t0)}ms)`
|
||||
);
|
||||
|
||||
channel.send(
|
||||
JSON.stringify({
|
||||
headers: resHeaders,
|
||||
id,
|
||||
status: response.status,
|
||||
type: 'res_start'
|
||||
} satisfies ResStartMsg)
|
||||
);
|
||||
|
||||
const reader = response.body?.getReader();
|
||||
|
||||
if (reader) {
|
||||
while (true) {
|
||||
const { done, value } = await reader.read();
|
||||
|
||||
if (done) break;
|
||||
|
||||
for (let i = 0; i < value.length; i += CHUNK_BYTES) {
|
||||
const slice = value.subarray(i, i + CHUNK_BYTES);
|
||||
|
||||
channel.send(
|
||||
JSON.stringify({
|
||||
data: uint8ToBase64(slice),
|
||||
id,
|
||||
type: 'res_chunk'
|
||||
} satisfies ResChunkMsg)
|
||||
);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
channel.send(JSON.stringify({ id, type: 'res_end' } satisfies ResEndMsg));
|
||||
} catch (e) {
|
||||
// AbortError means the client cancelled — no need to send an error back.
|
||||
if (!(e instanceof DOMException && e.name === 'AbortError')) {
|
||||
const errMsg = e instanceof Error ? e.message : String(e);
|
||||
|
||||
console.error(`[rtc] ${method} ${path} -> error: ${errMsg}`);
|
||||
channel.send(JSON.stringify({ id, message: errMsg, type: 'res_err' } satisfies ResErrMsg));
|
||||
}
|
||||
} finally {
|
||||
this.activeRequests.delete(id);
|
||||
}
|
||||
}
|
||||
|
||||
stop(): void {
|
||||
this.stopped = true;
|
||||
|
||||
if (this.announceTimer) clearInterval(this.announceTimer);
|
||||
|
||||
for (const t of this.trackers) t.close();
|
||||
for (const { channel, pc } of this.peers.values()) {
|
||||
channel.close();
|
||||
pc.close();
|
||||
}
|
||||
for (const ac of this.activeRequests.values()) ac.abort();
|
||||
this.trackers = [];
|
||||
this.peers.clear();
|
||||
this.activeRequests.clear();
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// ClientTunnel
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
export type ClientCallbacks = {
|
||||
onConnected?: () => void;
|
||||
onDisconnected?: () => void;
|
||||
};
|
||||
|
||||
type PendingReq = {
|
||||
onStart: (status: number, headers: Record<string, string>) => void;
|
||||
onChunk: (data: string) => void;
|
||||
onEnd: () => void;
|
||||
onError: (message: string) => void;
|
||||
};
|
||||
|
||||
export class ClientTunnel {
|
||||
private readonly passCode: string;
|
||||
private readonly infoHash: string;
|
||||
private readonly peerId: string;
|
||||
private readonly callbacks: ClientCallbacks;
|
||||
|
||||
private pc: RTCPeerConnection | null = null;
|
||||
private channel: RTCDataChannel | null = null;
|
||||
private tracker: Tracker | null = null;
|
||||
private pending = new Map<string, PendingReq>();
|
||||
|
||||
constructor(roomCode: string, passCode: string, callbacks: ClientCallbacks = {}) {
|
||||
this.passCode = passCode;
|
||||
this.infoHash = roomToInfoHash(roomCode);
|
||||
this.peerId = randomStr(20);
|
||||
this.callbacks = callbacks;
|
||||
}
|
||||
|
||||
get isConnected(): boolean {
|
||||
return this.channel?.readyState === 'open';
|
||||
}
|
||||
|
||||
async connect(): Promise<void> {
|
||||
let lastError: Error = new Error('no trackers available');
|
||||
|
||||
for (const url of TRACKER_URLS) {
|
||||
try {
|
||||
await this.connectViaTracker(url);
|
||||
|
||||
return;
|
||||
} catch (e) {
|
||||
lastError = e instanceof Error ? e : new Error(String(e));
|
||||
this.cleanupConnection();
|
||||
}
|
||||
}
|
||||
|
||||
throw lastError;
|
||||
}
|
||||
|
||||
private async connectViaTracker(trackerUrl: string): Promise<void> {
|
||||
const offerId = randomStr(20);
|
||||
const pc = new RTCPeerConnection(STUN_CONFIG);
|
||||
|
||||
this.pc = pc;
|
||||
|
||||
const channel = pc.createDataChannel('tunnel', { ordered: true });
|
||||
|
||||
this.channel = channel;
|
||||
|
||||
const offer = await pc.createOffer();
|
||||
|
||||
await pc.setLocalDescription(offer);
|
||||
await waitForIceComplete(pc);
|
||||
|
||||
const tracker = new Tracker(this.infoHash, this.peerId);
|
||||
|
||||
this.tracker = tracker;
|
||||
await tracker.connect(trackerUrl);
|
||||
|
||||
return new Promise((resolve, reject) => {
|
||||
const timer = setTimeout(() => {
|
||||
reject(new Error('connection timed out waiting for host'));
|
||||
}, CONNECT_TIMEOUT_MS);
|
||||
|
||||
tracker.onAnswer = async (_offerId, answer) => {
|
||||
if (_offerId !== offerId) return;
|
||||
|
||||
try {
|
||||
await pc.setRemoteDescription(new RTCSessionDescription(answer));
|
||||
} catch (e) {
|
||||
clearTimeout(timer);
|
||||
reject(e);
|
||||
}
|
||||
};
|
||||
|
||||
channel.onopen = () => {
|
||||
channel.send(JSON.stringify({ pass: this.passCode, type: 'auth' } satisfies AuthMsg));
|
||||
};
|
||||
|
||||
channel.onmessage = (event) => {
|
||||
try {
|
||||
const msg = JSON.parse(event.data as string) as TunnelMsg;
|
||||
|
||||
if (msg.type === 'auth_ok') {
|
||||
clearTimeout(timer);
|
||||
this.callbacks.onConnected?.();
|
||||
resolve();
|
||||
} else if (msg.type === 'auth_fail') {
|
||||
clearTimeout(timer);
|
||||
reject(new Error('authentication failed: invalid passcode'));
|
||||
} else {
|
||||
this.routeResponseMsg(msg);
|
||||
}
|
||||
} catch {
|
||||
// ignore
|
||||
}
|
||||
};
|
||||
|
||||
channel.onclose = () => {
|
||||
this.callbacks.onDisconnected?.();
|
||||
this.rejectAllPending('connection closed');
|
||||
};
|
||||
|
||||
channel.onerror = () => {
|
||||
clearTimeout(timer);
|
||||
reject(new Error('data channel error'));
|
||||
};
|
||||
|
||||
tracker.announce({
|
||||
numwant: 1,
|
||||
offers: [{ offer: pc.localDescription!, offer_id: offerId }]
|
||||
});
|
||||
});
|
||||
}
|
||||
|
||||
private routeResponseMsg(msg: TunnelMsg): void {
|
||||
if (
|
||||
msg.type !== 'res_start' &&
|
||||
msg.type !== 'res_chunk' &&
|
||||
msg.type !== 'res_end' &&
|
||||
msg.type !== 'res_err'
|
||||
)
|
||||
return;
|
||||
|
||||
const req = this.pending.get(msg.id);
|
||||
|
||||
if (!req) return;
|
||||
|
||||
if (msg.type === 'res_start') {
|
||||
req.onStart(msg.status, msg.headers);
|
||||
} else if (msg.type === 'res_chunk') {
|
||||
req.onChunk(msg.data);
|
||||
} else if (msg.type === 'res_end') {
|
||||
req.onEnd();
|
||||
this.pending.delete(msg.id);
|
||||
} else if (msg.type === 'res_err') {
|
||||
req.onError(msg.message);
|
||||
this.pending.delete(msg.id);
|
||||
}
|
||||
}
|
||||
|
||||
private rejectAllPending(reason: string): void {
|
||||
for (const req of this.pending.values()) req.onError(reason);
|
||||
this.pending.clear();
|
||||
}
|
||||
|
||||
private cleanupConnection(): void {
|
||||
this.channel?.close();
|
||||
this.pc?.close();
|
||||
this.tracker?.close();
|
||||
this.channel = null;
|
||||
this.pc = null;
|
||||
this.tracker = null;
|
||||
}
|
||||
|
||||
async fetch(input: RequestInfo | URL, init?: RequestInit): Promise<Response> {
|
||||
if (!this.channel || this.channel.readyState !== 'open') {
|
||||
throw new Error('tunnel not connected');
|
||||
}
|
||||
|
||||
const request = new Request(input, init);
|
||||
const signal = init?.signal ?? (input instanceof Request ? input.signal : undefined);
|
||||
const id = randomStr(16);
|
||||
|
||||
if (signal?.aborted) {
|
||||
return Promise.reject(new DOMException('Aborted', 'AbortError'));
|
||||
}
|
||||
|
||||
// Extract path+query so the host fetches relative to its own origin
|
||||
const reqUrl = new URL(request.url);
|
||||
const path = reqUrl.pathname + reqUrl.search;
|
||||
const headers: Record<string, string> = {};
|
||||
|
||||
request.headers.forEach((v, k) => {
|
||||
headers[k] = v;
|
||||
});
|
||||
|
||||
let bodyB64: string | null = null;
|
||||
|
||||
const bodyBytes = await request.arrayBuffer();
|
||||
|
||||
if (bodyBytes.byteLength > 0) {
|
||||
bodyB64 = uint8ToBase64(new Uint8Array(bodyBytes));
|
||||
}
|
||||
|
||||
return new Promise((resolve, reject) => {
|
||||
let streamController!: ReadableStreamDefaultController<Uint8Array>;
|
||||
|
||||
const stream = new ReadableStream<Uint8Array>({
|
||||
start(ctrl) {
|
||||
streamController = ctrl;
|
||||
}
|
||||
});
|
||||
const abortHandler = () => {
|
||||
this.pending.delete(id);
|
||||
try {
|
||||
streamController.error(new DOMException('Aborted', 'AbortError'));
|
||||
} catch {
|
||||
// stream may already be closed
|
||||
}
|
||||
reject(new DOMException('Aborted', 'AbortError'));
|
||||
|
||||
// Tell the host to stop the in-flight fetch
|
||||
if (this.channel?.readyState === 'open') {
|
||||
this.channel.send(JSON.stringify({ id, type: 'cancel' } satisfies CancelMsg));
|
||||
}
|
||||
};
|
||||
|
||||
signal?.addEventListener('abort', abortHandler, { once: true });
|
||||
|
||||
this.pending.set(id, {
|
||||
onChunk: (data) => {
|
||||
streamController.enqueue(base64ToUint8(data));
|
||||
},
|
||||
onEnd: () => {
|
||||
signal?.removeEventListener('abort', abortHandler);
|
||||
streamController.close();
|
||||
},
|
||||
onError: (message) => {
|
||||
signal?.removeEventListener('abort', abortHandler);
|
||||
try {
|
||||
streamController.error(new Error(message));
|
||||
} catch {
|
||||
// stream may already be closed
|
||||
}
|
||||
reject(new Error(message));
|
||||
this.pending.delete(id);
|
||||
},
|
||||
onStart: (status, resHeaders) => {
|
||||
resolve(new Response(stream, { headers: resHeaders, status }));
|
||||
}
|
||||
});
|
||||
|
||||
this.channel!.send(
|
||||
JSON.stringify({
|
||||
body: bodyB64,
|
||||
headers,
|
||||
id,
|
||||
method: request.method,
|
||||
path,
|
||||
type: 'req'
|
||||
} satisfies ReqMsg)
|
||||
);
|
||||
});
|
||||
}
|
||||
|
||||
disconnect(): void {
|
||||
this.rejectAllPending('disconnected');
|
||||
this.cleanupConnection();
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user