diff --git a/.devops/intel.Dockerfile b/.devops/intel.Dockerfile index b4bcd94b9264..db46fd868d9b 100644 --- a/.devops/intel.Dockerfile +++ b/.devops/intel.Dockerfile @@ -1,4 +1,4 @@ -ARG ONEAPI_VERSION=2025.3.3-0-devel-ubuntu24.04 +ARG ONEAPI_VERSION=2026.1.1-devel-ubuntu24.04 ARG BUILD_DATE=N/A ARG APP_VERSION=N/A ARG APP_REVISION=N/A @@ -19,7 +19,7 @@ RUN npm ci COPY tools/ui/ ./ RUN LLAMA_BUILD_NUMBER="$APP_VERSION" npm run build -FROM docker.io/intel/deep-learning-essentials:$ONEAPI_VERSION AS build +FROM docker.io/intel/oneapi-toolkit:$ONEAPI_VERSION AS build ARG GGML_SYCL_F16=ON ARG LEVEL_ZERO_VERSION=1.28.2 @@ -59,7 +59,7 @@ RUN mkdir -p /app/full \ && cp requirements.txt /app/full \ && cp .devops/tools.sh /app/full/tools.sh -FROM docker.io/intel/deep-learning-essentials:$ONEAPI_VERSION AS base +FROM docker.io/intel/oneapi-toolkit:$ONEAPI_VERSION AS base ARG BUILD_DATE=N/A ARG APP_VERSION=N/A diff --git a/.devops/musa.Dockerfile b/.devops/musa.Dockerfile index 33ac6bb82d49..0e3f6335972d 100644 --- a/.devops/musa.Dockerfile +++ b/.devops/musa.Dockerfile @@ -1,9 +1,9 @@ ARG UBUNTU_VERSION=22.04 # This needs to generally match the container host's environment. # Target the MUSA build image -ARG BASE_MUSA_DEV_CONTAINER=registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu${UBUNTU_VERSION}-amd64 +ARG BASE_MUSA_DEV_CONTAINER=registry.mthreads.com/mcconline/musa_sdk:5.2.0-devel-ubuntu${UBUNTU_VERSION}-s5000 -ARG BASE_MUSA_RUN_CONTAINER=${BASE_MUSA_DEV_CONTAINER} +ARG BASE_MUSA_RUN_CONTAINER=registry.mthreads.com/mcconline/musa_sdk:5.2.0-runtime-ubuntu${UBUNTU_VERSION}-s5000 ARG BUILD_DATE=N/A ARG APP_VERSION=N/A @@ -36,7 +36,10 @@ RUN apt-get update && \ python3-pip \ git \ libssl-dev \ - libgomp1 + libgomp1 \ + musa-mualg-5-2 \ + musa-muthrust-5-2 \ + libmthreads-compute WORKDIR /app @@ -79,13 +82,16 @@ LABEL org.opencontainers.image.created=$BUILD_DATE \ org.opencontainers.image.source=$IMAGE_SOURCE RUN apt-get update \ - && apt-get install -y libgomp1 curl ffmpeg \ + && apt-get install -y libgomp1 curl ffmpeg libmthreads-compute \ && apt autoremove -y \ && apt clean -y \ && rm -rf /tmp/* /var/tmp/* \ && find /var/cache/apt/archives /var/lib/apt/lists -not -name lock -type f -delete \ && find /var/cache -type f -delete +# The MUSA runtime image does not register its library directory +RUN echo "/usr/local/musa/lib" > /etc/ld.so.conf.d/musa-runtime.conf && ldconfig + COPY --from=build /app/lib/ /app ### Full diff --git a/.github/workflows/build-cuda-ubuntu.yml b/.github/workflows/build-cuda-ubuntu.yml index 3dc9255be327..68b6c01900bf 100644 --- a/.github/workflows/build-cuda-ubuntu.yml +++ b/.github/workflows/build-cuda-ubuntu.yml @@ -145,7 +145,7 @@ jobs: musa: runs-on: ubuntu-22.04 - container: registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64 + container: registry.mthreads.com/mcconline/musa_sdk:5.2.0-devel-ubuntu22.04-s5000 steps: - name: Clone @@ -156,7 +156,7 @@ jobs: id: depends run: | apt-get update - apt-get install -y build-essential git cmake libssl-dev jq python3-venv + apt-get install -y build-essential git cmake libssl-dev jq python3-venv musa-mualg-5-2 musa-muthrust-5-2 libmthreads-compute - name: ccache uses: ggml-org/ccache-action@v1.2.24 diff --git a/.github/workflows/build-sycl.yml b/.github/workflows/build-sycl.yml index 5424a8607e58..ddc4e7f2b1a1 100644 --- a/.github/workflows/build-sycl.yml +++ b/.github/workflows/build-sycl.yml @@ -48,7 +48,7 @@ jobs: env: ONEAPI_ROOT: /opt/intel/oneapi/ - ONEAPI_INSTALLER_VERSION: "2025.3.3" + ONEAPI_INSTALLER_VERSION: "2026.1" LEVEL_ZERO_VERSION: "1.33.1" LEVEL_ZERO_UBUNTU_VERSION: "u24.04" @@ -63,8 +63,8 @@ jobs: shell: bash run: | cd /tmp - wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/56f7923a-adb8-43f3-8b02-2b60fcac8cab/intel-deep-learning-essentials-2025.3.3.16_offline.sh -O intel-deep-learning-essentials_offline.sh - sudo bash intel-deep-learning-essentials_offline.sh -s -a --silent --eula accept + wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/5996e26b-f48a-42b1-8db0-b002ad0bd8d7/intel-oneapi-toolkit-2026.1.1.33_offline.sh -O intel-oneapi-toolkit_offline.sh + sudo bash intel-oneapi-toolkit_offline.sh -s -a --silent --eula accept - name: Install Level Zero SDK shell: bash @@ -129,11 +129,11 @@ jobs: shell: bash env: - WINDOWS_BASEKIT_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b60765d1-2b85-4e85-86b6-cb0e9563a699/intel-deep-learning-essentials-2025.3.3.18_offline.exe + WINDOWS_BASEKIT_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/0cb67a0d-67f6-410b-868b-f4a0a17ff0cf/intel-oneapi-toolkit-2026.1.1.32_offline.exe WINDOWS_DPCPP_MKL: intel.oneapi.win.cpp-dpcpp-common:intel.oneapi.win.mkl.devel:intel.oneapi.win.dnnl:intel.oneapi.win.tbb.devel LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.33.1/level-zero-win-sdk-1.33.1.zip ONEAPI_ROOT: "C:/Program Files (x86)/Intel/oneAPI" - ONEAPI_INSTALLER_VERSION: "2025.3.3" + ONEAPI_INSTALLER_VERSION: "2026.1" steps: - name: Clone id: checkout diff --git a/.github/workflows/release.yml b/.github/workflows/release.yml index be13d9118a46..b5dba37de8fb 100644 --- a/.github/workflows/release.yml +++ b/.github/workflows/release.yml @@ -1344,11 +1344,11 @@ jobs: shell: bash env: - WINDOWS_BASEKIT_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/b60765d1-2b85-4e85-86b6-cb0e9563a699/intel-deep-learning-essentials-2025.3.3.18_offline.exe + WINDOWS_BASEKIT_URL: https://registrationcenter-download.intel.com/akdlm/IRC_NAS/0cb67a0d-67f6-410b-868b-f4a0a17ff0cf/intel-oneapi-toolkit-2026.1.1.32_offline.exe WINDOWS_DPCPP_MKL: intel.oneapi.win.cpp-dpcpp-common:intel.oneapi.win.mkl.devel:intel.oneapi.win.dnnl:intel.oneapi.win.tbb.devel - LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.28.2/level-zero-win-sdk-1.28.2.zip + LEVEL_ZERO_SDK_URL: https://github.com/oneapi-src/level-zero/releases/download/v1.33.1/level-zero-win-sdk-1.33.1.zip ONEAPI_ROOT: "C:/Program Files (x86)/Intel/oneAPI" - ONEAPI_INSTALLER_VERSION: "2025.3.3" + ONEAPI_INSTALLER_VERSION: "2026.1" steps: - name: Clone @@ -1391,9 +1391,11 @@ jobs: run: | echo "cp oneAPI running time dll files in ${{ env.ONEAPI_ROOT }} to ./build/bin" - cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_sycl_blas.5.dll" ./build/bin - cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_core.2.dll" ./build/bin - cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_tbb_thread.2.dll" ./build/bin + cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_sycl_blas.6.dll" ./build/bin + cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_core.3.dll" ./build/bin + cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_def.3.dll" ./build/bin + cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_avx2.3.dll" ./build/bin + cp "${{ env.ONEAPI_ROOT }}/mkl/latest/bin/mkl_tbb_thread.3.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_adapter_level_zero.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/ur_adapter_level_zero_v2.dll" ./build/bin @@ -1408,13 +1410,11 @@ jobs: echo "Level Zero loader DLL not found in oneAPI or SDK; relying on system driver/runtime" fi - cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/sycl8.dll" ./build/bin + cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/sycl9.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/svml_dispmd.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/libmmd.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/libiomp5md.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/sycl-ls.exe" ./build/bin - cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/libsycl-fallback-bfloat16.spv" ./build/bin - cp "${{ env.ONEAPI_ROOT }}/compiler/latest/bin/libsycl-native-bfloat16.spv" ./build/bin cp "${{ env.ONEAPI_ROOT }}/dnnl/latest/bin/dnnl.dll" ./build/bin cp "${{ env.ONEAPI_ROOT }}/tbb/latest/bin/tbb12.dll" ./build/bin @@ -1454,8 +1454,8 @@ jobs: env: ONEAPI_ROOT: /opt/intel/oneapi/ - ONEAPI_INSTALLER_VERSION: "2025.3.3" - LEVEL_ZERO_VERSION: "1.28.2" + ONEAPI_INSTALLER_VERSION: "2026.1" + LEVEL_ZERO_VERSION: "1.33.1" LEVEL_ZERO_UBUNTU_VERSION: "u24.04" steps: @@ -1469,16 +1469,16 @@ jobs: shell: bash run: | cd /tmp - wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/56f7923a-adb8-43f3-8b02-2b60fcac8cab/intel-deep-learning-essentials-2025.3.3.16_offline.sh -O intel-deep-learning-essentials_offline.sh - sudo bash intel-deep-learning-essentials_offline.sh -s -a --silent --eula accept + wget https://registrationcenter-download.intel.com/akdlm/IRC_NAS/5996e26b-f48a-42b1-8db0-b002ad0bd8d7/intel-oneapi-toolkit-2026.1.1.33_offline.sh -O intel-oneapi-toolkit_offline.sh + sudo bash intel-oneapi-toolkit_offline.sh -s -a --silent --eula accept - name: Install Level Zero SDK shell: bash run: | cd /tmp - wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/level-zero_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O level-zero.deb - wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/level-zero-devel_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O level-zero-devel.deb - sudo apt-get install -y ./level-zero.deb ./level-zero-devel.deb + wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/libze1_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O libze1.deb + wget -q "https://github.com/oneapi-src/level-zero/releases/download/v${LEVEL_ZERO_VERSION}/libze-dev_${LEVEL_ZERO_VERSION}%2B${LEVEL_ZERO_UBUNTU_VERSION}_amd64.deb" -O libze-dev.deb + sudo apt-get install -y ./libze1.deb ./libze-dev.deb - name: Download UI build uses: actions/download-artifact@v7 diff --git a/ci/README-MUSA.md b/ci/README-MUSA.md index 2101dfa2c09b..40e9ea5af5ed 100644 --- a/ci/README-MUSA.md +++ b/ci/README-MUSA.md @@ -21,13 +21,13 @@ docker run --privileged -it \ -v $HOME/llama.cpp/ci-cache:/ci-cache \ -v $HOME/llama.cpp/ci-results:/ci-results \ -v $PWD:/ws -w /ws \ - registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64 + registry.mthreads.com/mcconline/musa_sdk:5.2.0-devel-ubuntu22.04-s5000 ``` Inside the container, execute the following commands: ```bash -apt update -y && apt install -y bc cmake ccache git python3.10-venv time unzip wget +apt update -y && apt install -y bc cmake ccache git python3.10-venv time unzip wget musa-mualg-5-2 musa-muthrust-5-2 libmthreads-compute git config --global --add safe.directory /ws GG_BUILD_MUSA=1 bash ./ci/run.sh /ci-results /ci-cache ``` diff --git a/ci/run.sh b/ci/run.sh index fc661f9767b9..ccfc0562f4ab 100755 --- a/ci/run.sh +++ b/ci/run.sh @@ -49,14 +49,6 @@ mkdir -p "$2" OUT=$(realpath "$1") MNT=$(realpath "$2") -# gpu-rocm self-hosted runner can't upload logs to blob; keep each run's logs in -# their own dir keyed by the GitHub run id so an Actions run URL maps to its logs. -if [ -n "${GG_BUILD_ROCM}" ] && [ -n "${GITHUB_RUN_ID}" ]; then - OUT="$OUT/run-${GITHUB_RUN_ID}-${GITHUB_RUN_ATTEMPT:-1}" - mkdir -p "$OUT" - echo "ci results dir: $OUT" -fi - rm -f $OUT/*.log sd=`dirname $0` diff --git a/common/arg.cpp b/common/arg.cpp index 2acac3b5cfbd..b70da5e9e6e0 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -351,7 +351,7 @@ static bool parse_bool_value(const std::string & value) { static std::string get_default_local_path(const std::string & url) { auto f = string_split(url, '#').front(); f = string_split(f, '?').front(); - return fs_get_cache_file(string_split(f, '/').back()); + return fs_path_to_utf8(fs_get_cache_file(string_split(f, '/').back())); } static bool spec_types_is_default(const common_params & params) { diff --git a/common/chat.cpp b/common/chat.cpp index ed1942e15349..3502b96eb34f 100644 --- a/common/chat.cpp +++ b/common/chat.cpp @@ -1099,6 +1099,12 @@ std::optional common_chat_try_specialized_template( return common_chat_params_init_ministral_3(tmpl, params); } + // LLM-jp-4.1 - GPT-OSS dialect (spaces after special tokens, <|end|>-separated parallel calls) + if (src.find("chat_format=llm-jp-harmony-v1") != std::string::npos) { + LOG_DBG("Using specialized template: LLM-jp Harmony v1\n"); + return common_chat_params_init_llm_jp_harmony(tmpl, params); + } + // GPT-OSS - has unique channel-based structure that needs dedicated handler if (src.find("<|channel|>") != std::string::npos) { LOG_DBG("Using specialized template: GPT-OSS\n"); diff --git a/common/common.cpp b/common/common.cpp index ae44f054ab51..d1e7b7e5dd11 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -49,7 +49,7 @@ #include #endif -#if defined(__linux__) +#if !defined(_WIN32) && !defined(__APPLE__) #include #include #endif @@ -613,34 +613,6 @@ std::string string_from(const struct llama_context * ctx, const std::vectorpw_dir)) { + const struct passwd * pw = getpwuid(getuid()); + if (!pw || !pw->pw_dir || !*pw->pw_dir) { throw std::runtime_error("Failed to find $HOME directory"); } - - cache_directory = std::string(pw->pw_dir) + std::string("/.cache/"); -#else /* defined(__linux__) */ - throw std::runtime_error("Failed to find $HOME directory"); -#endif /* defined(__linux__) */ - } -#elif defined(__APPLE__) - cache_directory = common_get_env("HOME"); - if (cache_directory.empty()) { - throw std::runtime_error("Failed to find $HOME directory"); - } - cache_directory += "/Library/Caches/"; -#elif defined(_WIN32) - cache_directory = common_get_env("LOCALAPPDATA"); - if (cache_directory.empty()) { - throw std::runtime_error("Failed to find %LOCALAPPDATA% directory"); + cache_directory = pw->pw_dir; + cache_directory /= ".cache"; } -#elif defined(__EMSCRIPTEN__) - GGML_ABORT("not implemented on this platform"); -#else -# error Unknown architecture -#endif - cache_directory = ensure_trailing_slash(cache_directory); - cache_directory += "llama.cpp"; } - return ensure_trailing_slash(cache_directory); +#endif + return cache_directory / "llama.cpp"; } std::string fs_get_config_directory() { @@ -1068,14 +1050,15 @@ std::string fs_get_config_directory() { return ensure_trailing_slash(config_directory); } -std::string fs_get_cache_file(const std::string & filename) { +std::filesystem::path fs_get_cache_file(const std::string & filename) { GGML_ASSERT(filename.find(DIRECTORY_SEPARATOR) == std::string::npos); - std::string cache_directory = fs_get_cache_directory(); - const bool success = fs_create_directory_with_parents(cache_directory); - if (!success) { - throw std::runtime_error("failed to create cache directory: " + cache_directory); + const std::filesystem::path cache_directory = fs_get_cache_directory(); + std::error_code ec; + std::filesystem::create_directories(cache_directory, ec); + if (ec) { + throw std::runtime_error("failed to create cache directory: " + fs_path_to_utf8(cache_directory)); } - return cache_directory + filename; + return cache_directory / std::filesystem::u8path(filename); } std::vector fs_list(const std::string & path, bool include_directories) { @@ -1480,7 +1463,8 @@ common_init_result_ptr common_init_from_params(common_params & params, bool mode } if (llama_model_has_encoder(model)) { - llama_encode(lctx, llama_batch_get_one(tmp.data(), tmp.size())); + common_batch batch = common_batch_get_one(lctx, tmp); + llama_process(lctx, LLAMA_PROCESS_TYPE_ENCODE, batch.get()); llama_token decoder_start_token_id = llama_model_decoder_start_token(model); if (decoder_start_token_id == LLAMA_TOKEN_NULL) { decoder_start_token_id = bos; @@ -1489,7 +1473,9 @@ common_init_result_ptr common_init_from_params(common_params & params, bool mode tmp.push_back(decoder_start_token_id); } if (llama_model_has_decoder(model)) { - llama_decode(lctx, llama_batch_get_one(tmp.data(), std::min(tmp.size(), (size_t) params.n_batch))); + tmp.resize(std::min(tmp.size(), (size_t) params.n_batch)); + common_batch batch = common_batch_get_one(lctx, tmp); + llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch.get()); } llama_memory_clear(llama_get_memory(lctx), true); llama_synchronize(lctx); @@ -1553,9 +1539,13 @@ common_context_seq_rm_type common_context_can_seq_rm(llama_context * ctx) { tmp.push_back(0); tmp.push_back(0); - int ret = llama_decode(ctx, llama_batch_get_one(tmp.data(), tmp.size())); + int ret; + { + common_batch batch = common_batch_get_one(ctx, tmp); + ret = llama_process(ctx, LLAMA_PROCESS_TYPE_DECODE, batch.get()); + } if (ret != 0) { - COM_ERR("llama_decode() failed: %d\n", ret); + COM_ERR("llama_process() failed: %d\n", ret); res = COMMON_CONTEXT_SEQ_RM_TYPE_NO; goto done; } @@ -2142,31 +2132,140 @@ float lr_opt::get_lr(float epoch) const { } bool common_replay_last_token(struct llama_context * ctx, llama_token last_token, int32_t pos) { - llama_batch batch = llama_batch_get_one(&last_token, 1); - batch.pos = &pos; - if (llama_decode(ctx, batch)) { + common_batch batch(ctx); + batch.add(last_token, pos, 0, true); + + if (llama_process(ctx, LLAMA_PROCESS_TYPE_DECODE, batch.get())) { LOG_ERR("%s: failed to replay last token\n", __func__); return false; } return true; } -llama_batch_ext_ptr common_batch_ext_get_one(llama_context * ctx, const llama_tokens & tokens) { - llama_batch_ext_ptr batch(llama_batch_ext_init(ctx)); +common_batch::common_batch(llama_context * ctx) : batch(llama_batch_ext_init(ctx)) { + const auto rope_type = llama_model_rope_type(llama_get_model(ctx)); + n_pos = rope_type == LLAMA_ROPE_TYPE_MROPE || rope_type == LLAMA_ROPE_TYPE_IMROPE ? GGML_MROPE_SECTIONS : 1; +} + +void common_batch::clear() { + tokens.clear(); + llama_batch_ext_clear(batch.get()); +} + +int32_t common_batch::add(llama_token id, llama_pos pos, llama_seq_id seq_id, bool output) { + const int32_t idx = llama_batch_ext_add_token(batch.get(), seq_id, id); + if (idx < 0) { + GGML_ABORT("%s: failed to add token %d to the batch (error %d, n_tokens = %d)\n", __func__, id, idx, size()); + } + llama_batch_ext_set_pos(batch.get(), idx, &pos); + if (output) { + llama_batch_ext_set_output_logits(batch.get(), idx, true); + } + tokens.push_back({ id, { pos, 0, 0, 0 }, seq_id, output, { nullptr, 0, 0 } }); + return idx; +} + +bool common_batch::set_output(int32_t idx, bool value) { + if (idx < 0 || idx >= (int32_t) tokens.size()) { + return false; + } + tokens[idx].output = value; + return llama_batch_ext_set_output_logits(batch.get(), idx, value); +} + +bool common_batch::set_embd(int32_t idx, llama_embd embd) { + if (idx < 0 || idx >= (int32_t) tokens.size()) { + return false; + } + if (!llama_batch_ext_set_embd_token(batch.get(), idx, embd)) { + return false; + } + tokens[idx].embd = embd; + return true; +} + +int32_t common_batch::add_embd(llama_embd embd, const llama_pos * pos, llama_seq_id seq_id, bool output) { + const int32_t idx = llama_batch_ext_add_embd(batch.get(), seq_id, embd); + if (idx < 0) { + GGML_ABORT("%s: failed to add embedding to the batch (error %d, n_tokens = %d)\n", __func__, idx, size()); + } + llama_batch_ext_set_pos(batch.get(), idx, pos); + if (output) { + llama_batch_ext_set_output_logits(batch.get(), idx, true); + } + token t = { LLAMA_TOKEN_NULL, { 0, 0, 0, 0 }, seq_id, output, embd }; + for (int32_t j = 0; j < n_pos; ++j) { + t.pos[j] = pos[j]; + } + tokens.push_back(t); + return idx; +} + +common_batch common_batch_from_llama_batch(llama_context * ctx, const llama_batch & batch) { + common_batch res(ctx); + + const bool has_token = batch.token != nullptr; + const bool has_embd = batch.embd != nullptr; + + const size_t n_embd = llama_model_n_embd_inp(llama_get_model(ctx)); + + // positions continue from the memory when none are given + auto * mem = llama_get_memory(ctx); + std::vector pos_next(llama_n_seq_max(ctx)); + for (llama_seq_id s = 0; s < (llama_seq_id) pos_next.size(); ++s) { + pos_next[s] = llama_memory_seq_pos_max(mem, s) + 1; + } + + for (int32_t i = 0; i < batch.n_tokens; ++i) { + const int32_t n_sid = batch.n_seq_id ? batch.n_seq_id[i] : 1; + const llama_seq_id seq_id = batch.seq_id ? batch.seq_id[i][0] : 0; + + llama_pos pos[GGML_MROPE_SECTIONS] = { 0, 0, 0, 0 }; + if (!batch.pos) { + pos[0] = pos_next[seq_id]++; + } else if (has_token) { + pos[0] = batch.pos[i]; + } else { + // embedding batch: section-major layout pos[j*n_tokens + i] + for (int32_t j = 0; j < res.n_pos; ++j) { + pos[j] = batch.pos[j * batch.n_tokens + i]; + } + } + + const bool output = batch.logits ? batch.logits[i] != 0 : i == batch.n_tokens - 1; + + const llama_embd embd = { has_embd ? batch.embd + (size_t) i * n_embd : nullptr, 1, n_embd }; + + int32_t idx; + if (has_token) { + idx = res.add(batch.token[i], pos[0], seq_id, output); + if (has_embd) { + res.set_embd(idx, embd); + } + } else { + idx = res.add_embd(embd, pos, seq_id, output); + } + + for (int32_t s = 1; s < n_sid; ++s) { + llama_batch_ext_add_seq(res.get(), idx, batch.seq_id[i][s]); + } + } + + return res; +} + +common_batch common_batch_get_one(llama_context * ctx, const llama_tokens & tokens) { + common_batch batch(ctx); auto mem = llama_get_memory(ctx); - llama_pos pos = mem ? llama_memory_seq_pos_max(mem, 0) + 1 : 0; + llama_pos pos = llama_memory_seq_pos_max(mem, 0) + 1; // -1 + 1 == 0 when the memory is empty for (size_t i = 0; i < tokens.size(); ++i) { - const int32_t idx = llama_batch_ext_add_token(batch.get(), 0, tokens[i]); - llama_batch_ext_set_pos(batch.get(), idx, &pos); + const bool output = i == tokens.size() - 1; + batch.add(tokens[i], pos, 0, output); pos++; } - if (!tokens.empty()) { - llama_batch_ext_set_output_logits(batch.get(), (int32_t) tokens.size() - 1, true); - } - return batch; } @@ -2194,7 +2293,7 @@ bool common_prompt_batch_decode( // memory, so we can't just remove the last token from the memory and replay the last token which // is the reason for this logic. llama_tokens prefix_tokens(all_tokens.begin() + offset, all_tokens.begin() + offset + n_tokens_before_last); - llama_batch_ext_ptr batch_prefix = common_batch_ext_get_one(ctx, prefix_tokens); + common_batch batch_prefix = common_batch_get_one(ctx, prefix_tokens); if (llama_process(ctx, LLAMA_PROCESS_TYPE_DECODE, batch_prefix.get())) { COM_ERR("%s", "failed to eval\n"); return false; @@ -2204,10 +2303,8 @@ bool common_prompt_batch_decode( llama_state_save_file(ctx, state_path.data(), all_tokens.data(), all_tokens.size()); COM_INF("saved session before last token to %s, n_new = %zu\n", state_path.data(), all_tokens.size()); - llama_token last_token = all_tokens.back(); - llama_batch_ext_ptr batch_last = common_batch_ext_get_one(ctx, { last_token }); - llama_pos pos = n_past; - llama_batch_ext_set_pos(batch_last.get(), 0, &pos); + common_batch batch_last(ctx); + batch_last.add(all_tokens.back(), n_past, 0, true); if (llama_process(ctx, LLAMA_PROCESS_TYPE_DECODE, batch_last.get())) { COM_ERR("%s", "failed to eval last token\n"); @@ -2216,7 +2313,7 @@ bool common_prompt_batch_decode( n_past++; } else { llama_tokens new_tokens(all_tokens.begin() + offset, all_tokens.begin() + offset + n_new); - llama_batch_ext_ptr batch = common_batch_ext_get_one(ctx, new_tokens); + common_batch batch = common_batch_get_one(ctx, new_tokens); if (llama_process(ctx, LLAMA_PROCESS_TYPE_DECODE, batch.get())) { COM_ERR("%s", "failed to eval\n"); return false; diff --git a/common/common.h b/common/common.h index 527f4be7d899..01726247b76e 100644 --- a/common/common.h +++ b/common/common.h @@ -8,6 +8,7 @@ #include "ggml.h" #include "llama.h" +#include #include #include #include @@ -879,7 +880,6 @@ void string_process_escapes(std::string & input); std::string string_from(bool value); std::string string_from(const std::vector & values); std::string string_from(const struct llama_context * ctx, const std::vector & tokens); -std::string string_from(const struct llama_context * ctx, const struct llama_batch & batch); bool glob_match(const std::string & pattern, const std::string & str); @@ -904,16 +904,18 @@ std::string fs_path_to_utf8(const std::filesystem::path & path); std::string common_get_env(const std::string & name); void common_set_env(const std::string & name, const std::string & value); +// reads a path from the environment, an unset variable gives an empty path +std::filesystem::path common_get_path_from_env(const std::string & name); + // // Filesystem utils // bool fs_validate_filename(const std::string & filename, bool allow_subdirs = false); -bool fs_create_directory_with_parents(const std::string & path); bool fs_is_directory(const std::string & path); -std::string fs_get_cache_directory(); -std::string fs_get_cache_file(const std::string & filename); +std::filesystem::path fs_get_cache_directory(); +std::filesystem::path fs_get_cache_file(const std::string & filename); std::string fs_get_config_directory(); struct common_file_info { @@ -927,6 +929,8 @@ std::vector fs_list(const std::string & path, bool include_dir // fs open, also handle UTF8 on Windows std::ifstream fs_open_ifstream(const std::string & fname, std::ios_base::openmode mode); +void fs_write_atomic(const std::filesystem::path & path, const std::string & data); + // // TTY utils // @@ -1036,9 +1040,54 @@ void common_batch_add( const std::vector & seq_ids, bool logits); +// wrapper around llama_batch_ext that provide getter functions for downstream code +struct common_batch { + struct token { + llama_token id; + std::array pos; // only pos[0] is used for text tokens + llama_seq_id seq_id; + bool output; + llama_embd embd; // non-owning view of the data passed to add_embd()/set_embd(), data == NULL if none + }; + + std::vector tokens; // mirror of the entries, tokens[i] describes batch index i + llama_batch_ext_ptr batch; + + int32_t n_pos = 1; // positions per embedding entry, GGML_MROPE_SECTIONS for MROPE/IMROPE + + common_batch() = default; + common_batch(struct llama_context * ctx); + + llama_batch_ext * get() const { return batch.get(); } + + // content type of the batch, all entries carry the same combination + bool has_token() const { return !tokens.empty() && tokens[0].id != LLAMA_TOKEN_NULL; } + bool has_embd () const { return !tokens.empty() && tokens[0].embd.data != nullptr; } + + void clear(); + + // returns the batch index (>= 0), aborts if the entry cannot be added (batch full, invalid token or seq id) + int32_t add(llama_token id, llama_pos pos, llama_seq_id seq_id, bool output); + + bool set_output(int32_t idx, bool value); + + // attach a token embedding to the entry at idx, can only be set once per entry + bool set_embd(int32_t idx, llama_embd embd); + + // add an embedding-only entry (no token id), aborts like add() on failure + // pos points to n_pos positions + int32_t add_embd(llama_embd embd, const llama_pos * pos, llama_seq_id seq_id, bool output); + + int32_t size() const { return (int32_t) tokens.size(); } +}; + // create a single-sequence batch from a list of tokens // last token always have output_logits set to true -llama_batch_ext_ptr common_batch_ext_get_one(struct llama_context * ctx, const llama_tokens & tokens); +common_batch common_batch_get_one(struct llama_context * ctx, const llama_tokens & tokens); + +// convert a legacy llama_batch, applying its defaults: seq 0, positions continue from memory, last token is output +// the embd rows are read at the model input width +common_batch common_batch_from_llama_batch(struct llama_context * ctx, const llama_batch & batch); // decodes a single batch of tokens for a prompt and manages session tokens // diff --git a/common/download.cpp b/common/download.cpp index 4b28a708c86e..6b7123e47671 100644 --- a/common/download.cpp +++ b/common/download.cpp @@ -46,39 +46,9 @@ // downloader // -// validate repo name format: owner/repo -static void write_file(const std::string & fname, const std::string & content) { - const std::string fname_tmp = fname + ".tmp"; - std::ofstream file(fname_tmp); - if (!file) { - throw std::runtime_error(string_format("error: failed to open file '%s'\n", fname.c_str())); - } - - try { - file << content; - file.close(); - - // Makes write atomic - if (rename(fname_tmp.c_str(), fname.c_str()) != 0) { - LOG_ERR("%s: unable to rename file: %s to %s\n", __func__, fname_tmp.c_str(), fname.c_str()); - // If rename fails, try to delete the temporary file - if (remove(fname_tmp.c_str()) != 0) { - LOG_ERR("%s: unable to delete temporary file: %s\n", __func__, fname_tmp.c_str()); - } - } - } catch (...) { - // If anything fails, try to delete the temporary file - if (remove(fname_tmp.c_str()) != 0) { - LOG_ERR("%s: unable to delete temporary file: %s\n", __func__, fname_tmp.c_str()); - } - - throw std::runtime_error(string_format("error: failed to write file '%s'\n", fname.c_str())); - } -} - static void write_etag(const std::string & path, const std::string & etag) { const std::string etag_path = path + ".etag"; - write_file(etag_path, etag); + fs_write_atomic(std::filesystem::u8path(etag_path), etag); LOG_DBG("%s: file etag saved: %s\n", __func__, etag_path.c_str()); } @@ -274,6 +244,12 @@ static bool common_pull_file(httplib::Client & cli, return false; } + ofs.close(); + if (!ofs) { + LOG_ERR("%s: error closing file: %s\n", __func__, path_tmp.c_str()); + return false; + } + return true; } @@ -286,7 +262,7 @@ static int common_download_file_single_online(const std::string & url, static const int max_attempts = 3; static const int retry_delay_seconds = 2; - const bool file_exists = std::filesystem::exists(path); + const bool file_exists = std::filesystem::exists(std::filesystem::u8path(path)); if (file_exists && skip_etag) { LOG_DBG("%s: using cached file: %s\n", __func__, path.c_str()); @@ -477,7 +453,7 @@ int common_download_file_single(const std::string & url, return common_download_file_single_online(url, path, online_opts, skip_etag); } - if (!std::filesystem::exists(path)) { + if (!std::filesystem::exists(std::filesystem::u8path(path))) { LOG_ERR("%s: required file is not available in cache (offline mode): %s\n", __func__, path.c_str()); return -1; } @@ -943,7 +919,7 @@ std::string common_docker_resolve_model(const std::string & docker) { std::string model_filename = repo; std::replace(model_filename.begin(), model_filename.end(), '/', '_'); model_filename += "_" + tag + ".gguf"; - std::string local_path = fs_get_cache_file(model_filename); + std::string local_path = fs_path_to_utf8(fs_get_cache_file(model_filename)); const std::string blob_url = url_prefix + "/blobs/" + gguf_digest; common_download_opts opts; diff --git a/common/hf-cache.cpp b/common/hf-cache.cpp index 65c9722f3322..4241d005af27 100644 --- a/common/hf-cache.cpp +++ b/common/hf-cache.cpp @@ -44,8 +44,7 @@ static fs::path get_cache_directory() { {HOME_DIR, fs::path(".cache") / "huggingface" / "hub"} }; for (const auto & entry : entries) { - if (auto * p = std::getenv(entry.var); p && *p) { - fs::path base(p); + if (fs::path base = common_get_path_from_env(entry.var); !base.empty()) { return entry.path.empty() ? base : base / entry.path; } } @@ -173,28 +172,6 @@ static bool is_valid_subpath(const fs::path & path, const fs::path & subpath) { return b_end == b.end(); } -static void safe_write_file(const fs::path & path, const std::string & data) { - fs::path path_tmp = path.string() + ".tmp"; - - if (path.has_parent_path()) { - fs::create_directories(path.parent_path()); - } - - std::ofstream file(path_tmp); - file << data; - file.close(); - - std::error_code ec; - - if (!file.fail()) { - fs::rename(path_tmp, path, ec); - } - if (file.fail() || ec) { - fs::remove(path_tmp, ec); - throw std::runtime_error("failed to write file: " + path.string()); - } -} - static common_json api_get(const std::string & url, const std::string & token) { auto [cli, parts] = common_http_client(url); @@ -241,6 +218,7 @@ static std::string get_repo_commit(const std::string & repo_id, fs::path refs_path = get_repo_path(repo_id) / "refs"; std::string name; std::string commit; + fs::path name_path; for (const auto & branch : json["branches"]) { if (!branch.is_object() || @@ -251,24 +229,28 @@ static std::string get_repo_commit(const std::string & repo_id, std::string _name = branch["name"].get(); std::string _commit = branch["targetCommit"].get(); - if (!is_valid_subpath(refs_path, _name)) { - LOG_WRN("%s: skip invalid branch: %s\n", __func__, _name.c_str()); - continue; - } if (!is_valid_commit(_commit)) { LOG_WRN("%s: skip invalid commit: %s\n", __func__, _commit.c_str()); continue; } + const fs::path candidate = fs::u8path(_name); + + if (!is_valid_subpath(refs_path, candidate)) { + LOG_WRN("%s: skip invalid branch: %s\n", __func__, _name.c_str()); + continue; + } if (_name == "main") { name = _name; commit = _commit; + name_path = candidate; break; } if (name.empty() || commit.empty()) { name = _name; commit = _commit; + name_path = candidate; } } @@ -277,7 +259,7 @@ static std::string get_repo_commit(const std::string & repo_id, return {}; } - safe_write_file(refs_path / name, commit); + fs_write_atomic(refs_path / name_path, commit); return commit; } catch (const common_json_error & e) { @@ -326,7 +308,9 @@ hf_files get_repo_files(const std::string & repo_id, file.repo_id = repo_id; file.path = item["path"].get(); - if (!is_valid_subpath(commit_path, file.path)) { + const fs::path subpath = fs::u8path(file.path); + + if (!is_valid_subpath(commit_path, subpath)) { LOG_WRN("%s: skip invalid path: %s\n", __func__, file.path.c_str()); continue; } @@ -346,12 +330,12 @@ hf_files get_repo_files(const std::string & repo_id, file.url = endpoint + repo_id + "/resolve/" + commit + "/" + file.path; - fs::path final_path = commit_path / file.path; - file.final_path = final_path.string(); + fs::path final_path = commit_path / subpath; + file.final_path = fs_path_to_utf8(final_path); if (!file.oid.empty() && !fs::exists(final_path)) { fs::path local_path = blobs_path / file.oid; - file.local_path = local_path.string(); + file.local_path = fs_path_to_utf8(local_path); } else { file.local_path = file.final_path; } @@ -418,7 +402,7 @@ hf_files get_cached_files(const std::string & repo_id) { if (!fs::exists(snapshots_path)) { continue; } - std::string _repo_id = folder_name_to_repo(repo.path().filename().string()); + std::string _repo_id = folder_name_to_repo(fs_path_to_utf8(repo.path().filename())); if (!is_valid_repo_id(_repo_id)) { continue; @@ -441,8 +425,9 @@ hf_files get_cached_files(const std::string & repo_id) { if (!path.empty()) { hf_file file; file.repo_id = _repo_id; - file.path = path.generic_string(); - file.local_path = entry.path().string(); + const auto generic_path = path.generic_u8string(); + file.path = std::string(generic_path.begin(), generic_path.end()); + file.local_path = fs_path_to_utf8(entry.path()); file.final_path = file.local_path; files.push_back(std::move(file)); } @@ -456,8 +441,8 @@ std::string finalize_file(const hf_file & file) { static std::atomic symlinks_disabled{false}; std::error_code ec; - fs::path local_path(file.local_path); - fs::path final_path(file.final_path); + fs::path local_path = fs::u8path(file.local_path); + fs::path final_path = fs::u8path(file.final_path); if (local_path == final_path || fs::exists(final_path, ec)) { return file.final_path; @@ -504,7 +489,7 @@ bool remove_cached_repo(const std::string & repo_id) { std::error_code ec; auto removed = fs::remove_all(repo_path, ec); if (ec) { - LOG_ERR("%s: failed to remove repo cache %s: %s\n", __func__, repo_path.string().c_str(), ec.message().c_str()); + LOG_ERR("%s: failed to remove repo cache %s: %s\n", __func__, fs_path_to_utf8(repo_path).c_str(), ec.message().c_str()); return false; } return removed > 0; diff --git a/common/parsers/llm-jp-harmony.cpp b/common/parsers/llm-jp-harmony.cpp new file mode 100644 index 000000000000..b012f9a05b94 --- /dev/null +++ b/common/parsers/llm-jp-harmony.cpp @@ -0,0 +1,164 @@ +#include "parsers.h" + +// LLM-jp-4.1: the GPT-OSS (Harmony) format with two differences +// - the tokenizer emits a space after every special token: "<|channel|> analysis<|message|> ..." +// - parallel tool calls are consecutive assistant messages, all but the last closed by <|end|> +common_chat_params common_chat_params_init_llm_jp_harmony(const common_chat_template & tmpl, + const autoparser::generation_params & inputs) { + common_chat_params data; + + // Copy reasoning to the "thinking" field as expected by the template + auto adjusted_messages = json::array(); + for (auto msg : inputs.messages) { + if (msg.contains("reasoning_content") && msg.at("reasoning_content").is_string()) { + msg["thinking"] = msg.at("reasoning_content"); + if (msg.contains("tool_calls") && msg.at("tool_calls").is_array() && !msg.at("tool_calls").empty()) { + msg.erase("content"); + } + } + adjusted_messages.push_back(msg); + } + + auto prompt = common_chat_template_direct_apply_impl(tmpl, inputs, /* messages_override= */ adjusted_messages); + + // Check if we need to replace the return token with end token during + // inference and without generation prompt. For more details see: + // https://github.com/ggml-org/llama.cpp/issues/15417 + if (inputs.is_inference && !inputs.add_generation_prompt) { + static constexpr std::string_view return_token = "<|return|>"; + static constexpr std::string_view end_token = "<|end|>"; + if (size_t pos = prompt.rfind(return_token); pos != std::string::npos) { + prompt.replace(pos, return_token.length(), end_token); + } + } + + data.prompt = prompt; + data.generation_prompt = common_chat_template_generation_prompt_impl(tmpl, inputs, /* messages_override= */ adjusted_messages); + data.message_delimiters = { + { COMMON_CHAT_ROLE_ASSISTANT, "<|start|>assistant" }, + { COMMON_CHAT_ROLE_USER, "<|start|>user" }, + { COMMON_CHAT_ROLE_SYSTEM, "<|start|>developer" }, + { COMMON_CHAT_ROLE_SYSTEM, "<|start|>system" }, + { COMMON_CHAT_ROLE_TOOL, "<|start|>functions" }, + }; + + data.format = COMMON_CHAT_FORMAT_PEG_NATIVE; + data.supports_thinking = true; + + data.thinking_start_tag = "<|channel|>analysis<|message|>"; + data.thinking_end_tags = {"<|end|>"}; + + // These special tokens are required to parse properly, so we include them + // even if parse_tool_calls is false. + data.preserved_tokens = { + "<|channel|>", "<|constrain|>", "<|message|>", "<|start|>", "<|end|>", + }; + + // Adjust prompt for continuation + if (inputs.has_continuation()) { + const auto & msg = inputs.continue_msg; + + data.generation_prompt = "<|start|>assistant<|channel|>analysis<|message|>" + msg.reasoning_content; + if (inputs.continue_final_message == COMMON_CHAT_CONTINUATION_CONTENT) { + data.generation_prompt += "<|end|><|start|>assistant<|channel|>final<|message|>" + msg.render_content(); + } + + data.prompt += data.generation_prompt; + } + + auto has_tools = inputs.tools.is_array() && !inputs.tools.empty(); + auto has_response_format = !inputs.json_schema.is_null() && inputs.json_schema.is_object(); + auto include_grammar = has_response_format || (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE); + auto extract_reasoning = inputs.reasoning_format != COMMON_REASONING_FORMAT_NONE; + + auto parser = build_chat_peg_parser([&](common_chat_peg_builder & p) { + // tokenizer space after special tokens; not p.space() since GBNF `space` allows one space only + auto sp = p.chars("[ ]", 0, -1); + auto channel_tag = p.literal("<|channel|>") + sp; + // one space only: keep an intentional leading space in the body + auto message = p.literal("<|message|>") + p.optional(p.literal(" ")); + + auto start = p.rule("start", p.literal("<|start|>") + sp + p.literal("assistant")); + auto end = p.rule("end", p.literal("<|end|>")); + auto content = p.rule("message-content", p.until("<|end|>")); + auto channel = channel_tag + (p.literal("commentary") | p.literal("analysis")); + auto constrain_type = p.chars("[A-Za-z0-9_-]", 1, -1); + auto constraint = p.optional(p.space() + p.optional(p.literal("<|constrain|>") + sp) + constrain_type); + + auto start_analysis = channel_tag + p.literal("analysis") + message; + if (extract_reasoning) { + p.rule("analysis", start_analysis + p.reasoning(content) + end); + } else { + p.rule("analysis", p.content(start_analysis + content + end)); + } + + auto analysis = p.ref("analysis"); + auto preamble = p.rule("preamble", channel_tag + p.literal("commentary") + message + p.content(content) + end); + auto final_msg = p.rule("final", channel_tag + p.literal("final") + message + p.content(content)); + + auto any = p.rule("any", preamble | analysis); + + if (has_response_format) { + auto response_format = p.rule("response-format", + channel_tag + p.literal("final") + constraint + message + + p.content(p.schema(p.json(), "response-format-schema", inputs.json_schema))); + + return p.zero_or_more(start + analysis) + start + response_format; + } + + if (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE) { + auto tool_choice = p.choice(); + + foreach_function(inputs.tools, [&](const json & tool) { + const auto & function = tool.at("function"); + std::string name = function.at("name"); + const auto params = common_chat_tool_parameters(function); + + auto func_name = p.literal(" to=functions.") + p.tool_name(p.literal(name)); + auto args = p.tool_args(p.schema(p.json(), "tool-" + name + "-schema", params)); + + // recipient in role header + // <|start|>assistant to=functions.NAME<|channel|>(commentary|analysis)[constraint]<|message|>ARGS + auto tool_in_role = p.tool(p.tool_open(func_name + channel + constraint + message) + args); + + // recipient in channel header + // <|channel|>(commentary|analysis) to=functions.NAME[constraint]<|message|>ARGS + auto tool_in_channel = p.tool(p.tool_open(channel + func_name + constraint + message) + args); + + tool_choice |= p.rule("tool-" + name, tool_in_role | tool_in_channel); + }); + + // parallel calls are separated by <|end|>; inside the trigger rule so the lazy grammar covers all of them + auto tool_calls = inputs.parallel_tool_calls + ? tool_choice + p.zero_or_more(end + start + tool_choice) + : tool_choice; + auto tool_call = p.trigger_rule("tool-call", tool_calls); + + if (inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_REQUIRED) { + return p.zero_or_more(start + any) + start + tool_call; + } + + return p.zero_or_more(start + any) + start + (tool_call | final_msg); + } + + return p.zero_or_more(start + any) + start + final_msg; + }); + + data.parser = parser.save(); + + if (include_grammar) { + data.grammar_lazy = !(has_response_format || (has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_REQUIRED)); + data.grammar = build_grammar([&](const common_grammar_builder & builder) { + parser.build_grammar(builder, data.grammar_lazy); + }); + + data.grammar_triggers = { + { COMMON_GRAMMAR_TRIGGER_TYPE_PATTERN, "^\\s+to$" }, + { COMMON_GRAMMAR_TRIGGER_TYPE_PATTERN, "^<\\|channel\\|>\\s*(?:commentary|analysis)\\s+to=functions$" }, + { COMMON_GRAMMAR_TRIGGER_TYPE_PATTERN, "<\\|start\\|>\\s*assistant(\\s+to)" }, + { COMMON_GRAMMAR_TRIGGER_TYPE_PATTERN, "<\\|start\\|>\\s*assistant(<\\|channel\\|>\\s*(?:commentary|analysis)\\s+to)" } + }; + } + + return data; +} diff --git a/common/parsers/muse-glimmer.cpp b/common/parsers/muse-glimmer.cpp index d03bf2d58f73..784928269b84 100644 --- a/common/parsers/muse-glimmer.cpp +++ b/common/parsers/muse-glimmer.cpp @@ -43,9 +43,10 @@ common_chat_params common_chat_params_init_muse_glimmer(const common_chat_templa auto extract_reasoning = inputs.reasoning_format != COMMON_REASONING_FORMAT_NONE; - auto has_tools = inputs.tools.is_array() && !inputs.tools.empty(); - // Constrained grammar whenever tools are offered. - auto include_grammar = has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE; + auto has_tools = inputs.tools.is_array() && !inputs.tools.empty(); + auto has_response_format = !inputs.json_schema.is_null() && inputs.json_schema.is_object(); + // Constrained grammar whenever tools are offered or a response format is requested. + auto include_grammar = has_response_format || (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE); auto parser = build_chat_peg_parser([&](common_chat_peg_builder & p) { auto start = p.rule("start", p.literal("<|start|>assistant")); @@ -65,6 +66,15 @@ common_chat_params common_chat_params_init_muse_glimmer(const common_chat_templa auto final_msg = p.rule("final", recipient + p.literal("<|message|>") + p.content(p.until_one_of({ "<|eot|>", "<|eom|>" }))); + if (has_response_format) { + auto response_json = p.content(p.schema(p.json(), "response-format-schema", inputs.json_schema)); + auto response_format = p.rule("response-format", + recipient + p.literal("<|message|>") + + ((p.literal("```json") + p.space() + response_json + p.space() + p.literal("```")) | response_json)); + + return p.zero_or_more(start + analysis) + start + response_format; + } + if (has_tools && inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_NONE) { auto string_value = p.ac( p.tool_arg_string_value(p.until("")) + p.tool_arg_close(p.literal("")), @@ -124,7 +134,7 @@ common_chat_params common_chat_params_init_muse_glimmer(const common_chat_templa data.parser = parser.save(); if (include_grammar) { - data.grammar_lazy = inputs.tool_choice != COMMON_CHAT_TOOL_CHOICE_REQUIRED; + data.grammar_lazy = !(has_response_format || (has_tools && inputs.tool_choice == COMMON_CHAT_TOOL_CHOICE_REQUIRED)); data.grammar = build_grammar([&](const common_grammar_builder & builder) { parser.build_grammar(builder, data.grammar_lazy); }); diff --git a/common/parsers/parsers.h b/common/parsers/parsers.h index f866360073bb..0caf62d8c055 100644 --- a/common/parsers/parsers.h +++ b/common/parsers/parsers.h @@ -68,6 +68,8 @@ common_chat_params common_chat_params_init_ling3(const common_chat_template & tm // tool_list_tokens preserves the LFM2 system tool-list markers; LFM2.5 renders without them common_chat_params common_chat_params_init_lfm2(const common_chat_template & tmpl, const autoparser::generation_params & inputs, bool tool_list_tokens); +common_chat_params common_chat_params_init_llm_jp_harmony(const common_chat_template & tmpl, const autoparser::generation_params & inputs); + common_chat_params common_chat_params_init_minicpm5(const common_chat_template & tmpl, const autoparser::generation_params & inputs); common_chat_params common_chat_params_init_minimax_m3(const common_chat_template & tmpl, const autoparser::generation_params & inputs); diff --git a/common/parsers/sources.cmake b/common/parsers/sources.cmake index 70af84e25110..5815939d6aeb 100644 --- a/common/parsers/sources.cmake +++ b/common/parsers/sources.cmake @@ -13,6 +13,7 @@ set(LLAMA_CHAT_PARSERS_SOURCES ${CMAKE_CURRENT_LIST_DIR}/kimi-k3.cpp ${CMAKE_CURRENT_LIST_DIR}/ling3.cpp ${CMAKE_CURRENT_LIST_DIR}/lfm2.cpp + ${CMAKE_CURRENT_LIST_DIR}/llm-jp-harmony.cpp ${CMAKE_CURRENT_LIST_DIR}/minicpm5.cpp ${CMAKE_CURRENT_LIST_DIR}/minimax-m3.cpp ${CMAKE_CURRENT_LIST_DIR}/ministral3.cpp diff --git a/common/speculative.cpp b/common/speculative.cpp index 6fdfa4dc33f0..82e9e92238b9 100644 --- a/common/speculative.cpp +++ b/common/speculative.cpp @@ -165,7 +165,7 @@ struct common_speculative_impl { virtual void begin(llama_seq_id seq_id, const llama_tokens & prompt) = 0; - virtual bool process(const llama_batch & batch) = 0; + virtual bool process(const common_batch & batch) = 0; virtual void draft(common_speculative_draft_params_vec & dparams) = 0; @@ -179,7 +179,11 @@ struct common_speculative_impl { struct common_speculative_impl_draft_simple : public common_speculative_impl { common_params_speculative_draft params; - llama_batch batch; + common_batch batch; + + // zero row at the draft input width, stands in for target embeddings the draft cannot read + std::vector zeros; + bool zeros_warned = false; // the substitution is reported once std::vector smpls; @@ -194,6 +198,8 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { throw std::runtime_error("draft-simple requires a draft context"); } + zeros.assign(llama_model_n_embd_inp(llama_get_model(ctx_dft)), 0.0f); + SPC_TRC("%s", "adding speculative implementation 'draft-simple'\n"); SPC_TRC("- n_max=%d, n_min=%d, p_min=%f\n", this->params.n_max, this->params.n_min, this->params.p_min); SPC_TRC("- gpu_layers=%d, cache_k=%s, cache_v=%s, ctx_tgt=%s, ctx_dft=%s, devices=[%s]\n", @@ -204,7 +210,7 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { ctx_dft ? "yes" : "no", common_speculative_get_devices_str(this->params.devices).c_str()); - batch = llama_batch_init(llama_n_batch(ctx_dft), 0, 1); + batch = common_batch(ctx_dft); // TODO: optimize or pass from outside? // { @@ -249,21 +255,46 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { } } - ~common_speculative_impl_draft_simple() override { - llama_batch_free(batch); - } - void begin(llama_seq_id /*seq_id*/, const llama_tokens & /*prompt*/) override { // noop } - bool process(const llama_batch & batch) override { + bool process(const common_batch & batch_in) override { auto * ctx_dft = params.ctx_dft; - llama_batch batch_dft = batch; - batch_dft.logits = nullptr; + // copy the entries to a batch owned by the draft context, only the last token is output + batch.clear(); + const int32_t n_tokens = batch_in.size(); + for (int32_t k = 0; k < n_tokens; ++k) { + const auto & t = batch_in.tokens[k]; + const bool output = k == n_tokens - 1; + if (t.id != LLAMA_TOKEN_NULL) { + const int32_t idx = batch.add(t.id, t.pos[0], t.seq_id, output); + if (t.embd.data) { + batch.set_embd(idx, t.embd); + } + } else { + // mtmd input is projected by the target encoder, a draft with a different width cannot read it + // it gets zeros instead, keeping its positions contiguous + // ref: https://github.com/ggml-org/llama.cpp/pull/29385#discussion_r4124743243 + const size_t n_embd = t.embd.n_rows * t.embd.n_embd; + const bool same_width = n_embd == zeros.size(); + if (!same_width && !zeros_warned) { + SPC_WRN("target embeddings of size %zu do not fit the draft input width %zu, " + "the draft receives zero rows for them and drafts after multimodal input will be poor\n", + n_embd, zeros.size()); + zeros_warned = true; + } + const llama_embd embd = same_width ? t.embd : llama_embd{ zeros.data(), 1, zeros.size() }; + batch.add_embd(embd, t.pos.data(), t.seq_id, output); + } + } - const int ret = llama_decode(ctx_dft, batch_dft); + if (batch.size() == 0) { + return true; + } + + const int ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { SPC_ERR("failed to decode draft batch, ret = %d\n", ret); @@ -277,7 +308,7 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { void draft(common_speculative_draft_params_vec & dparams) override { auto & ctx_dft = params.ctx_dft; - common_batch_clear(batch); + batch.clear(); // keep track of which sequences are still drafting int n_drafting = 0; @@ -294,12 +325,12 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { drafting[seq_id] = true; common_sampler_reset(smpls[seq_id].get()); - common_batch_add(batch, dp.id_last, dp.pos0, { seq_id }, true); + batch.add(dp.id_last, dp.pos0, seq_id, true); } - int ret = llama_decode(ctx_dft, batch); + int ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { - SPC_ERR("llama_decode returned %d\n", ret); + SPC_ERR("llama_process returned %d\n", ret); return; } @@ -308,7 +339,7 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { while (n_drafting > 0) { int i_batch = 0; - common_batch_clear(batch); + batch.clear(); for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { if (!drafting[seq_id]) { @@ -353,17 +384,17 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { continue; } - common_batch_add(batch, id, dp.pos0 + i + 1, { seq_id }, true); + batch.add(id, dp.pos0 + i + 1, seq_id, true); } - if (batch.n_tokens == 0) { + if (batch.size() == 0) { break; } // evaluate the drafted tokens on the draft model - ret = llama_decode(ctx_dft, batch); + ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { - SPC_ERR("llama_decode[%d] returned %d\n", i, ret); + SPC_ERR("llama_process[%d] returned %d\n", i, ret); break; } @@ -423,7 +454,8 @@ struct common_speculative_impl_draft_simple : public common_speculative_impl { // encoder+decoder on n_accepted+1 rows). struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { common_params_speculative_draft params; - llama_batch batch; + common_batch batch; // decoder input, (token, g_embd) pairs + common_batch batch_enc; // encoder input, built from the extracted target features std::vector smpls; @@ -477,11 +509,8 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { n_embd_enc = (int32_t) target_layer_ids_n * n_embd_tgt; n_layer_tgt = llama_model_n_layer(model_tgt); - const int32_t n_b = (int32_t) llama_n_batch(ctx_dft); - batch = llama_batch_init(/*n_tokens=*/ n_b, /*embd=*/ n_embd_dec, /*n_seq_max=*/ 1); - // llama_batch_init allocates only one of token/embd; eagle3 decoder needs both. - // TODO: fix, how to call without malloc - batch.token = (llama_token *) malloc(sizeof(llama_token) * n_b); + batch = common_batch(ctx_dft); + batch_enc = common_batch(ctx_dft); smpls.resize(n_seq); for (auto & s : smpls) { @@ -543,12 +572,6 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { llama_sampler_free(backend_chains[seq_id]); } backend_chains.clear(); - - if (batch.token != nullptr) { - free(batch.token); - batch.token = nullptr; - } - llama_batch_free(batch); } void begin(llama_seq_id seq_id, const llama_tokens & prompt) override { @@ -567,16 +590,16 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { } } - bool process(const llama_batch & batch_in) override { - if (batch_in.n_tokens <= 0) { + bool process(const common_batch & batch_in) override { + if (batch_in.size() <= 0) { return true; } - if (batch_in.token == nullptr || batch_in.embd != nullptr) { + if (!batch_in.has_token() || batch_in.has_embd()) { return true; } - const int32_t n_tokens = batch_in.n_tokens; + const int32_t n_tokens = batch_in.size(); // i_batch_beg[seq] / i_batch_end[seq]: inclusive batch indices of this seq's // first/last token in batch_in. Assumes per-seq tokens are contiguous within @@ -584,8 +607,7 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { std::vector i_batch_beg(n_seq, -1); std::vector i_batch_end(n_seq, -1); for (int k = 0; k < n_tokens; ++k) { - GGML_ASSERT(batch_in.n_seq_id[k] == 1); - const llama_seq_id seq_id = batch_in.seq_id[k][0]; + const llama_seq_id seq_id = batch_in.tokens[k].seq_id; if (seq_id < 0 || seq_id >= (llama_seq_id) n_seq) { continue; } @@ -619,24 +641,23 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { g_embd_buf.resize((size_t) n_tokens * n_embd_dec); - // llama_encode() requires the full encoder batch to fit in n_ubatch. + // llama_process() requires the full encoder batch to fit in n_ubatch. // Allow batch > ubatch: eagle3's per-token encoder can be chunked safely. const int32_t n_ubatch_dft = (int32_t) llama_n_ubatch(ctx_dft); for (int32_t i = 0; i < n_tokens; i += n_ubatch_dft) { const int32_t n_chunk = std::min(n_ubatch_dft, n_tokens - i); - llama_batch enc_batch = { - /*.n_tokens =*/ n_chunk, - /*.token =*/ nullptr, - /*.embd =*/ features_buf.data() + (size_t) i * n_embd_enc, - /*.pos =*/ nullptr, - /*.n_seq_id =*/ nullptr, - /*.seq_id =*/ nullptr, - /*.logits =*/ nullptr, - }; - const int32_t rc = llama_encode(ctx_dft, enc_batch); + // the per-token encoder does not use positions, generate placeholder ones from the memory state + batch_enc.clear(); + llama_pos pos = llama_memory_seq_pos_max(llama_get_memory(ctx_dft), 0) + 1; + for (int32_t j = 0; j < n_chunk; ++j) { + batch_enc.add_embd({ features_buf.data() + (size_t) (i + j) * n_embd_enc, 1, (size_t) n_embd_enc }, &pos, 0, true); + pos++; + } + + const int32_t rc = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_ENCODE, batch_enc.get()); if (rc != 0) { - SPC_ERR("llama_encode(ctx_dft) failed rc=%d (n_tokens=%d, offset=%d)\n", + SPC_ERR("llama_process(ctx_dft) failed rc=%d (n_tokens=%d, offset=%d)\n", rc, (int) n_chunk, (int) i); return false; } @@ -664,7 +685,7 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { // deferred boundary, completed by the next process() or draft() call. // (c) refresh deferred state — stash this ubatch's full g_embd into verify_g, // update pending_g_last / pending_pos_last to the last row. - common_batch_clear(batch); + batch.clear(); for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { const int32_t beg = i_batch_beg[seq_id]; @@ -679,36 +700,34 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { // 2) pending_pos_last + 1 == pos[beg] // 3) pending_pos_last > dft_pos_max // TODO: is this check needed? const llama_pos pending_pos = pending_pos_last[seq_id]; - if (pending_pos >= 0 && pending_pos + 1 == batch_in.pos[beg]) { + if (pending_pos >= 0 && pending_pos + 1 == batch_in.tokens[beg].pos[0]) { const llama_pos dft_pos_max = llama_memory_seq_pos_max(llama_get_memory(ctx_dft), seq_id); if (pending_pos > dft_pos_max) { - common_batch_add(batch, batch_in.token[beg], pending_pos, { seq_id }, /*logits=*/ false); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd_dec, - pending_g_last[seq_id].data(), row_bytes); + const int32_t idx = batch.add(batch_in.tokens[beg].id, pending_pos, seq_id, /*output=*/ false); + batch.set_embd(idx, { pending_g_last[seq_id].data(), 1, (size_t) n_embd_dec }); } } for (int32_t k = beg; k < end; ++k) { - common_batch_add(batch, batch_in.token[k + 1], batch_in.pos[k], { seq_id }, /*logits=*/ false); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd_dec, - g_embd + (size_t) k * n_embd_dec, row_bytes); + const int32_t idx = batch.add(batch_in.tokens[k + 1].id, batch_in.tokens[k].pos[0], seq_id, /*output=*/ false); + batch.set_embd(idx, { g_embd + (size_t) k * n_embd_dec, 1, (size_t) n_embd_dec }); } // refresh deferred state const int32_t n_rows = end - beg + 1; - verify_pos_first[seq_id] = batch_in.pos[beg]; - pending_pos_last[seq_id] = batch_in.pos[end]; + verify_pos_first[seq_id] = batch_in.tokens[beg].pos[0]; + pending_pos_last[seq_id] = batch_in.tokens[end].pos[0]; verify_g_rows[seq_id] = n_rows; verify_g[seq_id].resize((size_t) n_rows * n_embd_dec, 0.0f); std::memcpy(verify_g[seq_id].data(), g_embd + (size_t) beg * n_embd_dec, row_bytes * n_rows); std::memcpy(pending_g_last[seq_id].data(), g_embd + (size_t) end * n_embd_dec, row_bytes); } - if (batch.n_tokens > 0) { - const int32_t rc = llama_decode(ctx_dft, batch); + if (batch.size() > 0) { + const int32_t rc = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (rc != 0) { - SPC_ERR("llama_decode(ctx_dft) failed rc=%d (n_tokens=%d, ubatch_pos[0]=%d)\n", - rc, (int) batch.n_tokens, (int) batch_in.pos[0]); + SPC_ERR("llama_process(ctx_dft) failed rc=%d (n_tokens=%d, ubatch_pos[0]=%d)\n", + rc, (int) batch.size(), (int) batch_in.tokens[0].pos[0]); return false; } } @@ -719,14 +738,12 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { void draft(common_speculative_draft_params_vec & dparams) override { auto & ctx_dft = params.ctx_dft; - common_batch_clear(batch); + batch.clear(); // keep track of which sequences are still drafting int n_drafting = 0; std::vector drafting(n_seq); - const size_t row_bytes = (size_t) n_embd_dec * sizeof(float); - // Complete the deferred boundary pair (dp.id_last, pending_g_last) at memory // pos pending_pos_last. dp.id_last is target's freshest sample (= corrected // token after verify, or first generated token after prefill), matching the @@ -747,19 +764,17 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { llama_memory_seq_rm(llama_get_memory(ctx_dft), seq_id, pending_pos_last[seq_id], -1); - common_batch_add(batch, dp.id_last, pending_pos_last[seq_id], { seq_id }, true); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd_dec, - pending_g_last[seq_id].data(), - row_bytes); + const int32_t idx = batch.add(dp.id_last, pending_pos_last[seq_id], seq_id, true); + batch.set_embd(idx, { pending_g_last[seq_id].data(), 1, (size_t) n_embd_dec }); } - if (batch.n_tokens == 0) { + if (batch.size() == 0) { return; } - int ret = llama_decode(ctx_dft, batch); + int ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { - SPC_ERR("llama_decode returned %d\n", ret); + SPC_ERR("llama_process returned %d\n", ret); return; } @@ -768,7 +783,7 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { while (n_drafting > 0) { int i_batch = 0; - common_batch_clear(batch); + batch.clear(); for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { if (!drafting[seq_id]) { @@ -814,17 +829,17 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { continue; } - common_batch_add(batch, id, pending_pos_last[seq_id] + (i + 1), { seq_id }, true); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd_dec, prenorm, row_bytes); + const int32_t idx = batch.add(id, pending_pos_last[seq_id] + (i + 1), seq_id, true); + batch.set_embd(idx, { prenorm, 1, (size_t) n_embd_dec }); } - if (batch.n_tokens == 0) { + if (batch.size() == 0) { break; } - ret = llama_decode(ctx_dft, batch); + ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { - SPC_ERR("llama_decode[%d] returned %d\n", i, ret); + SPC_ERR("llama_process[%d] returned %d\n", i, ret); break; } @@ -908,8 +923,10 @@ struct common_speculative_impl_draft_eagle3 : public common_speculative_impl { struct common_speculative_impl_draft_dflash : public common_speculative_impl { common_params_speculative_draft params; - llama_batch batch; // noise tokens - llama_batch batch_inject; // target features for KV cache injection + common_batch batch; // noise tokens + common_batch batch_inject; // target features for KV cache injection + + std::vector features_buf; // [n_chunk, n_embd_enc] gathered target features std::vector smpls; @@ -1005,15 +1022,11 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { } this->n_max = this->params.n_max; - batch = llama_batch_init(llama_n_batch(ctx_dft), 0, n_seq); - batch_inject = llama_batch_init(llama_n_ubatch(ctx_dft), n_embd_enc, n_seq); + batch = common_batch(ctx_dft); + batch_inject = common_batch(ctx_dft); - // embd batches on an M-RoPE draft need 4 position rows per token + // embd batches on an M-RoPE draft carry 4 position rows per token is_mrope = llama_model_rope_type(model_dft) == LLAMA_ROPE_TYPE_MROPE; - if (is_mrope) { - free(batch_inject.pos); - batch_inject.pos = (llama_pos *) malloc(sizeof(llama_pos) * 4 * llama_n_batch(ctx_dft)); - } smpls.resize(n_seq); for (auto & s : smpls) { @@ -1062,9 +1075,6 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { llama_sampler_free(backend_chains[seq_id]); } backend_chains.clear(); - - llama_batch_free(batch); - llama_batch_free(batch_inject); } void begin(llama_seq_id seq_id, const llama_tokens & prompt) override { @@ -1085,8 +1095,8 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { } } - bool process(const llama_batch & batch_in) override { - if (batch_in.n_tokens <= 0) { + bool process(const common_batch & batch_in) override { + if (batch_in.size() <= 0) { return true; } @@ -1094,20 +1104,19 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { // produce the target-layer features used to seed the draft KV cache, so // embeddings are injected too, except the pinned ones skipped below. // TODO: revisit after https://github.com/ggml-org/llama.cpp/pull/24669 is merged - const bool has_tokens = batch_in.token != nullptr; - const bool has_embeddings = batch_in.embd != nullptr; + const bool has_tokens = batch_in.has_token(); + const bool has_embeddings = batch_in.has_embd(); if (has_tokens == has_embeddings) { return true; } - const int32_t n_tokens = batch_in.n_tokens; + const int32_t n_tokens = batch_in.size(); // per-seq inclusive batch range (assumes each seq's tokens are contiguous in the batch) std::vector i_batch_beg(n_seq, -1); std::vector i_batch_end(n_seq, -1); for (int32_t k = 0; k < n_tokens; ++k) { - GGML_ASSERT(batch_in.n_seq_id[k] == 1); - const llama_seq_id seq_id = batch_in.seq_id[k][0]; + const llama_seq_id seq_id = batch_in.tokens[k].seq_id; if (seq_id < 0 || seq_id >= (llama_seq_id) n_seq) { continue; } @@ -1130,7 +1139,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { // an M-RoPE image pins all its rows to one position, so a windowed draft // cache cannot free cells for it - skip it, the draft can jump over the gap - const bool pos_pinned = batch_in.pos[i_batch_beg[seq_id]] == batch_in.pos[i_batch_end[seq_id]]; + const bool pos_pinned = batch_in.tokens[i_batch_beg[seq_id]].pos[0] == batch_in.tokens[i_batch_end[seq_id]].pos[0]; if (has_embeddings && n_rows > 1 && pos_pinned) { continue; } @@ -1140,34 +1149,28 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { // gather target features per extract layer; the fused decode encodes and // injects them into the K/V cache at the target positions - batch_inject.n_tokens = n_chunk; + features_buf.resize((size_t) n_chunk * n_embd_enc); for (uint32_t k = 0; k < target_layer_ids_n; ++k) { const float * layer = llama_get_embeddings_layer_inp(ctx_tgt, (uint32_t) target_layer_ids[k]); if (!layer) { GGML_ABORT("DFlash: target layer %d input not extracted.", target_layer_ids[k]); } for (int32_t i = 0; i < n_chunk; ++i) { - float * dst = batch_inject.embd + (size_t) i * n_embd_enc + k * (size_t) n_embd_tgt; + float * dst = features_buf.data() + (size_t) i * n_embd_enc + k * (size_t) n_embd_tgt; const float * src = layer + (size_t) (i_batch_beg[seq_id] + offset + i) * n_embd_tgt; std::memcpy(dst, src, (size_t) n_embd_tgt * sizeof(float)); } } + batch_inject.clear(); for (int32_t i = 0; i < n_chunk; ++i) { - const llama_pos p = batch_in.pos[i_batch_beg[seq_id] + offset + i]; - batch_inject.pos[i] = p; - if (is_mrope) { - batch_inject.pos[1 * n_chunk + i] = p; - batch_inject.pos[2 * n_chunk + i] = p; - batch_inject.pos[3 * n_chunk + i] = 0; - } - batch_inject.n_seq_id[i] = 1; - batch_inject.seq_id[i][0] = seq_id; - batch_inject.logits[i] = false; + const llama_pos p = batch_in.tokens[i_batch_beg[seq_id] + offset + i].pos[0]; + const llama_pos pos_arr[4] = { p, p, p, 0 }; + batch_inject.add_embd({ features_buf.data() + (size_t) i * n_embd_enc, 1, (size_t) n_embd_enc }, pos_arr, seq_id, false); } - const int32_t rc = llama_decode(ctx_dft, batch_inject); + const int32_t rc = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch_inject.get()); if (rc != 0) { - LOG_ERR("%s: llama_decode(ctx_dft) failed rc=%d (n_tokens=%d, offset=%d)\n", + LOG_ERR("%s: llama_process(ctx_dft) failed rc=%d (n_tokens=%d, offset=%d)\n", __func__, rc, (int) n_chunk, (int) offset); return false; } @@ -1180,7 +1183,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { void draft(common_speculative_draft_params_vec & dparams) override { auto & ctx_dft = params.ctx_dft; - common_batch_clear(batch); + batch.clear(); // build one batch holding every drafting sequence's noise block into a single decode) // record where each block starts and its size @@ -1200,21 +1203,21 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { const int32_t n_draft = params.n_max; const int32_t n_block_tokens = n_draft + (is_dspark && sample_from_anchor ? 0 : 1); - i_block_beg[seq_id] = batch.n_tokens; + i_block_beg[seq_id] = batch.size(); n_block [seq_id] = n_block_tokens; for (int32_t i = 0; i < n_block_tokens; ++i) { - common_batch_add(batch, i == 0 ? dp.id_last : mask_token_id, n + i, { seq_id }, !is_dflash2); + batch.add(i == 0 ? dp.id_last : mask_token_id, n + i, seq_id, !is_dflash2); } } - if (batch.n_tokens == 0) { + if (batch.size() == 0) { return; } // decode all sequence's noise block in a single batch - int ret = llama_decode(ctx_dft, batch); + int ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { - LOG_WRN("%s: llama_decode returned %d\n", __func__, ret); + LOG_WRN("%s: llama_process returned %d\n", __func__, ret); return; } @@ -1328,7 +1331,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl { struct common_speculative_impl_draft_mtp : public common_speculative_impl { common_params_speculative_draft params; // reuses the draft-model params slot (ctx_tgt/ctx_dft) - llama_batch batch; + common_batch batch; std::vector smpls; @@ -1384,11 +1387,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { ctx_dft ? "yes" : "no", common_speculative_get_devices_str(this->params.devices).c_str()); - const int32_t n_b = (int32_t) llama_n_batch(ctx_dft); - batch = llama_batch_init(/*n_tokens=*/ n_b, /*embd=*/ n_embd, /*n_seq_max=*/ 1); - // llama_batch_init allocates only one of token/embd; MTP needs both. - // TODO: fix, how to call without malloc - batch.token = (llama_token *) malloc(sizeof(llama_token) * n_b); + batch = common_batch(ctx_dft); smpls.resize(n_seq); for (auto & s : smpls) { @@ -1453,12 +1452,6 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { llama_sampler_free(backend_chains[seq_id]); } backend_chains.clear(); - - if (batch.token != nullptr) { - free(batch.token); - batch.token = nullptr; - } - llama_batch_free(batch); } void begin(llama_seq_id seq_id, const llama_tokens & prompt) override { @@ -1473,23 +1466,23 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { if (pos_max < N - 1 && !is_mem_shared) { SPC_WRN("ctx_dft pos_max=%d < N-1=%d - " "process() hook may not have run on every prefill ubatch " - "(need_embd / logits=1 on every prompt position?). " + "(need_embd / output flag on every prompt position?). " "Drafts may degrade.\n", (int) pos_max, N - 1); } } - bool process(const llama_batch & batch_in) override { - if (batch_in.n_tokens <= 0) { + bool process(const common_batch & batch_in) override { + if (batch_in.size() <= 0) { return true; } // TODO: how to make it work with vision tokens? - if (batch_in.token == nullptr || batch_in.embd != nullptr) { + if (!batch_in.has_token() || batch_in.has_embd()) { return true; } - const int32_t n_tokens = batch_in.n_tokens; + const int32_t n_tokens = batch_in.size(); // remember the first and last batch index for each sequence std::fill(i_batch_beg.begin(), i_batch_beg.end(), -1); @@ -1497,9 +1490,7 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { for (int k = 0; k < n_tokens; ++k) { for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { - GGML_ASSERT(batch_in.n_seq_id[k] == 1); - - if (batch_in.seq_id[k][0] == seq_id) { + if (batch_in.tokens[k].seq_id == seq_id) { i_batch_end[seq_id] = k; if (i_batch_beg[seq_id] < 0) { i_batch_beg[seq_id] = k; @@ -1515,33 +1506,26 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { // if kv is shared with target (e.g Gemma4), then we can skip this catch-up decode if (!is_mem_shared) { - common_batch_clear(batch); - - for (int k = 0; k < n_tokens; ++k) { - common_batch_add(batch, batch_in.token[k], batch_in.pos[k], { batch_in.seq_id[k][0] }, 0); - } + batch.clear(); - // shift the tgt embeddings to the right by one position + // pair each token with the tgt embedding shifted right by one position, and + // the first token of each sequence with the pending embedding from a previous run // assumes that the tokens in the batch are sequential for each sequence // i.e. we cannot have seq_id like this: [0, 0, 0, 1, 1, 0, 1, 1] // ^--- this is a problem // TODO:this is generally true, but would be nice to assert it - { - const float * h_tgt = llama_get_embeddings_nextn(ctx_tgt); - std::memcpy(batch.embd + (size_t) 1 * n_embd, h_tgt, row_bytes * (n_tokens-1)); - } + const float * h_tgt = llama_get_embeddings_nextn(ctx_tgt); - // fill the pending embeddings from a previous run - auto set_h = [&](int idx, const float * h_row) { - std::memcpy(batch.embd + (size_t) idx * n_embd, h_row, row_bytes); - }; + for (int k = 0; k < n_tokens; ++k) { + const llama_seq_id seq_id = batch_in.tokens[k].seq_id; - for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { - if (i_batch_beg[seq_id] < 0) { - continue; - } + const int32_t idx = batch.add(batch_in.tokens[k].id, batch_in.tokens[k].pos[0], seq_id, false); + + const float * h_row = k == i_batch_beg[seq_id] + ? pending_h[seq_id].data() + : h_tgt + (size_t) (k - 1) * n_embd; - set_h(i_batch_beg[seq_id], pending_h[seq_id].data()); + batch.set_embd(idx, { h_row, 1, (size_t) n_embd }); } auto * mem_dft = llama_get_memory(ctx_dft); @@ -1554,15 +1538,15 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { if (i_batch_beg[seq_id] < 0) { continue; } - llama_memory_seq_rm(mem_dft, seq_id, batch_in.pos[i_batch_beg[seq_id]], -1); + llama_memory_seq_rm(mem_dft, seq_id, batch_in.tokens[i_batch_beg[seq_id]].pos[0], -1); } llama_set_nextn_layer_offset(ctx_dft, head); } - const int32_t rc = llama_decode(ctx_dft, batch); + const int32_t rc = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (rc != 0) { - SPC_ERR("llama_decode(ctx_dft) head=%d failed rc=%d (pos=%d)\n", - head, (int) rc, (int) batch_in.pos[0]); + SPC_ERR("llama_process(ctx_dft) head=%d failed rc=%d (pos=%d)\n", + head, (int) rc, (int) batch_in.tokens[0].pos[0]); ok = false; break; } @@ -1600,14 +1584,12 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { void draft(common_speculative_draft_params_vec & dparams) override { auto & ctx_dft = params.ctx_dft; - common_batch_clear(batch); + batch.clear(); // keep track of which sequences are still drafting int n_drafting = 0; std::vector drafting(n_seq); - const size_t row_bytes = (size_t) n_embd * sizeof(float); - for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { auto & dp = dparams[seq_id]; @@ -1619,10 +1601,10 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { drafting[seq_id] = true; common_sampler_reset(smpls[seq_id].get()); - common_batch_add(batch, dp.id_last, dp.pos0, { seq_id }, true); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd, pending_h[seq_id].data(), row_bytes); + const int32_t idx = batch.add(dp.id_last, dp.pos0, seq_id, true); + batch.set_embd(idx, { pending_h[seq_id].data(), 1, (size_t) n_embd }); - i_last[seq_id] = batch.n_tokens - 1; + i_last[seq_id] = idx; if (chain_heads) { chain_h[seq_id].assign(pending_h[seq_id].begin(), pending_h[seq_id].end()); @@ -1648,16 +1630,16 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { llama_set_nextn_layer_offset(ctx_dft, i); } - int ret = llama_decode(ctx_dft, batch); + int ret = llama_process(ctx_dft, LLAMA_PROCESS_TYPE_DECODE, batch.get()); if (ret != 0) { - SPC_ERR("llama_decode[%d] returned %d\n", i, ret); + SPC_ERR("llama_process[%d] returned %d\n", i, ret); break; } // rebuild the batch for the next step: the growing-KV paths re-add only the // new token (the KV already holds the prefix), while chained heads re-add the // whole prefix at the next head. dropped sequences are simply not re-added. - common_batch_clear(batch); + batch.clear(); for (llama_seq_id seq_id = 0; seq_id < (llama_seq_id) n_seq; ++seq_id) { if (!drafting[seq_id]) { @@ -1708,24 +1690,24 @@ struct common_speculative_impl_draft_mtp : public common_speculative_impl { const int n_rows = (int) result.size() + 1; // id_last + tokens drafted so far for (int t = 0; t < n_rows; ++t) { const llama_token tok = (t == 0) ? dp.id_last : result[t - 1]; - common_batch_add(batch, tok, dp.pos0 + t, { seq_id }, t == n_rows - 1); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd, - chain_h[seq_id].data() + (size_t) t * n_embd, row_bytes); + const int32_t idx = batch.add(tok, dp.pos0 + t, seq_id, t == n_rows - 1); + batch.set_embd(idx, { chain_h[seq_id].data() + (size_t) t * n_embd, 1, (size_t) n_embd }); + i_last[seq_id] = idx; } } else if (is_mem_shared) { // note: with shared memory (e.g. Gemma4 assistants) we use the same position for all draft tokens // ref: https://github.com/huggingface/transformers/blob/effde20942e3f82a1b97449f60b3a48c5ff96145/docs/source/en/model_doc/gemma4_assistant.md?plain=1#L36-L37 - common_batch_add(batch, id, dp.pos0, { seq_id }, true); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd, h_row, row_bytes); + const int32_t idx = batch.add(id, dp.pos0, seq_id, true); + batch.set_embd(idx, { h_row, 1, (size_t) n_embd }); + i_last[seq_id] = idx; } else { - common_batch_add(batch, id, dp.pos0 + i + 1, { seq_id }, true); - std::memcpy(batch.embd + (size_t) (batch.n_tokens - 1) * n_embd, h_row, row_bytes); + const int32_t idx = batch.add(id, dp.pos0 + i + 1, seq_id, true); + batch.set_embd(idx, { h_row, 1, (size_t) n_embd }); + i_last[seq_id] = idx; } - - i_last[seq_id] = batch.n_tokens - 1; } - if (batch.n_tokens == 0) { + if (batch.size() == 0) { break; } @@ -1787,7 +1769,7 @@ struct common_speculative_impl_ngram_simple : public common_speculative_impl { // noop } - bool process(const llama_batch & /*batch*/) override { + bool process(const common_batch & /*batch*/) override { // TODO: implement return true; } @@ -1835,7 +1817,7 @@ struct common_speculative_impl_ngram_map_k : public common_speculative_impl { common_ngram_map_begin(config[seq_id], prompt); } - bool process(const llama_batch & /*batch*/) override { + bool process(const common_batch & /*batch*/) override { // TODO: implement return true; } @@ -1993,7 +1975,7 @@ struct common_speculative_impl_ngram_mod : public common_speculative_impl { sinfo.n_draft_last = result.size(); } - bool process(const llama_batch & /*batch*/) override { + bool process(const common_batch & /*batch*/) override { // TODO: implement return true; } @@ -2155,7 +2137,7 @@ struct common_speculative_impl_ngram_cache : public common_speculative_impl { } } - bool process(const llama_batch & /*batch*/) override { + bool process(const common_batch & /*batch*/) override { // TODO: implement return true; } @@ -2181,6 +2163,9 @@ struct common_speculative_impl_ngram_cache : public common_speculative_impl { struct common_speculative { common_speculative_draft_params_vec dparams; + // the target context, used to convert legacy llama_batch inputs + llama_context * ctx_tgt = nullptr; + // list of implementations to use and their states std::vector> impls; @@ -2726,6 +2711,7 @@ common_speculative * common_speculative_init(common_params_speculative & params, common_speculative_ptr result(new common_speculative { /* .dparams = */ common_speculative_draft_params_vec(n_seq), + /* .ctx_tgt = */ params.draft.ctx_tgt, /* .impls = */ std::move(impls), /* .impl_last = */ std::vector(n_seq, nullptr), /* .synth_probs = */ {}, @@ -2789,6 +2775,17 @@ void common_speculative_begin(common_speculative * spec, llama_seq_id seq_id, co } bool common_speculative_process(common_speculative * spec, const llama_batch & batch) { + if (spec == nullptr) { + return true; + } + + // ngram-only setups have no target context, they do not read the batch anyway + const common_batch tmp = spec->ctx_tgt ? common_batch_from_llama_batch(spec->ctx_tgt, batch) : common_batch(); + + return common_speculative_process(spec, tmp); +} + +bool common_speculative_process(common_speculative * spec, const common_batch & batch) { bool result = true; if (spec == nullptr) { diff --git a/common/speculative.h b/common/speculative.h index c968750e2d80..211fcdabd1c9 100644 --- a/common/speculative.h +++ b/common/speculative.h @@ -77,6 +77,9 @@ common_speculative_draft_params & common_speculative_get_draft_params(common_spe void common_speculative_begin(common_speculative * spec, llama_seq_id seq_id, const llama_tokens & prompt); // process the batch and update the internal state of the speculative context +bool common_speculative_process(common_speculative * spec, const common_batch & batch); + +// legacy llama_batch input, converted with common_batch_from_llama_batch() bool common_speculative_process(common_speculative * spec, const llama_batch & batch); // generate drafts for the sequences specified with `common_speculative_get_draft_params` diff --git a/docs/autoparser.md b/docs/autoparser.md index 2a7ea00b4f03..0f8b70171114 100644 --- a/docs/autoparser.md +++ b/docs/autoparser.md @@ -509,6 +509,7 @@ The following templates have active tests in `tests/test-chat.cpp`: | Kimi-K2 / Kimi-K2-Instruct | JSON_NATIVE | JSON tools with special markers | | Llama 3.1/3.2/3.3 | JSON_NATIVE | Standard Llama tool format | | OpenAI GPT-OSS | Specialized | Channel-based (dedicated handler) | +| LLM-jp-4.1 | Specialized | GPT-OSS dialect (dedicated handler) | | Apriel 1.5 | JSON_NATIVE | `` wrapper with JSON array | | Apriel 1.6 Thinker | Reasoning | Implicit reasoning start | | Mistral Small 3.2 | JSON_NATIVE | `[TOOL_CALLS]func[ARGS]{...}` with call ID | diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index f9da1ab3edbb..7f2b551829fb 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -52,6 +52,10 @@ The packages for FP32 and FP16 would have different accuracy and performance on ## News +- 2026.09 + - Update the CI build environment for oneAPI 2026.1 (unified oneAPI Toolkit). oneDNN is removed from the Deep Learning Essentials package in 2026.0, so the CI now uses the oneAPI Toolkit installer which still includes oneDNN. + - oneAPI 2026.1 improves the SYCL build performance: measured with the same code on Arc B570, prompt processing 1331 vs 434 t/s (3.1x) vs the 2025.3-based release build. + - 2026.04-05 - Optimize mul_mat by reorder feature for data type: Q4_K, Q5_K, Q6_K, Q8_0. - Fused MoE. @@ -257,7 +261,7 @@ Platform #0: Intel(R) OpenCL HD Graphics `-- Device #0: Intel(R) Iris(R) Xe Graphics [0x9a49] ``` -2. **Install Intel® oneAPI Base toolkit** +2. **Install Intel® oneAPI Toolkit** SYCL backend depends on: - Intel® oneAPI DPC++/C++ compiler/running-time. @@ -267,11 +271,11 @@ SYCL backend depends on: - **For Intel GPU** -All above are included in both **Intel® oneAPI Base toolkit** and **Intel® Deep Learning Essentials** packages. +With the 2026.0 release, the Intel® oneAPI Base toolkit and the HPC toolkit are combined into the **Intel® oneAPI Toolkit**, and **oneDNN is removed from the Intel® Deep Learning Essentials** package (oneDNN is distributed separately since then). The **Intel® oneAPI Toolkit** includes oneDNN until 2027.0. -It's recommended to install **Intel® Deep Learning Essentials** which only provides the necessary libraries with less size. +It's recommended to install the **Intel® oneAPI Toolkit**. -The **Intel® oneAPI Base toolkit** and **Intel® Deep Learning Essentials** can be obtained from the official [Intel® oneAPI Base Toolkit](https://www.intel.com/content/www/us/en/developer/tools/oneapi/base-toolkit.html) page. +The **Intel® oneAPI Toolkit** can be obtained from the official [Intel® oneAPI Toolkit](https://www.intel.com/content/www/us/en/developer/tools/oneapi/base-toolkit-download.html) page. Please follow the instructions for downloading and installing the Toolkit for Linux, and preferably keep the default installation values unchanged, notably the installation path *(`/opt/intel/oneapi` by default)*. @@ -281,6 +285,7 @@ Upon a successful installation, SYCL is enabled for the available Intel devices, |Verified release| |-| +|2026.1 | |2025.3.3 | |2025.2.1| |2025.1| diff --git a/docs/function-calling.md b/docs/function-calling.md index 850b59ce7aa7..28eecbe25749 100644 --- a/docs/function-calling.md +++ b/docs/function-calling.md @@ -16,6 +16,7 @@ Function calling is supported for all models (see https://github.com/ggml-org/ll - Firefunction v2 - Command R7B - DeepSeek R1 (WIP / seems reluctant to call any tools?) + - GPT-OSS (Harmony), LLM-jp-4.1 (Harmony dialect) - Generic tool call is supported when the template isn't recognized by native format handlers (you'll see `Chat format: Generic` in the logs). - Use `--chat-template-file` to override the template when appropriate (see examples below) diff --git a/examples/simple-chat/simple-chat.cpp b/examples/simple-chat/simple-chat.cpp index 97e9dc9842f5..30a0966e07e1 100644 --- a/examples/simple-chat/simple-chat.cpp +++ b/examples/simple-chat/simple-chat.cpp @@ -65,8 +65,7 @@ int main(int argc, char ** argv) { } }, nullptr); - // load dynamic backends - ggml_backend_load_all(); + llama_backend_init(); // initialize the model llama_model_params model_params = llama_model_default_params(); diff --git a/examples/simple/simple.cpp b/examples/simple/simple.cpp index 9f0a25d713f4..982a4d860084 100644 --- a/examples/simple/simple.cpp +++ b/examples/simple/simple.cpp @@ -77,9 +77,7 @@ int main(int argc, char ** argv) { } } - // load dynamic backends - - ggml_backend_load_all(); + llama_backend_init(); // initialize the model diff --git a/examples/speculative-simple/speculative-simple.cpp b/examples/speculative-simple/speculative-simple.cpp index 863af5a2c71a..81aa106f14fd 100644 --- a/examples/speculative-simple/speculative-simple.cpp +++ b/examples/speculative-simple/speculative-simple.cpp @@ -228,7 +228,6 @@ int main(int argc, char ** argv) { common_batch_add(batch_tgt, draft[i], n_past + i, { seq_id }, true); } - //LOG_DBG("target batch: %s\n", string_from(ctx_tgt, batch_tgt).c_str()); llama_decode(ctx_tgt, batch_tgt); } diff --git a/ggml/src/ggml-backend.cpp b/ggml/src/ggml-backend.cpp index 20bf965017e3..273dc92b2571 100644 --- a/ggml/src/ggml-backend.cpp +++ b/ggml/src/ggml-backend.cpp @@ -1372,30 +1372,6 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra const int src_backend_id = sched->hv_tensor_backend_ids[src_id]; GGML_ASSERT(src_backend_id != -1); // all inputs should be assigned by now - if (src->flags & GGML_TENSOR_FLAG_INPUT && sched->n_copies > 1) { - if (tensor_id_copy(src_id, src_backend_id, 0) == NULL) { - ggml_backend_t backend = sched->backends[src_backend_id]; - for (int c = 0; c < sched->n_copies; c++) { - struct ggml_tensor * tensor_copy; - if (c == sched->cur_copy) { - tensor_copy = src; // use the original tensor as the current copy - } else { - tensor_copy = ggml_dup_tensor_layout(sched->ctx, src); - ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c); - } - ggml_set_input(tensor_copy); - ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor - tensor_id_copy(src_id, src_backend_id, c) = tensor_copy; - SET_CAUSE(tensor_copy, "4.cpy"); - } - int n_graph_inputs = sched->n_graph_inputs++; - if (n_graph_inputs >= sched->graph_inputs_capacity) { - ggml_backend_sched_graph_inputs_grow(sched); - } - sched->graph_inputs[n_graph_inputs] = src; - } - } - if (src_backend_id != cur_backend_id && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) { // create a copy of the input in the split's backend if (tensor_id_copy(src_id, cur_backend_id, 0) == NULL) { @@ -1428,6 +1404,46 @@ void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgra ggml_backend_sched_print_assignments(sched, graph); } + // pass 6: collect all input tensors into graph_inputs + // this includes inputs not consumed by any node (e.g. the embeddings input of a text-only batch) so that + // the graph composition does not depend on which inputs are used, which would otherwise cause graph + // reallocations when switching between different types of batches [GGML_SCHED_DEBUG_REALLOC] + if (sched->n_copies > 1) { + for (int i = 0; i < graph->n_leafs; i++) { + struct ggml_tensor * leaf = graph->leafs[i]; + if ((leaf->flags & GGML_TENSOR_FLAG_INPUT) == 0) { + continue; + } + + const size_t leaf_id = hash_id(leaf); + const int leaf_backend_id = tensor_backend_id(leaf); + GGML_ASSERT(leaf_backend_id != -1); // all leafs should be assigned by now + + if (tensor_id_copy(leaf_id, leaf_backend_id, 0) == NULL) { + ggml_backend_t backend = sched->backends[leaf_backend_id]; + for (int c = 0; c < sched->n_copies; c++) { + struct ggml_tensor * tensor_copy; + if (c == sched->cur_copy) { + tensor_copy = leaf; // use the original tensor as the current copy + } else { + tensor_copy = ggml_dup_tensor_layout(sched->ctx, leaf); + ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), leaf->name, c); + } + ggml_set_input(tensor_copy); + ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor + tensor_id_copy(leaf_id, leaf_backend_id, c) = tensor_copy; + SET_CAUSE(tensor_copy, "6.cpy"); + } + } + + int n_graph_inputs = sched->n_graph_inputs++; + if (n_graph_inputs >= sched->graph_inputs_capacity) { + ggml_backend_sched_graph_inputs_grow(sched); + } + sched->graph_inputs[n_graph_inputs] = leaf; + } + } + // swap node_backend_ids and leaf _backend_ids with prevs { int * tmp = sched->node_backend_ids; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index ba00a0a73ed8..a07e1f96373c 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -9037,6 +9037,11 @@ static void ggml_compute_forward_flash_attn_ext_tiled( simd_gemm(KQ, (const float *)Q_q, K_f32, Q_TILE_SZ, DK, KV_TILE_SZ); ggml_vec_scale_f32(Q_TILE_SZ * KV_TILE_SZ, KQ, scale); + if (logit_softcap != 0.0f) { + ggml_vec_tanh_f32(Q_TILE_SZ * KV_TILE_SZ, KQ, KQ); + ggml_vec_scale_f32(Q_TILE_SZ * KV_TILE_SZ, KQ, logit_softcap); + } + // Set padded KQ entries to -inf so softmax gives them zero weight if (kv_tile < KV_TILE_SZ) { for (int tq = 0; tq < Q_TILE_SZ; tq++) { @@ -9046,11 +9051,6 @@ static void ggml_compute_forward_flash_attn_ext_tiled( } } - if (logit_softcap != 0.0f) { - ggml_vec_tanh_f32(Q_TILE_SZ * KV_TILE_SZ, KQ, KQ); - ggml_vec_scale_f32(Q_TILE_SZ * KV_TILE_SZ, KQ, logit_softcap); - } - if (mask) { ggml_vec_add_f32(tile_rows * KV_TILE_SZ, KQ, KQ, mask32); } @@ -9320,7 +9320,7 @@ static void ggml_compute_forward_flash_attn_ext_f16( kv_is_f32_or_f16 && k->type == v->type && neq1 >= Q_TILE_SZ); -#ifdef GGML_SIMD +#if defined(GGML_SIMD) && !defined(__x86_64__) && !defined(_M_X64) #if defined(__ARM_FEATURE_SVE) const int64_t f32_epr = svcntw(); #else diff --git a/ggml/src/ggml-cpu/simd-gemm.h b/ggml/src/ggml-cpu/simd-gemm.h index 2ebd10051ed8..4b9396d54066 100644 --- a/ggml/src/ggml-cpu/simd-gemm.h +++ b/ggml/src/ggml-cpu/simd-gemm.h @@ -56,6 +56,56 @@ static inline void simd_gemm_ukernel( } } +template +static inline void simd_gemm_ukernel_tail( + float * GGML_RESTRICT C, + const float * GGML_RESTRICT A, + const float * GGML_RESTRICT B, + int K, int N, int cols) +{ +#if defined(__AVX512F__) + const __mmask16 mask = (1u << cols) - 1; + __m512 acc[RM]; + for (int64_t i = 0; i < RM; i++) { + acc[i] = _mm512_maskz_loadu_ps(mask, C + i * N); + } + for (int64_t kk = 0; kk < K; kk++) { + const __m512 b = _mm512_maskz_loadu_ps(mask, B + kk * N); + for (int64_t i = 0; i < RM; i++) { + acc[i] = _mm512_mask3_fmadd_ps(_mm512_set1_ps(A[i * K + kk]), b, acc[i], mask); + } + } + for (int64_t i = 0; i < RM; i++) { + _mm512_mask_storeu_ps(C + i * N, mask, acc[i]); + } +#elif defined(__AVX2__) + const __m256i mask = _mm256_cmpgt_epi32(_mm256_set1_epi32(cols), _mm256_setr_epi32(0, 1, 2, 3, 4, 5, 6, 7)); + __m256 acc[RM]; + for (int64_t i = 0; i < RM; i++) { + acc[i] = _mm256_maskload_ps(C + i * N, mask); + } + for (int64_t kk = 0; kk < K; kk++) { + const __m256 b = _mm256_maskload_ps(B + kk * N, mask); + for (int64_t i = 0; i < RM; i++) { + acc[i] = GGML_F32_VEC_FMA(acc[i], b, _mm256_set1_ps(A[i * K + kk])); + } + } + for (int64_t i = 0; i < RM; i++) { + _mm256_maskstore_ps(C + i * N, mask, acc[i]); + } +#else + for (int64_t j = 0; j < cols; j++) { + for (int64_t i = 0; i < RM; i++) { + float a = C[i * N + j]; + for (int64_t kk = 0; kk < K; kk++) { + a += A[i * K + kk] * B[kk * N + j]; + } + C[i * N + j] = a; + } + } +#endif +} + // C[M x N] += A[M x K] * B[K x N] static void simd_gemm( float * GGML_RESTRICT C, @@ -74,14 +124,8 @@ static void simd_gemm( for (; jj + KN <= N; jj += KN) { simd_gemm_ukernel(C + jj, A, B + jj, K, N); } - for (; jj < N; jj++) { - for (int64_t i = 0; i < GEMM_RM; i++) { - float a = C[i * N + jj]; - for (int64_t kk = 0; kk < K; kk++) { - a += A[i * K + kk] * B[kk * N + jj]; - } - C[i * N + jj] = a; - } + if (jj < N) { + simd_gemm_ukernel_tail(C + jj, A, B + jj, K, N, N - jj); } A += GEMM_RM * K; @@ -97,12 +141,8 @@ static void simd_gemm( for (; jj + KN <= N; jj += KN) { simd_gemm_ukernel<1, 1>(C + jj, A, B + jj, K, N); } - for (; jj < N; jj++) { - float a = C[jj]; - for (int64_t kk = 0; kk < K; kk++) { - a += A[kk] * B[kk * N + jj]; - } - C[jj] = a; + if (jj < N) { + simd_gemm_ukernel_tail<1>(C + jj, A, B + jj, K, N, N - jj); } A += K; diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh index fa0d347afb67..083d3228abdd 100644 --- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh +++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh @@ -1875,7 +1875,7 @@ static __global__ void flash_attn_ext_f16( #endif // defined(AMD_WMMA_AVAILABLE) #if defined(AMD_MFMA_AVAILABLE) - if (ncols1*ncols2 < 16 || (DKQ > 256 && ncols1*ncols2 < 64)) { + if (ncols1*ncols2 < 16 || (DKQ > 256 && ncols1*ncols2 < 32)) { NO_DEVICE_CODE; return; } diff --git a/ggml/src/ggml-metal/ggml-metal-fusion.cpp b/ggml/src/ggml-metal/ggml-metal-fusion.cpp index eac3bd6fef5a..e55b01503da7 100644 --- a/ggml/src/ggml-metal/ggml-metal-fusion.cpp +++ b/ggml/src/ggml-metal/ggml-metal-fusion.cpp @@ -418,7 +418,8 @@ static bool ggml_metal_fusion_check_topk_moe( const int64_t n_tokens = logits->ne[1]; const int64_t n_expert_used = ids->ne[0]; - if (n_expert <= 0 || n_tokens <= 0 || n_expert_used <= 0 || n_expert_used > n_expert || + // note: n_tokens == 0 (no-output batch) must match so that the packing stays shape-independent + if (n_expert <= 0 || n_expert_used <= 0 || n_expert_used > n_expert || n_expert > GGML_METAL_TOPK_MOE_MAX_EXPERTS || n_expert_used > GGML_METAL_TOPK_MOE_MAX_EXPERTS) { return false; } @@ -545,7 +546,8 @@ static bool ggml_metal_fusion_match_moe_reduce( const int64_t n_embd = experts->ne[0]; const int64_t n_tokens = experts->ne[2]; - if (n_embd <= 0 || n_tokens <= 0 || experts->ne[1] != n_expert_used || experts->ne[3] != 1 || + // note: n_tokens == 0 (no-output batch) must match so that the packing stays shape-independent + if (n_embd <= 0 || experts->ne[1] != n_expert_used || experts->ne[3] != 1 || weights->ne[0] != 1 || weights->ne[1] != n_expert_used || weights->ne[2] != n_tokens || weights->ne[3] != 1 || dst->ne[0] != n_embd || dst->ne[1] != n_tokens || dst->ne[2] != 1 || dst->ne[3] != 1) { return false; @@ -1079,16 +1081,17 @@ const ggml_metal_fusion * ggml_metal_fusion_next( // transparent) node sequence that the compute phase uses, so the returned count is the raw index // span from idx to the last matched node (intermediate views are packed along). int ggml_metal_fusion_max(const ggml_cgraph * gf, int idx) { - // an empty/view node cannot start a pattern - pack it alone - if (ggml_op_is_empty(gf->nodes[idx]->op) || ggml_is_empty(gf->nodes[idx])) { + // a view node cannot start a pattern - pack it alone + if (ggml_op_is_empty(gf->nodes[idx]->op)) { return 1; } - // collect the non-empty node indices starting at idx + // collect the non-view node indices starting at idx; 0-element tensors are included so + // that empty graphs pack like their non-empty counterparts (see ggml_metal_fusion_filter_ops) int idxs[GGML_METAL_FUSION_MAX]; int n_idxs = 0; for (int i = idx; i < gf->n_nodes && n_idxs < GGML_METAL_FUSION_MAX; i++) { - if (!ggml_op_is_empty(gf->nodes[i]->op) && !ggml_is_empty(gf->nodes[i])) { + if (!ggml_op_is_empty(gf->nodes[i]->op)) { idxs[n_idxs++] = i; } } diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 8bb70d07778d..eb1868f62e84 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -1242,7 +1242,7 @@ typedef struct { // widths at or above this use the threadgroup FWHT kernel, one row per threadgroup // with GGML_METAL_FWHT_TG_NT threads, instead of one row per simdgroup -#define GGML_METAL_FWHT_TG_MIN_N 1024 +#define GGML_METAL_FWHT_TG_MIN_N 512 #define GGML_METAL_FWHT_TG_NT 256 typedef struct { diff --git a/ggml/src/ggml-metal/kernels/misc.metal b/ggml/src/ggml-metal/kernels/misc.metal index d29786e1c634..fb27983db3e6 100644 --- a/ggml/src/ggml-metal/kernels/misc.metal +++ b/ggml/src/ggml-metal/kernels/misc.metal @@ -541,18 +541,18 @@ typedef decltype(kernel_fwht<64, half>) kernel_fwht_f16_t; template [[host_name("kernel_fwht_f32_64")]] kernel kernel_fwht_f32_t kernel_fwht<64, float>; template [[host_name("kernel_fwht_f32_128")]] kernel kernel_fwht_f32_t kernel_fwht<128, float>; template [[host_name("kernel_fwht_f32_256")]] kernel kernel_fwht_f32_t kernel_fwht<256, float>; -template [[host_name("kernel_fwht_f32_512")]] kernel kernel_fwht_f32_t kernel_fwht<512, float>; template [[host_name("kernel_fwht_f16_64")]] kernel kernel_fwht_f16_t kernel_fwht<64, half>; template [[host_name("kernel_fwht_f16_128")]] kernel kernel_fwht_f16_t kernel_fwht<128, half>; template [[host_name("kernel_fwht_f16_256")]] kernel kernel_fwht_f16_t kernel_fwht<256, half>; -template [[host_name("kernel_fwht_f16_512")]] kernel kernel_fwht_f16_t kernel_fwht<512, half>; +template [[host_name("kernel_fwht_f32_512")]] kernel kernel_fwht_f32_t kernel_fwht_tg<512, GGML_METAL_FWHT_TG_NT, float>; template [[host_name("kernel_fwht_f32_1024")]] kernel kernel_fwht_f32_t kernel_fwht_tg<1024, GGML_METAL_FWHT_TG_NT, float>; template [[host_name("kernel_fwht_f32_2048")]] kernel kernel_fwht_f32_t kernel_fwht_tg<2048, GGML_METAL_FWHT_TG_NT, float>; template [[host_name("kernel_fwht_f32_4096")]] kernel kernel_fwht_f32_t kernel_fwht_tg<4096, GGML_METAL_FWHT_TG_NT, float>; template [[host_name("kernel_fwht_f32_8192")]] kernel kernel_fwht_f32_t kernel_fwht_tg<8192, GGML_METAL_FWHT_TG_NT, float>; +template [[host_name("kernel_fwht_f16_512")]] kernel kernel_fwht_f16_t kernel_fwht_tg<512, GGML_METAL_FWHT_TG_NT, half>; template [[host_name("kernel_fwht_f16_1024")]] kernel kernel_fwht_f16_t kernel_fwht_tg<1024, GGML_METAL_FWHT_TG_NT, half>; template [[host_name("kernel_fwht_f16_2048")]] kernel kernel_fwht_f16_t kernel_fwht_tg<2048, GGML_METAL_FWHT_TG_NT, half>; template [[host_name("kernel_fwht_f16_4096")]] kernel kernel_fwht_f16_t kernel_fwht_tg<4096, GGML_METAL_FWHT_TG_NT, half>; diff --git a/ggml/src/ggml-openvino/ggml-openvino.cpp b/ggml/src/ggml-openvino/ggml-openvino.cpp index 02c5962238a7..b2d81df97d92 100644 --- a/ggml/src/ggml-openvino/ggml-openvino.cpp +++ b/ggml/src/ggml-openvino/ggml-openvino.cpp @@ -9,6 +9,7 @@ #include "ggml-quants.h" #include "ggml.h" +#include #include #include #include @@ -960,6 +961,24 @@ static bool has_view_op_input(const ggml_tensor * op) { return false; } +// OV slices whole elements per axis, so each stride must be a multiple of the next smaller one +// (e.g. a batch stride of m*nb[1] + pad bytes cannot be expressed and would be read wrongly). +static bool has_strides_on_element_grid(const ggml_tensor * t) { + std::vector strides; + for (int i = 0; i < GGML_MAX_DIMS; i++) { + if (t->ne[i] > 1) { + strides.push_back(t->nb[i]); + } + } + std::sort(strides.begin(), strides.end()); + for (size_t i = 1; i < strides.size(); i++) { + if (strides[i - 1] == 0 || strides[i] % strides[i - 1] != 0) { + return false; + } + } + return true; +} + static bool has_non_contiguous_view_input(const ggml_tensor * op) { for (int i = 0; i < GGML_MAX_SRC; i++) { if (op->src[i] == nullptr) { @@ -1540,6 +1559,9 @@ static ggml_openvino_op_support ggml_backend_openvino_device_supports_op_impl(gg if (supported_types.find(src->type) == supported_types.end()) { return {false, "src[" + std::to_string(i) + "] type " + std::string(ggml_type_name(src->type)) + " is not supported"}; } + if (!has_strides_on_element_grid(src)) { + return {false, "src[" + std::to_string(i) + "] strides are not multiples of each other"}; + } const bool is_supported_3d_moe_expert = op->op == GGML_OP_MUL_MAT_ID && i == 0 && (src->type == GGML_TYPE_MXFP4 || src->ne[3] == 1); if (ggml_is_quantized(src->type) && src->ne[2] != 1 && !is_supported_3d_moe_expert) { diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-common.h b/ggml/src/ggml-vulkan/ggml-vulkan-common.h index 4ae5fea7a856..125e7608a3dc 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan-common.h +++ b/ggml/src/ggml-vulkan/ggml-vulkan-common.h @@ -97,7 +97,7 @@ vk_pipeline ggml_vk_get_quantize_pipeline(ggml_backend_vk_context * ctx, ggml_ty void ggml_vk_quantize_q8_1(ggml_backend_vk_context * ctx, vk_context& subctx, const vk_subbuffer & in, const vk_subbuffer & out, uint32_t ne); void ggml_vk_dsv4_hc_comb(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * mixes, const ggml_tensor * scale, const ggml_tensor * base, ggml_tensor * dst); void ggml_vk_dsv4_hc_pre(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * weights, ggml_tensor * dst); -void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst); +void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst, const ggml_tensor * gate_scale_in = nullptr); void ggml_vk_mul_mat(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx); bool ggml_vk_use_mul_mat_vec_id(const struct ggml_cgraph * cgraph, int node_idx); void ggml_vk_mul_mat_id(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx); @@ -263,9 +263,36 @@ inline void ggml_vk_dispatch_pipeline(ggml_backend_vk_context* ctx, vk_context& GGML_ASSERT(pipeline->parameter_count == descriptor_buffer_infos.size()); GGML_ASSERT(pipeline->push_constant_size == push_constant_size(push_constants)); - vk::DescriptorSet& descriptor_set = ctx->descriptor_sets[ctx->descriptor_set_idx++]; - vk::WriteDescriptorSet write_descriptor_set{ descriptor_set, 0, 0, pipeline->parameter_count, vk::DescriptorType::eStorageBuffer, nullptr, descriptor_buffer_infos.begin() }; - ctx->device->device.updateDescriptorSets({ write_descriptor_set }, {}); + const uint32_t descriptor_set_idx = ctx->descriptor_set_idx++; + vk::DescriptorSet& descriptor_set = ctx->descriptor_sets[descriptor_set_idx]; + + // a new buffer can get the handle of a destroyed one, so drop all cached bindings after any destroy + const uint64_t destroy_count = ctx->device->buffer_destroy_count.load(std::memory_order_acquire); + if (ctx->descriptor_set_bindings_destroy_count != destroy_count) { + for (auto & b : ctx->descriptor_set_bindings) { + b.clear(); + } + ctx->descriptor_set_bindings_destroy_count = destroy_count; + } + + // skip the write if this set already holds these bindings from the last graph + std::vector & bindings = ctx->descriptor_set_bindings[descriptor_set_idx]; + bool same = !ctx->device->disable_descriptor_reuse && bindings.size() == descriptor_buffer_infos.size(); + if (same) { + size_t i = 0; + for (const vk::DescriptorBufferInfo & info : descriptor_buffer_infos) { + const vk::DescriptorBufferInfo & prev = bindings[i++]; + if (prev.buffer != info.buffer || prev.offset != info.offset || prev.range != info.range) { + same = false; + break; + } + } + } + if (!same) { + vk::WriteDescriptorSet write_descriptor_set{ descriptor_set, 0, 0, pipeline->parameter_count, vk::DescriptorType::eStorageBuffer, nullptr, descriptor_buffer_infos.begin() }; + ctx->device->device.updateDescriptorSets({ write_descriptor_set }, {}); + bindings.assign(descriptor_buffer_infos.begin(), descriptor_buffer_infos.end()); + } subctx->s->buffer->buf.pushConstants(pipeline->layout, vk::ShaderStageFlagBits::eCompute, 0, push_constant_size(push_constants), push_constant_data(push_constants)); subctx->s->buffer->buf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline->pipeline); diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h b/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h index f066d106444a..037bcee825c3 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h +++ b/ggml/src/ggml-vulkan/ggml-vulkan-push-constants.h @@ -205,6 +205,10 @@ struct vk_op_dsv4_hc_post_push_constants { uint32_t p_offset; uint32_t c_offset; uint32_t d_offset; + + uint32_t gate; + float gate_scale_in; + float gate_scale_out; }; struct vk_op_count_experts_push_constants { diff --git a/ggml/src/ggml-vulkan/ggml-vulkan-types.h b/ggml/src/ggml-vulkan/ggml-vulkan-types.h index 252359bf891c..ee50f51a3836 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan-types.h +++ b/ggml/src/ggml-vulkan/ggml-vulkan-types.h @@ -51,6 +51,8 @@ typedef struct VkPhysicalDeviceCooperativeMatrixDecodeVectorFeaturesNV { #include +#include + #include #include @@ -553,6 +555,15 @@ static constexpr std::initializer_list rms_norm_view_set_rows_pattern { static constexpr std::initializer_list rope_view_set_rows_pattern { GGML_OP_ROPE, GGML_OP_VIEW, GGML_OP_SET_ROWS }; +// scale_out*sigmoid(scale_in*x) as the hc_post weights (qwen4exp hc_combine) +static constexpr std::initializer_list hc_post_gate_pattern { GGML_OP_SCALE, GGML_OP_UNARY, GGML_OP_SCALE, GGML_OP_DSV4_HC_POST }; + +static constexpr std::initializer_list> hc_post_gate_edges { + { 1, 0, 0 }, // sigmoid->src[0] == scale + { 2, 0, 1 }, // scale->src[0] == sigmoid + { 3, 2, 2 }, // hc_post->src[2] == scale (post) +}; + static constexpr std::initializer_list> topk_moe_early_softmax_norm_edges { { 1, 0, 0 }, // reshape->src[0] == softmax { 2, 0, 0 }, // argsort->src[0] == softmax @@ -1014,6 +1025,8 @@ struct vk_device_struct { ggml_backend_buffer_type buffer_type; bool disable_fusion; + bool disable_descriptor_reuse; + std::atomic buffer_destroy_count {}; bool disable_host_visible_vidmem; bool allow_sysmem_fallback; bool disable_graph_optimize; @@ -1056,6 +1069,8 @@ struct vk_buffer_struct { } VK_LOG_DEBUG("~vk_buffer_struct(" << buffer << ", " << size << ")"); + // bump before destroying, so a thread that sees the buffer gone also sees the new count + device->buffer_destroy_count.fetch_add(1, std::memory_order_release); device->device.freeMemory(device_memory); device->device.destroyBuffer(buffer); } @@ -1267,6 +1282,9 @@ struct ggml_backend_vk_context { std::vector descriptor_pools; std::vector descriptor_sets; + // last bindings written to each set; descriptor_sets is append-only so an index always names the same set + std::vector> descriptor_set_bindings; + uint64_t descriptor_set_bindings_destroy_count {}; uint32_t descriptor_set_idx {}; uint32_t pipeline_descriptor_set_requirements {}; @@ -1284,6 +1302,7 @@ struct ggml_backend_vk_context { bool fused_topk_moe_scale {}; // QSA indexer gather+add+top_k fused into one radix-select bool fused_topk_qsa {}; + bool fused_hc_post_gate {}; rms_norm_mode fused_rms_norm_mode {RMS_NORM_COUNT}; // for GGML_VK_PERF_LOGGER diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 1521e508c97d..38ec1f9e82ab 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -881,6 +881,7 @@ void ggml_pipeline_allocate_descriptor_sets(ggml_backend_vk_context * ctx) { pool_idx++; } + ctx->descriptor_set_bindings.resize(ctx->descriptor_sets.size()); } static vk_command_buffer* ggml_vk_create_cmd_buffer(vk_device& device, vk_command_pool& p) { @@ -4885,6 +4886,8 @@ vk_device ggml_vk_get_device(size_t idx) { device->disable_fusion = getenv("GGML_VK_DISABLE_FUSION") != nullptr; + device->disable_descriptor_reuse = getenv("GGML_VK_DISABLE_DESCRIPTOR_REUSE") != nullptr; + device->add_rms_fusion = !device->disable_fusion && device->subgroup_arithmetic && device->vendor_id != VK_VENDOR_ID_INTEL; @@ -7167,7 +7170,7 @@ void ggml_vk_dsv4_hc_pre(ggml_backend_vk_context * ctx, vk_context& subctx, cons ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { x_buf, w_buf, d_buf }, pc, { n_embd, n_tokens, 1 }); } -void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst) { +void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, const ggml_tensor * x, const ggml_tensor * residual, const ggml_tensor * post, const ggml_tensor * comb, ggml_tensor * dst, const ggml_tensor * gate_scale_in) { VK_LOG_DEBUG("ggml_vk_dsv4_hc_post(" << x << ", " << residual << ", " << post << ", " << comb << ", " << dst << ")"); vk_pipeline pipeline = comb ? ctx->device->pipeline_dsv4_hc_post_f32 : ctx->device->pipeline_dsv4_hc_post_nocomb_f32; @@ -7180,7 +7183,9 @@ void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, con const vk_subbuffer x_buf = ggml_vk_tensor_subbuffer(ctx, x, true); const vk_subbuffer r_buf = ggml_vk_tensor_subbuffer(ctx, residual, true); - const vk_subbuffer p_buf = ggml_vk_tensor_subbuffer(ctx, post, true); + // with a fused gate, post is scale(sigmoid(scale(p_src))) and the shader applies it to p_src + const ggml_tensor * p_src = gate_scale_in ? gate_scale_in->src[0] : post; + const vk_subbuffer p_buf = ggml_vk_tensor_subbuffer(ctx, p_src, true); const vk_subbuffer c_buf = comb ? ggml_vk_tensor_subbuffer(ctx, comb, true) : x_buf; const vk_subbuffer d_buf = ggml_vk_tensor_subbuffer(ctx, dst, true); @@ -7188,12 +7193,15 @@ void ggml_vk_dsv4_hc_post(ggml_backend_vk_context * ctx, vk_context& subctx, con n_embd, n_tokens, ggml_vk_nb_elem(x, 0), ggml_vk_nb_elem(x, 1), ggml_vk_nb_elem(residual, 0), ggml_vk_nb_elem(residual, 1), ggml_vk_nb_elem(residual, 2), - ggml_vk_nb_elem(post, 0), ggml_vk_nb_elem(post, 1), + ggml_vk_nb_elem(p_src, 0), ggml_vk_nb_elem(p_src, 1), comb ? ggml_vk_nb_elem(comb, 0) : 0, comb ? ggml_vk_nb_elem(comb, 1) : 0, comb ? ggml_vk_nb_elem(comb, 2) : 0, ggml_vk_nb_elem(dst, 0), ggml_vk_nb_elem(dst, 1), ggml_vk_nb_elem(dst, 2), 0, 0, 0, 0, 0, + gate_scale_in ? 1u : 0u, + gate_scale_in ? ggml_get_op_params_f32(gate_scale_in, 0) : 1.0f, + gate_scale_in ? ggml_get_op_params_f32(post, 0) : 1.0f, }; - init_pushconst_tensor_offsets(ctx, pc, x, residual, post, comb, dst); + init_pushconst_tensor_offsets(ctx, pc, x, residual, p_src, comb, dst); ggml_vk_dispatch_pipeline(ctx, subctx, pipeline, { x_buf, r_buf, p_buf, c_buf, d_buf }, pc, { n_embd, n_tokens, 1 }); } @@ -12356,7 +12364,12 @@ bool ggml_vk_build_graph(ggml_backend_vk_context * ctx, ggml_cgraph * cgraph, in break; case GGML_OP_SCALE: - ggml_vk_scale(ctx, compute_ctx, src0, node); + if (ctx->fused_hc_post_gate) { + ggml_tensor * hc_post = cgraph->nodes[node_idx + ctx->num_additional_fused_ops]; + ggml_vk_dsv4_hc_post(ctx, compute_ctx, hc_post->src[0], hc_post->src[1], hc_post->src[2], hc_post->src[3], hc_post, node); + } else { + ggml_vk_scale(ctx, compute_ctx, src0, node); + } break; case GGML_OP_SQR: @@ -12836,6 +12849,7 @@ void ggml_vk_cleanup(ggml_backend_vk_context * ctx) { } ctx->descriptor_pools.clear(); ctx->descriptor_sets.clear(); + ctx->descriptor_set_bindings.clear(); ctx->compute_cmd_pool.destroy(ctx->device->device); if (ctx->device->async_use_transfer_queue) { @@ -13431,6 +13445,19 @@ static bool ggml_vk_can_fuse_unary_mul_pair(const struct ggml_cgraph * cgraph, i ggml_vk_can_fuse_unary_mul(cgraph, node_idx, node_idx + 1); } +static bool ggml_vk_can_fuse_hc_post_gate(const struct ggml_cgraph * cgraph, int node_idx) { + const ggml_tensor * scale_in = cgraph->nodes[node_idx]; + const ggml_tensor * sigmoid = cgraph->nodes[node_idx + 1]; + const ggml_tensor * scale_out = cgraph->nodes[node_idx + 2]; + + // the shader folds scale -> sigmoid -> scale; a bias on either scale is not handled + return ggml_get_unary_op(sigmoid) == GGML_UNARY_OP_SIGMOID && + ggml_get_op_params_f32(scale_in, 1) == 0.0f && + ggml_get_op_params_f32(scale_out, 1) == 0.0f && + scale_in->src[0]->type == GGML_TYPE_F32 && + ggml_are_same_shape(scale_in->src[0], scale_out); +} + bool ggml_vk_can_fuse(const ggml_backend_vk_context * ctx, const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list ops) { if (ops.size() == 2 && ops.begin()[0] == GGML_OP_UNARY && ops.begin()[1] == GGML_OP_MUL) { return ggml_vk_can_fuse_unary_mul_pair(cgraph, node_idx); @@ -14278,6 +14305,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ctx->fused_topk_moe_mode = TOPK_MOE_COUNT; ctx->fused_topk_moe_scale = false; ctx->fused_topk_qsa = false; + ctx->fused_hc_post_gate = false; ctx->fused_rms_norm_mode = RMS_NORM_COUNT; const char *fusion_string {}; if (!ctx->device->disable_fusion) { @@ -14334,6 +14362,13 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg op_srcs_fused_elementwise[0] = false; op_srcs_fused_elementwise[1] = true; op_srcs_fused_elementwise[2] = true; + } else if (ggml_can_fuse_subgraph(cgraph, i, hc_post_gate_pattern, { i + 3 }) && + ggml_check_edges(cgraph, i, hc_post_gate_edges) && + ggml_vk_can_fuse_hc_post_gate(cgraph, i)) { + ctx->num_additional_fused_ops = hc_post_gate_pattern.size() - 1; + ctx->fused_hc_post_gate = true; + fusion_string = "HC_POST_GATE"; + std::fill_n(op_srcs_fused_elementwise, ctx->num_additional_fused_ops + 1, false); } else if (ggml_vk_can_fuse(ctx, cgraph, i, rms_norm_mul_add_mul_pattern)) { ctx->num_additional_fused_ops = 3; ctx->fused_rms_norm_mode = RMS_NORM_MUL_ADD_MUL; @@ -14512,6 +14547,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg ctx->fused_topk_moe_mode = TOPK_MOE_COUNT; ctx->fused_topk_moe_scale = false; ctx->fused_topk_qsa = false; + ctx->fused_hc_post_gate = false; ctx->fused_rms_norm_mode = RMS_NORM_COUNT; fusion_string = nullptr; } @@ -14768,6 +14804,11 @@ void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * graph, if (keep_pattern(rope_view_set_rows_pattern)) { continue; } + if (match_pattern(hc_post_gate_pattern, first_unused)) { + add_pattern_alloc_deps(hc_post_gate_pattern, first_unused + (int) hc_post_gate_pattern.size() - 1); + keep_pattern(hc_post_gate_pattern); + continue; + } // First, grab the next unused node. current_set.push_back(first_unused); @@ -14807,7 +14848,8 @@ void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * graph, match_pattern(rms_norm_mul_add_pattern, j) || match_pattern(rms_norm_mul_rope_view_set_rows_pattern, j) || match_pattern(rms_norm_view_set_rows_pattern, j) || - match_pattern(rope_view_set_rows_pattern, j)) { + match_pattern(rope_view_set_rows_pattern, j) || + match_pattern(hc_post_gate_pattern, j)) { continue; } bool ok = true; diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dsv4_hc_post.comp b/ggml/src/ggml-vulkan/vulkan-shaders/dsv4_hc_post.comp index e521fd9d45da..b80c077259a6 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/dsv4_hc_post.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dsv4_hc_post.comp @@ -33,6 +33,10 @@ layout(push_constant) uniform parameter uint p_offset; uint c_offset; uint d_offset; + + uint gate; // post = gate_scale_out*sigmoid(gate_scale_in*p) + float gate_scale_in; + float gate_scale_out; }; layout(binding = 0, std430) readonly buffer X { float data_x[]; }; @@ -51,7 +55,8 @@ void main() { const uint it = gl_WorkGroupID.y; if (tid < hc) { - post_s[tid] = data_p[p_offset + tid * nbp0 + it * nbp1]; + const float p = data_p[p_offset + tid * nbp0 + it * nbp1]; + post_s[tid] = gate != 0 ? (1.0f / (1.0f + exp(-(p * gate_scale_in)))) * gate_scale_out : p; } if (HAS_COMB == 1 && tid < hc * hc) { const uint idst = tid & 3; diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp index 9c5dc768efe7..1ebff43f3814 100644 --- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp +++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp @@ -3774,7 +3774,28 @@ static void ggml_backend_webgpu_buffer_set_tensor(ggml_backend_buffer_t buffer, size_t total_offset = ggml_webgpu_tensor_offset(tensor) + offset; - buf_ctx->global_ctx->queue.WriteBuffer(buf_ctx->buffer, total_offset, data, (size / 4) * 4); + // WriteBuffer needs the offset and the size to be multiples of 4. + // Write the misaligned head bytes using compute memset, then increment total_offset + // and data pointer so that they are 4-aligned. + if (total_offset % 4 != 0) { + size_t lane = total_offset % 4; // in-word lane the head starts at (the tail below always starts at 0) + size_t head = std::min(4 - lane, size); + + // Pack head bytes into a uint32_t + uint32_t head_val = 0; + for (size_t i = 0; i < head; i++) { + ((uint8_t *) &head_val)[lane + i] = ((const uint8_t *) data)[i]; + } + ggml_backend_webgpu_buffer_memset(buf_ctx->global_ctx, buf_ctx->buffer, head_val, total_offset, head); + + total_offset += head; + size -= head; + data = (const uint8_t *) data + head; + } + + if (size > 0) { + buf_ctx->global_ctx->queue.WriteBuffer(buf_ctx->buffer, total_offset, data, (size / 4) * 4); + } if (size % 4 != 0) { // If size is not a multiple of 4, we need to memset the remaining bytes diff --git a/ggml/src/ggml-zdnn/ggml-zdnn.cpp b/ggml/src/ggml-zdnn/ggml-zdnn.cpp index 4007ac9dfc7d..3b45370a8a64 100644 --- a/ggml/src/ggml-zdnn/ggml-zdnn.cpp +++ b/ggml/src/ggml-zdnn/ggml-zdnn.cpp @@ -225,6 +225,11 @@ static enum ggml_status ggml_backend_zdnn_buffer_init_tensor(ggml_backend_buffer return GGML_STATUS_SUCCESS; } + // reject empty tensors to avoid zDNN crash + if (ggml_is_empty(tensor)) { + return GGML_STATUS_SUCCESS; + } + ggml_backend_zdnn_buffer_context * ctx = (ggml_backend_zdnn_buffer_context *)buffer->context; const int64_t tsize = ggml_nbytes(tensor); diff --git a/ggml/src/gguf.cpp b/ggml/src/gguf.cpp index 0eb9fb744d28..5c26f751fd7d 100644 --- a/ggml/src/gguf.cpp +++ b/ggml/src/gguf.cpp @@ -14,6 +14,7 @@ #include #include #include +#include #include #define GGUF_MAX_STRING_LENGTH (1024*1024*1024) @@ -550,6 +551,8 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr // KV pairs { + std::unordered_set seen_keys; + for (int64_t i = 0; ok && i < n_kv; ++i) { std::string key; gguf_type type = gguf_type(-1); @@ -569,11 +572,9 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr GGML_LOG_ERROR("%s: key %" PRIi64 " is empty\n", __func__, i); ok = false; } - for (size_t j = 0; ok && j < ctx->kv.size(); ++j) { - if (key == ctx->kv[j].key) { - GGML_LOG_ERROR("%s: duplicate key '%s' for tensors %zu and %" PRIi64 " \n", __func__, key.c_str(), j, i); - ok = false; - } + if (ok && !seen_keys.insert(key).second) { + GGML_LOG_ERROR("%s: duplicate key '%s' for KV pair %" PRIi64 "\n", __func__, key.c_str(), i); + ok = false; } if (!ok) { break; @@ -636,6 +637,8 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr } // read the tensor info + std::unordered_set seen_tensor_names; + for (int64_t i = 0; ok && i < n_tensors; ++i) { struct gguf_tensor_info info; @@ -659,12 +662,10 @@ static struct gguf_context * gguf_init_from_reader(const struct gguf_reader & gr ggml_set_name(&info.t, name.c_str()); // make sure there are no duplicate tensor names - for (int64_t j = 0; ok && j < i; ++j) { - if (strcmp(info.t.name, ctx->info[j].t.name) == 0) { - GGML_LOG_ERROR("%s: duplicate tensor name '%s' for tensors %" PRIi64 " and %" PRIi64 "\n", __func__, info.t.name, j, i); - ok = false; - break; - } + if (ok && !seen_tensor_names.insert(name).second) { + GGML_LOG_ERROR("%s: duplicate tensor name '%s' for tensor %" PRIi64 "\n", __func__, info.t.name, i); + ok = false; + break; } } if (!ok) { diff --git a/include/llama.h b/include/llama.h index e808fa9ec126..dea4d46062b7 100644 --- a/include/llama.h +++ b/include/llama.h @@ -489,6 +489,7 @@ extern "C" { LLAMA_API void llama_backend_free(void); // Optional: enable numa optimizations + // TODO: deprecate and make part of llama_backend_init() LLAMA_API void llama_numa_init(enum ggml_numa_strategy numa); // Optional: an auto threadpool gets created in ggml if not passed explicitly diff --git a/models/templates/llm-jp-llm-jp-4.1-8b-thinking.jinja b/models/templates/llm-jp-llm-jp-4.1-8b-thinking.jinja new file mode 100644 index 000000000000..ccf4ddf1cade --- /dev/null +++ b/models/templates/llm-jp-llm-jp-4.1-8b-thinking.jinja @@ -0,0 +1,352 @@ +{#- chat_format=llm-jp-harmony-v1 -#} +{#- + In addition to the normal inputs of `messages` and `tools`, this template also accepts the + following kwargs: + - "builtin_tools": A list, can contain "browser" and/or "python". + - "model_identity": A string that optionally describes the model identity. + - "reasoning_effort": A string that describes the reasoning effort, defaults to "medium". + #} + +{#- Tool Definition Rendering ============================================== #} +{%- macro render_typescript_type(param_spec, required_params, is_nullable=false) -%} + {%- if param_spec.type == "array" -%} + {%- if param_spec['items'] -%} + {%- if param_spec['items']['type'] == "string" -%} + {{- "string[]" }} + {%- elif param_spec['items']['type'] == "number" -%} + {{- "number[]" }} + {%- elif param_spec['items']['type'] == "integer" -%} + {{- "number[]" }} + {%- elif param_spec['items']['type'] == "boolean" -%} + {{- "boolean[]" }} + {%- else -%} + {%- set inner_type = render_typescript_type(param_spec['items'], required_params) -%} + {%- if inner_type == "object | object" or inner_type|length > 50 -%} + {{- "any[]" }} + {%- else -%} + {{- inner_type + "[]" }} + {%- endif -%} + {%- endif -%} + {%- if param_spec.nullable -%} + {{- " | null" }} + {%- endif -%} + {%- else -%} + {{- "any[]" }} + {%- if param_spec.nullable -%} + {{- " | null" }} + {%- endif -%} + {%- endif -%} + {%- elif param_spec.type is defined and param_spec.type is iterable and param_spec.type is not string and param_spec.type is not mapping and param_spec.type[0] is defined -%} + {#- Handle array of types like ["object", "object"] from Union[dict, list] #} + {%- if param_spec.type | length > 1 -%} + {{- param_spec.type | join(" | ") }} + {%- else -%} + {{- param_spec.type[0] }} + {%- endif -%} + {%- elif param_spec.oneOf -%} + {#- Handle oneOf schemas - check for complex unions and fallback to any #} + {%- set has_object_variants = false -%} + {%- for variant in param_spec.oneOf -%} + {%- if variant.type == "object" -%} + {%- set has_object_variants = true -%} + {%- endif -%} + {%- endfor -%} + {%- if has_object_variants and param_spec.oneOf|length > 1 -%} + {{- "any" }} + {%- else -%} + {%- for variant in param_spec.oneOf -%} + {{- render_typescript_type(variant, required_params) -}} + {%- if variant.description %} + {{- "// " + variant.description }} + {%- endif -%} + {%- if variant.default is defined %} + {{ "// default: " + variant.default|tojson }} + {%- endif -%} + {%- if not loop.last %} + {{- " | " }} + {% endif -%} + {%- endfor -%} + {%- endif -%} + {%- elif param_spec.type == "string" -%} + {%- if param_spec.enum -%} + {{- '"' + param_spec.enum|join('" | "') + '"' -}} + {%- else -%} + {{- "string" }} + {%- if param_spec.nullable %} + {{- " | null" }} + {%- endif -%} + {%- endif -%} + {%- elif param_spec.type == "number" -%} + {{- "number" }} + {%- elif param_spec.type == "integer" -%} + {{- "number" }} + {%- elif param_spec.type == "boolean" -%} + {{- "boolean" }} + + {%- elif param_spec.type == "object" -%} + {%- if param_spec.properties -%} + {{- "{\n" }} + {%- for prop_name, prop_spec in param_spec.properties.items() -%} + {{- prop_name -}} + {%- if prop_name not in (param_spec.required or []) -%} + {{- "?" }} + {%- endif -%} + {{- ": " }} + {{ render_typescript_type(prop_spec, param_spec.required or []) }} + {%- if not loop.last -%} + {{-", " }} + {%- endif -%} + {%- endfor -%} + {{- "}" }} + {%- else -%} + {{- "object" }} + {%- endif -%} + {%- else -%} + {{- "any" }} + {%- endif -%} +{%- endmacro -%} + +{%- macro render_tool_namespace(namespace_name, tools) -%} + {{- "## " + namespace_name + "\n\n" }} + {{- "namespace " + namespace_name + " {\n\n" }} + {%- for tool in tools %} + {%- set tool = tool.function %} + {{- "// " + tool.description + "\n" }} + {{- "type "+ tool.name + " = " }} + {%- if tool.parameters and tool.parameters.properties %} + {{- "(_: {\n" }} + {%- for param_name, param_spec in tool.parameters.properties.items() %} + {%- if param_spec.description %} + {{- "// " + param_spec.description + "\n" }} + {%- endif %} + {{- param_name }} + {%- if param_name not in (tool.parameters.required or []) -%} + {{- "?" }} + {%- endif -%} + {{- ": " }} + {{- render_typescript_type(param_spec, tool.parameters.required or []) }} + {%- if param_spec.default is defined -%} + {%- if param_spec.enum %} + {{- ", // default: " + param_spec.default|tojson }} + {%- elif param_spec.oneOf %} + {{- "// default: " + param_spec.default|tojson }} + {%- else %} + {{- ", // default: " + param_spec.default|tojson }} + {%- endif -%} + {%- endif -%} + {%- if not loop.last %} + {{- ",\n" }} + {%- else %} + {{- ",\n" }} + {%- endif -%} + {%- endfor %} + {{- "}) => any;\n\n" }} + {%- else -%} + {{- "() => any;\n\n" }} + {%- endif -%} + {%- endfor %} + {{- "} // namespace " + namespace_name }} +{%- endmacro -%} + +{%- macro render_builtin_tools(browser_tool, python_tool) -%} + {%- if browser_tool %} + {{- "## browser\n\n" }} + {{- "// Tool for browsing.\n" }} + {{- "// The `cursor` appears in brackets before each browsing display: `[{cursor}]`.\n" }} + {{- "// Cite information from the tool using the following format:\n" }} + {{- "// `【{cursor}†L{line_start}(-L{line_end})?】`, for example: `【6†L9-L11】` or `【8†L3】`.\n" }} + {{- "// Do not quote more than 10 words directly from the tool output.\n" }} + {{- "// sources=web (default: web)\n" }} + {{- "namespace browser {\n\n" }} + {{- "// Searches for information related to `query` and displays `topn` results.\n" }} + {{- "type search = (_: {\n" }} + {{- "query: string,\n" }} + {{- "topn?: number, // default: 10\n" }} + {{- "source?: string,\n" }} + {{- "}) => any;\n\n" }} + {{- "// Opens the link `id` from the page indicated by `cursor` starting at line number `loc`, showing `num_lines` lines.\n" }} + {{- "// Valid link ids are displayed with the formatting: `【{id}†.*】`.\n" }} + {{- "// If `cursor` is not provided, the most recent page is implied.\n" }} + {{- "// If `id` is a string, it is treated as a fully qualified URL associated with `source`.\n" }} + {{- "// If `loc` is not provided, the viewport will be positioned at the beginning of the document or centered on the most relevant passage, if available.\n" }} + {{- "// Use this function without `id` to scroll to a new location of an opened page.\n" }} + {{- "type open = (_: {\n" }} + {{- "id?: number | string, // default: -1\n" }} + {{- "cursor?: number, // default: -1\n" }} + {{- "loc?: number, // default: -1\n" }} + {{- "num_lines?: number, // default: -1\n" }} + {{- "view_source?: boolean, // default: false\n" }} + {{- "source?: string,\n" }} + {{- "}) => any;\n\n" }} + {{- "// Finds exact matches of `pattern` in the current page, or the page given by `cursor`.\n" }} + {{- "type find = (_: {\n" }} + {{- "pattern: string,\n" }} + {{- "cursor?: number, // default: -1\n" }} + {{- "}) => any;\n\n" }} + {{- "} // namespace browser\n\n" }} + {%- endif -%} + + {%- if python_tool %} + {{- "## python\n\n" }} + {{- "Use this tool to execute Python code in your chain of thought. The code will not be shown to the user. This tool should be used for internal reasoning, but not for code that is intended to be visible to the user (e.g. when creating plots, tables, or files).\n\n" }} + {{- "When you send a message containing Python code to python, it will be executed in a stateful Jupyter notebook environment. python will respond with the output of the execution or time out after 120.0 seconds. The drive at '/mnt/data' can be used to save and persist user files. Internet access for this session is UNKNOWN. Depends on the cluster.\n\n" }} + {%- endif -%} +{%- endmacro -%} + +{#- System Message Construction ============================================ #} +{%- macro build_system_message() -%} + {%- if model_identity is not defined %} + {%- set model_identity = "You are LLM-jp-4, a large language model trained by LLM-jp." %} + {%- endif %} + {{- model_identity + "\n" -}} + {% if knowledge_cutoff is not defined %} + {%- set knowledge_cutoff = "2025-12" %} + {%- endif %} + {{- "Knowledge cutoff: " + knowledge_cutoff + "\n" -}} + {% if conversation_start_date is not defined %} + {%- set conversation_start_date = strftime_now("%Y-%m-%d") %} + {%- endif %} + {{- "Current date: " + conversation_start_date + "\n\n" }} + {%- if reasoning_effort is not defined %} + {%- set reasoning_effort = "medium" %} + {%- endif %} + {{- "Reasoning: " + reasoning_effort + "\n\n" }} + {%- if builtin_tools %} + {{- "# Tools\n\n" }} + {%- set available_builtin_tools = namespace(browser=false, python=false) %} + {%- for tool in builtin_tools %} + {%- if tool == "browser" %} + {%- set available_builtin_tools.browser = true %} + {%- elif tool == "python" %} + {%- set available_builtin_tools.python = true %} + {%- endif %} + {%- endfor %} + {{- render_builtin_tools(available_builtin_tools.browser, available_builtin_tools.python) }} + {%- endif -%} + {{- "# Valid channels: analysis, commentary, final. Channel must be included for every message." }} + {%- if tools -%} + {{- "\nCalls to these tools must go to the commentary channel: 'functions'." }} + {%- endif -%} +{%- endmacro -%} + +{#- Main Template Logic ================================================= #} +{#- Set defaults #} + +{#- Render system message #} +{{- "<|start|>system<|message|>" }} +{{- build_system_message() }} +{{- "<|end|>" }} + +{#- Extract developer message #} +{%- if messages[0].role == "developer" or messages[0].role == "system" %} + {%- set developer_message = messages[0].content %} + {%- set loop_messages = messages[1:] %} +{%- else %} + {%- set developer_message = "" %} + {%- set loop_messages = messages %} +{%- endif %} + +{#- Render developer message #} +{%- if developer_message or tools %} + {{- "<|start|>developer<|message|>" }} + {%- if developer_message %} + {{- "# Instructions\n\n" }} + {{- developer_message }} + {{- "\n\n" }} + {%- endif %} + {%- if tools -%} + {{- "# Tools\n\n" }} + {{- render_tool_namespace("functions", tools) }} + {%- endif -%} + {{- "<|end|>" }} +{%- endif %} + +{#- Render messages #} +{#- tool_group: call names of the latest assistant tool_calls message, in call order, #} +{#- plus how many of their results have been rendered (nameless results are attributed by order) #} +{%- set last_tool_call = namespace(name=none) %} +{%- set tool_group = namespace(names=[], result_idx=0) %} +{%- for message in loop_messages -%} + {#- At this point only assistant/user/tool messages should remain #} + {%- if message.role == 'assistant' -%} + {#- Checks to ensure the messages are being passed in the format we expect #} + {%- if "content" in message %} + {%- if "<|channel|>analysis<|message|>" in message.content or "<|channel|>final<|message|>" in message.content %} + {{- raise_exception("You have passed a message containing <|channel|> tags in the content field. Instead of doing this, you should pass analysis messages (the string between '<|message|>' and '<|end|>') in the 'thinking' field, and final messages (the string between '<|message|>' and '<|end|>') in the 'content' field.") }} + {%- endif %} + {%- endif %} + {%- if "thinking" in message %} + {%- if "<|channel|>analysis<|message|>" in message.thinking or "<|channel|>final<|message|>" in message.thinking %} + {{- raise_exception("You have passed a message containing <|channel|> tags in the thinking field. Instead of doing this, you should pass analysis messages (the string between '<|message|>' and '<|end|>') in the 'thinking' field, and final messages (the string between '<|message|>' and '<|end|>') in the 'content' field.") }} + {%- endif %} + {%- endif %} + {%- if "tool_calls" in message %} + {#- Drop-all: previous-turn analysis/CoT is never rendered. It is emitted only for the #} + {#- training target (last message, add_generation_prompt=false), once before the calls. #} + {%- if message.content and message.thinking %} + {{- raise_exception("Cannot pass both content and thinking in an assistant message with tool calls! Put the analysis message in one or the other, but not both.") }} + {%- elif message.content and loop.last and not add_generation_prompt %} + {{- "<|start|>assistant<|channel|>analysis<|message|>" + message.content + "<|end|>" }} + {%- elif message.thinking and loop.last and not add_generation_prompt %} + {{- "<|start|>assistant<|channel|>analysis<|message|>" + message.thinking + "<|end|>" }} + {%- endif %} + {#- Parallel tool calls: every call is rendered; all but the last end with <|end|>, #} + {#- the last with <|call|>. A single call renders as in the non-parallel template. #} + {%- set tool_group.names = [] %} + {%- set tool_group.result_idx = 0 %} + {%- for raw_tool_call in message.tool_calls %} + {%- set tool_call = raw_tool_call.function if raw_tool_call.function else raw_tool_call %} + {{- "<|start|>assistant to=" }} + {{- "functions." + tool_call.name + "<|channel|>commentary " }} + {#- content_type default must match the training data ("<|constrain|> json") #} + {{- (tool_call.content_type if tool_call.content_type is defined else "<|constrain|> json") + "<|message|>" }} + {#- String arguments are already JSON: emit verbatim; only dict arguments need |tojson #} + {{- tool_call.arguments if tool_call.arguments is string else tool_call.arguments|tojson }} + {%- if loop.last %} + {{- "<|call|>" }} + {%- else %} + {{- "<|end|>" }} + {%- endif %} + {%- set last_tool_call.name = tool_call.name %} + {%- set tool_group.names = tool_group.names + [tool_call.name] %} + {%- endfor %} + {%- elif loop.last and not add_generation_prompt %} + {#- Only render the CoT if the final turn is an assistant turn and add_generation_prompt is false #} + {#- This is a situation that should only occur in training, never in inference. #} + {%- if "thinking" in message %} + {{- "<|start|>assistant<|channel|>analysis<|message|>" + message.thinking + "<|end|>" }} + {%- endif %} + {#- <|return|> indicates the end of generation, but <|end|> does not #} + {#- <|return|> should never be an input to the model, but we include it as the final token #} + {#- when training, so the model learns to emit it. #} + {{- "<|start|>assistant<|channel|>final<|message|>" + message.content + "<|return|>" }} + {%- else %} + {#- CoT is dropped during all previous turns, so we never render it for inference #} + {{- "<|start|>assistant<|channel|>final<|message|>" + message.content + "<|end|>" }} + {%- set last_tool_call.name = none %} + {%- endif %} + {%- elif message.role == 'tool' -%} + {#- Tool name: message.name if given, else by position in the current call group #} + {#- (call order == result order), else the last call (legacy fallback). #} + {%- if message.name is defined and message.name %} + {%- set tool_name = message.name %} + {%- elif tool_group.result_idx < tool_group.names|length %} + {%- set tool_name = tool_group.names[tool_group.result_idx] %} + {%- else %} + {%- set tool_name = last_tool_call.name %} + {%- endif %} + {%- set tool_group.result_idx = tool_group.result_idx + 1 %} + {%- if tool_name is none %} + {{- raise_exception("Message has tool role, but could not resolve its source function name (no 'name' field and no preceding assistant tool call).") }} + {%- endif %} + {{- "<|start|>functions." + tool_name }} + {{- " to=assistant<|channel|>commentary<|message|>" + message.content|tojson + "<|end|>" }} + {%- elif message.role == 'user' -%} + {{- "<|start|>user<|message|>" + message.content + "<|end|>" }} + {%- endif -%} +{%- endfor -%} + +{#- Generation prompt #} +{%- if add_generation_prompt -%} +<|start|>assistant +{%- endif -%} \ No newline at end of file diff --git a/scripts/hip/gcn-cdna-vgpr-check.py b/scripts/hip/gcn-cdna-vgpr-check.py index bc4eba687ece..238a9665c329 100644 --- a/scripts/hip/gcn-cdna-vgpr-check.py +++ b/scripts/hip/gcn-cdna-vgpr-check.py @@ -73,6 +73,15 @@ def main(): '_ZL18flash_attn_ext_f16ILi512ELi512ELi8ELi8ELb1ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', '_ZL18flash_attn_ext_f16ILi576ELi512ELi16ELi4ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', '_ZL18flash_attn_ext_f16ILi576ELi512ELi4ELi16ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi512ELi512ELi16ELi2ELb0ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi512ELi512ELi16ELi2ELb1ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi512ELi512ELi4ELi8ELb0ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi512ELi512ELi4ELi8ELb1ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi512ELi512ELi8ELi4ELb0ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi512ELi512ELi8ELi4ELb1ELb0ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi576ELi512ELi1ELi32ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi576ELi512ELi2ELi16ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', + '_ZL18flash_attn_ext_f16ILi576ELi512ELi8ELi4ELb0ELb1ELb0EEvPKcS1_S1_S1_S1_PKiPfP15HIP_vector_typeIfLj2EEffffjfiS5_IjLj3EEiiiiiiiiiiiliiliiiiil', } functions = parse_log_file(log_file) diff --git a/scripts/snapdragon/ggml-hexagon-trace.py b/scripts/snapdragon/ggml-hexagon-trace.py index 760eb57d9c10..5241dca15311 100755 --- a/scripts/snapdragon/ggml-hexagon-trace.py +++ b/scripts/snapdragon/ggml-hexagon-trace.py @@ -514,7 +514,7 @@ def generate_perfetto_trace(filtered_ops, trace_events, output_path): tm = time_mappers[dev] e['ts_ns'] = tm.cycle_to_ns(e['start_cyc']) dur_ns = tm.dur_cycles_to_ns(e['start_cyc'], e['end_cyc'] - e['start_cyc']) - e['dur_ns'] = max(dur_ns, 100) + e['dur_ns'] = max(dur_ns, 1) # Allocate slots (sub-tracks) to prevent overlaps on same virtual track active_slots = defaultdict(list) diff --git a/src/models/dflash.cpp b/src/models/dflash.cpp index 9b56ac9eca53..1e8881c0c7c8 100644 --- a/src/models/dflash.cpp +++ b/src/models/dflash.cpp @@ -446,19 +446,16 @@ static ggml_tensor * build_dflash2_conv( ggml_tensor * weight_all = ggml_add(ctx0, coeff_all, base_side); + // taps at or past block_size only read the left padding and add nothing + const int64_t n_taps = std::min(kernel_size, block_size); + ggml_tensor * result = nullptr; - for (int64_t tap = 0; tap < kernel_size; ++tap) { + for (int64_t tap = 0; tap < n_taps; ++tap) { ggml_tensor * values = blocks; if (tap > 0) { - ggml_tensor * zeros = ggml_fill(ctx0, - ggml_new_tensor_3d(ctx0, hidden->type, hidden_size, std::min(tap, block_size), n_blocks), 0.0f); - if (tap < block_size) { - ggml_tensor * previous = ggml_view_3d(ctx0, blocks, hidden_size, block_size - tap, n_blocks, - blocks->nb[1], blocks->nb[2], 0); - values = ggml_concat(ctx0, zeros, previous, 1); - } else { - values = zeros; - } + ggml_tensor * previous = ggml_view_3d(ctx0, blocks, hidden_size, block_size - tap, n_blocks, + blocks->nb[1], blocks->nb[2], 0); + values = ggml_pad_ext(ctx0, previous, 0, 0, tap, 0, 0, 0, 0, 0); } values = ggml_reshape_2d(ctx0, values, hidden_size, n_tokens); diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 01af4ae37368..96d910a752b7 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -207,42 +207,13 @@ if (NOT WIN32 OR NOT BUILD_SHARED_LIBS) FIXTURES_SETUP generate-models ) + # Test recurrent-state rollback across all architectures, using the generated dummy models llama_test( test-recurrent-state-rollback LABEL main - ARGS -m "${MODEL_DIR}/qwen35-dense.gguf" - ) - set_tests_properties(test-recurrent-state-rollback PROPERTIES - FIXTURES_REQUIRED generate-models - ) - - llama_test( - test-recurrent-state-rollback - NAME test-recurrent-state-rollback-nemotron-h - LABEL main - ARGS -m "${MODEL_DIR}/nemotron_h-dense.gguf" - ) - set_tests_properties(test-recurrent-state-rollback-nemotron-h PROPERTIES - FIXTURES_REQUIRED generate-models - ) - llama_test( - test-recurrent-state-rollback - NAME test-recurrent-state-rollback-dsv4 - LABEL main - ARGS -m "${MODEL_DIR}/deepseek4-moe.gguf" - ) - set_tests_properties(test-recurrent-state-rollback-dsv4 PROPERTIES - FIXTURES_REQUIRED generate-models - ) - llama_test( - test-recurrent-state-rollback - NAME test-recurrent-state-rollback-kimi-k3 - LABEL main - ARGS -m "${MODEL_DIR}/kimi-k3-moe.gguf" - ) - set_tests_properties(test-recurrent-state-rollback-kimi-k3 PROPERTIES - FIXTURES_REQUIRED generate-models + ARGS --models "${MODEL_DIR}" ) + set_tests_properties(test-recurrent-state-rollback PROPERTIES FIXTURES_REQUIRED generate-models) # Test state save/load functionality across all architectures, using the generated dummy models llama_test( diff --git a/tests/gguf-model-data.cpp b/tests/gguf-model-data.cpp index fe8b4ca76e7f..3e2147dd2b13 100644 --- a/tests/gguf-model-data.cpp +++ b/tests/gguf-model-data.cpp @@ -371,8 +371,8 @@ static std::optional gguf_parse_meta(const std::vector } // cache handling for local download -static std::string get_default_cache_dir() { - return fs_get_cache_directory() + "gguf-headers/"; +static std::filesystem::path get_default_cache_dir() { + return fs_get_cache_directory() / "gguf-headers"; } static std::string sanitize_for_path(const std::string & s) { @@ -385,7 +385,7 @@ static std::string sanitize_for_path(const std::string & s) { return out; } -static bool read_file(const std::string & path, std::vector & out) { +static bool read_file(const std::filesystem::path & path, std::vector & out) { std::ifstream f(path, std::ios::binary | std::ios::ate); if (!f.good()) { return false; @@ -400,7 +400,7 @@ static bool read_file(const std::string & path, std::vector & out) { return f.good(); } -static bool write_file(const std::string & path, const std::vector & data) { +static bool write_file(const std::filesystem::path & path, const std::vector & data) { std::ofstream f(path, std::ios::binary | std::ios::trunc); if (!f.good()) { return false; @@ -516,7 +516,7 @@ static std::string detect_gguf_filename(const std::string & repo, const std::str static std::optional fetch_and_parse( const std::string & repo, const std::string & filename, - const std::string & cache_path, + const std::filesystem::path & cache_path, bool verbose) { std::string url = "https://huggingface.co/" + repo + "/resolve/main/" + filename; @@ -564,19 +564,19 @@ static std::optional fetch_and_parse( return std::nullopt; } -static std::string get_cache_file_path(const std::string& cdir, const std::string& repo_part, const std::string& filename) { +static std::filesystem::path get_cache_file_path(const std::filesystem::path & cdir, const std::string & repo_part, const std::string & filename) { std::string fname_part = sanitize_for_path(filename); - return cdir + "/" + repo_part + "--" + fname_part + ".partial"; + return cdir / std::filesystem::u8path(repo_part + "--" + fname_part + ".partial"); } // Try cache first, then fetch and parse a single GGUF shard. static std::optional fetch_or_cached( const std::string & repo, const std::string & filename, - const std::string & cdir, + const std::filesystem::path & cdir, const std::string & repo_part, bool verbose) { - std::string cache_path = get_cache_file_path(cdir, repo_part, filename); + std::filesystem::path cache_path = get_cache_file_path(cdir, repo_part, filename); { std::vector cached; @@ -584,14 +584,15 @@ static std::optional fetch_or_cached( auto result = gguf_parse_meta(cached); if (result.has_value()) { if (verbose) { - fprintf(stderr, "gguf_fetch: loaded from cache: %s\n", cache_path.c_str()); + fprintf(stderr, "gguf_fetch: loaded from cache: %s\n", fs_path_to_utf8(cache_path).c_str()); } return result; } } } - fs_create_directory_with_parents(cdir); + std::error_code ec; + std::filesystem::create_directories(cdir, ec); return fetch_and_parse(repo, filename, cache_path, verbose); } @@ -600,7 +601,7 @@ std::optional gguf_fetch_model_meta( const std::string & quant, const std::string & cache_dir, bool verbose) { - std::string cdir = cache_dir.empty() ? get_default_cache_dir() : cache_dir; + const std::filesystem::path cdir = cache_dir.empty() ? get_default_cache_dir() : std::filesystem::u8path(cache_dir); std::string repo_part = sanitize_for_path(repo); std::string split_prefix; @@ -661,7 +662,7 @@ gguf_context_ptr gguf_fetch_gguf_ctx( const std::string & quant, const std::string & cache_dir, bool verbose) { - std::string cdir = cache_dir.empty() ? get_default_cache_dir() : cache_dir; + const std::filesystem::path cdir = cache_dir.empty() ? get_default_cache_dir() : std::filesystem::u8path(cache_dir); std::string repo_part = sanitize_for_path(repo); std::string split_prefix; @@ -679,12 +680,12 @@ gguf_context_ptr gguf_fetch_gguf_ctx( auto & model = model_opt.value(); - const std::string cache_path = get_cache_file_path(cdir, repo_part, filename); + const std::filesystem::path cache_path = get_cache_file_path(cdir, repo_part, filename); ggml_context_ptr ggml_ctx_ptr; ggml_context * ggml_ctx{}; gguf_init_params params{true, &ggml_ctx}; - gguf_context_ptr ctx{gguf_init_from_file(cache_path.c_str(), params)}; + gguf_context_ptr ctx{gguf_init_from_file(fs_path_to_utf8(cache_path).c_str(), params)}; ggml_ctx_ptr.reset(ggml_ctx); if (ctx == nullptr) { @@ -718,11 +719,11 @@ gguf_context_ptr gguf_fetch_gguf_ctx( } // Load tensors from shard and add to main gguf_context - const std::string shard_path = get_cache_file_path(cdir, repo_part, shard_name); + const std::filesystem::path shard_path = get_cache_file_path(cdir, repo_part, shard_name); ggml_context_ptr shard_ggml_ctx_ptr; ggml_context * shard_ggml_ctx{}; gguf_init_params shard_params{true, &shard_ggml_ctx}; - gguf_context_ptr shard_ctx{gguf_init_from_file(shard_path.c_str(), shard_params)}; + gguf_context_ptr shard_ctx{gguf_init_from_file(fs_path_to_utf8(shard_path).c_str(), shard_params)}; shard_ggml_ctx_ptr.reset(shard_ggml_ctx); if (shard_ctx == nullptr) { diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index ff5a83295737..17ac9dafef28 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -4342,6 +4342,7 @@ struct test_dsv4_hc_post : public test_dsv4_hc { const int64_t n_embd; const int64_t n_tokens; const bool identity; + const bool gated; std::string op_desc(ggml_tensor * t) override { GGML_UNUSED(t); @@ -4349,11 +4350,14 @@ struct test_dsv4_hc_post : public test_dsv4_hc { } std::string vars() override { - return VARS_TO_STR3(n_embd, n_tokens, identity); + return VARS_TO_STR4(n_embd, n_tokens, identity, gated); } - test_dsv4_hc_post(int64_t n_embd = 31, int64_t n_tokens = 17, bool identity = false) - : n_embd(n_embd), n_tokens(n_tokens), identity(identity) {} + // gated: post = 2*sigmoid(post/hc), as qwen4exp builds it, so backends can fuse the chain + bool run_whole_graph() override { return gated; } + + test_dsv4_hc_post(int64_t n_embd = 31, int64_t n_tokens = 17, bool identity = false, bool gated = false) + : n_embd(n_embd), n_tokens(n_tokens), identity(identity), gated(gated) {} ggml_tensor * build_graph(ggml_context * ctx) override { ggml_tensor * x = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, n_embd, n_tokens); @@ -4365,6 +4369,10 @@ struct test_dsv4_hc_post : public test_dsv4_hc { ggml_tensor * post = ggml_new_tensor_2d(ctx, GGML_TYPE_F32, hc, n_tokens); ggml_set_name(post, "post"); + if (gated) { + post = ggml_scale(ctx, ggml_sigmoid(ctx, ggml_scale(ctx, post, 1.0f / (float) hc)), 2.0f); + } + ggml_tensor * comb = nullptr; if (!identity) { comb = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, hc, hc, n_tokens); @@ -9198,6 +9206,9 @@ static std::vector> make_test_cases_eval() { test_cases.emplace_back(new test_dsv4_hc_post(4096, 21)); test_cases.emplace_back(new test_dsv4_hc_post(31, 17, true)); test_cases.emplace_back(new test_dsv4_hc_post(4096, 21, true)); + test_cases.emplace_back(new test_dsv4_hc_post(31, 17, true, true)); + test_cases.emplace_back(new test_dsv4_hc_post(2560, 21, true, true)); + test_cases.emplace_back(new test_dsv4_hc_post(31, 17, false, true)); // glu ops for (ggml_type type : {GGML_TYPE_F16, GGML_TYPE_F32}) { @@ -10998,6 +11009,9 @@ static std::vector> make_test_cases_eval() { // asymmetric head_dim (hsk != hsv) with one or both sides not 64-aligned test_cases.emplace_back(new test_flash_attn_ext(72, 64, 4, {1, 1}, 256, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); test_cases.emplace_back(new test_flash_attn_ext(64, 72, 4, {1, 1}, 256, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + test_cases.emplace_back(new test_flash_attn_ext(65, 67, 4, {1, 1}, 113, 75, true, true, 8.0f, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + test_cases.emplace_back(new test_flash_attn_ext(65, 67, 4, {1, 1}, 17, 75, false, false, 0, 1.0f, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); + test_cases.emplace_back(new test_flash_attn_ext(65, 67, 4, {1, 1}, 113, 75, false, false, 0, 1.0f, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16)); // mixed quant and Q1_0 test cases test_cases.emplace_back(new test_flash_attn_ext(64, 64, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q8_0, GGML_TYPE_Q4_0)); diff --git a/tests/test-chat.cpp b/tests/test-chat.cpp index f2728f7ccadd..d0fe08585bfe 100644 --- a/tests/test-chat.cpp +++ b/tests/test-chat.cpp @@ -1217,6 +1217,16 @@ static void test_peg_parser(common_chat_templates * tmpls, } assert_msg_equals(tc.expect, msg_accum, true); + // A response format must be enforced by an eager grammar + if (!tc.params.json_schema.empty()) { + if (parser.params_.grammar.empty()) { + throw std::runtime_error("json_schema is set but no grammar was produced"); + } + if (parser.params_.grammar_lazy) { + throw std::runtime_error("json_schema is set but the grammar is lazy"); + } + } + // Test grammar if present in params if (!parser.params_.grammar.empty()) { auto grammar = build_grammar(parser.params_.grammar); @@ -6384,6 +6394,22 @@ static void test_template_output_peg_parsers(bool detailed_debug) { .expect_content("You invoke it like this:\n" + call_markup) .run(); + // Structured output, straight to the final answer + tst.test(" to=user<|message|>" R"({"amount": 123.45, "date": "2025-12-03"})") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .json_schema(invoice_schema) + .expect_content(R"({"amount": 123.45, "date": "2025-12-03"})") + .run(); + + // Structured output after a reasoning message: reasoning stays free-form + tst.test(" to=self<|message|>I need to output the invoice details in JSON<|eom|>" + "<|start|>assistant to=user<|message|>" R"({"amount": 123.45, "date": "2025-12-03"})") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .json_schema(invoice_schema) + .expect_reasoning("I need to output the invoice details in JSON") + .expect_content(R"({"amount": 123.45, "date": "2025-12-03"})") + .run(); + // Tool markup inside the analysis channel is reasoning, not a call tst.test(" to=self<|message|>I could use " + call_markup + " here<|eom|>" "<|start|>assistant to=user<|message|>Hello!<|eot|>") @@ -6575,6 +6601,76 @@ static void test_template_output_peg_parsers(bool detailed_debug) { .run(); } + // LLM-jp-4.1: GPT-OSS dialect with spaces after special tokens and <|end|>-separated parallel calls + { + auto tst = peg_tester("models/templates/llm-jp-llm-jp-4.1-8b-thinking.jinja", detailed_debug); + + // Final channel + tst.test("<|channel|> final<|message|> Hello, world!\nWhat's up?").expect(message_assist).run(); + + // One space rule: an intentional leading space survives + tst.test("<|channel|> final<|message|> padded").expect_content(" padded").run(); + + // Reasoning + content + tst.test( + "<|channel|> analysis<|message|> I'm\nthinking<|end|><|start|> assistant<|channel|> final<|message|> Hello, world!\nWhat's " + "up?") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .expect(message_assist_thoughts) + .run(); + + // Partial reasoning + tst.test("<|channel|> analysis<|message|> I'm\nthinking") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .is_partial(true) + .expect_reasoning("I'm\nthinking") + .run(); + + // Tool call, recipient in role header + tst.test( + "<|channel|> analysis<|message|> I'm\nthinking<|end|>" + "<|start|> assistant to=functions.special_function<|channel|> commentary <|constrain|> json<|message|> {\"arg1\": 1}") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .tools({ special_function_tool }) + .expect(message_assist_call_thoughts) + .run(); + + // Tool call, recipient in channel header + tst.test("<|channel|> commentary to=functions.special_function<|message|> {\"arg1\": 1}") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .tools({ special_function_tool }) + .expect(message_assist_call) + .run(); + + // Parallel tool calls separated by <|end|> + tst.test( + "<|channel|> analysis<|message|> I'm\nthinking<|end|>" + "<|start|> assistant to=functions.special_function<|channel|> commentary <|constrain|> json<|message|> {\"arg1\": 1}<|end|>" + "<|start|> assistant to=functions.special_function<|channel|> commentary <|constrain|> json<|message|> {\"arg1\": 2}") + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .tools({ special_function_tool }) + .parallel_tool_calls(true) + .expect_reasoning("I'm\nthinking") + .expect_tool_calls({ + { "special_function", "{\"arg1\": 1}", {} }, + { "special_function", "{\"arg1\": 2}", {} }, + }) + .run(); + + // Structured output + tst.test( + "<|channel|> analysis<|message|> I need to output the invoice details in JSON<|end|>" + "<|start|> assistant<|channel|> final <|constrain|> json" + "<|message|> " + R"({"amount": 123.45, "date": "2025-12-03"})" + ) + .reasoning_format(COMMON_REASONING_FORMAT_AUTO) + .json_schema(invoice_schema) + .expect_reasoning("I need to output the invoice details in JSON") + .expect_content(R"({"amount": 123.45, "date": "2025-12-03"})") + .run(); + } + { auto tst = peg_tester("models/templates/StepFun3.5-Flash.jinja", detailed_debug); diff --git a/tests/test-fusion.cpp b/tests/test-fusion.cpp index 467248f0f101..65d444ecf46f 100644 --- a/tests/test-fusion.cpp +++ b/tests/test-fusion.cpp @@ -297,7 +297,7 @@ int main(int argc, char ** argv) { } common_init(); - ggml_backend_load_all(); + llama_backend_init(); ggml_backend_dev_t dev = ggml_backend_dev_by_name(device_name.c_str()); if (!dev) { diff --git a/tests/test-recurrent-state-rollback.cpp b/tests/test-recurrent-state-rollback.cpp index 1944a441cb13..4ad0d6f9efca 100644 --- a/tests/test-recurrent-state-rollback.cpp +++ b/tests/test-recurrent-state-rollback.cpp @@ -1,6 +1,8 @@ #include "arg.h" #include "common.h" #include "ggml-backend.h" +#include "log.h" +#include "llama-cpp.h" #include "llama.h" #include "../src/llama-io.h" @@ -10,10 +12,28 @@ #include #include #include +#include +#include #include #include +#include #include +enum class test_status { + PASS, + FAIL, + SKIP, +}; + +static const char * test_status_str(test_status status) { + switch (status) { + case test_status::PASS: return "\033[1;32mPASS\033[m"; + case test_status::FAIL: return "\033[1;31mFAIL\033[m"; + case test_status::SKIP: return "\033[1;33mSKIP\033[m"; + } + return ""; +} + static bool decode_tokens(llama_context * ctx, const std::vector & tokens, uint32_t count) { llama_batch batch = llama_batch_init(count, 0, 1); for (uint32_t pos = 0; pos < count; ++pos) { @@ -50,25 +70,23 @@ struct cache_buffer_collector : llama_io_write_i { } }; -static llama_context * init_ctx(llama_model * model, llama_context_params cparams, uint8_t fill) { - llama_context * ctx = llama_init_from_model(model, cparams); - if (ctx == nullptr || fill == 0) { +static llama_context_ptr init_ctx(llama_model * model, llama_context_params cparams, uint8_t fill) { + llama_context_ptr ctx{llama_init_from_model(model, cparams)}; + if (!ctx || fill == 0) { return ctx; } // Use a full ubatch so buffer discovery preserves prefill allocation sizes. - const uint32_t n_tokens = llama_n_ubatch(ctx); - if (!decode_tokens(ctx, std::vector(n_tokens, 0), n_tokens)) { - llama_free(ctx); + const uint32_t n_tokens = llama_n_ubatch(ctx.get()); + if (!decode_tokens(ctx.get(), std::vector(n_tokens, 0), n_tokens)) { return nullptr; } - llama_synchronize(ctx); + llama_synchronize(ctx.get()); cache_buffer_collector collector; - llama_get_memory(ctx)->state_write(collector); - llama_memory_clear(llama_get_memory(ctx), true); + llama_get_memory(ctx.get())->state_write(collector); + llama_memory_clear(llama_get_memory(ctx.get()), true); if (collector.buffers.empty()) { - fprintf(stderr, "%s : no cache buffers found\n", __func__); - llama_free(ctx); + LOG_ERR("%s: no cache buffers found\n", __func__); return nullptr; } for (auto * buffer : collector.buffers) { @@ -77,7 +95,7 @@ static llama_context * init_ctx(llama_model * model, llama_context_params cparam return ctx; } -static llama_context * make_ctx(const common_params & params, llama_model * model, uint8_t fill) { +static llama_context_ptr make_ctx(const common_params & params, llama_model * model, uint8_t fill) { auto cparams = common_context_params_to_llama(params); cparams.n_seq_max = 1; cparams.n_rs_seq = 8; @@ -109,7 +127,9 @@ static double nmse(const float * a, const float * b, int n) { // ubatches while its rollback restore is still pending. Compared against a // reference context that never advanced past the rollback point and decodes // the identical replay batch. -static bool test_multi_seq_split_replay(const common_params & params, llama_model * model, const int n_vocab, uint8_t fill) { +static test_status test_multi_seq_split_replay(const common_params & params, llama_model * model, uint8_t fill) { + const int n_vocab = llama_vocab_n_tokens(llama_model_get_vocab(model)); + constexpr uint32_t n_seqs = 2; constexpr uint32_t n_ubatch = 16; constexpr uint32_t n_prompt = 19; @@ -128,22 +148,16 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode return init_ctx(model, cparams, fill); }; - llama_context * ctx_roll = make_ctx_multi(); - llama_context * ctx_ref = make_ctx_multi(); - if (ctx_roll == nullptr || ctx_ref == nullptr) { - fprintf(stderr, "%s : failed to init multi-seq contexts\n", __func__); - return false; + llama_context_ptr ctx_roll = make_ctx_multi(); + llama_context_ptr ctx_ref = make_ctx_multi(); + if (!ctx_roll || !ctx_ref) { + LOG_ERR("%s: failed to init multi-seq contexts\n", __func__); + return test_status::FAIL; } - const auto cleanup = [&]() { - llama_free(ctx_roll); - llama_free(ctx_ref); - }; - - if (llama_n_rs_seq(ctx_roll) < n_rollback) { - fprintf(stderr, "%s : skipping because n_rs_seq is too small\n", __func__); - cleanup(); - return true; + if (llama_n_rs_seq(ctx_roll.get()) < n_rollback) { + LOG_INF("%s: skipping because n_rs_seq is too small\n", __func__); + return test_status::SKIP; } const auto tok = [&](uint32_t seq, llama_pos pos) { @@ -159,25 +173,24 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode for (llama_pos pos = 0; pos < (llama_pos) p0; ++pos) { common_batch_add(batch, tok(s, pos), pos, { (llama_seq_id) s }, false); } - ok = ok && llama_decode(ctx_roll, batch) == 0; - ok = ok && llama_decode(ctx_ref, batch) == 0; + ok = ok && llama_decode(ctx_roll.get(), batch) == 0; + ok = ok && llama_decode(ctx_ref.get(), batch) == 0; common_batch_clear(batch); for (llama_pos pos = p0; pos < (llama_pos) n_prompt; ++pos) { common_batch_add(batch, tok(s, pos), pos, { (llama_seq_id) s }, false); } - ok = ok && llama_decode(ctx_roll, batch) == 0; + ok = ok && llama_decode(ctx_roll.get(), batch) == 0; llama_batch_free(batch); - ok = ok && llama_memory_seq_rm(llama_get_memory(ctx_roll), (llama_seq_id) s, p0, -1); + ok = ok && llama_memory_seq_rm(llama_get_memory(ctx_roll.get()), (llama_seq_id) s, p0, -1); // a second partial removal while one is pending must be refused - ok = ok && !llama_memory_seq_rm(llama_get_memory(ctx_roll), (llama_seq_id) s, p0 - 1, -1); + ok = ok && !llama_memory_seq_rm(llama_get_memory(ctx_roll.get()), (llama_seq_id) s, p0 - 1, -1); } if (!ok) { - fprintf(stderr, "%s : multi-seq prefill/rollback failed\n", __func__); - cleanup(); - return false; + LOG_ERR("%s: multi-seq prefill/rollback failed\n", __func__); + return test_status::FAIL; } llama_batch batch = llama_batch_init(n_seqs*n_replay, 0, 1); @@ -187,18 +200,17 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode common_batch_add(batch, tok(s, pos), pos, { (llama_seq_id) s }, true); } } - ok = llama_decode(ctx_roll, batch) == 0; - ok = ok && llama_decode(ctx_ref, batch) == 0; + ok = llama_decode(ctx_roll.get(), batch) == 0; + ok = ok && llama_decode(ctx_ref.get(), batch) == 0; llama_batch_free(batch); if (!ok) { - fprintf(stderr, "%s : multi-seq replay decode failed\n", __func__); - cleanup(); - return false; + LOG_ERR("%s: multi-seq replay decode failed\n", __func__); + return test_status::FAIL; } - // identical ubatch shapes should produce identical states, but the larger - // stdev makes the model sensitive to backend scheduling/rounding noise - constexpr float nmse_eps = 1e-5f; + // both contexts decode identical batches, so the logits should match; + // random dummy models can still drift up to ~1.7e-5, so the bound is 1e-4 + constexpr float nmse_eps = 1e-4f; float diff_max = 0.0f; uint32_t seq_first = 0; @@ -206,12 +218,11 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode double nmse_ab = 0.0; double nmse_a0 = 0.0; for (uint32_t i = 0; i < n_seqs*n_replay; ++i) { - const float * l_roll = llama_get_logits_ith(ctx_roll, i); - const float * l_ref = llama_get_logits_ith(ctx_ref, i); + const float * l_roll = llama_get_logits_ith(ctx_roll.get(), i); + const float * l_ref = llama_get_logits_ith(ctx_ref.get(), i); if (l_roll == nullptr || l_ref == nullptr) { - fprintf(stderr, "%s : missing multi-seq logits at index %u\n", __func__, i); - cleanup(); - return false; + LOG_ERR("%s: missing multi-seq logits at index %u\n", __func__, i); + return test_status::FAIL; } for (int t = 0; t < n_vocab; ++t) { const float r = l_roll[t]; @@ -235,13 +246,12 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode const double nmse_val = nmse_a0 == 0.0 ? (nmse_ab == 0.0 ? 0.0 : std::numeric_limits::infinity()) : nmse_ab/nmse_a0; if (nmse_val > nmse_eps) { - fprintf(stderr, "%s : multi-seq split replay logits mismatch (max diff %g, nmse %g, first at seq %u pos %d)\n", + LOG_ERR("%s: multi-seq split replay logits mismatch (max diff %g, nmse %g, first at seq %u pos %d)\n", __func__, (double) diff_max, nmse_val, seq_first, pos_first); - cleanup(); - return false; + return test_status::FAIL; } - fprintf(stderr, "%s : multi-seq split replay matched (max diff %g, nmse %g)\n", __func__, (double) diff_max, nmse_val); + LOG_INF("%s: multi-seq split replay matched (max diff %g, nmse %g)\n", __func__, (double) diff_max, nmse_val); // seq-1-only decodes must be independent of seq 0's content: diverge seq 0 // in ctx_ref only, then compare identical seq-1-only continuations bitwise @@ -253,7 +263,7 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode const llama_pos pos = p0 + (llama_pos) (n_replay + i); common_batch_add(batch_tail, tok(0, pos + 7), pos, { 0 }, false); } - ok = llama_decode(ctx_ref, batch_tail) == 0; + ok = llama_decode(ctx_ref.get(), batch_tail) == 0; llama_batch_free(batch_tail); } @@ -264,15 +274,15 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode const llama_pos pos = p0 + (llama_pos) (n_replay + i); llama_batch batch_one = llama_batch_init(1, 0, 1); common_batch_add(batch_one, tok(1, pos), pos, { 1 }, true); - ok = llama_decode(ctx_roll, batch_one) == 0; - ok = ok && llama_decode(ctx_ref, batch_one) == 0; + ok = llama_decode(ctx_roll.get(), batch_one) == 0; + ok = ok && llama_decode(ctx_ref.get(), batch_one) == 0; llama_batch_free(batch_one); if (!ok) { break; } - const float * l_roll = llama_get_logits_ith(ctx_roll, 0); - const float * l_ref = llama_get_logits_ith(ctx_ref, 0); + const float * l_roll = llama_get_logits_ith(ctx_roll.get(), 0); + const float * l_ref = llama_get_logits_ith(ctx_ref.get(), 0); ok = l_roll != nullptr && l_ref != nullptr; for (int t = 0; ok && t < n_vocab; ++t) { const float r = l_roll[t]; @@ -291,53 +301,48 @@ static bool test_multi_seq_split_replay(const common_params & params, llama_mode const double nmse_tail = nmse_tail_a0 == 0.0 ? (nmse_tail_ab == 0.0 ? 0.0 : std::numeric_limits::infinity()) : nmse_tail_ab/nmse_tail_a0; if (!ok || nmse_tail > nmse_eps) { - fprintf(stderr, "%s : seq-1-only decode leaked seq 0 state (ok=%d, max diff %g, nmse %g)\n", + LOG_ERR("%s: seq-1-only decode leaked seq 0 state (ok=%d, max diff %g, nmse %g)\n", __func__, ok ? 1 : 0, (double) diff_tail, nmse_tail); - cleanup(); - return false; + return test_status::FAIL; } - fprintf(stderr, "%s : seq-1-only decode independent of seq 0 (max diff %g, nmse %g)\n", __func__, (double) diff_tail, nmse_tail); - cleanup(); - return true; + LOG_INF("%s: seq-1-only decode independent of seq 0 (max diff %g, nmse %g)\n", __func__, (double) diff_tail, nmse_tail); + return test_status::PASS; } -static int test_rollback(const common_params & params, llama_model * model, uint8_t fill) { +// Save a rolled-back single-seq state, restore it into fresh and dirty +// contexts, and verify exact logit matches on replay. +static test_status test_rollback(const common_params & params, llama_model * model, uint8_t fill) { const llama_vocab * vocab = llama_model_get_vocab(model); const int n_vocab = llama_vocab_n_tokens(vocab); - // TODO: use smart pointers - llama_context * ctx_src = make_ctx(params, model, fill); - llama_context * ctx_dst = make_ctx(params, model, fill); - if (ctx_src == nullptr || ctx_dst == nullptr) { - fprintf(stderr, "%s : failed to init contexts\n", __func__); - return 1; + llama_context_ptr ctx_src = make_ctx(params, model, fill); + llama_context_ptr ctx_dst = make_ctx(params, model, fill); + if (!ctx_src || !ctx_dst) { + LOG_ERR("%s: failed to init contexts\n", __func__); + return test_status::FAIL; } - if (llama_n_rs_seq(ctx_src) == 0) { - fprintf(stderr, "%s : skipping because n_rs_seq is disabled\n", __func__); - llama_free(ctx_src); - llama_free(ctx_dst); - return 0; + if (llama_n_rs_seq(ctx_src.get()) == 0) { + LOG_INF("%s: skipping because n_rs_seq is disabled\n", __func__); + return test_status::SKIP; } std::vector tokens; if (llama_vocab_type(vocab) == LLAMA_VOCAB_TYPE_NONE) { tokens = { 1, 2, 3, 4, 5, 6, 7, 8, 9 }; } else { - tokens = common_tokenize(ctx_src, "The quick brown fox jumps over the lazy dog", true); + tokens = common_tokenize(ctx_src.get(), "The quick brown fox jumps over the lazy dog", true); } - const uint32_t n_rs_seq = llama_n_rs_seq(ctx_src); + const uint32_t n_rs_seq = llama_n_rs_seq(ctx_src.get()); constexpr uint32_t n_rollback = 3; if (n_rs_seq < n_rollback) { - fprintf(stderr, "%s : skipping because n_rs_seq is too small\n", __func__); - llama_free(ctx_src); - llama_free(ctx_dst); - return 0; + LOG_INF("%s: skipping because n_rs_seq is too small\n", __func__); + return test_status::SKIP; } if (tokens.empty()) { - fprintf(stderr, "%s : not enough prompt tokens\n", __func__); - return 1; + LOG_ERR("%s: not enough prompt tokens\n", __func__); + return test_status::FAIL; } tokens.resize(n_rs_seq + 1, tokens.back()); @@ -347,35 +352,35 @@ static int test_rollback(const common_params & params, llama_model * model, uint // Decode the full prompt on the source, then roll back three positions. // Replaying them crosses DSV4's ratio-4 compressor boundary. // Rollback leaves the recurrent memory in a snapshot state (rs_idx != 0). - if (!decode_tokens(ctx_src, tokens, n_tokens)) { - fprintf(stderr, "%s : failed to decode prompt\n", __func__); - return 1; + if (!decode_tokens(ctx_src.get(), tokens, n_tokens)) { + LOG_ERR("%s: failed to decode prompt\n", __func__); + return test_status::FAIL; } - if (!llama_memory_seq_rm(llama_get_memory(ctx_src), 0, rollback_pos, -1)) { - fprintf(stderr, "%s : rollback failed\n", __func__); - return 1; + if (!llama_memory_seq_rm(llama_get_memory(ctx_src.get()), 0, rollback_pos, -1)) { + LOG_ERR("%s: rollback failed\n", __func__); + return test_status::FAIL; } // Save the rolled-back state and restore it into a fresh context. common_prompt_checkpoint ckpt; - ckpt.update_tgt(ctx_src, 0, 0); - ckpt.load_tgt(ctx_dst, 0, 0); + ckpt.update_tgt(ctx_src.get(), 0, 0); + ckpt.load_tgt(ctx_dst.get(), 0, 0); constexpr float nmse_eps = 0.0; std::vector> logits_src_replay(n_rollback); const auto replay_and_compare = [&](const char * mode) { for (uint32_t i = 0; i < n_rollback; ++i) { const llama_pos pos = rollback_pos + i; - if (!decode_one(ctx_src, tokens[pos], pos) || - !decode_one(ctx_dst, tokens[pos], pos)) { - fprintf(stderr, "%s : %s replay failed at position %d\n", __func__, mode, pos); + if (!decode_one(ctx_src.get(), tokens[pos], pos) || + !decode_one(ctx_dst.get(), tokens[pos], pos)) { + LOG_ERR("%s: %s replay failed at position %d\n", __func__, mode, pos); return false; } - const float * logits_src = llama_get_logits_ith(ctx_src, 0); - const float * logits_dst = llama_get_logits_ith(ctx_dst, 0); + const float * logits_src = llama_get_logits_ith(ctx_src.get(), 0); + const float * logits_dst = llama_get_logits_ith(ctx_dst.get(), 0); if (logits_src == nullptr || logits_dst == nullptr) { - fprintf(stderr, "%s : missing %s logits at position %d\n", __func__, mode, pos); + LOG_ERR("%s: missing %s logits at position %d\n", __func__, mode, pos); return false; } @@ -388,7 +393,7 @@ static int test_rollback(const common_params & params, llama_model * model, uint } } if (nmse_val > nmse_eps) { - fprintf(stderr, "%s : %s logits mismatch at position %d, first token %d, nmse %g\n", + LOG_ERR("%s: %s logits mismatch at position %d, first token %d, nmse %g\n", __func__, mode, pos, token_first, nmse_val); return false; } @@ -396,7 +401,7 @@ static int test_rollback(const common_params & params, llama_model * model, uint return true; }; if (!replay_and_compare("full")) { - return 1; + return test_status::FAIL; } // TODO: this test is invalid because RS rollback is only correct once after a ubatch with more than n_rs_seq tokens @@ -419,10 +424,10 @@ static int test_rollback(const common_params & params, llama_model * model, uint // Repeat the load into a context that already has its own rollback state: // groups 1..n_rs_seq hold a different prompt's history, and rs_idx[0] is // non-zero at load time. The restore must wipe that state and still match. - llama_context * ctx_dirty = make_ctx(params, model, fill); - if (ctx_dirty == nullptr) { - fprintf(stderr, "%s : failed to init dirty ctx\n", __func__); - return 1; + llama_context_ptr ctx_dirty = make_ctx(params, model, fill); + if (!ctx_dirty) { + LOG_ERR("%s: failed to init dirty ctx\n", __func__); + return test_status::FAIL; } std::vector noise = tokens; @@ -432,28 +437,28 @@ static int test_rollback(const common_params & params, llama_model * model, uint t = 0; } } - if (!decode_tokens(ctx_dirty, noise, n_tokens)) { - fprintf(stderr, "%s : dirty prompt decode failed\n", __func__); - return 1; + if (!decode_tokens(ctx_dirty.get(), noise, n_tokens)) { + LOG_ERR("%s: dirty prompt decode failed\n", __func__); + return test_status::FAIL; } - if (!llama_memory_seq_rm(llama_get_memory(ctx_dirty), 0, rollback_pos, -1)) { - fprintf(stderr, "%s : dirty rollback failed\n", __func__); - return 1; + if (!llama_memory_seq_rm(llama_get_memory(ctx_dirty.get()), 0, rollback_pos, -1)) { + LOG_ERR("%s: dirty rollback failed\n", __func__); + return test_status::FAIL; } - ckpt.load_tgt(ctx_dirty, 0, 0); + ckpt.load_tgt(ctx_dirty.get(), 0, 0); for (uint32_t i = 0; i < n_rollback; ++i) { const llama_pos pos = rollback_pos + i; - if (!decode_one(ctx_dirty, tokens[pos], pos)) { - fprintf(stderr, "%s : dirty replay failed at position %d\n", __func__, pos); - return 1; + if (!decode_one(ctx_dirty.get(), tokens[pos], pos)) { + LOG_ERR("%s: dirty replay failed at position %d\n", __func__, pos); + return test_status::FAIL; } - const float * logits_dirty = llama_get_logits_ith(ctx_dirty, 0); + const float * logits_dirty = llama_get_logits_ith(ctx_dirty.get(), 0); if (logits_dirty == nullptr) { - fprintf(stderr, "%s : missing dirty logits at position %d\n", __func__, pos); - return 1; + LOG_ERR("%s: missing dirty logits at position %d\n", __func__, pos); + return test_status::FAIL; } const double nmse_dirty = nmse(logits_src_replay[i].data(), logits_dirty, n_vocab); @@ -464,22 +469,75 @@ static int test_rollback(const common_params & params, llama_model * model, uint } } if (nmse_dirty > nmse_eps) { - fprintf(stderr, "%s : dirty-ctx logits mismatch at position %d, first token %d, nmse %g\n", + LOG_ERR("%s: dirty-ctx logits mismatch at position %d, first token %d, nmse %g\n", __func__, pos, token_first, nmse_dirty); - return 1; + return test_status::FAIL; } } - fprintf(stderr, "%s : recurrent rollback checkpoint restored successfully\n", __func__); - llama_free(ctx_src); - llama_free(ctx_dst); - llama_free(ctx_dirty); + LOG_INF("%s: recurrent rollback checkpoint restored successfully\n", __func__); + return test_status::PASS; +} - if (!test_multi_seq_split_replay(params, model, n_vocab, fill)) { - return 1; +static test_status merge_status(test_status a, test_status b) { + if (a == test_status::FAIL || b == test_status::FAIL) { + return test_status::FAIL; + } + if (a == test_status::PASS || b == test_status::PASS) { + return test_status::PASS; + } + return test_status::SKIP; +} + +struct test_results { + test_status rollback = test_status::SKIP; + test_status replay = test_status::SKIP; +}; + +// Run every test for an initialized model over both cache fills. +static test_results run_tests(const common_params & params, llama_model * model) { + test_results res; + for (uint8_t fill : { 0, 0x3e }) { + LOG_INF("%s: testing with cache fill 0x%02x\n", __func__, fill); + const test_status rb = test_rollback(params, model, fill); + const test_status rp = test_multi_seq_split_replay(params, model, fill); + res.rollback = merge_status(res.rollback, rb); + res.replay = merge_status(res.replay, rp); + if (rb == test_status::FAIL || rp == test_status::FAIL) { + break; + } + } + return res; +} + +// Run the tests for a single model file. +// Returns the per-test statuses. +static test_results run_tests_for_model(const std::string & model_path, const struct common_params & base_params) { + struct common_params params = base_params; + params.model.path = model_path; + + auto llama_init = common_init_from_params(params, true); + auto * model = llama_init->model(); + + if (model == nullptr) { + LOG_ERR("%s: failed to init model '%s'\n", __func__, model_path.c_str()); + // a model that cannot be loaded is a failure, not a skip + return { test_status::FAIL, test_status::FAIL }; } - return 0; + if (!llama_model_is_recurrent(model) && !llama_model_is_hybrid(model)) { + LOG_INF("%s: skipping for non-recurrent model\n", __func__); + return {}; + } + + return run_tests(params, model); +} + +static void print_usage(int /* argc */, char ** argv) { + LOG("\nexample usage:\n"); + LOG("\n %s -m your_model.gguf\n", argv[0]); + LOG("\n %s --models tests/test-models\n", argv[0]); + LOG("\n"); } int main(int argc, char ** argv) { @@ -491,30 +549,107 @@ int main(int argc, char ** argv) { common_init(); - if (!common_params_parse(argc, argv, params, LLAMA_EXAMPLE_COMMON)) { - return 1; + // extract our own --models DIR option before handing the rest to the common arg parser + std::string models_dir; + std::vector filtered_argv; + filtered_argv.push_back(argv[0]); + for (int i = 1; i < argc; i++) { + if (strcmp(argv[i], "--models") == 0) { + if (i + 1 >= argc) { + LOG_ERR("%s: --models requires a directory argument\n", __func__); + return 1; + } + models_dir = argv[i + 1]; + i++; + } else { + filtered_argv.push_back(argv[i]); + } } + filtered_argv.push_back(nullptr); + const int fargc = (int)filtered_argv.size() - 1; - llama_backend_init(); + // in --models mode there is no single model; set a placeholder so the common parser's + // "--model is required" check passes (each model is set individually inside the loop) + if (!models_dir.empty()) { + params.model.path = models_dir; + } - common_init_result_ptr llama_init = common_init_from_params(params); - llama_model * model = llama_init->model(); - if (model == nullptr) { - fprintf(stderr, "%s : failed to init model\n", __func__); + if (!common_params_parse(fargc, filtered_argv.data(), params, LLAMA_EXAMPLE_COMMON, print_usage)) { return 1; } - if (!llama_model_is_recurrent(model) && !llama_model_is_hybrid(model)) { - fprintf(stderr, "%s : skipping for non-recurrent model\n", __func__); - return 0; - } + llama_backend_init(); - for (uint8_t fill : { 0, 0x3e }) { - fprintf(stderr, "%s : testing with cache fill 0x%02x\n", __func__, fill); - if (test_rollback(params, model, fill) != 0) { + if (!models_dir.empty()) { + // run every test over each dummy model in the directory + if (!std::filesystem::exists(models_dir) || !std::filesystem::is_directory(models_dir)) { + LOG_ERR("%s: models directory '%s' does not exist\n", __func__, models_dir.c_str()); return 1; } + + std::vector models; + for (const auto & entry : std::filesystem::directory_iterator(models_dir)) { + if (entry.is_regular_file() && entry.path().extension() == ".gguf") { + models.push_back(entry.path().string()); + } + } + std::sort(models.begin(), models.end()); + + if (models.empty()) { + LOG_ERR("%s: no .gguf models found in '%s'\n", __func__, models_dir.c_str()); + return 1; + } + + size_t name_width = 5; // "Model" + for (const auto & model_path : models) { + name_width = std::max(name_width, std::filesystem::path(model_path).filename().string().size()); + } + + // silence everything but the table itself (LOG has verbosity LOG_LEVEL_OUTPUT = 0) + common_log_set_verbosity_thold(0); + + LOG("%-*s %-8s %s\n", (int) name_width, "Model", "rollback", "split replay"); + common_log_flush(common_log_main()); + + size_t n_pass[2] = { 0, 0 }; + size_t n_skip[2] = { 0, 0 }; + size_t n_fail[2] = { 0, 0 }; + for (const auto & model_path : models) { + const auto name = std::filesystem::path(model_path).filename().string(); + + LOG("%-*s", (int) name_width, name.c_str()); + + const test_results res = run_tests_for_model(model_path, params); + + // all status strings have the same raw length, so the columns line up; + // pad the first status to the width of the "rollback" header + separator + LOG(" %s %s", test_status_str(res.rollback), test_status_str(res.replay)); + LOG("\n"); + common_log_flush(common_log_main()); + + const test_status all[2] = { res.rollback, res.replay }; + for (int t = 0; t < 2; ++t) { + switch (all[t]) { + case test_status::PASS: n_pass[t]++; break; + case test_status::FAIL: n_fail[t]++; break; + case test_status::SKIP: n_skip[t]++; break; + } + } + } + + common_log_set_verbosity_thold(LOG_DEFAULT_LLAMA); + common_log_flush(common_log_main()); + + LOG_INF("%s: rollback: %zu passed, %zu skipped, %zu failed (of %zu)\n", + __func__, n_pass[0], n_skip[0], n_fail[0], models.size()); + LOG_INF("%s: split replay: %zu passed, %zu skipped, %zu failed (of %zu)\n", + __func__, n_pass[1], n_skip[1], n_fail[1], models.size()); + + return (n_fail[0] + n_fail[1]) == 0 ? 0 : 1; } - return 0; + // single-model mode + const test_results res = run_tests_for_model(params.model.path, params); + + return (res.rollback == test_status::FAIL || res.replay == test_status::FAIL) ? 1 : 0; } diff --git a/tools/cli/cli.cpp b/tools/cli/cli.cpp index dcdb6aeac21d..92c86614f820 100644 --- a/tools/cli/cli.cpp +++ b/tools/cli/cli.cpp @@ -41,6 +41,9 @@ int llama_cli(int argc, char ** argv) { return 1; } + llama_backend_init(); + llama_numa_init(params.numa); + #if defined (__unix__) || (defined (__APPLE__) && defined (__MACH__)) struct sigaction sigint_action; sigint_action.sa_handler = signal_handler; diff --git a/tools/mtmd/models/conformer.cpp b/tools/mtmd/models/conformer.cpp index 5f2c7b97314e..18c3d27bcd7b 100644 --- a/tools/mtmd/models/conformer.cpp +++ b/tools/mtmd/models/conformer.cpp @@ -124,8 +124,7 @@ ggml_cgraph * clip_graph_conformer::build() { const auto pos_len = matrix_bd->ne[0]; const auto q_len = matrix_bd->ne[1]; const auto h = matrix_bd->ne[2]; - matrix_bd = ggml_pad(ctx0, matrix_bd, 1, 0, 0, 0); - matrix_bd = ggml_roll(ctx0, matrix_bd, 1, 0, 0, 0); + matrix_bd = ggml_pad_ext(ctx0, matrix_bd, 1, 0, 0, 0, 0, 0, 0, 0); matrix_bd = ggml_reshape_3d(ctx0, matrix_bd, q_len, pos_len + 1, h); matrix_bd = ggml_view_3d(ctx0, matrix_bd, q_len, pos_len, h, matrix_bd->nb[1], matrix_bd->nb[2], matrix_bd->nb[0] * q_len); diff --git a/tools/mtmd/models/gemma4a.cpp b/tools/mtmd/models/gemma4a.cpp index 5dd64b7833b0..f98a8b6fc198 100644 --- a/tools/mtmd/models/gemma4a.cpp +++ b/tools/mtmd/models/gemma4a.cpp @@ -117,14 +117,12 @@ ggml_cgraph * clip_graph_gemma4a::build() { Qcur = ggml_cont(ctx0, ggml_permute(ctx0, Qcur, 0, 3, 1, 2)); // [D, C, B, H] // K/V block context extraction via overlapping view: - // Pad to S*B elements, roll right by P to create left-padding, + // Left pad by P and right pad to S*B elements, // then view with stride C in the block dimension (overlapping windows). auto extract_blocks = [&](ggml_tensor * t) -> ggml_tensor * { - // [D, H, N] -> pad to S*B -> roll right by P -> cont (materialize) + // [D, H, N] -> left pad by P, right pad to S*B const int64_t pad_kv = S * B - n_pos; - t = ggml_pad(ctx0, t, 0, 0, pad_kv, 0); // [D, H, S*B] - t = ggml_roll(ctx0, t, 0, 0, P, 0); // left-pad by P - t = ggml_cont(ctx0, t); // materialize roll (removes view offset) + t = ggml_pad_ext(ctx0, t, 0, 0, 0, 0, P, pad_kv - P, 0, 0); // [D, H, S*B] // Overlapping view: stride for B dim is C positions, not S // ne = [D, H, S, B], data_size = D*H*S*B*sizeof = source_nbytes (exact fit) // nb1=D*sizeof, nb2=D*H*sizeof, nb3=C*D*H*sizeof (overlap: C < S) @@ -219,9 +217,8 @@ ggml_cgraph * clip_graph_gemma4a::build() { x = ggml_cont(ctx0, ggml_transpose(ctx0, x)); } - // Causal depthwise Conv1D via ggml_ssm_conv (pad+roll for left-only padding). - x = ggml_pad(ctx0, x, 4, 0, 0, 0); - x = ggml_roll(ctx0, x, 4, 0, 0, 0); + // Causal depthwise Conv1D via ggml_ssm_conv, left padded only. + x = ggml_pad_ext(ctx0, x, 4, 0, 0, 0, 0, 0, 0, 0); x = ggml_ssm_conv(ctx0, x, layer.conv_dw_w); if (layer.conv_dw_b) { x = ggml_add(ctx0, x, layer.conv_dw_b); diff --git a/tools/mtmd/models/granite-speech.cpp b/tools/mtmd/models/granite-speech.cpp index a158a59ce9ac..9725def82733 100644 --- a/tools/mtmd/models/granite-speech.cpp +++ b/tools/mtmd/models/granite-speech.cpp @@ -143,9 +143,7 @@ ggml_cgraph * clip_graph_granite_speech::build() { } cb(x, "conv_glu", il); - x = ggml_pad(ctx0, x, conv_pad, 0, 0, 0); - x = ggml_roll(ctx0, x, conv_pad, 0, 0, 0); - x = ggml_pad(ctx0, x, conv_pad, 0, 0, 0); + x = ggml_pad_ext(ctx0, x, conv_pad, conv_pad, 0, 0, 0, 0, 0, 0); x = ggml_ssm_conv(ctx0, x, layer.conv_dw_w); cb(x, "conv_dw", il); diff --git a/tools/mtmd/models/parakeet.cpp b/tools/mtmd/models/parakeet.cpp index 8be141d93b37..a6d7f0739943 100644 --- a/tools/mtmd/models/parakeet.cpp +++ b/tools/mtmd/models/parakeet.cpp @@ -287,8 +287,7 @@ ggml_cgraph * clip_graph_parakeet::build() { const auto n_frame = rel_pos_scores->ne[1]; const auto n_head = rel_pos_scores->ne[2]; - rel_pos_scores = ggml_pad(ctx0, rel_pos_scores, 1, 0, 0, 0); - rel_pos_scores = ggml_roll(ctx0, rel_pos_scores, 1, 0, 0, 0); + rel_pos_scores = ggml_pad_ext(ctx0, rel_pos_scores, 1, 0, 0, 0, 0, 0, 0, 0); rel_pos_scores = ggml_reshape_3d(ctx0, rel_pos_scores, n_frame, pos_window + 1, n_head); rel_pos_scores = ggml_cont(ctx0, rel_pos_scores); @@ -366,9 +365,7 @@ ggml_cgraph * clip_graph_parakeet::build() { // use ggml_ssm_conv for f32 precision const int dw_pad = (hparams.audio_conv_kernel_size - 1) / 2; - cur = ggml_pad(ctx0, cur, dw_pad, 0, 0, 0); - cur = ggml_roll(ctx0, cur, dw_pad, 0, 0, 0); - cur = ggml_pad(ctx0, cur, dw_pad, 0, 0, 0); + cur = ggml_pad_ext(ctx0, cur, dw_pad, dw_pad, 0, 0, 0, 0, 0, 0); ggml_format_name(cur, "enc_%d_conv_dw_pad", il); cur = ggml_ssm_conv(ctx0, cur, layer.conv_dw_w); diff --git a/tools/mtmd/mtmd-cli.cpp b/tools/mtmd/mtmd-cli.cpp index ba18b3e32b5f..6fe058fd8cd0 100644 --- a/tools/mtmd/mtmd-cli.cpp +++ b/tools/mtmd/mtmd-cli.cpp @@ -81,7 +81,7 @@ struct mtmd_cli_context { llama_context * lctx; const llama_vocab * vocab; common_sampler * smpl; - llama_batch batch; + common_batch batch; int n_batch; mtmd::bitmaps bitmaps; @@ -115,7 +115,7 @@ struct mtmd_cli_context { vocab = llama_model_get_vocab(model); smpl = common_sampler_init(model, params.sampling); n_threads = params.cpuparams.n_threads; - batch = llama_batch_init(1, 0, 1); // batch for next token generation + batch = common_batch(lctx); // batch for next token generation n_batch = params.n_batch; init_vision_context(params); @@ -148,7 +148,6 @@ struct mtmd_cli_context { } ~mtmd_cli_context() { - llama_batch_free(batch); common_sampler_free(smpl); } @@ -230,9 +229,9 @@ static int generate_response(mtmd_cli_context & ctx, int n_predict) { } // eval the token - common_batch_clear(ctx.batch); - common_batch_add(ctx.batch, token_id, ctx.n_past++, {0}, true); - if (llama_decode(ctx.lctx, ctx.batch)) { + ctx.batch.clear(); + ctx.batch.add(token_id, ctx.n_past++, 0, true); + if (llama_process(ctx.lctx, LLAMA_PROCESS_TYPE_DECODE, ctx.batch.get())) { LOG_ERR("failed to decode token\n"); return 1; } diff --git a/tools/mtmd/mtmd-helper-common.h b/tools/mtmd/mtmd-helper-common.h index f907346c7b58..e347868ef214 100644 --- a/tools/mtmd/mtmd-helper-common.h +++ b/tools/mtmd/mtmd-helper-common.h @@ -6,6 +6,7 @@ #include "ggml.h" #include "llama.h" +#include "llama-cpp.h" #include "mtmd.h" #include @@ -73,112 +74,99 @@ inline mtmd_helper_logger g_logger; struct decode_embd_batch { int n_pos_per_embd; int n_mmproj_embd; - std::vector pos; - std::vector pos_view; // used by mrope - std::vector n_seq_id; - std::vector seq_id_0; - std::vector seq_ids; - std::vector logits; - llama_batch batch; - decode_embd_batch(float * embd, int32_t n_tokens, int n_pos_per_embd, int n_mmproj_embd) : n_pos_per_embd(n_pos_per_embd), n_mmproj_embd(n_mmproj_embd) { + int32_t n_tokens; + const float * embd; // [n_tokens, n_mmproj_embd], not owned + std::vector pos; // [n_pos_per_embd, n_tokens], section-major + std::vector pos_view; // sliced positions of the last get_view() + std::vector logits; + llama_seq_id seq_id = 0; + + llama_batch_ext_ptr batch; // rendered sub-batch, see render() + + decode_embd_batch(const float * embd, int32_t n_tokens, int n_pos_per_embd, int n_mmproj_embd) + : n_pos_per_embd(n_pos_per_embd), n_mmproj_embd(n_mmproj_embd), n_tokens(n_tokens), embd(embd) { GGML_ASSERT(n_tokens > 0 && n_pos_per_embd > 0 && n_mmproj_embd > 0); - pos .resize((size_t) n_tokens * (size_t) n_pos_per_embd); - n_seq_id.resize(n_tokens); - seq_ids .resize(n_tokens + 1); - logits .resize(n_tokens); - seq_id_0.resize(1); - seq_ids [n_tokens] = nullptr; - batch = { - /*n_tokens =*/ n_tokens, - /*tokens =*/ nullptr, - /*embd =*/ embd, - /*pos =*/ pos.data(), - /*n_seq_id =*/ n_seq_id.data(), - /*seq_id =*/ seq_ids.data(), - /*logits =*/ logits.data(), - }; + pos .resize((size_t) n_tokens * (size_t) n_pos_per_embd); + logits.resize(n_tokens); } void set_position_normal(llama_pos pos_0, llama_seq_id seq_id) { - seq_id_0[0] = seq_id; - for (int i = 0; i < batch.n_tokens; i++) { - batch.pos [i] = pos_0 + i; - batch.n_seq_id[i] = 1; - batch.seq_id [i] = seq_id_0.data(); - batch.logits [i] = false; + this->seq_id = seq_id; + for (int i = 0; i < n_tokens; i++) { + pos[i] = pos_0 + i; } } // M-RoPE for image void set_position_mrope_2d(const std::vector & rel_pos, llama_seq_id seq_id) { GGML_ASSERT(n_pos_per_embd == 4); - GGML_ASSERT(!rel_pos.empty() && (int32_t)rel_pos.size() == batch.n_tokens); - seq_id_0[0] = seq_id; - for (int32_t i = 0; i < batch.n_tokens; i++) { + GGML_ASSERT(!rel_pos.empty() && (int32_t)rel_pos.size() == n_tokens); + this->seq_id = seq_id; + for (int32_t i = 0; i < n_tokens; i++) { const size_t idx = (size_t) i; - const size_t n_tokens = (size_t) batch.n_tokens; - pos[idx ] = rel_pos[i].t; - pos[idx + n_tokens ] = rel_pos[i].y; - pos[idx + n_tokens * 2 ] = rel_pos[i].x; - pos[idx + n_tokens * 3 ] = rel_pos[i].z; - } - for (int i = 0; i < batch.n_tokens; i++) { - batch.n_seq_id[i] = 1; - batch.seq_id [i] = seq_id_0.data(); - batch.logits [i] = false; + const size_t n = (size_t) n_tokens; + pos[idx ] = rel_pos[i].t; + pos[idx + n ] = rel_pos[i].y; + pos[idx + n * 2] = rel_pos[i].x; + pos[idx + n * 3] = rel_pos[i].z; } } // M-RoPE for audio void set_position_mrope_1d(llama_pos pos_0, llama_seq_id seq_id) { GGML_ASSERT(n_pos_per_embd == 4); - seq_id_0[0] = seq_id; - for (int i = 0; i < batch.n_tokens; i++) { + this->seq_id = seq_id; + for (int i = 0; i < n_tokens; i++) { const size_t idx = (size_t) i; - const size_t n_tokens = (size_t) batch.n_tokens; - pos[idx ] = pos_0 + i; - pos[idx + n_tokens ] = pos_0 + i; - pos[idx + n_tokens * 2 ] = pos_0 + i; - pos[idx + n_tokens * 3 ] = pos_0 + i; - } - for (int i = 0; i < batch.n_tokens; i++) { - batch.n_seq_id[i] = 1; - batch.seq_id [i] = seq_id_0.data(); - batch.logits [i] = false; + const size_t n = (size_t) n_tokens; + pos[idx ] = pos_0 + i; + pos[idx + n ] = pos_0 + i; + pos[idx + n * 2] = pos_0 + i; + pos[idx + n * 3] = pos_0 + i; } } - llama_batch get_view(int offset, int n_tokens) { - GGML_ASSERT(offset >= 0 && n_tokens > 0 && offset + n_tokens <= batch.n_tokens); - llama_pos * pos_ptr; + // describe the entries [offset, offset + n) with section-major positions + mtmd_helper_embd_batch get_view(int offset, int n) { + GGML_ASSERT(offset >= 0 && n > 0 && offset + n <= n_tokens); pos_view.clear(); - pos_view.reserve((size_t) n_tokens * (size_t) n_pos_per_embd); - if (n_pos_per_embd > 1) { - // mrope - // for example, with layout of src: 1234...1234...1234...1234... - // offset 2 will give us dst: 34...34...34...34... - for (int i = 0; i < n_pos_per_embd; i++) { - // assume n_tokens is less than or equal to batch.n_tokens - // batch.n_tokens is number of **total** tokens - // n_tokens is number of viewed token - size_t src_idx = (size_t) i * (size_t) batch.n_tokens + (size_t) offset; - pos_view.insert(pos_view.end(), - pos.data() + src_idx, - pos.data() + src_idx + n_tokens); - } - pos_ptr = pos_view.data(); - } else { - // normal - pos_ptr = pos.data() + offset; + pos_view.reserve((size_t) n * (size_t) n_pos_per_embd); + for (int j = 0; j < n_pos_per_embd; j++) { + const size_t src = (size_t) j * (size_t) n_tokens + (size_t) offset; + pos_view.insert(pos_view.end(), pos.data() + src, pos.data() + src + n); } return { - /*n_tokens =*/ n_tokens, - /*tokens =*/ nullptr, - /*embd =*/ batch.embd + offset * n_mmproj_embd, - /*pos =*/ pos_ptr, - /*n_seq_id =*/ batch.n_seq_id + offset, - /*seq_id =*/ batch.seq_id + offset, - /*logits =*/ batch.logits + offset, + /*n_tokens =*/ n, + /*embd =*/ embd + (size_t) offset * n_mmproj_embd, + /*n_embd =*/ n_mmproj_embd, + /*pos =*/ pos_view.data(), + /*n_pos =*/ n_pos_per_embd, + /*seq_id =*/ seq_id, }; } + + // render the entries [offset, offset + n) into a batch owned by this object, ready for llama_process() + llama_batch_ext * render(llama_context * lctx, int offset, int n) { + GGML_ASSERT(offset >= 0 && n > 0 && offset + n <= n_tokens); + if (!batch) { + batch.reset(llama_batch_ext_init(lctx)); + } + llama_batch_ext_clear(batch.get()); + for (int i = offset; i < offset + n; i++) { + const llama_embd e = { embd + (size_t) i * n_mmproj_embd, 1, (size_t) n_mmproj_embd }; + const int32_t idx = llama_batch_ext_add_embd(batch.get(), seq_id, e); + GGML_ASSERT(idx >= 0); + + llama_pos p[GGML_MROPE_SECTIONS] = { 0, 0, 0, 0 }; + for (int j = 0; j < n_pos_per_embd; j++) { + p[j] = pos[(size_t) j * (size_t) n_tokens + (size_t) i]; + } + llama_batch_ext_set_pos(batch.get(), idx, p); + + if (logits[i]) { + llama_batch_ext_set_output_logits(batch.get(), idx, true); + } + } + return batch.get(); + } }; diff --git a/tools/mtmd/mtmd-helper-gen.cpp b/tools/mtmd/mtmd-helper-gen.cpp index 1c58d3ae1959..5fb7ea9a0de8 100644 --- a/tools/mtmd/mtmd-helper-gen.cpp +++ b/tools/mtmd/mtmd-helper-gen.cpp @@ -222,14 +222,13 @@ class qwen3tts_gen_audio_pipeline : public mtmd_gen_audio_pipeline { return 0; } const int32_t n_tokens_batch = std::min(n_batch, n_prompt - prompt_pos); - llama_batch batch_view = prompt_batch->get_view(prompt_pos, n_tokens_batch); const bool is_last_batch = (prompt_pos + n_tokens_batch) == n_prompt; if (is_last_batch) { - batch_view.logits[n_tokens_batch - 1] = 1; + prompt_batch->logits[prompt_pos + n_tokens_batch - 1] = 1; } - if (llama_decode(lctx, batch_view) != 0) { + if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, prompt_batch->render(lctx, prompt_pos, n_tokens_batch)) != 0) { LOG_ERR("mtmd_helper_gen_audio: prompt decode failed\n"); return -1; } @@ -286,10 +285,10 @@ class qwen3tts_gen_audio_pipeline : public mtmd_gen_audio_pipeline { decode_embd_batch batch_embd(fb.data(), 1, n_pos_per_embd, n_embd); if (mrope) batch_embd.set_position_mrope_1d(pos, seq_id); else batch_embd.set_position_normal (pos, seq_id); - batch_embd.batch.logits[0] = 1; + batch_embd.logits[0] = 1; pos++; - if (llama_decode(lctx, batch_embd.batch) != 0) { + if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch_embd.render(lctx, 0, 1)) != 0) { LOG_ERR("mtmd_helper_gen_audio: decode failed\n"); return 1; } @@ -586,13 +585,12 @@ class pockettts_gen_audio_pipeline : public mtmd_gen_audio_pipeline { return 0; } const int32_t n_tokens_batch = std::min(n_batch, n_prompt - prompt_pos); - llama_batch batch_view = prompt_batch->get_view(prompt_pos, n_tokens_batch); if ((prompt_pos + n_tokens_batch) == n_prompt) { - batch_view.logits[n_tokens_batch - 1] = 1; + prompt_batch->logits[prompt_pos + n_tokens_batch - 1] = 1; } - if (llama_decode(lctx, batch_view) != 0) { + if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, prompt_batch->render(lctx, prompt_pos, n_tokens_batch)) != 0) { LOG_ERR("mtmd_helper_gen_audio: prompt decode failed\n"); return -1; } @@ -646,12 +644,12 @@ class pockettts_gen_audio_pipeline : public mtmd_gen_audio_pipeline { } } - decode_embd_batch batch_embd(const_cast(out.embd), 1, 1, n_embd); + decode_embd_batch batch_embd(out.embd, 1, 1, n_embd); batch_embd.set_position_normal(pos, seq_id); - batch_embd.batch.logits[0] = 1; + batch_embd.logits[0] = 1; pos++; - if (llama_decode(lctx, batch_embd.batch) != 0) { + if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch_embd.render(lctx, 0, 1)) != 0) { LOG_ERR("mtmd_helper_gen_audio: decode failed\n"); return 1; } @@ -842,8 +840,8 @@ class pockettts_gen_audio_pipeline : public mtmd_gen_audio_pipeline { GGML_ASSERT(n_rows > 0); decode_embd_batch batch(prompt_embd_buf.data(), n_rows, 1, n_e); batch.set_position_normal(pos, seq_id); - batch.batch.logits[n_rows - 1] = 1; - if (llama_decode(lctx, batch.batch) != 0) { + batch.logits[n_rows - 1] = 1; + if (llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch.render(lctx, 0, n_rows)) != 0) { LOG_ERR("mtmd_helper_gen_audio: chunk prompt decode failed\n"); return 1; } diff --git a/tools/mtmd/mtmd-helper.cpp b/tools/mtmd/mtmd-helper.cpp index bdf8bf6fe450..cd05ade2d7ac 100644 --- a/tools/mtmd/mtmd-helper.cpp +++ b/tools/mtmd/mtmd-helper.cpp @@ -169,19 +169,19 @@ int32_t mtmd_helper_decode_image_chunk( while (i_batch < n_img_batches) { // split into batches int pos_offset = i_batch*n_batch; int n_tokens_batch = std::min(n_batch, n_tokens - pos_offset); - llama_batch batch_embd_view = batch_embd.get_view(pos_offset, n_tokens_batch); LOG_INF("decoding %s batch %d/%d, n_tokens_batch = %d\n", name, i_batch+1, n_img_batches, n_tokens_batch); int64_t t1 = ggml_time_ms(); - int32_t ret = llama_decode(lctx, batch_embd_view); + int32_t ret = llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, batch_embd.render(lctx, pos_offset, n_tokens_batch)); if (ret != 0) { LOG_ERR("failed to decode %s\n", name); return ret; } if (callback != nullptr) { - ret = callback(batch_embd_view, user_data); + const mtmd_helper_embd_batch view = batch_embd.get_view(pos_offset, n_tokens_batch); + ret = callback(&view, user_data); if (ret != 0) { LOG_ERR("post-decode callback failed\n"); return ret; @@ -209,37 +209,35 @@ int32_t mtmd_helper_eval_chunk_single(mtmd_context * ctx, llama_pos * new_n_past) { GGML_ASSERT(n_batch > 0); int32_t ret; - llama_batch text_batch = llama_batch_init(n_batch, 0, 1); auto chunk_type = mtmd_input_chunk_get_type(chunk); if (chunk_type == MTMD_INPUT_CHUNK_TYPE_TEXT) { size_t n_tokens; const auto tokens = mtmd_input_chunk_get_tokens_text(chunk, &n_tokens); // LOG_INF("decoding text chunk, n_tokens = %zu\n", n_tokens); + llama_batch_ext_ptr text_batch(llama_batch_ext_init(lctx)); size_t i = 0; while (i < n_tokens) { // split into batches - text_batch.n_tokens = 0; // clear the batch - for (; i < n_tokens && text_batch.n_tokens < n_batch; i++) { - int32_t j = text_batch.n_tokens; - text_batch.token [j] = tokens[i]; - text_batch.pos [j] = n_past++; - text_batch.n_seq_id[j] = 1; - text_batch.seq_id [j][0] = seq_id; - text_batch.logits [j] = false; - - text_batch.n_tokens++; + llama_batch_ext_clear(text_batch.get()); + int32_t n_added = 0; + int32_t idx = -1; + for (; i < n_tokens && n_added < n_batch; i++) { + idx = llama_batch_ext_add_token(text_batch.get(), seq_id, tokens[i]); + GGML_ASSERT(idx >= 0); + llama_pos pos = n_past++; + llama_batch_ext_set_pos(text_batch.get(), idx, &pos); + n_added++; } bool is_last_token = (i == n_tokens); if (logits_last && is_last_token) { - text_batch.logits[text_batch.n_tokens - 1] = true; + llama_batch_ext_set_output_logits(text_batch.get(), idx, true); } - ret = llama_decode(lctx, text_batch); + ret = llama_process(lctx, LLAMA_PROCESS_TYPE_DECODE, text_batch.get()); if (ret != 0) { LOG_ERR("failed to decode text\n"); - llama_batch_free(text_batch); return ret; } - *new_n_past += text_batch.n_tokens; + *new_n_past += n_added; } } else if (chunk_type == MTMD_INPUT_CHUNK_TYPE_IMAGE || chunk_type == MTMD_INPUT_CHUNK_TYPE_AUDIO) { @@ -251,7 +249,6 @@ int32_t mtmd_helper_eval_chunk_single(mtmd_context * ctx, ret = mtmd_encode_chunk(ctx, chunk); if (ret != 0) { LOG_ERR("failed to encode %s slice\n", name); - llama_batch_free(text_batch); return ret; } @@ -261,14 +258,12 @@ int32_t mtmd_helper_eval_chunk_single(mtmd_context * ctx, ret = mtmd_helper_decode_image_chunk(ctx, lctx, chunk, embd, n_past, seq_id, n_batch, new_n_past, nullptr, nullptr); if (ret != 0) { LOG_ERR("failed to decode %s\n", name); - llama_batch_free(text_batch); return ret; } } else { GGML_ABORT("chunk type not supported"); } - llama_batch_free(text_batch); return 0; } diff --git a/tools/mtmd/mtmd-helper.h b/tools/mtmd/mtmd-helper.h index 10f2171c0fdd..7436230f0dca 100644 --- a/tools/mtmd/mtmd-helper.h +++ b/tools/mtmd/mtmd-helper.h @@ -92,9 +92,9 @@ MTMD_API llama_pos mtmd_helper_get_n_pos(const mtmd_input_chunks * chunks); MTMD_API void mtmd_helper_image_get_decoder_pos(const mtmd_image_tokens * image, llama_pos pos_0, struct mtmd_decoder_pos * out_pos); // helper function that automatically: -// 1. run llama_decode() on text chunks -// 2. run mtmd_encode_chunk() on image chunks, then mtmd_get_output_embd() and then llama_decode() -// if any of the mtmd_encode_chunk() or llama_decode() calls return non-zero, stop and forward the error +// 1. decode text chunks +// 2. run mtmd_encode_chunk() on image chunks, then mtmd_get_output_embd() and then decode the embeddings +// if any of the mtmd_encode_chunk() or decode calls return non-zero, stop and forward the error // otherwise, returns 0 on success // this function is NOT thread-safe MTMD_API int32_t mtmd_helper_eval_chunks(mtmd_context * ctx, @@ -117,7 +117,17 @@ MTMD_API int32_t mtmd_helper_eval_chunk_single(mtmd_context * ctx, bool logits_last, llama_pos * new_n_past); -typedef int32_t (*mtmd_helper_post_decode_callback)(struct llama_batch batch, void * user_data); +// one decoded sub-batch of embeddings, passed to mtmd_helper_post_decode_callback +struct mtmd_helper_embd_batch { + int32_t n_tokens; + const float * embd; // [n_tokens, n_embd] + int32_t n_embd; + const llama_pos * pos; // [n_pos, n_tokens], section-major + int32_t n_pos; // 4 for M-RoPE models, 1 otherwise + llama_seq_id seq_id; +}; + +typedef int32_t (*mtmd_helper_post_decode_callback)(const struct mtmd_helper_embd_batch * batch, void * user_data); // helper function to decode an image whose embeddings have already been calculated // this helper will handle batching and pre/post decoding setup (for ex. gemma 3 requires non-causal attention) diff --git a/tools/rpc/rpc-server.cpp b/tools/rpc/rpc-server.cpp index 100df3c40900..5c30e6fe4716 100644 --- a/tools/rpc/rpc-server.cpp +++ b/tools/rpc/rpc-server.cpp @@ -2,12 +2,10 @@ #include "ggml-rpc.h" #ifdef _WIN32 # define NOMINMAX -# define DIRECTORY_SEPARATOR '\\' # include # include # include #else -# define DIRECTORY_SEPARATOR '/' # include #endif #include @@ -20,64 +18,68 @@ #include #include -#if defined(__linux__) +#if !defined(_WIN32) && !defined(__APPLE__) #include #include #endif -// NOTE: this is copied from common.cpp to avoid linking with libcommon -// returns true if successful, false otherwise -static bool fs_create_directory_with_parents(const std::string & path) { - std::error_code ec; - std::filesystem::create_directories(std::filesystem::u8path(path), ec); - return !ec; + +static std::string fs_path_to_utf8(const std::filesystem::path & path) { + const auto value = path.u8string(); + return std::string(value.begin(), value.end()); +} + +// common_get_path_from_env() is adapted to avoid utf8_to_wstring +static std::filesystem::path common_get_path_from_env(const std::string & name) { +#ifdef _WIN32 + std::wstring wname; + for (const char * p = name.c_str(); *p; ++p) { + wname.push_back((wchar_t)*p); + } + const wchar_t * wvalue = _wgetenv(wname.c_str()); + return wvalue ? std::filesystem::path(wvalue) : std::filesystem::path(); +#else + const char * value = std::getenv(name.c_str()); + return value ? std::filesystem::path(value) : std::filesystem::path(); +#endif } // NOTE: this is copied from common.cpp to avoid linking with libcommon -static std::string fs_get_cache_directory() { - std::string cache_directory = ""; - auto ensure_trailing_slash = [](std::string p) { - // Make sure to add trailing slash - if (p.back() != DIRECTORY_SEPARATOR) { - p += DIRECTORY_SEPARATOR; - } - return p; - }; - if (getenv("LLAMA_CACHE")) { - cache_directory = std::getenv("LLAMA_CACHE"); - } else { -#if defined(__linux__) || defined(__FreeBSD__) || defined(_AIX) || \ - defined(__OpenBSD__) || defined(__NetBSD__) - if (std::getenv("XDG_CACHE_HOME")) { - cache_directory = std::getenv("XDG_CACHE_HOME"); - } else if (std::getenv("HOME")) { - cache_directory = std::getenv("HOME") + std::string("/.cache/"); +static std::filesystem::path fs_get_cache_directory() { + std::filesystem::path cache_directory = common_get_path_from_env("LLAMA_CACHE"); + if (!cache_directory.empty()) { + return cache_directory; + } + +#if defined(_WIN32) + cache_directory = common_get_path_from_env("LOCALAPPDATA"); + if (cache_directory.empty()) { + throw std::runtime_error("Failed to find %LOCALAPPDATA% directory"); + } +#elif defined(__APPLE__) + cache_directory = common_get_path_from_env("HOME"); + if (cache_directory.empty()) { + throw std::runtime_error("Failed to find $HOME directory"); + } + cache_directory /= "Library/Caches"; +#else + cache_directory = common_get_path_from_env("XDG_CACHE_HOME"); + if (cache_directory.empty()) { + cache_directory = common_get_path_from_env("HOME"); + if (!cache_directory.empty()) { + cache_directory /= ".cache"; } else { -#if defined(__linux__) /* no $HOME is defined, fallback to getpwuid */ - struct passwd *pw = getpwuid(getuid()); - if ((!pw) || (!pw->pw_dir)) { + const struct passwd * pw = getpwuid(getuid()); + if (!pw || !pw->pw_dir || !*pw->pw_dir) { throw std::runtime_error("Failed to find $HOME directory"); } - - cache_directory = std::string(pw->pw_dir) + std::string("/.cache/"); -#else /* defined(__linux__) */ - throw std::runtime_error("Failed to find $HOME directory"); -#endif /* defined(__linux__) */ + cache_directory = pw->pw_dir; + cache_directory /= ".cache"; } -#elif defined(__APPLE__) - cache_directory = std::getenv("HOME") + std::string("/Library/Caches/"); -#elif defined(_WIN32) - cache_directory = std::getenv("LOCALAPPDATA"); -#elif defined(__EMSCRIPTEN__) - GGML_ABORT("not implemented on this platform"); -#else -# error Unknown architecture -#endif - cache_directory = ensure_trailing_slash(cache_directory); - cache_directory += "llama.cpp"; } - return ensure_trailing_slash(cache_directory); +#endif + return cache_directory / "llama.cpp"; } struct rpc_server_params { @@ -229,11 +231,14 @@ int main(int argc, char * argv[]) { const char * cache_dir = nullptr; std::string cache_dir_str; if (params.use_cache) { - cache_dir_str = fs_get_cache_directory() + "rpc" + DIRECTORY_SEPARATOR; - if (!fs_create_directory_with_parents(cache_dir_str)) { - fprintf(stderr, "Failed to create cache directory: %s\n", cache_dir_str.c_str()); + const std::filesystem::path cache_dir_path = fs_get_cache_directory() / "rpc"; + std::error_code ec; + std::filesystem::create_directories(cache_dir_path, ec); + if (ec) { + fprintf(stderr, "Failed to create cache directory: %s\n", fs_path_to_utf8(cache_dir_path).c_str()); return 1; } + cache_dir_str = fs_path_to_utf8(cache_dir_path); cache_dir = cache_dir_str.c_str(); } diff --git a/tools/server/README.md b/tools/server/README.md index e665904ba2c1..09fdb87e5daf 100644 --- a/tools/server/README.md +++ b/tools/server/README.md @@ -1509,6 +1509,13 @@ This endpoint requires that the model uses a pooling different than type `none`. See [OpenAI Embeddings API documentation](https://platform.openai.com/docs/api-reference/embeddings). +For multimodal models (loaded with `--mmproj`), each element of `input` can also be an object with a `content` array, using the same parts as `/v1/chat/completions`: +- `{ "type": "text", "text": "..." }`: text is added to the prompt as-is +- `{ "type": "image_url", "image_url": { "url": "..." } }`: remote URL, base64 data URI, or local file (`file://`, requires `--media-path`) +- `{ "type": "input_audio", "input_audio": { "data": "..." } }` and `{ "type": "input_video", "input_video": { "url": "..." } }`: same as `/v1/chat/completions`, requires a model with audio or video support + +Each object gives one embedding. This input shape is not part of the OpenAI Embeddings API; it follows the shape used by providers like OpenRouter for vision embedding models. + *Examples:* - input as string @@ -1537,6 +1544,26 @@ See [OpenAI Embeddings API documentation](https://platform.openai.com/docs/api-r }' ``` +- `input` as multimodal content + + ```shell + curl http://localhost:8080/v1/embeddings \ + -H "Content-Type: application/json" \ + -H "Authorization: Bearer no-key" \ + -d '{ + "input": [ + { "content": [ + { "type": "image_url", "image_url": { "url": "data:image/jpeg;base64,/9j/4AAQSkZJRg..." } }, + { "type": "text", "text": "Describe this image" } + ] }, + { "content": [ + { "type": "text", "text": "hello" } + ] } + ], + "encoding_format": "float" + }' + ``` + ### POST `/v1/responses/input_tokens`: Token Counting Similar to [Response input token counts API](https://developers.openai.com/api/reference/python/resources/responses/subresources/input_tokens/methods/count). diff --git a/tools/server/server-common.cpp b/tools/server/server-common.cpp index 2ff790a5de31..8cf794f1e24f 100644 --- a/tools/server/server-common.cpp +++ b/tools/server/server-common.cpp @@ -970,7 +970,7 @@ server_tokens process_mtmd_prompt( } /** - * break the input "prompt" object into multiple prompt if needed, then tokenize them + * tokenize a single input "prompt" object * use tokenize_input_prompts() if the input could be an array. * this supports these cases: * - "prompt": "string" @@ -978,7 +978,7 @@ server_tokens process_mtmd_prompt( * - "prompt": [12, 34, "string", 56, 78] * - "prompt": { "prompt_string": "string", "multimodal_data": [ "base64" ] } */ -static server_tokens tokenize_input_subprompt(const llama_vocab * vocab, mtmd_context * mctx, const json & json_prompt, bool add_special, bool parse_special, const mtmd_helper_init_opt & init_opt) { +server_tokens tokenize_input_subprompt(const llama_vocab * vocab, mtmd_context * mctx, const json & json_prompt, bool add_special, bool parse_special, const mtmd_helper_init_opt & init_opt) { constexpr char JSON_STRING_PROMPT_KEY[] = "prompt_string"; constexpr char JSON_MTMD_DATA_KEY[] = "multimodal_data"; const bool has_mtmd = mctx != nullptr; @@ -1147,6 +1147,79 @@ static void handle_media( } } +// load media files from an OAI content array, then replace each media part with a media marker text part +static void oaicompat_content_load_media(json & content, const server_chat_params & opt, std::vector & out_files) { + for (auto & p : content) { + std::string type = json_value(p, "type", std::string()); + if (type == "image_url") { + if (!opt.allow_image) { + throw std::runtime_error("image input is not supported - hint: if this is unexpected, you may need to provide the mmproj"); + } + + json image_url = json_value(p, "image_url", json::object()); + std::string url = json_value(image_url, "url", std::string()); + handle_media(out_files, url, opt.media_path); + + p["type"] = "media_marker"; + p["text"] = get_media_marker(); + p.erase("image_url"); + + } else if (type == "input_audio") { + if (!opt.allow_audio) { + throw std::runtime_error("audio input is not supported - hint: if this is unexpected, you may need to provide the mmproj"); + } + + // note: don't need to validate "format", it's redundant + json input_audio = json_value(p, "input_audio", json::object()); + std::string url = json_value(input_audio, "data", + json_value(input_audio, "url", std::string())); + handle_media(out_files, url, opt.media_path); + + p["type"] = "media_marker"; + p["text"] = get_media_marker(); + p.erase("input_audio"); + + } else if (type == "input_video" || type == "video_url") { + if (!opt.allow_video) { + throw std::runtime_error("video input is not supported - hint: if this is unexpected, you may need to provide the mmproj"); + } + + // accept the OpenAI-style "video_url" key as an alias of "input_video" + json input_video = json_value(p, type, json::object()); + std::string url = json_value(input_video, "data", + json_value(input_video, "url", std::string())); + handle_media(out_files, url, opt.media_path); + + p["type"] = "media_marker"; + p["text"] = get_media_marker(); + p.erase("input_video"); + p.erase("video_url"); + + } else if (type != "text") { + throw std::invalid_argument("unsupported content[].type"); + } + } +} + +server_tokens tokenize_oai_content_array(const llama_vocab * vocab, mtmd_context * mctx, const server_chat_params & opt, json content, bool add_special, bool parse_special, const mtmd_helper_init_opt & init_opt) { + if (!content.is_array()) { + throw std::invalid_argument("\"content\" must be an array"); + } + + std::vector files; + oaicompat_content_load_media(content, opt, files); + + std::string prompt; + for (const auto & p : content) { + prompt += json_value(p, "text", std::string()); + } + + if (files.empty()) { + return server_tokens(common_tokenize(vocab, prompt, add_special, parse_special), false); + } + return process_mtmd_prompt(mctx, prompt, files, init_opt); +} + // used by /chat/completions endpoint json oaicompat_chat_params_parse( json & body, /* openai api json semantics */ @@ -1233,56 +1306,7 @@ json oaicompat_chat_params_parse( throw std::invalid_argument("Expected 'content' to be a string or an array"); } - for (auto & p : content) { - std::string type = json_value(p, "type", std::string()); - if (type == "image_url") { - if (!opt.allow_image) { - throw std::runtime_error("image input is not supported - hint: if this is unexpected, you may need to provide the mmproj"); - } - - json image_url = json_value(p, "image_url", json::object()); - std::string url = json_value(image_url, "url", std::string()); - handle_media(out_files, url, opt.media_path); - - p["type"] = "media_marker"; - p["text"] = get_media_marker(); - p.erase("image_url"); - - } else if (type == "input_audio") { - if (!opt.allow_audio) { - throw std::runtime_error("audio input is not supported - hint: if this is unexpected, you may need to provide the mmproj"); - } - - // note: don't need to validate "format", it's redundant - json input_audio = json_value(p, "input_audio", json::object()); - std::string url = json_value(input_audio, "data", - json_value(input_audio, "url", std::string())); - handle_media(out_files, url, opt.media_path); - - p["type"] = "media_marker"; - p["text"] = get_media_marker(); - p.erase("input_audio"); - - } else if (type == "input_video" || type == "video_url") { - if (!opt.allow_video) { - throw std::runtime_error("video input is not supported - hint: if this is unexpected, you may need to provide the mmproj"); - } - - // accept the OpenAI-style "video_url" key as an alias of "input_video" - json input_video = json_value(p, type, json::object()); - std::string url = json_value(input_video, "data", - json_value(input_video, "url", std::string())); - handle_media(out_files, url, opt.media_path); - - p["type"] = "media_marker"; - p["text"] = get_media_marker(); - p.erase("input_video"); - p.erase("video_url"); - - } else if (type != "text") { - throw std::invalid_argument("unsupported content[].type"); - } - } + oaicompat_content_load_media(content, opt, out_files); } auto caps = common_chat_templates_get_caps(opt.tmpls.get()); diff --git a/tools/server/server-common.h b/tools/server/server-common.h index e00408cd4f9e..8cb6b90da62f 100644 --- a/tools/server/server-common.h +++ b/tools/server/server-common.h @@ -300,6 +300,15 @@ std::vector tokenize_input_prompts( bool parse_special, const mtmd_helper_init_opt & init_opt); +// tokenize a single prompt, see tokenize_input_prompts() for the supported shapes +server_tokens tokenize_input_subprompt( + const llama_vocab * vocab, + mtmd_context * mctx, + const json & json_prompt, + bool add_special, + bool parse_special, + const mtmd_helper_init_opt & init_opt); + // // OAI utils // @@ -330,6 +339,16 @@ json oaicompat_chat_params_parse( const server_chat_params & opt, std::vector & out_files); +// used by /embeddings endpoint, content has the same format as a chat message content array +server_tokens tokenize_oai_content_array( + const llama_vocab * vocab, + mtmd_context * mctx, + const server_chat_params & opt, + json content, + bool add_special, + bool parse_special, + const mtmd_helper_init_opt & init_opt); + // TODO: move it to server-task.cpp json format_embeddings_response_oaicompat( const json & request, diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 611e82a6afdf..2036319c1381 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -109,8 +109,7 @@ enum slot_state { struct server_slot; // forward declaration struct server_batch { - llama_batch batch; - bool batch_rendered = false; + common_batch view; // the rendered sub-batch [off, off + n_tokens), see render() struct token { int32_t id_slot; @@ -126,36 +125,21 @@ struct server_batch { // track if given slot can be batched with slots already in the batch server_slot * slot_batched = nullptr; - // in embd mode, we temporarily swap out the tokens arr and restore it on clear() bool has_embd = false; - llama_token * tokens_ptr = nullptr; std::vector embd; float alora_scale = -1.0f; size_t alora_disabled_id = 0; - server_batch() { - batch.pos = nullptr; // sentinel: uninitialized batch - } - - ~server_batch() { - if (batch.pos != nullptr) { - clear(); - llama_batch_free(batch); - } - } - - void init(int32_t n_tokens_alloc, int32_t n_embd) { + void init(llama_context * ctx, int32_t n_tokens_alloc, int32_t n_embd) { this->n_tokens_alloc = n_tokens_alloc; this->n_embd = n_embd; - batch = llama_batch_init(n_tokens_alloc, 0, 1); - tokens_ptr = batch.token; + view = common_batch(ctx); tokens.reserve(n_tokens_alloc); } bool add(int32_t id_slot, llama_token token, llama_pos pos, bool output, bool is_prompt) { GGML_ASSERT(!has_embd); // cannot mix tokens + embd in same batch - GGML_ASSERT(batch.pos != nullptr); if ((int32_t)tokens.size() >= n_tokens_alloc) { return false; } @@ -164,7 +148,6 @@ struct server_batch { } bool add(int32_t id_slot, const std::vector & embd_in, llama_pos pos, bool output, bool is_prompt) { - GGML_ASSERT(batch.pos != nullptr); if ((int32_t)tokens.size() >= n_tokens_alloc) { return false; } @@ -177,16 +160,11 @@ struct server_batch { void clear() { tokens.clear(); embd.clear(); - common_batch_clear(batch); + view.clear(); slot_batched = nullptr; alora_scale = -1.0f; alora_disabled_id = 0; - batch_rendered = false; has_embd = false; - if (batch.token == nullptr) { - batch.token = tokens_ptr; - batch.embd = nullptr; - } } int32_t size() const { @@ -198,41 +176,22 @@ struct server_batch { tokens[idx].output = output; } - void render() { - GGML_ASSERT(!batch_rendered); - GGML_ASSERT(batch.pos != nullptr); - common_batch_clear(batch); - for (int32_t i = 0; i < size(); i++) { - const auto & t = tokens[i]; - common_batch_add(batch, t.token, t.pos, { t.id_slot }, t.output); - } - if (has_embd) { - batch.token = nullptr; // will be restored on clear() - batch.embd = embd.data(); - } - batch_rendered = true; - } - - llama_batch get_view(int32_t off, int32_t n_tokens) const { - GGML_ASSERT(batch.pos != nullptr); - GGML_ASSERT(batch_rendered); + // render the sub-batch [off, off + n_tokens) into view, index i in view is index off + i here + void render(int32_t off, int32_t n_tokens) { GGML_ASSERT(off >= 0 && off < size()); GGML_ASSERT(n_tokens > 0 && off + n_tokens <= size()); - auto * token = batch.token ? batch.token + off : nullptr; - auto * embd = batch.embd ? batch.embd + off * n_embd : nullptr; - - llama_batch view = { - n_tokens, - token, - embd, - batch.pos + off, - batch.n_seq_id + off, - batch.seq_id + off, - batch.logits + off, - }; - - return view; + view.clear(); + for (int32_t i = off; i < off + n_tokens; i++) { + const auto & t = tokens[i]; + if (has_embd) { + // text embeddings broadcast the same position across the M-RoPE sections + const llama_pos pos[GGML_MROPE_SECTIONS] = { t.pos, t.pos, t.pos, 0 }; + view.add_embd({ embd.data() + (size_t) i * n_embd, 1, (size_t) n_embd }, pos, t.id_slot, t.output); + } else { + view.add(t.token, t.pos, t.id_slot, t.output); + } + } } }; @@ -761,13 +720,24 @@ static int process_mtmd_chunk(const server_slot & slot, mtmd::batch_ptr & mbatch if (mbatch) { float * embd = mtmd_batch_get_output_embd(mbatch.get(), chunk.get()); if (embd) { - void * cb_data = slot.spec; - static auto cb = [](llama_batch batch, void * user_data) { - common_speculative * spec = static_cast(user_data); - if (!common_speculative_process(spec, batch)) { - return 1; + struct cb_data_t { + common_speculative * spec; + llama_context * ctx; + } cb_data = { slot.spec, slot.ctx_tgt }; + + static auto cb = [](const mtmd_helper_embd_batch * b, void * user_data) { + const auto * data = static_cast(user_data); + + common_batch batch(data->ctx); + for (int32_t i = 0; i < b->n_tokens; ++i) { + llama_pos pos[GGML_MROPE_SECTIONS] = { 0, 0, 0, 0 }; + for (int32_t j = 0; j < b->n_pos; ++j) { + pos[j] = b->pos[j * b->n_tokens + i]; + } + batch.add_embd({ b->embd + (size_t) i * b->n_embd, 1, (size_t) b->n_embd }, pos, b->seq_id, false); } - return 0; + + return common_speculative_process(data->spec, batch) ? 0 : 1; }; llama_pos new_n_past; // unused for now @@ -781,7 +751,7 @@ static int process_mtmd_chunk(const server_slot & slot, mtmd::batch_ptr & mbatch llama_n_batch(slot.ctx_tgt), &new_n_past, cb, - cb_data + &cb_data ); if (res != 0) { SLT_ERR(slot, "failed to decode mtmd chunk, idx = %zu, res = %d\n", idx, res); @@ -1356,7 +1326,7 @@ struct server_context_impl { { const int32_t n_batch = llama_n_batch(ctx_tgt); const int32_t n_embd = llama_model_n_embd_inp(model_tgt); - batch.init(std::max(n_batch, params_base.n_parallel), n_embd); + batch.init(ctx_tgt, std::max(n_batch, params_base.n_parallel), n_embd); } if (params_base.cache_ram_mib != 0) { @@ -2160,7 +2130,7 @@ struct server_context_impl { queue_results.send(std::move(res)); } - void send_embedding(const server_slot & slot, const llama_batch & batch) { + void send_embedding(const server_slot & slot, const common_batch & batch) { auto res = std::make_unique(); res->id = slot.task->id; res->index = slot.task->index; @@ -2171,8 +2141,8 @@ struct server_context_impl { std::vector embd_res(n_embd_out, 0.0f); - for (int i = 0; i < batch.n_tokens; ++i) { - if (!batch.logits[i] || batch.seq_id[i][0] != slot.id) { + for (int i = 0; i < batch.size(); ++i) { + if (!batch.tokens[i].output || batch.tokens[i].seq_id != slot.id) { continue; } @@ -2180,11 +2150,11 @@ struct server_context_impl { if (llama_pooling_type(slot.ctx_tgt) == LLAMA_POOLING_TYPE_NONE) { embd = llama_get_embeddings_ith(slot.ctx_tgt, i); } else { - embd = llama_get_embeddings_seq(slot.ctx_tgt, batch.seq_id[i][0]); + embd = llama_get_embeddings_seq(slot.ctx_tgt, batch.tokens[i].seq_id); } if (embd == nullptr) { - SLT_ERR(slot, "failed to get embeddings, token = %d, seq_id = %d\n", batch.token[i], batch.seq_id[i][0]); + SLT_ERR(slot, "failed to get embeddings, token = %d, seq_id = %d\n", batch.tokens[i].id, batch.tokens[i].seq_id); res->embedding.push_back(std::vector(n_embd_out, 0.0f)); continue; @@ -2205,24 +2175,24 @@ struct server_context_impl { queue_results.send(std::move(res)); } - void send_rerank(const server_slot & slot, const llama_batch & batch) { + void send_rerank(const server_slot & slot, const common_batch & batch) { auto res = std::make_unique(); res->id = slot.task->id; res->index = slot.task->index; res->n_tokens = slot.task->n_tokens(); - for (int i = 0; i < batch.n_tokens; ++i) { - if (!batch.logits[i] || batch.seq_id[i][0] != slot.id) { + for (int i = 0; i < batch.size(); ++i) { + if (!batch.tokens[i].output || batch.tokens[i].seq_id != slot.id) { continue; } - const float * embd = llama_get_embeddings_seq(ctx_tgt, batch.seq_id[i][0]); + const float * embd = llama_get_embeddings_seq(ctx_tgt, batch.tokens[i].seq_id); if (embd == NULL) { embd = llama_get_embeddings_ith(ctx_tgt, i); } if (embd == NULL) { - SLT_ERR(slot, "failed to get embeddings, token = %d, seq_id = %d\n", batch.token[i], batch.seq_id[i][0]); + SLT_ERR(slot, "failed to get embeddings, token = %d, seq_id = %d\n", batch.tokens[i].id, batch.tokens[i].seq_id); res->score = -1e6; continue; @@ -2845,7 +2815,6 @@ struct server_context_impl { try { scoped_timer t(t_pre_decode, n_pre_decode); pre_decode(); - batch.render(); } catch (const std::exception & e) { SRV_ERR("pre_decode() failed: %s\n", e.what()); abort_all_slots("pre_decode() failed: " + std::string(e.what())); @@ -2875,7 +2844,6 @@ struct server_context_impl { llama_set_embeddings(ctx_tgt, slot_batched->need_embd()); } - llama_batch batch_view; int32_t off_next = 0; int32_t n_batch = llama_n_batch(ctx_tgt); for (int32_t off = 0; off < batch.size(); off = off_next) { @@ -2884,8 +2852,8 @@ struct server_context_impl { scoped_timer t(t_decode, n_decode); // TODO @ngxson : maybe handle n_batch == 1 here instead of inside decode() - batch_view = batch.get_view(off, n_tokens); - bool ok = decode(n_batch, off, batch_view); + batch.render(off, n_tokens); + bool ok = decode(n_batch, off); #ifdef DEBUG_TIMINGS llama_synchronize(ctx_tgt); #endif @@ -2908,7 +2876,7 @@ struct server_context_impl { try { scoped_timer t(t_post_decode, n_post_decode); - post_decode(n_tokens, off, batch_view); + post_decode(n_tokens, off); } catch (const std::exception & e) { SRV_ERR("post_decode() failed: %s\n", e.what()); abort_all_slots("post_decode() failed: " + std::string(e.what())); @@ -3225,7 +3193,9 @@ struct server_context_impl { return; } - if (slot.task->params.cache_prompt) { + const bool is_stateless_task = slot.task->type == SERVER_TASK_TYPE_EMBEDDING || slot.task->type == SERVER_TASK_TYPE_RERANK; + + if (slot.task->params.cache_prompt && !is_stateless_task) { // reuse any previously computed tokens that are common with the new prompt n_past = slot.prompt.tokens.get_common_prefix(input_tokens); @@ -3655,7 +3625,7 @@ struct server_context_impl { // returns true = success ; false = retry with smaller batch size // throw std::runtime_error on fatal error - bool decode(int32_t & n_batch, int32_t off, llama_batch & batch_view) { + bool decode(int32_t & n_batch, int32_t off) { SRV_DBG("n_batch (effective) = %d, off = %d\n", n_batch, off); metrics_pre_decode(); @@ -3682,7 +3652,7 @@ struct server_context_impl { } bool has_output = false; - for (int i = off; i < off + batch_view.n_tokens; ++i) { + for (int i = off; i < off + batch.view.size(); ++i) { has_output |= batch.tokens[i].output; } @@ -3690,7 +3660,7 @@ struct server_context_impl { // note: the sync is done here too, so that the wait is also covered by the yield int ret = 0; queue_tasks.yield_to_queue([&]() { - ret = llama_decode(ctx_tgt, batch_view); + ret = llama_process(ctx_tgt, LLAMA_PROCESS_TYPE_DECODE, batch.view.get()); if (ret == 0 && has_output) { llama_synchronize(ctx_tgt); } @@ -3746,7 +3716,7 @@ struct server_context_impl { return false; // retry with the updated n_batch } else { // success, apply batch metrics - metrics_post_decode(off, batch_view.n_tokens, has_output); + metrics_post_decode(off, batch.view.size(), has_output); } // TODO: avoid restoring the draft context and re-evaluating the drafted tokens when not needed [TAG_SPEC_AVOID_DRAFT_REEVAL] @@ -3755,7 +3725,7 @@ struct server_context_impl { if (spec) { bool ok = true; queue_tasks.yield_to_queue([&]() { - ok = common_speculative_process(spec.get(), batch_view); + ok = common_speculative_process(spec.get(), batch.view); }); if (!ok) { @@ -3792,8 +3762,8 @@ struct server_context_impl { return true; } - void post_decode(int32_t n_batch_tokens, int32_t off, llama_batch & batch_view) { - // for checking if a given batch index is inside batch_view + void post_decode(int32_t n_batch_tokens, int32_t off) { + // for checking if a given batch index is inside the current sub-batch auto is_inside_view = [&](int32_t idx) { return idx >= off && idx < off + n_batch_tokens; }; @@ -3829,14 +3799,14 @@ struct server_context_impl { if (slot.state == SLOT_STATE_DONE_PROMPT) { if (slot.task->type == SERVER_TASK_TYPE_EMBEDDING) { // prompt evaluated for embedding - send_embedding(slot, batch_view); + send_embedding(slot, batch.view); slot.release(); slot.i_batch = -1; return; } if (slot.task->type == SERVER_TASK_TYPE_RERANK) { - send_rerank(slot, batch_view); + send_rerank(slot, batch.view); slot.release(); slot.i_batch = -1; return; @@ -5435,7 +5405,27 @@ std::unique_ptr server_routes::handle_embeddings_impl(cons } } - auto tokenized_prompts = tokenize_input_prompts(ctx_server.vocab, ctx_server.mctx, prompt, true, true, ctx_server.init_opt); + // same shapes as tokenize_input_prompts(), plus OAI content: { "content": [ { "type": "text"|"image_url"|"input_audio"|"input_video", ... } ] } + auto tokenize_entry = [&](const json & p) { + if (p.is_object() && p.contains("content")) { + return tokenize_oai_content_array(ctx_server.vocab, ctx_server.mctx, meta->chat_params, p.at("content"), true, true, ctx_server.init_opt); + } + return tokenize_input_subprompt(ctx_server.vocab, ctx_server.mctx, p, true, true, ctx_server.init_opt); + }; + + std::vector tokenized_prompts; + if (prompt.is_array() && !json_is_array_and_contains_numbers(prompt)) { + for (const auto & p : prompt) { + tokenized_prompts.push_back(tokenize_entry(p)); + } + } else { + tokenized_prompts.push_back(tokenize_entry(prompt)); + } + if (tokenized_prompts.empty()) { + res->error(format_error_response("\"input\" must not be empty", ERROR_TYPE_INVALID_REQUEST)); + return res; + } + for (const auto & tokens : tokenized_prompts) { // this check is necessary for models that do not add BOS token to the input if (tokens.empty()) { diff --git a/tools/server/tests/tests.sh b/tools/server/tests/tests.sh index c7d5d7805993..fd83b7fd5491 100755 --- a/tools/server/tests/tests.sh +++ b/tools/server/tests/tests.sh @@ -8,13 +8,19 @@ set -eu WORKERS="${PYTEST_WORKERS:-4}" +if [ "${WORKERS}" -eq 1 ]; then + WORKER_FLAGS="" +else + WORKER_FLAGS="-n ${WORKERS} --dist=worksteal" +fi + if [ $# -lt 1 ] then if [[ "${SLOW_TESTS:-0}" == 1 ]]; then - pytest --durations=30 -v -x -n "${WORKERS}" --dist=worksteal + pytest --durations=30 -v -x ${WORKER_FLAGS} else - pytest --durations=30 -v -x -n "${WORKERS}" --dist=worksteal -m "not slow" + pytest --durations=30 -v -x ${WORKER_FLAGS} -m "not slow" fi else - pytest --durations=30 -n "${WORKERS}" --dist=worksteal "$@" + pytest --durations=30 ${WORKER_FLAGS} "$@" fi diff --git a/tools/server/tests/unit/test_embedding.py b/tools/server/tests/unit/test_embedding.py index 17ba09554bb7..c4a7d35fefb5 100644 --- a/tools/server/tests/unit/test_embedding.py +++ b/tools/server/tests/unit/test_embedding.py @@ -83,6 +83,10 @@ def test_embedding_multiple_with_fa(): (["string1", [12, 34, 56]], True), ([[12, 34, 56], [12, 34, 56]], True), ([[12, 34, 56], [12, "string", 34, 56]], True), + # object entries + ({"prompt_string": "string"}, False), + ({"content": [{"type": "text", "text": "string"}]}, False), + (["string1", {"prompt_string": "string2"}, {"content": [{"type": "text", "text": "string3"}]}], True), ] ) def test_embedding_mixed_input(input, is_multi_prompt: bool): @@ -101,6 +105,40 @@ def test_embedding_mixed_input(input, is_multi_prompt: bool): assert len(data[0]['embedding']) > 1 +def test_embedding_content_text_same_as_string(): + global server + server.pooling = 'last' + server.start() + res = server.make_request("POST", "/v1/embeddings", data={ + "input": [ + "hello world", + {"content": [{"type": "text", "text": "hello "}, {"type": "text", "text": "world"}]}, + ], + }) + assert res.status_code == 200 + data = res.body['data'] + assert data[0]['embedding'] == data[1]['embedding'] + + +@pytest.mark.parametrize( + "input", + [ + [], + {"content": "string"}, + {"content": [{"type": "unknown"}]}, + # model is not multimodal + {"content": [{"type": "image_url", "image_url": {"url": "data:image/png;base64,AAAA"}}]}, + {"content": [{"type": "input_audio", "input_audio": {"data": "AAAA", "format": "wav"}}]}, + {"content": [{"type": "input_video", "input_video": {"url": "data:video/mp4;base64,AAAA"}}]}, + ] +) +def test_embedding_invalid_input(input): + global server + server.start() + res = server.make_request("POST", "/v1/embeddings", data={"input": input}) + assert res.status_code != 200 + + def test_embedding_pooling_mean(): global server server.pooling = 'mean' diff --git a/tools/server/tests/unit/test_vision_api.py b/tools/server/tests/unit/test_vision_api.py index 3bf868e66207..dd04e2e2b4c2 100644 --- a/tools/server/tests/unit/test_vision_api.py +++ b/tools/server/tests/unit/test_vision_api.py @@ -179,3 +179,28 @@ def test_vision_embeddings(prompt, image_data, success): assert content[0]['embedding'] != content[2]['embedding'] else: assert res.status_code != 200 + + +def test_vision_embeddings_oai_content(): + global server + server.server_embeddings = True + server.pooling = 'mean' + server.n_batch = 512 + server.start() + res = server.make_request("POST", "/v1/embeddings", data={ + "input": [ + {"content": [ + {"type": "text", "text": "What is this: "}, + {"type": "image_url", "image_url": {"url": get_img_url("IMG_BASE64_URI_0")}}, + {"type": "text", "text": "\n"}, + ]}, + {JSON_PROMPT_STRING_KEY: "What is this: <__media__>\n", JSON_MULTIMODAL_KEY: [get_img_url("IMG_BASE64_0")]}, + "What is this: \n", + ], + }) + assert res.status_code == 200 + data = res.body["data"] + assert len(data) == 3 + # same prompt and image in both formats + assert data[0]["embedding"] == data[1]["embedding"] + assert data[0]["embedding"] != data[2]["embedding"]